← Back to branch summary

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

« back to all changes in this revision

Viewing changes to lib/cuda/pair_sw_cuda_kernel_nc.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/pair_sw_cuda_kernel_nc.cu
27
27
 
28
28
 
29
29
 
30
 
__device__ void twobody(int iparam, F_FLOAT rsq, F_FLOAT &fforce,
31
 
                        int eflag, ENERGY_FLOAT &eng)
 
30
__device__ void twobody(int iparam, F_CFLOAT rsq, F_CFLOAT &fforce,
 
31
                        int eflag, ENERGY_CFLOAT &eng)
32
32
{
33
 
  F_FLOAT r, rp, rq, rainv, expsrainv;
 
33
  F_CFLOAT r, rp, rq, rainv, expsrainv;
34
34
 
35
35
  r = sqrt(rsq);
36
36
  rp = pow(r, -params_sw[iparam].powerp);
44
44
}
45
45
 
46
46
__device__ void threebody(int paramij, int paramik, int paramijk,
47
 
                          F_FLOAT4 &delr1,
48
 
                          F_FLOAT4 &delr2,
49
 
                          F_FLOAT3 &fj, F_FLOAT3 &fk, int eflag, ENERGY_FLOAT &eng)
 
47
                          F_CFLOAT4 &delr1,
 
48
                          F_CFLOAT4 &delr2,
 
49
                          F_CFLOAT3 &fj, F_CFLOAT3 &fk, int eflag, ENERGY_CFLOAT &eng)
50
50
{
51
 
  F_FLOAT r1, rinvsq1, rainv1, gsrainv1, gsrainvsq1, expgsrainv1;
52
 
  F_FLOAT r2, rinvsq2, rainv2, gsrainv2, gsrainvsq2, expgsrainv2;
53
 
  F_FLOAT rinv12, cs, delcs, delcssq, facexp, facrad, frad1, frad2;
54
 
  F_FLOAT facang, facang12, csfacang, csfac1, csfac2;
 
51
  F_CFLOAT r1, rinvsq1, rainv1, gsrainv1, gsrainvsq1, expgsrainv1;
 
52
  F_CFLOAT r2, rinvsq2, rainv2, gsrainv2, gsrainvsq2, expgsrainv2;
 
53
  F_CFLOAT rinv12, cs, delcs, delcssq, facexp, facrad, frad1, frad2;
 
54
  F_CFLOAT facang, facang12, csfacang, csfac1, csfac2;
55
55
 
56
56
  r1 = sqrt(delr1.w);
57
57
  rinvsq1 = F_F(1.0) / delr1.w;
99
99
}
100
100
 
101
101
__device__ void threebody_fj(int paramij, int paramik, int paramijk,
102
 
                             F_FLOAT4 &delr1,
103
 
                             F_FLOAT4 &delr2,
104
 
                             F_FLOAT3 &fj)
 
102
                             F_CFLOAT4 &delr1,
 
103
                             F_CFLOAT4 &delr2,
 
104
                             F_CFLOAT3 &fj)
105
105
{
106
 
  F_FLOAT r1, rinvsq1, rainv1, gsrainv1, gsrainvsq1, expgsrainv1;
107
 
  F_FLOAT r2, rainv2, gsrainv2, expgsrainv2;
108
 
  F_FLOAT rinv12, cs, delcs, delcssq, facexp, facrad, frad1;
109
 
  F_FLOAT facang, facang12, csfacang, csfac1;
 
106
  F_CFLOAT r1, rinvsq1, rainv1, gsrainv1, gsrainvsq1, expgsrainv1;
 
107
  F_CFLOAT r2, rainv2, gsrainv2, expgsrainv2;
 
108
  F_CFLOAT rinv12, cs, delcs, delcssq, facexp, facrad, frad1;
 
109
  F_CFLOAT facang, facang12, csfacang, csfac1;
110
110
 
111
111
  r1 = sqrt(delr1.w);
112
112
  rinvsq1 = F_F(1.0) / delr1.w;
143
143
}
144
144
 
145
145
 
146
 
__global__ void Pair_SW_Kernel_TpA_RIJ()//F_FLOAT4* _glob_r_ij,int* _glob_numneigh_red,int* _glob_neighbors_red,int* _glob_neightype_red)
 
