NAMD
Functions
CudaNonbondedTables.C File Reference
#include "charm++.h"
#include "NamdTypes.h"
#include "ComputeNonbondedUtil.h"
#include "LJTable.h"
#include "CudaUtils.h"
#include "CudaNonbondedTables.h"

Go to the source code of this file.

Functions

template<typename T >
void copyTable (int size, T *h_table, T *&d_table, bool update=false)
 
template<typename T >
void bindTextureObject (int size, T *h_table, T *&d_table, cudaTextureObject_t &tex, bool update=false)
 
template<typename T >
void buildForceAndEnergyTable (const int tableSize, const double *r2list, const BigReal *src_table, const bool flip, const BigReal prefac, const int dst_stride, T *dst_force, T *dst_energy)
 
template<typename T >
void bindTextureObject (int tableSize, int tableWidth, T *h_table, cudaArray_t &array, cudaTextureObject_t &tableTex, T **d_table)
 

Function Documentation

template<typename T >
void bindTextureObject ( int  size,
T *  h_table,
T *&  d_table,
cudaTextureObject_t &  tex,
bool  update = false 
)

Definition at line 67 of file CudaNonbondedTables.C.

References cudaCheck.

67  {
68  // Copy to device
69  if ( ! update) {
70  allocate_device<T>(&d_table, size);
71  }
72  else {
73  cudaCheck(cudaDestroyTextureObject(tex));
74  }
75  copy_HtoD_sync<T>(h_table, d_table, size);
76 
77  // Create texture object
78  cudaResourceDesc resDesc;
79  memset(&resDesc, 0, sizeof(resDesc));
80  resDesc.resType = cudaResourceTypeLinear;
81  resDesc.res.linear.devPtr = d_table;
82  resDesc.res.linear.desc.f = cudaChannelFormatKindFloat;
83  resDesc.res.linear.desc.x = sizeof(float)*8; // bits per channel
84  if (sizeof(T) >= sizeof(float)*2) resDesc.res.linear.desc.y = sizeof(float)*8; // bits per channel
85  if (sizeof(T) >= sizeof(float)*3) resDesc.res.linear.desc.z = sizeof(float)*8; // bits per channel
86  if (sizeof(T) >= sizeof(float)*4) resDesc.res.linear.desc.w = sizeof(float)*8; // bits per channel
87  resDesc.res.linear.sizeInBytes = size*sizeof(T);
88 
89  cudaTextureDesc texDesc;
90  memset(&texDesc, 0, sizeof(texDesc));
91  texDesc.readMode = cudaReadModeElementType;
92  //texDesc.normalizedCoords = 0;
93 
94  cudaCheck(cudaCreateTextureObject(&tex, &resDesc, &texDesc, NULL));
95 
96 }
#define cudaCheck(stmt)
Definition: CudaUtils.h:95
template<typename T >
void bindTextureObject ( int  tableSize,
int  tableWidth,
T *  h_table,
cudaArray_t &  array,
cudaTextureObject_t &  tableTex,
T **  d_table 
)

Definition at line 194 of file CudaNonbondedTables.C.

References allocate_device_T(), copy_HtoD_T(), and cudaCheck.

