← Back to branch summary

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

« back to all changes in this revision

Viewing changes to lib/cuda/atom_vec_cuda.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/atom_vec_cuda.cu
85
85
{
86
86
  cudaMemcpyToSymbol(MY_AP(nlocal)  , & sdata->atom.nlocal        , sizeof(int));
87
87
  cudaMemcpyToSymbol(MY_AP(nmax)    , & sdata->atom.nmax          , sizeof(int));
88
 
  cudaMemcpyToSymbol(MY_AP(x)       , & sdata->atom.x    .dev_data, sizeof(X_FLOAT*));
89
 
  cudaMemcpyToSymbol(MY_AP(v)       , & sdata->atom.v    .dev_data, sizeof(V_FLOAT*));
90
 
  cudaMemcpyToSymbol(MY_AP(f)       , & sdata->atom.f    .dev_data, sizeof(F_FLOAT*));
 
88
  cudaMemcpyToSymbol(MY_AP(x)       , & sdata->atom.x    .dev_data, sizeof(X_CFLOAT*));
 
89
  cudaMemcpyToSymbol(MY_AP(v)       , & sdata->atom.v    .dev_data, sizeof(V_CFLOAT*));
 
90
  cudaMemcpyToSymbol(MY_AP(f)       , & sdata->atom.f    .dev_data, sizeof(F_CFLOAT*));
91
91
  cudaMemcpyToSymbol(MY_AP(tag)     , & sdata->atom.tag  .dev_data, sizeof(int*));
92
92
  cudaMemcpyToSymbol(MY_AP(type)    , & sdata->atom.type .dev_data, sizeof(int*));
93
93
  cudaMemcpyToSymbol(MY_AP(mask)    , & sdata->atom.mask .dev_data, sizeof(int*));
94
94
  cudaMemcpyToSymbol(MY_AP(image)   , & sdata->atom.image.dev_data, sizeof(int*));
95
95
 
96
 
  if(data_mask & Q_MASK) cudaMemcpyToSymbol(MY_AP(q)       , & sdata->atom.q    .dev_data, sizeof(F_FLOAT*));
 
96
  if(data_mask & Q_MASK) cudaMemcpyToSymbol(MY_AP(q)       , & sdata->atom.q    .dev_data, sizeof(F_CFLOAT*));
97
97
 
98
98
  if(data_mask & MOLECULE_MASK) cudaMemcpyToSymbol(MY_AP(molecule)   , & sdata->atom.molecule.dev_data, sizeof(int*));
99
99
 
121
121
    cudaMemcpyToSymbol(MY_AP(nlocal)  , & sdata->atom.nlocal        , sizeof(int));
122
122
 
123
123
  MYDBG(printf("# CUDA: Cuda_AtomVecCuda_Init ... post Nmax\n");)
124
 
  cudaMemcpyToSymbol(MY_AP(prd)   , sdata->domain.prd, 3 * sizeof(X_FLOAT));
125
 
  cudaMemcpyToSymbol(MY_AP(sublo)   , & sdata->domain.sublo, 3 * sizeof(X_FLOAT));
126
 
  cudaMemcpyToSymbol(MY_AP(subhi)   , & sdata->domain.subhi, 3 * sizeof(X_FLOAT));
 
124
  cudaMemcpyToSymbol(MY_AP(prd)   , sdata->domain.prd, 3 * sizeof(X_CFLOAT));
 
125
  cudaMemcpyToSymbol(MY_AP(sublo)   , & sdata->domain.sublo, 3 * sizeof(X_CFLOAT));
 
126
  cudaMemcpyToSymbol(MY_AP(subhi)   , & sdata->domain.subhi, 3 * sizeof(X_CFLOAT));
127
127
  cudaMemcpyToSymbol(MY_AP(flag)   , & sdata->flag, sizeof(int*));
128
128
  cudaThreadSynchronize();
129
129
  MYDBG(printf("# CUDA: Cuda_AtomVecCuda_Init ... end\n");)
143
143
    cudaMemcpyToSymbol(MY_AP(nlocal)  , & sdata->atom.nlocal        , sizeof(int));
144
144
 
145
145
  int n_data_items = AtomVecCuda_CountDataItems(data_mask);
146
 
  int size = (n * n_data_items) * sizeof(X_FLOAT);
 
146
  int size = (n * n_data_items) * sizeof(X_CFLOAT);
147
147
 
148
148
  if(sdata->buffer_new or (size > sdata->buffersize))
149
149
    Cuda_AtomVecCuda_UpdateBuffer(sdata, size);
150
150
 
151
 
  X_FLOAT dx = 0.0;
152
 
  X_FLOAT dy = 0.0;
153
 
  X_FLOAT dz = 0.0;
 
151
  X_CFLOAT dx = 0.0;
 
152
  X_CFLOAT dy = 0.0;
 
153
  X_CFLOAT dz = 0.0;
154
154
 
155
155
  if(pbc_flag != 0) {
156
156
    if(sdata->domain.triclinic == 0) {
185
185
    CUT_CHECK_ERROR("Cuda_AtomVecCuda_PackComm: Kernel execution failed");
186
186
 
187
187
    if(not sdata->overlap_comm)
188
 
      cudaMemcpy(buf_send, sdata->buffer, n* n_data_items* sizeof(X_FLOAT), cudaMemcpyDeviceToHost);
189
 
    //cudaMemcpy(buf_send, sdata->comm.buf_send_dev[iswap], n*3*sizeof(X_FLOAT), cudaMemcpyDeviceToHost);
 
188
      cudaMemcpy(buf_send, sdata->buffer, n* n_data_items* sizeof(X_CFLOAT), cudaMemcpyDeviceToHost);
 
189
    //cudaMemcpy(buf_send, sdata->comm.buf_send_dev[iswap], n*3*sizeof(X_CFLOAT), cudaMemcpyDeviceToHost);
190
190
 
191
191
    my_gettime(CLOCK_REALTIME, &time1);
192
192
    sdata->cuda_timings.comm_forward_download +=
216
216
    cudaMemcpyToSymbol(MY_AP(nlocal)  , & sdata->atom.nlocal        , sizeof(int));
217
217
 
218
218
  int n_data_items = AtomVecCuda_CountDataItems(data_mask);
219
 
  int size = (n * n_data_items) * sizeof(X_FLOAT);
 
219
  int size = (n * n_data_items) * sizeof(X_CFLOAT);
220
220
 
221
221
  if(sdata->buffer_new or (size > sdata->buffersize))
222
222
    Cuda_AtomVecCuda_UpdateBuffer(sdata, size);
223
223
 
224
224
  static int count = -1;
225
225
  count++;
226
 
  X_FLOAT dx = 0.0;
227
 
  X_FLOAT dy = 0.0;
228
 
  X_FLOAT dz = 0.0;
 
226
  X_CFLOAT dx = 0.0;
 
227
  X_CFLOAT dy = 0.0;
 
228
  X_CFLOAT dz = 0.0;
229
229
 
230
230
  if(pbc_flag != 0) {
231
231
    if(sdata->domain.triclinic == 0) {
276
276
    cudaMemcpyToSymbol(MY_AP(nlocal)  , & sdata->atom.nlocal        , sizeof(int));
277
277
 
278
278
  int n_data_items = AtomVecCuda_CountDataItems(data_mask);
279
 
  int size = (n * n_data_items) * sizeof(X_FLOAT);
 
279
  int size = (n * n_data_items) * sizeof(X_CFLOAT);
280
280
 
281
281
  if(sdata->buffer_new or (size > sdata->buffersize))
282
282
    Cuda_AtomVecCuda_UpdateBuffer(sdata, size);
289
289
    my_gettime(CLOCK_REALTIME, &time1);
290
290
 
291
291
    if(not sdata->overlap_comm || iswap < 0)
292
 
      cudaMemcpy(sdata->buffer, (void*)buf_recv, n_data_items * n * sizeof(X_FLOAT), cudaMemcpyHostToDevice);
 
292
      cudaMemcpy(sdata->buffer, (void*)buf_recv, n_data_items * n * sizeof(X_CFLOAT), cudaMemcpyHostToDevice);
293
293
 
294
294
    my_gettime(CLOCK_REALTIME, &time2);
295
295
    sdata->cuda_timings.comm_forward_upload +=
463
463
 
464
464
  int n_data_items = AtomVecCuda_CountDataItems(data_mask);
465
465
 
466
 
  int size = nsend * n_data_items * sizeof(X_FLOAT);
 
466
  int size = nsend * n_data_items * sizeof(X_CFLOAT);
467
467
 
468
468
  if(sdata->buffer_new or (size > sdata->buffersize))
469
469
    Cuda_AtomVecCuda_UpdateBuffer(sdata, size);
470
470
 
471
 
  X_FLOAT dx = 0.0;
472
 
  X_FLOAT dy = 0.0;
473
 
  X_FLOAT dz = 0.0;
 
471
  X_CFLOAT dx = 0.0;
 
472
  X_CFLOAT dy = 0.0;
 
473
  X_CFLOAT dz = 0.0;
474
474
 
475
475
  if(pbc_flag != 0) {
476
476
    if(sdata->domain.triclinic == 0) {
522
522
 
523
523
  int n_data_items = AtomVecCuda_CountDataItems(data_mask);
524
524
 
525
 
  int size = n * n_data_items * sizeof(X_FLOAT);
 
525
  int size = n * n_data_items * sizeof(X_CFLOAT);
526
526
 
527
527
  if(sdata->buffer_new or (size > sdata->buffersize))
528
528
    Cuda_AtomVecCuda_UpdateBuffer(sdata, size);
529
529
 
530
 
  X_FLOAT dx = 0.0;
531
 
  X_FLOAT dy = 0.0;
532
 
  X_FLOAT dz = 0.0;
 
530
  X_CFLOAT dx = 0.0;
 
531
  X_CFLOAT dy = 0.0;
 
532
  X_CFLOAT dz = 0.0;
533
533
 
534
534
  if(pbc_flag != 0) {
535
535
    if(sdata->domain.triclinic == 0) {
584
584
 
585
585
  int n_data_items = AtomVecCuda_CountDataItems(data_mask);
586
586
 
587
 
  int size = n * n_data_items * sizeof(X_FLOAT);
 
587
  int size = n * n_data_items * sizeof(X_CFLOAT);
588
588
 
589
589
  if(sdata->buffer_new or (size > sdata->buffersize))
590
590
    Cuda_AtomVecCuda_UpdateBuffer(sdata, size);