146
__global__ void Pair_SW_Kernel_TpA_RIJ()//F_CFLOAT4* _glob_r_ij,int* _glob_numneigh_red,int* _glob_neighbors_red,int* _glob_neightype_red)
147
147
{
148
148
  int ii = (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x;
149
149
 
150
150
  if(ii >= _nall) return;
151
151
 
152
 
  X_FLOAT4 myxtype;
153
 
  F_FLOAT4 delij;
154
 
  F_FLOAT xtmp, ytmp, ztmp;
 
152
  X_CFLOAT4 myxtype;
 
153
  F_CFLOAT4 delij;
 
154
  F_CFLOAT xtmp, ytmp, ztmp;
155
155
  int itype, jnum, i, j;
156
156
  int* jlist;
157
157
  int neigh_red = 0;
195
195
 
196
196
 
197
197
template <int eflag, int vflagm>
198
 
__global__ void Pair_SW_Kernel_TpA(int eflag_atom, int vflag_atom) //,F_FLOAT* _glob_zeta_ij,F_FLOAT4* _glob_r_ij,int* _glob_numneigh_red,int* _glob_neighbors_red,int* _glob_neightype_red)
 
198
__global__ void Pair_SW_Kernel_TpA(int eflag_atom, int vflag_atom) //,F_CFLOAT* _glob_zeta_ij,F_CFLOAT4* _glob_r_ij,int* _glob_numneigh_red,int* _glob_neighbors_red,int* _glob_neightype_red)
199
199
{
200
 
  ENERGY_FLOAT evdwl = ENERGY_F(0.0);
201
 
 
202
 
  ENERGY_FLOAT* sharedE = &sharedmem[threadIdx.x];
203
 
  ENERGY_FLOAT* sharedV = &sharedmem[threadIdx.x];
204
 
 
205
 
  F_FLOAT* shared_F_F = (F_FLOAT*) sharedmem;
206
 
 
207
 
  if((eflag || eflag_atom) && (vflagm || vflag_atom)) shared_F_F = (F_FLOAT*) &sharedmem[7 * blockDim.x];
208
 
  else if(eflag) shared_F_F = (F_FLOAT*) &sharedmem[blockDim.x];
209
 
  else if(vflagm) shared_F_F = (F_FLOAT*) &sharedmem[6 * blockDim.x];
 
200
  ENERGY_CFLOAT evdwl = ENERGY_F(0.0);
 
201
 
 
202
  ENERGY_CFLOAT* sharedE = &sharedmem[threadIdx.x];
 
203
  ENERGY_CFLOAT* sharedV = &sharedmem[threadIdx.x];
 
204
 
 
205
  F_CFLOAT* shared_F_F = (F_CFLOAT*) sharedmem;
 
206
 
 
207
  if((eflag || eflag_atom) && (vflagm || vflag_atom)) shared_F_F = (F_CFLOAT*) &sharedmem[7 * blockDim.x];
 
208
  else if(eflag) shared_F_F = (F_CFLOAT*) &sharedmem[blockDim.x];
 
209
  else if(vflagm) shared_F_F = (F_CFLOAT*) &sharedmem[6 * blockDim.x];
210
210
 
211
211
  shared_F_F += threadIdx.x;
212
212
 
231
231
  //#define jnum_red (static_cast <int> (shared_F_F[3*blockDim.x]))
232
232
 
233
233
  int ii = (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x;
234
 
  X_FLOAT4 myxtype_i, myxtype_j, myxtype_k;
235
 
  F_FLOAT4 delij, delik, deljk;
236
 
  F_FLOAT fpair;
 
234
  X_CFLOAT4 myxtype_i, myxtype_j, myxtype_k;
 
235
  F_CFLOAT4 delij, delik, deljk;
 
236
  F_CFLOAT fpair;
237
237
 
238
238
  int itype, i, j;
239
239
  int* jlist_red;
277
277
      volatile int iparam_ji = elem2param[(jtype * nelements + itype) * nelements + itype];
278
278
 
279
279
      if(delij.w < params_sw[iparam_ij].cutsq) {
280
 
        F_FLOAT dxfp, dyfp, dzfp;
 
280
        F_CFLOAT dxfp, dyfp, dzfp;
281
281
        twobody(iparam_ij, delij.w, fpair, eflag, evdwl);
282
282
        fxtmp += dxfp = delij.x * fpair;
283
283
        fytmp += dyfp = delij.y * fpair;
316
316
          vec3_scale(F_F(-1.0), delik, delik);
317
317
 
318
318
          if(delik.w <= params_sw[iparam_ijk].cutsq) {
319
 
            F_FLOAT3 fj, fk;
 
319
            F_CFLOAT3 fj, fk;
320
320
            threebody(iparam_ij, iparam_ik, iparam_ijk,
321
321
                      delij, delik, fj, fk, eflag, evdwl);
322
322
            fxtmp -= fj.x + fk.x;
377
377
          vec3_scale(F_F(-1.0), delij, delij);
378
378
 
379
379
          if(deljk.w <= params_sw[iparam_jik].cutsq) {
380
 
            F_FLOAT3 fj;
 
380
            F_CFLOAT3 fj;
381
381
 
382
382
            threebody_fj(iparam_ji, iparam_jk, iparam_jik,
383
383
                         delij, deljk, fj);
397
397
  __syncthreads();
398
398
 
399
399
  if(ii < _inum) {
400
 
    F_FLOAT* my_f;
 
400
    F_CFLOAT* my_f;
401
401
 
402
402
    if(_collect_forces_later) {
403
 
      ENERGY_FLOAT* buffer = (ENERGY_FLOAT*) _buffer;
 
403
      ENERGY_CFLOAT* buffer = (ENERGY_CFLOAT*) _buffer;
404
404
 
405
405
      if(eflag) {
406
406
        buffer = &buffer[1 * gridDim.x * gridDim.y];
410
410
        buffer = &buffer[6 * gridDim.x * gridDim.y];
411
411
      }
412
412
 
413
 
      my_f = (F_FLOAT*) buffer;
 
413
      my_f = (F_CFLOAT*) buffer;
414
414
      my_f += i;
415
415
      *my_f = fxtmp;
416
416
      my_f += _nmax;