← Back to branch summary

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

« back to all changes in this revision

Viewing changes to lib/cuda/fix_freeze_cuda_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/fix_freeze_cuda_kernel.cu
21
21
   This software is distributed under the GNU General Public License.
22
22
------------------------------------------------------------------------- */
23
23
 
24
 
extern __shared__ F_FLOAT sharedmem[];
 
24
extern __shared__ F_CFLOAT sharedmem[];
25
25
 
26
26
 
27
27
__global__ void Cuda_FixFreezeCuda_PostForce_Kernel(int groupbit)
49
49
  reduceBlock(sharedmem);
50
50
  reduceBlock(&sharedmem[blockDim.x]);
51
51
  reduceBlock(&sharedmem[2 * blockDim.x]);
52
 
  F_FLOAT* buffer = (F_FLOAT*)_buffer;
 
52
  F_CFLOAT* buffer = (F_CFLOAT*)_buffer;
53
53
 
54
54
  if(threadIdx.x == 0) {
55
55
    buffer[blockIdx.x * gridDim.y + blockIdx.y] = sharedmem[0];
59
59
}
60
60
 
61
61
 
62
 
__global__ void Cuda_FixFreezeCuda_Reduce_FOriginal(int n, F_FLOAT* foriginal)
 
62
__global__ void Cuda_FixFreezeCuda_Reduce_FOriginal(int n, F_CFLOAT* foriginal)
63
63
{
64
64
  int i = 0;
65
65
  sharedmem[threadIdx.x] = 0;
66
 
  F_FLOAT myforig = 0.0;
67
 
  F_FLOAT* buf = (F_FLOAT*)_buffer;
 
66
  F_CFLOAT myforig = 0.0;
 
67
  F_CFLOAT* buf = (F_CFLOAT*)_buffer;
68
68
  buf = &buf[blockIdx.x * n];
69
69
 
70
70
  while(i < n) {