forked from lijiext/lammps
Getting rid of extra CR characters at ends of lines.
git-svn-id: svn://svn.icms.temple.edu/lammps-ro/trunk@5285 f3b2605a-c512-4ea7-a41b-209d697bcdaa
This commit is contained in:
pscrozi
parent
80a30880f8
commit
4366bfffd3
|
|
@ -10,11 +10,11 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
#include "atomic_gpu_memory.h"
|
||||
#define AtomicGPUMemoryT AtomicGPUMemory<numtyp, acctyp>
|
||||
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -1,337 +1,337 @@
|
|||
// -------------------------------------------------------------
|
||||
// CUDPP -- CUDA Data Parallel Primitives library
|
||||
// -------------------------------------------------------------
|
||||
// $Revision$
|
||||
// $Date$
|
||||
// -------------------------------------------------------------
|
||||
// This source code is distributed under the terms of license.txt
|
||||
// in the root directory of this source distribution.
|
||||
// -------------------------------------------------------------
|
||||
#include <cudpp_globals.h>
|
||||
#include "cudpp_radixsort.h"
|
||||
#include "cta/scan_cta.cu"
|
||||
#include <cudpp.h>
|
||||
#include <stdio.h>
|
||||
|
||||
#include <cudpp_util.h>
|
||||
#include <math.h>
|
||||
#include "sharedmem.h"
|
||||
|
||||
|
||||
#ifdef __DEVICE_EMULATION__
|
||||
#define __EMUSYNC __syncthreads()
|
||||
#else
|
||||
#define __EMUSYNC
|
||||
#endif
|
||||
|
||||
/**
|
||||
* @file
|
||||
* sort_cta.cu
|
||||
*
|
||||
* @brief CUDPP CTA-level sort routines
|
||||
*/
|
||||
|
||||
/** \addtogroup cudpp_cta
|
||||
* @{
|
||||
*/
|
||||
|
||||
/** @name Radix Sort Functions
|
||||
* @{
|
||||
*/
|
||||
|
||||
|
||||
typedef unsigned int uint;
|
||||
|
||||
/**
|
||||
* @brief Flips bits of single-precision floating-point number (parameterized by doFlip)
|
||||
*
|
||||
* flip a float for sorting
|
||||
* finds SIGN of fp number.
|
||||
* if it's 1 (negative float), it flips all bits
|
||||
* if it's 0 (positive float), it flips the sign only
|
||||
* @param[in] f floating-point input (passed as unsigned int)
|
||||
* @see floatUnflip
|
||||
**/
|
||||
|
||||
template <bool doFlip>
|
||||
__device__ uint floatFlip(uint f)
|
||||
{
|
||||
if (doFlip)
|
||||
{
|
||||
uint mask = -int(f >> 31) | 0x80000000;
|
||||
return f ^ mask;
|
||||
}
|
||||
else
|
||||
return f;
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Reverses bit-flip of single-precision floating-point number (parameterized by doFlip)
|
||||
*
|
||||
* flip a float back (invert FloatFlip)
|
||||
* signed was flipped from above, so:
|
||||
* if sign is 1 (negative), it flips the sign bit back
|
||||
* if sign is 0 (positive), it flips all bits back
|
||||
* @param[in] f floating-point input (passed as unsigned int)
|
||||
* @see floatFlip
|
||||
**/
|
||||
template <bool doFlip>
|
||||
__device__ uint floatUnflip(uint f)
|
||||
{
|
||||
if (doFlip)
|
||||
{
|
||||
uint mask = ((f >> 31) - 1) | 0x80000000;
|
||||
return f ^ mask;
|
||||
}
|
||||
else
|
||||
return f;
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Scans one warp quickly, optimized for 32-element warps, using shared memory
|
||||
*
|
||||
* Scans each warp in parallel ("warp-scan"), one element per thread.
|
||||
* uses 2 numElements of shared memory per thread (64 numElements per warp)
|
||||
*
|
||||
* @param[in] val Elements per thread to scan
|
||||
* @param[in,out] sData
|
||||
**/
|
||||
template<class T, int maxlevel>
|
||||
__device__ T scanwarp(T val, volatile T* sData)
|
||||
{
|
||||
// The following is the same as 2 * WARP_SIZE * warpId + threadInWarp =
|
||||
// 64*(threadIdx.x >> 5) + (threadIdx.x & (WARP_SIZE - 1))
|
||||
int idx = 2 * threadIdx.x - (threadIdx.x & (WARP_SIZE - 1));
|
||||
sData[idx] = 0;
|
||||
idx += WARP_SIZE;
|
||||
T t = sData[idx] = val; __EMUSYNC;
|
||||
|
||||
#ifdef __DEVICE_EMULATION__
|
||||
t = sData[idx - 1]; __EMUSYNC;
|
||||
sData[idx] += t; __EMUSYNC;
|
||||
t = sData[idx - 2]; __EMUSYNC;
|
||||
sData[idx] += t; __EMUSYNC;
|
||||
t = sData[idx - 4]; __EMUSYNC;
|
||||
sData[idx] += t; __EMUSYNC;
|
||||
t = sData[idx - 8]; __EMUSYNC;
|
||||
sData[idx] += t; __EMUSYNC;
|
||||
t = sData[idx - 16]; __EMUSYNC;
|
||||
sData[idx] += t; __EMUSYNC;
|
||||
#else
|
||||
if (0 <= maxlevel) { sData[idx] = t = t + sData[idx - 1]; } __EMUSYNC;
|
||||
if (1 <= maxlevel) { sData[idx] = t = t + sData[idx - 2]; } __EMUSYNC;
|
||||
if (2 <= maxlevel) { sData[idx] = t = t + sData[idx - 4]; } __EMUSYNC;
|
||||
if (3 <= maxlevel) { sData[idx] = t = t + sData[idx - 8]; } __EMUSYNC;
|
||||
if (4 <= maxlevel) { sData[idx] = t = t + sData[idx -16]; } __EMUSYNC;
|
||||
#endif
|
||||
return sData[idx] - val; // convert inclusive -> exclusive
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Scans 4*CTA_SIZE unsigned ints in a block
|
||||
*
|
||||
* scan4 scans 4*CTA_SIZE numElements in a block (4 per
|
||||
* thread), using a warp-scan algorithm
|
||||
*
|
||||
* @param[in] idata 4-vector of integers to scan
|
||||
**/
|
||||
__device__ uint4 scan4(uint4 idata)
|
||||
{
|
||||
extern __shared__ uint ptr[];
|
||||
|
||||
uint idx = threadIdx.x;
|
||||
|
||||
uint4 val4 = idata;
|
||||
uint sum[3];
|
||||
sum[0] = val4.x;
|
||||
sum[1] = val4.y + sum[0];
|
||||
sum[2] = val4.z + sum[1];
|
||||
|
||||
uint val = val4.w + sum[2];
|
||||
|
||||
val = scanwarp<uint, 4>(val, ptr);
|
||||
__syncthreads();
|
||||
|
||||
if ((idx & (WARP_SIZE - 1)) == WARP_SIZE - 1)
|
||||
{
|
||||
ptr[idx >> 5] = val + val4.w + sum[2];
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
#ifndef __DEVICE_EMULATION__
|
||||
if (idx < WARP_SIZE)
|
||||
#endif
|
||||
{
|
||||
ptr[idx] = scanwarp<uint, 2>(ptr[idx], ptr);
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
val += ptr[idx >> 5];
|
||||
|
||||
val4.x = val;
|
||||
val4.y = val + sum[0];
|
||||
val4.z = val + sum[1];
|
||||
val4.w = val + sum[2];
|
||||
|
||||
return val4;
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Computes output position for each thread given predicate; trues come first then falses
|
||||
*
|
||||
* Rank is the core of the radix sort loop. Given a predicate, it
|
||||
* computes the output position for each thread in an ordering where all
|
||||
* True threads come first, followed by all False threads.
|
||||
* This version handles 4 predicates per thread; hence, "rank4".
|
||||
*
|
||||
* @param[in] preds true/false values for each of the 4 elements in this thread
|
||||
*
|
||||
* @todo is the description of "preds" correct?
|
||||
**/
|
||||
template <int ctasize>
|
||||
__device__ uint4 rank4(uint4 preds)
|
||||
{
|
||||
uint4 address = scan4(preds);
|
||||
|
||||
__shared__ uint numtrue;
|
||||
if (threadIdx.x == ctasize-1)
|
||||
{
|
||||
numtrue = address.w + preds.w;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
uint4 rank;
|
||||
uint idx = threadIdx.x << 2;
|
||||
rank.x = (preds.x) ? address.x : numtrue + idx - address.x;
|
||||
rank.y = (preds.y) ? address.y : numtrue + idx + 1 - address.y;
|
||||
rank.z = (preds.z) ? address.z : numtrue + idx + 2 - address.z;
|
||||
rank.w = (preds.w) ? address.w : numtrue + idx + 3 - address.w;
|
||||
|
||||
return rank;
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Sorts one block
|
||||
*
|
||||
* Uses rank to sort one bit at a time: Sorts a block according
|
||||
* to bits startbit -> nbits + startbit
|
||||
* @param[in,out] key
|
||||
* @param[in,out] value
|
||||
**/
|
||||
template<uint nbits, uint startbit>
|
||||
__device__ void radixSortBlock(uint4 &key, uint4 &value)
|
||||
{
|
||||
extern __shared__ uint sMem1[];
|
||||
for(uint shift = startbit; shift < (startbit + nbits); ++shift)
|
||||
{
|
||||
uint4 lsb;
|
||||
lsb.x = !((key.x >> shift) & 0x1);
|
||||
lsb.y = !((key.y >> shift) & 0x1);
|
||||
lsb.z = !((key.z >> shift) & 0x1);
|
||||
lsb.w = !((key.w >> shift) & 0x1);
|
||||
|
||||
uint4 r = rank4<256>(lsb);
|
||||
|
||||
#if 1
|
||||
// This arithmetic strides the ranks across 4 SORT_CTA_SIZE regions
|
||||
sMem1[(r.x & 3) * SORT_CTA_SIZE + (r.x >> 2)] = key.x;
|
||||
sMem1[(r.y & 3) * SORT_CTA_SIZE + (r.y >> 2)] = key.y;
|
||||
sMem1[(r.z & 3) * SORT_CTA_SIZE + (r.z >> 2)] = key.z;
|
||||
sMem1[(r.w & 3) * SORT_CTA_SIZE + (r.w >> 2)] = key.w;
|
||||
__syncthreads();
|
||||
|
||||
// The above allows us to read without 4-way bank conflicts:
|
||||
key.x = sMem1[threadIdx.x];
|
||||
key.y = sMem1[threadIdx.x + SORT_CTA_SIZE];
|
||||
key.z = sMem1[threadIdx.x + 2 * SORT_CTA_SIZE];
|
||||
key.w = sMem1[threadIdx.x + 3 * SORT_CTA_SIZE];
|
||||
|
||||
__syncthreads();
|
||||
|
||||
sMem1[(r.x & 3) * SORT_CTA_SIZE + (r.x >> 2)] = value.x;
|
||||
sMem1[(r.y & 3) * SORT_CTA_SIZE + (r.y >> 2)] = value.y;
|
||||
sMem1[(r.z & 3) * SORT_CTA_SIZE + (r.z >> 2)] = value.z;
|
||||
sMem1[(r.w & 3) * SORT_CTA_SIZE + (r.w >> 2)] = value.w;
|
||||
__syncthreads();
|
||||
|
||||
value.x = sMem1[threadIdx.x];
|
||||
value.y = sMem1[threadIdx.x + SORT_CTA_SIZE];
|
||||
value.z = sMem1[threadIdx.x + 2 * SORT_CTA_SIZE];
|
||||
value.w = sMem1[threadIdx.x + 3 * SORT_CTA_SIZE];
|
||||
#else
|
||||
sMem1[r.x] = key.x;
|
||||
sMem1[r.y] = key.y;
|
||||
sMem1[r.z] = key.z;
|
||||
sMem1[r.w] = key.w;
|
||||
__syncthreads();
|
||||
|
||||
// This access has 4-way bank conflicts
|
||||
key = sMem[threadIdx.x];
|
||||
|
||||
__syncthreads();
|
||||
|
||||
sMem1[r.x] = value.x;
|
||||
sMem1[r.y] = value.y;
|
||||
sMem1[r.z] = value.z;
|
||||
sMem1[r.w] = value.w;
|
||||
__syncthreads();
|
||||
|
||||
value = sMem[threadIdx.x];
|
||||
#endif
|
||||
|
||||
__syncthreads();
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Sorts one block. Key-only version.
|
||||
*
|
||||
* Uses rank to sort one bit at a time: Sorts a block according
|
||||
* to bits startbit -> nbits + startbit
|
||||
* @param[in,out] key
|
||||
**/
|
||||
|
||||
template<uint nbits, uint startbit>
|
||||
__device__ void radixSortBlockKeysOnly(uint4 &key)
|
||||
{
|
||||
extern __shared__ uint sMem1[];
|
||||
for(uint shift = startbit; shift < (startbit + nbits); ++shift)
|
||||
{
|
||||
uint4 lsb;
|
||||
lsb.x = !((key.x >> shift) & 0x1);
|
||||
lsb.y = !((key.y >> shift) & 0x1);
|
||||
lsb.z = !((key.z >> shift) & 0x1);
|
||||
lsb.w = !((key.w >> shift) & 0x1);
|
||||
|
||||
uint4 r = rank4<256>(lsb);
|
||||
|
||||
#if 1
|
||||
// This arithmetic strides the ranks across 4 CTA_SIZE regions
|
||||
sMem1[(r.x & 3) * SORT_CTA_SIZE + (r.x >> 2)] = key.x;
|
||||
sMem1[(r.y & 3) * SORT_CTA_SIZE + (r.y >> 2)] = key.y;
|
||||
sMem1[(r.z & 3) * SORT_CTA_SIZE + (r.z >> 2)] = key.z;
|
||||
sMem1[(r.w & 3) * SORT_CTA_SIZE + (r.w >> 2)] = key.w;
|
||||
__syncthreads();
|
||||
|
||||
// The above allows us to read without 4-way bank conflicts:
|
||||
key.x = sMem1[threadIdx.x];
|
||||
key.y = sMem1[threadIdx.x + SORT_CTA_SIZE];
|
||||
key.z = sMem1[threadIdx.x + 2 * SORT_CTA_SIZE];
|
||||
key.w = sMem1[threadIdx.x + 3 * SORT_CTA_SIZE];
|
||||
#else
|
||||
sMem1[r.x] = key.x;
|
||||
sMem1[r.y] = key.y;
|
||||
sMem1[r.z] = key.z;
|
||||
sMem1[r.w] = key.w;
|
||||
__syncthreads();
|
||||
|
||||
// This access has 4-way bank conflicts
|
||||
key = sMem[threadIdx.x];
|
||||
#endif
|
||||
|
||||
__syncthreads();
|
||||
}
|
||||
}
|
||||
|
||||
/** @} */ // end radix sort functions
|
||||
/** @} */ // end cudpp_cta
|
||||
// -------------------------------------------------------------
|
||||
// CUDPP -- CUDA Data Parallel Primitives library
|
||||
// -------------------------------------------------------------
|
||||
// $Revision$
|
||||
// $Date$
|
||||
// -------------------------------------------------------------
|
||||
// This source code is distributed under the terms of license.txt
|
||||
// in the root directory of this source distribution.
|
||||
// -------------------------------------------------------------
|
||||
#include <cudpp_globals.h>
|
||||
#include "cudpp_radixsort.h"
|
||||
#include "cta/scan_cta.cu"
|
||||
#include <cudpp.h>
|
||||
#include <stdio.h>
|
||||
|
||||
#include <cudpp_util.h>
|
||||
#include <math.h>
|
||||
#include "sharedmem.h"
|
||||
|
||||
|
||||
#ifdef __DEVICE_EMULATION__
|
||||
#define __EMUSYNC __syncthreads()
|
||||
#else
|
||||
#define __EMUSYNC
|
||||
#endif
|
||||
|
||||
/**
|
||||
* @file
|
||||
* sort_cta.cu
|
||||
*
|
||||
* @brief CUDPP CTA-level sort routines
|
||||
*/
|
||||
|
||||
/** \addtogroup cudpp_cta
|
||||
* @{
|
||||
*/
|
||||
|
||||
/** @name Radix Sort Functions
|
||||
* @{
|
||||
*/
|
||||
|
||||
|
||||
typedef unsigned int uint;
|
||||
|
||||
/**
|
||||
* @brief Flips bits of single-precision floating-point number (parameterized by doFlip)
|
||||
*
|
||||
* flip a float for sorting
|
||||
* finds SIGN of fp number.
|
||||
* if it's 1 (negative float), it flips all bits
|
||||
* if it's 0 (positive float), it flips the sign only
|
||||
* @param[in] f floating-point input (passed as unsigned int)
|
||||
* @see floatUnflip
|
||||
**/
|
||||
|
||||
template <bool doFlip>
|
||||
__device__ uint floatFlip(uint f)
|
||||
{
|
||||
if (doFlip)
|
||||
{
|
||||
uint mask = -int(f >> 31) | 0x80000000;
|
||||
return f ^ mask;
|
||||
}
|
||||
else
|
||||
return f;
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Reverses bit-flip of single-precision floating-point number (parameterized by doFlip)
|
||||
*
|
||||
* flip a float back (invert FloatFlip)
|
||||
* signed was flipped from above, so:
|
||||
* if sign is 1 (negative), it flips the sign bit back
|
||||
* if sign is 0 (positive), it flips all bits back
|
||||
* @param[in] f floating-point input (passed as unsigned int)
|
||||
* @see floatFlip
|
||||
**/
|
||||
template <bool doFlip>
|
||||
__device__ uint floatUnflip(uint f)
|
||||
{
|
||||
if (doFlip)
|
||||
{
|
||||
uint mask = ((f >> 31) - 1) | 0x80000000;
|
||||
return f ^ mask;
|
||||
}
|
||||
else
|
||||
return f;
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Scans one warp quickly, optimized for 32-element warps, using shared memory
|
||||
*
|
||||
* Scans each warp in parallel ("warp-scan"), one element per thread.
|
||||
* uses 2 numElements of shared memory per thread (64 numElements per warp)
|
||||
*
|
||||
* @param[in] val Elements per thread to scan
|
||||
* @param[in,out] sData
|
||||
**/
|
||||
template<class T, int maxlevel>
|
||||
__device__ T scanwarp(T val, volatile T* sData)
|
||||
{
|
||||
// The following is the same as 2 * WARP_SIZE * warpId + threadInWarp =
|
||||
// 64*(threadIdx.x >> 5) + (threadIdx.x & (WARP_SIZE - 1))
|
||||
int idx = 2 * threadIdx.x - (threadIdx.x & (WARP_SIZE - 1));
|
||||
sData[idx] = 0;
|
||||
idx += WARP_SIZE;
|
||||
T t = sData[idx] = val; __EMUSYNC;
|
||||
|
||||
#ifdef __DEVICE_EMULATION__
|
||||
t = sData[idx - 1]; __EMUSYNC;
|
||||
sData[idx] += t; __EMUSYNC;
|
||||
t = sData[idx - 2]; __EMUSYNC;
|
||||
sData[idx] += t; __EMUSYNC;
|
||||
t = sData[idx - 4]; __EMUSYNC;
|
||||
sData[idx] += t; __EMUSYNC;
|
||||
t = sData[idx - 8]; __EMUSYNC;
|
||||
sData[idx] += t; __EMUSYNC;
|
||||
t = sData[idx - 16]; __EMUSYNC;
|
||||
sData[idx] += t; __EMUSYNC;
|
||||
#else
|
||||
if (0 <= maxlevel) { sData[idx] = t = t + sData[idx - 1]; } __EMUSYNC;
|
||||
if (1 <= maxlevel) { sData[idx] = t = t + sData[idx - 2]; } __EMUSYNC;
|
||||
if (2 <= maxlevel) { sData[idx] = t = t + sData[idx - 4]; } __EMUSYNC;
|
||||
if (3 <= maxlevel) { sData[idx] = t = t + sData[idx - 8]; } __EMUSYNC;
|
||||
if (4 <= maxlevel) { sData[idx] = t = t + sData[idx -16]; } __EMUSYNC;
|
||||
#endif
|
||||
return sData[idx] - val; // convert inclusive -> exclusive
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Scans 4*CTA_SIZE unsigned ints in a block
|
||||
*
|
||||
* scan4 scans 4*CTA_SIZE numElements in a block (4 per
|
||||
* thread), using a warp-scan algorithm
|
||||
*
|
||||
* @param[in] idata 4-vector of integers to scan
|
||||
**/
|
||||
__device__ uint4 scan4(uint4 idata)
|
||||
{
|
||||
extern __shared__ uint ptr[];
|
||||
|
||||
uint idx = threadIdx.x;
|
||||
|
||||
uint4 val4 = idata;
|
||||
uint sum[3];
|
||||
sum[0] = val4.x;
|
||||
sum[1] = val4.y + sum[0];
|
||||
sum[2] = val4.z + sum[1];
|
||||
|
||||
uint val = val4.w + sum[2];
|
||||
|
||||
val = scanwarp<uint, 4>(val, ptr);
|
||||
__syncthreads();
|
||||
|
||||
if ((idx & (WARP_SIZE - 1)) == WARP_SIZE - 1)
|
||||
{
|
||||
ptr[idx >> 5] = val + val4.w + sum[2];
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
#ifndef __DEVICE_EMULATION__
|
||||
if (idx < WARP_SIZE)
|
||||
#endif
|
||||
{
|
||||
ptr[idx] = scanwarp<uint, 2>(ptr[idx], ptr);
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
val += ptr[idx >> 5];
|
||||
|
||||
val4.x = val;
|
||||
val4.y = val + sum[0];
|
||||
val4.z = val + sum[1];
|
||||
val4.w = val + sum[2];
|
||||
|
||||
return val4;
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Computes output position for each thread given predicate; trues come first then falses
|
||||
*
|
||||
* Rank is the core of the radix sort loop. Given a predicate, it
|
||||
* computes the output position for each thread in an ordering where all
|
||||
* True threads come first, followed by all False threads.
|
||||
* This version handles 4 predicates per thread; hence, "rank4".
|
||||
*
|
||||
* @param[in] preds true/false values for each of the 4 elements in this thread
|
||||
*
|
||||
* @todo is the description of "preds" correct?
|
||||
**/
|
||||
template <int ctasize>
|
||||
__device__ uint4 rank4(uint4 preds)
|
||||
{
|
||||
uint4 address = scan4(preds);
|
||||
|
||||
__shared__ uint numtrue;
|
||||
if (threadIdx.x == ctasize-1)
|
||||
{
|
||||
numtrue = address.w + preds.w;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
uint4 rank;
|
||||
uint idx = threadIdx.x << 2;
|
||||
rank.x = (preds.x) ? address.x : numtrue + idx - address.x;
|
||||
rank.y = (preds.y) ? address.y : numtrue + idx + 1 - address.y;
|
||||
rank.z = (preds.z) ? address.z : numtrue + idx + 2 - address.z;
|
||||
rank.w = (preds.w) ? address.w : numtrue + idx + 3 - address.w;
|
||||
|
||||
return rank;
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Sorts one block
|
||||
*
|
||||
* Uses rank to sort one bit at a time: Sorts a block according
|
||||
* to bits startbit -> nbits + startbit
|
||||
* @param[in,out] key
|
||||
* @param[in,out] value
|
||||
**/
|
||||
template<uint nbits, uint startbit>
|
||||
__device__ void radixSortBlock(uint4 &key, uint4 &value)
|
||||
{
|
||||
extern __shared__ uint sMem1[];
|
||||
for(uint shift = startbit; shift < (startbit + nbits); ++shift)
|
||||
{
|
||||
uint4 lsb;
|
||||
lsb.x = !((key.x >> shift) & 0x1);
|
||||
lsb.y = !((key.y >> shift) & 0x1);
|
||||
lsb.z = !((key.z >> shift) & 0x1);
|
||||
lsb.w = !((key.w >> shift) & 0x1);
|
||||
|
||||
uint4 r = rank4<256>(lsb);
|
||||
|
||||
#if 1
|
||||
// This arithmetic strides the ranks across 4 SORT_CTA_SIZE regions
|
||||
sMem1[(r.x & 3) * SORT_CTA_SIZE + (r.x >> 2)] = key.x;
|
||||
sMem1[(r.y & 3) * SORT_CTA_SIZE + (r.y >> 2)] = key.y;
|
||||
sMem1[(r.z & 3) * SORT_CTA_SIZE + (r.z >> 2)] = key.z;
|
||||
sMem1[(r.w & 3) * SORT_CTA_SIZE + (r.w >> 2)] = key.w;
|
||||
__syncthreads();
|
||||
|
||||
// The above allows us to read without 4-way bank conflicts:
|
||||
key.x = sMem1[threadIdx.x];
|
||||
key.y = sMem1[threadIdx.x + SORT_CTA_SIZE];
|
||||
key.z = sMem1[threadIdx.x + 2 * SORT_CTA_SIZE];
|
||||
key.w = sMem1[threadIdx.x + 3 * SORT_CTA_SIZE];
|
||||
|
||||
__syncthreads();
|
||||
|
||||
sMem1[(r.x & 3) * SORT_CTA_SIZE + (r.x >> 2)] = value.x;
|
||||
sMem1[(r.y & 3) * SORT_CTA_SIZE + (r.y >> 2)] = value.y;
|
||||
sMem1[(r.z & 3) * SORT_CTA_SIZE + (r.z >> 2)] = value.z;
|
||||
sMem1[(r.w & 3) * SORT_CTA_SIZE + (r.w >> 2)] = value.w;
|
||||
__syncthreads();
|
||||
|
||||
value.x = sMem1[threadIdx.x];
|
||||
value.y = sMem1[threadIdx.x + SORT_CTA_SIZE];
|
||||
value.z = sMem1[threadIdx.x + 2 * SORT_CTA_SIZE];
|
||||
value.w = sMem1[threadIdx.x + 3 * SORT_CTA_SIZE];
|
||||
#else
|
||||
sMem1[r.x] = key.x;
|
||||
sMem1[r.y] = key.y;
|
||||
sMem1[r.z] = key.z;
|
||||
sMem1[r.w] = key.w;
|
||||
__syncthreads();
|
||||
|
||||
// This access has 4-way bank conflicts
|
||||
key = sMem[threadIdx.x];
|
||||
|
||||
__syncthreads();
|
||||
|
||||
sMem1[r.x] = value.x;
|
||||
sMem1[r.y] = value.y;
|
||||
sMem1[r.z] = value.z;
|
||||
sMem1[r.w] = value.w;
|
||||
__syncthreads();
|
||||
|
||||
value = sMem[threadIdx.x];
|
||||
#endif
|
||||
|
||||
__syncthreads();
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Sorts one block. Key-only version.
|
||||
*
|
||||
* Uses rank to sort one bit at a time: Sorts a block according
|
||||
* to bits startbit -> nbits + startbit
|
||||
* @param[in,out] key
|
||||
**/
|
||||
|
||||
template<uint nbits, uint startbit>
|
||||
__device__ void radixSortBlockKeysOnly(uint4 &key)
|
||||
{
|
||||
extern __shared__ uint sMem1[];
|
||||
for(uint shift = startbit; shift < (startbit + nbits); ++shift)
|
||||
{
|
||||
uint4 lsb;
|
||||
lsb.x = !((key.x >> shift) & 0x1);
|
||||
lsb.y = !((key.y >> shift) & 0x1);
|
||||
lsb.z = !((key.z >> shift) & 0x1);
|
||||
lsb.w = !((key.w >> shift) & 0x1);
|
||||
|
||||
uint4 r = rank4<256>(lsb);
|
||||
|
||||
#if 1
|
||||
// This arithmetic strides the ranks across 4 CTA_SIZE regions
|
||||
sMem1[(r.x & 3) * SORT_CTA_SIZE + (r.x >> 2)] = key.x;
|
||||
sMem1[(r.y & 3) * SORT_CTA_SIZE + (r.y >> 2)] = key.y;
|
||||
sMem1[(r.z & 3) * SORT_CTA_SIZE + (r.z >> 2)] = key.z;
|
||||
sMem1[(r.w & 3) * SORT_CTA_SIZE + (r.w >> 2)] = key.w;
|
||||
__syncthreads();
|
||||
|
||||
// The above allows us to read without 4-way bank conflicts:
|
||||
key.x = sMem1[threadIdx.x];
|
||||
key.y = sMem1[threadIdx.x + SORT_CTA_SIZE];
|
||||
key.z = sMem1[threadIdx.x + 2 * SORT_CTA_SIZE];
|
||||
key.w = sMem1[threadIdx.x + 3 * SORT_CTA_SIZE];
|
||||
#else
|
||||
sMem1[r.x] = key.x;
|
||||
sMem1[r.y] = key.y;
|
||||
sMem1[r.z] = key.z;
|
||||
sMem1[r.w] = key.w;
|
||||
__syncthreads();
|
||||
|
||||
// This access has 4-way bank conflicts
|
||||
key = sMem[threadIdx.x];
|
||||
#endif
|
||||
|
||||
__syncthreads();
|
||||
}
|
||||
}
|
||||
|
||||
/** @} */ // end radix sort functions
|
||||
/** @} */ // end cudpp_cta
|
||||
|
|
|
|||
File diff suppressed because it is too large
Load Diff
|
|
@ -7,7 +7,7 @@
|
|||
// This source code is distributed under the terms of license.txt in
|
||||
// the root directory of this source distribution.
|
||||
// -------------------------------------------------------------
|
||||
|
||||
|
||||
/**
|
||||
* @file
|
||||
* cudpp.cpp
|
||||
|
|
|
|||
|
|
@ -7,7 +7,7 @@
|
|||
// This source code is distributed under the terms of license.txt in
|
||||
// the root directory of this source distribution.
|
||||
// -------------------------------------------------------------
|
||||
|
||||
|
||||
/**
|
||||
* @file
|
||||
* cudpp.h
|
||||
|
|
|
|||
|
|
@ -7,7 +7,7 @@
|
|||
// This source code is distributed under the terms of license.txt in
|
||||
// the root directory of this source distribution.
|
||||
// -------------------------------------------------------------
|
||||
|
||||
|
||||
/**
|
||||
* @file
|
||||
* cudpp_globals.h
|
||||
|
|
|
|||
|
|
@ -8,7 +8,7 @@
|
|||
// in the root directory of this source distribution.
|
||||
// -------------------------------------------------------------
|
||||
#include "cudpp_maximal_launch.h"
|
||||
|
||||
|
||||
inline size_t min(size_t x, size_t y)
|
||||
{
|
||||
return (x <= y) ? x : y;
|
||||
|
|
|
|||
|
|
@ -9,7 +9,7 @@
|
|||
// -------------------------------------------------------------
|
||||
#ifndef _MAXIMAL_LAUNCH_H_
|
||||
#define _MAXIMAL_LAUNCH_H_
|
||||
|
||||
|
||||
#include "cuda_runtime.h"
|
||||
|
||||
extern "C"
|
||||
|
|
|
|||
|
|
@ -7,7 +7,7 @@
|
|||
// This source code is distributed under the terms of license.txt
|
||||
// in the root directory of this source distribution.
|
||||
// -------------------------------------------------------------
|
||||
|
||||
|
||||
#include "cudpp.h"
|
||||
#include "cudpp_plan_manager.h"
|
||||
#include "cudpp_scan.h"
|
||||
|
|
|
|||
|
|
@ -9,7 +9,7 @@
|
|||
// -------------------------------------------------------------
|
||||
#ifndef __CUDPP_PLAN_H__
|
||||
#define __CUDPP_PLAN_H__
|
||||
|
||||
|
||||
typedef void* KernelPointer;
|
||||
|
||||
extern "C" size_t getNumCTAs(KernelPointer kernel);
|
||||
|
|
|
|||
|
|
@ -11,7 +11,7 @@
|
|||
#include "cudpp_plan.h"
|
||||
#include "cudpp_plan_manager.h"
|
||||
#include "cudpp_maximal_launch.h"
|
||||
|
||||
|
||||
typedef void* KernelPointer;
|
||||
|
||||
extern "C" size_t getNumCTAs(KernelPointer kernel)
|
||||
|
|
|
|||
|
|
@ -9,7 +9,7 @@
|
|||
// -------------------------------------------------------------
|
||||
#ifndef __CUDPP_PLAN_MANAGER_H__
|
||||
#define __CUDPP_PLAN_MANAGER_H__
|
||||
|
||||
|
||||
#include <map>
|
||||
|
||||
class CUDPPPlan;
|
||||
|
|
|
|||
|
|
@ -9,7 +9,7 @@
|
|||
// -------------------------------------------------------------
|
||||
#ifndef __RADIXSORT_H__
|
||||
#define __RADIXSORT_H__
|
||||
|
||||
|
||||
#define SORT_CTA_SIZE 256 //This CTA_SIZE must equal 16 * number of radices
|
||||
|
||||
#include "cudpp_globals.h"
|
||||
|
|
|
|||
|
|
@ -7,7 +7,7 @@
|
|||
// This source code is distributed under the terms of license.txt
|
||||
// in the root directory of this source distribution.
|
||||
// -------------------------------------------------------------
|
||||
|
||||
|
||||
/**
|
||||
* @file
|
||||
* cudpp_scan.h
|
||||
|
|
|
|||
|
|
@ -7,7 +7,7 @@
|
|||
// This source code is distributed under the terms of license.txt in
|
||||
// the root directory of this source distribution.
|
||||
// -------------------------------------------------------------
|
||||
|
||||
|
||||
/**
|
||||
* @file
|
||||
* cudpp_util.h
|
||||
|
|
|
|||
|
|
@ -26,7 +26,7 @@
|
|||
* 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the
|
||||
* source code with only those rights set forth herein.
|
||||
*/
|
||||
|
||||
|
||||
|
||||
/* CUda UTility Library */
|
||||
|
||||
|
|
|
|||
File diff suppressed because it is too large
Load Diff
|
|
@ -23,3 +23,4 @@ DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
|
|||
LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE
|
||||
OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF
|
||||
ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
|
|
@ -7,7 +7,7 @@
|
|||
// This source code is distributed under the terms of license.txt
|
||||
// in the root directory of this source distribution.
|
||||
// -------------------------------------------------------------
|
||||
|
||||
|
||||
/**
|
||||
* @file
|
||||
* sharedmem.h
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -1 +1,2 @@
|
|||
Geryon Version 10.280
|
||||
|
||||
|
|
@ -6,7 +6,7 @@
|
|||
# requires: sed, rm, mv
|
||||
#
|
||||
# Author: Axel Kohlmeyer, Temple University
|
||||
|
||||
|
||||
num_args=$#
|
||||
|
||||
# we write to a scratch file, since
|
||||
|
|
|
|||
|
|
@ -13,7 +13,7 @@
|
|||
copyright : (C) 2009 by W. Michael Brown
|
||||
email : brownw@ornl.gov
|
||||
***************************************************************************/
|
||||
|
||||
|
||||
/* -----------------------------------------------------------------------
|
||||
Copyright (2009) Sandia Corporation. Under the terms of Contract
|
||||
DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
|
||||
|
|
|
|||
|
|
@ -6,7 +6,7 @@
|
|||
#undef _GLIBCXX_ATOMIC_BUILTINS
|
||||
#endif // _GLIBCXX_ATOMIC_BUILTINS
|
||||
#endif // __APPLE__
|
||||
|
||||
|
||||
#include <stdio.h>
|
||||
#include <cassert>
|
||||
#include <cuda_runtime.h>
|
||||
|
|
|
|||
|
|
@ -20,7 +20,7 @@
|
|||
certain rights in this software. This software is distributed under
|
||||
the Simplified BSD License.
|
||||
----------------------------------------------------------------------- */
|
||||
|
||||
|
||||
#ifndef NVC_TEXTURE
|
||||
#define NVC_TEXTURE
|
||||
|
||||
|
|
|
|||
|
|
@ -20,7 +20,7 @@
|
|||
certain rights in this software. This software is distributed under
|
||||
the Simplified BSD License.
|
||||
----------------------------------------------------------------------- */
|
||||
|
||||
|
||||
#ifndef NVD_DEVICE
|
||||
#define NVD_DEVICE
|
||||
|
||||
|
|
|
|||
|
|
@ -20,7 +20,7 @@
|
|||
certain rights in this software. This software is distributed under
|
||||
the Simplified BSD License.
|
||||
----------------------------------------------------------------------- */
|
||||
|
||||
|
||||
#ifndef NVD_KERNEL
|
||||
#define NVD_KERNEL
|
||||
|
||||
|
|
|
|||
|
|
@ -1,6 +1,6 @@
|
|||
#ifndef NVD_MACROS_H
|
||||
#define NVD_MACROS_H
|
||||
|
||||
|
||||
#include <stdio.h>
|
||||
#include <cassert>
|
||||
#include <cuda.h>
|
||||
|
|
|
|||
|
|
@ -20,7 +20,7 @@
|
|||
certain rights in this software. This software is distributed under
|
||||
the Simplified BSD License.
|
||||
----------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/*! \file */
|
||||
|
||||
#ifndef NVD_MAT_H
|
||||
|
|
|
|||
|
|
@ -20,7 +20,7 @@
|
|||
certain rights in this software. This software is distributed under
|
||||
the Simplified BSD License.
|
||||
----------------------------------------------------------------------- */
|
||||
|
||||
|
||||
#ifndef NVD_MEMORY_H
|
||||
#define NVD_MEMORY_H
|
||||
|
||||
|
|
|
|||
|
|
@ -20,7 +20,7 @@
|
|||
certain rights in this software. This software is distributed under
|
||||
the Simplified BSD License.
|
||||
----------------------------------------------------------------------- */
|
||||
|
||||
|
||||
#ifndef NVD_TEXTURE
|
||||
#define NVD_TEXTURE
|
||||
|
||||
|
|
|
|||
|
|
@ -20,7 +20,7 @@
|
|||
certain rights in this software. This software is distributed under
|
||||
the Simplified BSD License.
|
||||
----------------------------------------------------------------------- */
|
||||
|
||||
|
||||
#ifndef NVD_TIMER_H
|
||||
#define NVD_TIMER_H
|
||||
|
||||
|
|
|
|||
|
|
@ -20,7 +20,7 @@
|
|||
certain rights in this software. This software is distributed under
|
||||
the Simplified BSD License.
|
||||
----------------------------------------------------------------------- */
|
||||
|
||||
|
||||
#ifndef OCL_DEVICE
|
||||
#define OCL_DEVICE
|
||||
|
||||
|
|
|
|||
|
|
@ -13,7 +13,7 @@
|
|||
copyright : (C) 2010 by W. Michael Brown
|
||||
email : brownw@ornl.gov
|
||||
***************************************************************************/
|
||||
|
||||
|
||||
/* -----------------------------------------------------------------------
|
||||
Copyright (2010) Sandia Corporation. Under the terms of Contract
|
||||
DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
|
||||
|
|
|
|||
|
|
@ -20,7 +20,7 @@
|
|||
certain rights in this software. This software is distributed under
|
||||
the Simplified BSD License.
|
||||
----------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/*! \file */
|
||||
|
||||
#ifndef OCL_MAT_H
|
||||
|
|
|
|||
|
|
@ -20,7 +20,7 @@
|
|||
certain rights in this software. This software is distributed under
|
||||
the Simplified BSD License.
|
||||
----------------------------------------------------------------------- */
|
||||
|
||||
|
||||
#ifndef OCL_TEXTURE
|
||||
#define OCL_TEXTURE
|
||||
|
||||
|
|
|
|||
|
|
@ -20,7 +20,7 @@
|
|||
certain rights in this software. This software is distributed under
|
||||
the Simplified BSD License.
|
||||
----------------------------------------------------------------------- */
|
||||
|
||||
|
||||
#ifndef OCL_TIMER_H
|
||||
#define OCL_TIMER_H
|
||||
|
||||
|
|
|
|||
|
|
@ -13,7 +13,7 @@
|
|||
copyright : (C) 2010 by W. Michael Brown
|
||||
email : brownw@ornl.gov
|
||||
***************************************************************************/
|
||||
|
||||
|
||||
/* -----------------------------------------------------------------------
|
||||
Copyright (2010) Sandia Corporation. Under the terms of Contract
|
||||
DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
|
||||
|
|
|
|||
|
|
@ -13,7 +13,7 @@
|
|||
copyright : (C) 2009 by W. Michael Brown
|
||||
email : brownw@ornl.gov
|
||||
***************************************************************************/
|
||||
|
||||
|
||||
/* -----------------------------------------------------------------------
|
||||
Copyright (2009) Sandia Corporation. Under the terms of Contract
|
||||
DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
|
||||
|
|
|
|||
|
|
@ -13,7 +13,7 @@
|
|||
copyright : (C) 2010 by W. Michael Brown
|
||||
email : brownw@ornl.gov
|
||||
***************************************************************************/
|
||||
|
||||
|
||||
/* -----------------------------------------------------------------------
|
||||
Copyright (2010) Sandia Corporation. Under the terms of Contract
|
||||
DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
|
||||
|
|
|
|||
|
|
@ -13,7 +13,7 @@
|
|||
copyright : (C) 2009 by W. Michael Brown
|
||||
email : brownw@ornl.gov
|
||||
***************************************************************************/
|
||||
|
||||
|
||||
/* -----------------------------------------------------------------------
|
||||
Copyright (2009) Sandia Corporation. Under the terms of Contract
|
||||
DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
|
||||
|
|
|
|||
|
|
@ -13,7 +13,7 @@
|
|||
copyright : (C) 2009 by W. Michael Brown
|
||||
email : brownw@ornl.gov
|
||||
***************************************************************************/
|
||||
|
||||
|
||||
/* -----------------------------------------------------------------------
|
||||
Copyright (2009) Sandia Corporation. Under the terms of Contract
|
||||
DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
|
||||
|
|
|
|||
|
|
@ -13,7 +13,7 @@
|
|||
copyright : (C) 2009 by W. Michael Brown
|
||||
email : brownw@ornl.gov
|
||||
***************************************************************************/
|
||||
|
||||
|
||||
/* -----------------------------------------------------------------------
|
||||
Copyright (2009) Sandia Corporation. Under the terms of Contract
|
||||
DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
|
||||
|
|
|
|||
|
|
@ -13,7 +13,7 @@
|
|||
copyright : (C) 2009 by W. Michael Brown
|
||||
email : brownw@ornl.gov
|
||||
***************************************************************************/
|
||||
|
||||
|
||||
/* -----------------------------------------------------------------------
|
||||
Copyright (2009) Sandia Corporation. Under the terms of Contract
|
||||
DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
|
||||
|
|
|
|||
|
|
@ -13,7 +13,7 @@
|
|||
copyright : (C) 2009 by W. Michael Brown
|
||||
email : brownw@ornl.gov
|
||||
***************************************************************************/
|
||||
|
||||
|
||||
/* -----------------------------------------------------------------------
|
||||
Copyright (2009) Sandia Corporation. Under the terms of Contract
|
||||
DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
|
||||
|
|
|
|||
|
|
@ -13,7 +13,7 @@
|
|||
copyright : (C) 2010 by W. Michael Brown
|
||||
email : brownw@ornl.gov
|
||||
***************************************************************************/
|
||||
|
||||
|
||||
/* -----------------------------------------------------------------------
|
||||
Copyright (2010) Sandia Corporation. Under the terms of Contract
|
||||
DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
|
||||
|
|
|
|||
|
|
@ -13,7 +13,7 @@
|
|||
copyright : (C) 2010 by W. Michael Brown
|
||||
email : brownw@ornl.gov
|
||||
***************************************************************************/
|
||||
|
||||
|
||||
/* -----------------------------------------------------------------------
|
||||
Copyright (2010) Sandia Corporation. Under the terms of Contract
|
||||
DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
|
||||
|
|
|
|||
|
|
@ -13,7 +13,7 @@
|
|||
copyright : (C) 2010 by W. Michael Brown
|
||||
email : brownw@ornl.gov
|
||||
***************************************************************************/
|
||||
|
||||
|
||||
/* -----------------------------------------------------------------------
|
||||
Copyright (2010) Sandia Corporation. Under the terms of Contract
|
||||
DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
Peng Wang (Nvidia), penwang@nvidia.com
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
|
|
@ -10,7 +10,7 @@
|
|||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
|
||||
------------------------------------------------------------------------- */
|
||||
|
|
|
|||
Loading…
Reference in New Issue