← Back to branch summary

~ubuntu-branches/debian/sid/lammps/sid

« back to all changes in this revision

Viewing changes to lib/cuda/domain_kernel.cu

  • Committer: Package Import Robot
  • Author(s): Anton Gladky
  • Date: 2015-04-29 23:44:49 UTC
  • mfrom: (5.1.3 experimental)
  • Revision ID: package-import@ubuntu.com-20150429234449-mbhy9utku6hp6oq8
Tags: 0~20150313.gitfa668e1-1
Upload into unstable.

Show diffs side-by-side

added added

removed removed

Lines of Context:
lib/cuda/domain_kernel.cu
21
21
   This software is distributed under the GNU General Public License.
22
22
------------------------------------------------------------------------- */
23
23
 
24
 
extern __shared__ X_FLOAT sharedmem[];
 
24
extern __shared__ X_CFLOAT sharedmem[];
25
25
 
26
26
#define BIG 1e10
27
27
__global__ void Domain_PBC_Kernel(int deform_remap, int deform_groupbit, int box_change)
29
29
  int idim, otherdims;
30
30
  int i = (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x;
31
31
 
32
 
  X_FLOAT lo[3];
33
 
  X_FLOAT hi[3];
34
 
  X_FLOAT* period;
 
32
  X_CFLOAT lo[3];
 
33
  X_CFLOAT hi[3];
 
34
  X_CFLOAT* period;
35
35
 
36
36
  if(_triclinic == 0) {
37
37
    lo[0] = _boxlo[0];
54
54
  }
55
55
 
56
56
 
57
 
  X_FLOAT tmpx = X_F(0.5) * (hi[0] + lo[0]);
58
 
  X_FLOAT tmpy = X_F(0.5) * (hi[1] + lo[1]);
59
 
  X_FLOAT tmpz = X_F(0.5) * (hi[2] + lo[2]);
 
57
  X_CFLOAT tmpx = X_F(0.5) * (hi[0] + lo[0]);
 
58
  X_CFLOAT tmpy = X_F(0.5) * (hi[1] + lo[1]);
 
59
  X_CFLOAT tmpz = X_F(0.5) * (hi[2] + lo[2]);
60
60
 
61
 
  X_FLOAT* buf = (X_FLOAT*) _buffer;
 
61
  X_CFLOAT* buf = (X_CFLOAT*) _buffer;
62
62
  buf += blockIdx.x * gridDim.y + blockIdx.y;
63
63
  buf[0] = tmpx;
64
64
  buf += gridDim.x * gridDim.y;
181
181
  __syncthreads();
182
182
 
183
183
  if(box_change) {
184
 
    X_FLOAT minx = BIG;
185
 
    X_FLOAT maxx = -BIG;
186
 
    X_FLOAT miny = BIG;
187
 
    X_FLOAT maxy = -BIG;
188
 
    X_FLOAT minz = BIG;
189
 
    X_FLOAT maxz = -BIG;
 
184
    X_CFLOAT minx = BIG;
 
185
    X_CFLOAT maxx = -BIG;
 
186
    X_CFLOAT miny = BIG;
 
187
    X_CFLOAT maxy = -BIG;
 
188
    X_CFLOAT minz = BIG;
 
189
    X_CFLOAT maxz = -BIG;
190
190
 
191
191
    if(not _periodicity[0]) {
192
192
      sharedmem[threadIdx.x] = tmpx;
231
231
    }
232
232
 
233
233
    if(threadIdx.x == 0) {
234
 
      buf = (X_FLOAT*) _buffer;
 
234
      buf = (X_CFLOAT*) _buffer;
235
235
      buf += blockIdx.x * gridDim.y + blockIdx.y;
236
236
      buf[0] = minx;
237
237
      buf += gridDim.x * gridDim.y;
250
250
 
251
251
__global__ void Domain_reduceBoxExtent(double* extent, int n)
252
252
{
253
 
  X_FLOAT* buf = (X_FLOAT*) _buffer;
 
253
  X_CFLOAT* buf = (X_CFLOAT*) _buffer;
254
254
  buf += blockIdx.x * n;
255
255
  copyGlobToShared(buf, sharedmem, n);
256
256
 
267
267
  int i = (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x;
268
268
 
269
269
  if(i < n) {
270
 
    X_FLOAT ytmp = _x[i + _nmax];
271
 
    X_FLOAT ztmp = _x[i + 2 * _nmax];
 
270
    X_CFLOAT ytmp = _x[i + _nmax];
 
271
    X_CFLOAT ztmp = _x[i + 2 * _nmax];
272
272
    _x[i] = _h[0] * _x[i] + _h[5] * ytmp + _h[4] * ztmp + _boxlo[0];
273
273
    _x[i + _nmax] = _h[1] * ytmp + _h[3] * ztmp + _boxlo[1];
274
274
    _x[i + 2 * _nmax] = _h[2] * ztmp + _boxlo[2];
279
279
{
280
280
  int i = (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x;
281
281
 
282
 
  X_FLOAT delta[3];
 
282
  X_CFLOAT delta[3];
283
283
 
284
284
  if(i < n) {
285
285
    delta[0] = _x[i] - _boxlo[0];