195  {
196 
197 #if defined(NAMD_CUDA)
198  allocate_device_T((void **)d_table, tableSize, sizeof(T)*tableWidth);
199  copy_HtoD_T(h_table, *d_table, tableSize, sizeof(T)*tableWidth);
200 
201  cudaChannelFormatDesc desc;
202  memset(&desc, 0, sizeof(desc));
203  desc.x = sizeof(T)*8;
204  if (tableWidth >= 2) desc.y = sizeof(T)*8;
205  if (tableWidth >= 3) desc.z = sizeof(T)*8;
206  if (tableWidth >= 4) desc.w = sizeof(T)*8;
207  desc.f = cudaChannelFormatKindFloat;
208  cudaCheck(cudaMallocArray(&array, &desc, tableSize, 0));
209  cudaCheck(cudaMemcpyToArray(array, 0, 0, h_table, tableSize*sizeof(T)*tableWidth, cudaMemcpyHostToDevice));
210 
211  cudaResourceDesc resDesc;
212  memset(&resDesc, 0, sizeof(resDesc));
213  resDesc.resType = cudaResourceTypeArray;
214  resDesc.res.array.array = array;
215 
216  cudaTextureDesc texDesc;
217  memset(&texDesc, 0, sizeof(texDesc));
218  texDesc.addressMode[0] = cudaAddressModeClamp;
219  texDesc.filterMode = cudaFilterModeLinear;
220  texDesc.normalizedCoords = 1;
221 
222  cudaCheck(cudaCreateTextureObject(&tableTex, &resDesc, &texDesc, NULL));
223 #else
224  // tex1Dfetch is used in kernels, so create a linear texture
225  // The texture is 1 texel wider to simplify (and optimize) index clamping for tex1Dfetch
226  allocate_device_T((void **)d_table, tableSize + 1, sizeof(T)*tableWidth);
227  copy_HtoD_T(h_table, *d_table, tableSize, sizeof(T)*tableWidth);
228  copy_HtoD_T(h_table + tableWidth * (tableSize - 1), *d_table + tableWidth * tableSize, 1, sizeof(T)*tableWidth);
229 
230  cudaResourceDesc resDesc;
231  memset(&resDesc, 0, sizeof(resDesc));
232  resDesc.resType = cudaResourceTypeLinear;
233  resDesc.res.linear.devPtr = *d_table;
234  resDesc.res.linear.desc.f = cudaChannelFormatKindFloat;
235  resDesc.res.linear.desc.x = sizeof(T)*8;
236  if (tableWidth >= 2) resDesc.res.linear.desc.y = sizeof(T)*8;
237  if (tableWidth >= 3) resDesc.res.linear.desc.z = sizeof(T)*8;
238  if (tableWidth >= 4) resDesc.res.linear.desc.w = sizeof(T)*8;
239  resDesc.res.linear.sizeInBytes = (tableSize + 1)*sizeof(T)*tableWidth;
240 
241  cudaTextureDesc texDesc;
242  memset(&texDesc, 0, sizeof(texDesc));
243  texDesc.readMode = cudaReadModeElementType;
244 
245  cudaCheck(cudaCreateTextureObject(&tableTex, &resDesc, &texDesc, NULL));
246 #endif
247 }
void allocate_device_T(void **pp, const int len, const size_t sizeofT)
Definition: CudaUtils.C:75
void copy_HtoD_T(const void *h_array, void *d_array, int array_len, const size_t sizeofT)
Definition: CudaUtils.C:180
#define cudaCheck(stmt)
Definition: CudaUtils.h:95
template<typename T >
void buildForceAndEnergyTable ( const int  tableSize,
const double *  r2list,
const BigReal src_table,
const bool  flip,
const BigReal  prefac,
const int  dst_stride,
T *  dst_force,
T *  dst_energy 
)

Definition at line 147 of file CudaNonbondedTables.C.

References ComputeNonbondedUtil::cutoff, ComputeNonbondedUtil::r2_delta, ComputeNonbondedUtil::r2_delta_exp, and ComputeNonbondedUtil::r2_table.

148  {
149 
150  const BigReal r2_delta = ComputeNonbondedUtil:: r2_delta;
151  const int r2_delta_exp = ComputeNonbondedUtil:: r2_delta_exp;
152  const int r2_delta_expc = 64 * (r2_delta_exp - 1023);
153 
154  union { double f; int32 i[2]; } byte_order_test;
155  byte_order_test.f = 1.0; // should occupy high-order bits only
156  int32 *r2iilist = (int32*)r2list + ( byte_order_test.i[0] ? 0 : 1 );
157 
158  for ( int i=1; i<tableSize; ++i ) {
159  double r = ((double) tableSize) / ( (double) i + 0.5 );
160  int table_i = (r2iilist[2*i] >> 14) + r2_delta_expc; // table_i >= 0
161 
162  if ( r > ComputeNonbondedUtil::cutoff ) {
163  dst_force[i*dst_stride] = 0.;
164  dst_energy[i*dst_stride] = 0.;
165  continue;
166  }
167 
168  BigReal diffa = r2list[i] - ComputeNonbondedUtil::r2_table[table_i];
169 
170  BigReal table_a, table_b, table_c, table_d;
171  if (flip) {
172  table_a = src_table[4*table_i+3];
173  table_b = src_table[4*table_i+2];
174  table_c = src_table[4*table_i+1];
175  table_d = src_table[4*table_i];
176  } else {
177  table_a = src_table[4*table_i];
178  table_b = src_table[4*table_i+1];
179  table_c = src_table[4*table_i+2];
180  table_d = src_table[4*table_i+3];
181  }
182 
183  BigReal grad = ( 3. * table_d * diffa + 2. * table_c ) * diffa + table_b;
184  dst_force[i*dst_stride] = prefac * 2. * grad;
185  BigReal ener = table_a + diffa * ( ( table_d * diffa + table_c ) * diffa + table_b);
186  dst_energy[i*dst_stride] = prefac * ener;
187  }
188 
189  dst_force[0] = 0.;
190  dst_energy[0] = dst_energy[1*dst_stride];
191 }
short int32
Definition: dumpdcd.c:24
double BigReal
Definition: common.h:114
template<typename T >
void copyTable ( int  size,
T *  h_table,
T *&  d_table,
bool  update = false 
)

Definition at line 57 of file CudaNonbondedTables.C.

57  {
58  // Copy to device
59  if ( ! update) {
60  allocate_device<T>(&d_table, size);
61  }
62  copy_HtoD_sync<T>(h_table, d_table, size);
63 
64 }