← Back to branch summary

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

« back to all changes in this revision

Viewing changes to lib/cuda/fix_set_force_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_set_force_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[];
25
 
 
26
 
 
27
 
__global__ void Cuda_FixSetForceCuda_PostForce_Kernel(int groupbit, F_FLOAT xvalue, F_FLOAT yvalue, F_FLOAT zvalue, int flagx, int flagy, int flagz)
 
24
extern __shared__ F_CFLOAT sharedmem[];
 
25
 
 
26
 
 
27
__global__ void Cuda_FixSetForceCuda_PostForce_Kernel(int groupbit, F_CFLOAT xvalue, F_CFLOAT yvalue, F_CFLOAT zvalue, int flagx, int flagy, int flagz)
28
28
{
29
29
  int i = (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x;
30
30
  sharedmem[threadIdx.x] = 0;
48
48
  reduceBlock(sharedmem);
49
49
  reduceBlock(&sharedmem[blockDim.x]);
50
50
  reduceBlock(&sharedmem[2 * blockDim.x]);
51
 
  F_FLOAT* buffer = (F_FLOAT*)_buffer;
 
51
  F_CFLOAT* buffer = (F_CFLOAT*)_buffer;
52
52
 
53
53
  if(threadIdx.x == 0) {
54
54
    buffer[blockIdx.x * gridDim.y + blockIdx.y] = sharedmem[0];
58
58
}
59
59
 
60
60
 
61
 
__global__ void Cuda_FixSetForceCuda_Reduce_FOriginal(int n, F_FLOAT* foriginal)
 
61
__global__ void Cuda_FixSetForceCuda_Reduce_FOriginal(int n, F_CFLOAT* foriginal)
62
62
{
63
63
  int i = 0;
64
64
  sharedmem[threadIdx.x] = 0;
65
 
  F_FLOAT myforig = 0.0;
66
 
  F_FLOAT* buf = (F_FLOAT*)_buffer;
 
65
  F_CFLOAT myforig = 0.0;
 
66
  F_CFLOAT* buf = (F_CFLOAT*)_buffer;
67
67
  buf = &buf[blockIdx.x * n];
68
68
 
69
69
  while(i < n) {