← Back to branch summary

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

« back to all changes in this revision

Viewing changes to lib/cuda/compute_temp_partial_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/compute_temp_partial_cuda_kernel.cu
21
21
   This software is distributed under the GNU General Public License.
22
22
------------------------------------------------------------------------- */
23
23
 
24
 
extern __shared__ ENERGY_FLOAT sharedmem[];
 
24
extern __shared__ ENERGY_CFLOAT sharedmem[];
25
25
 
26
26
 
27
27
__global__ void Cuda_ComputeTempPartialCuda_Scalar_Kernel(int groupbit, int xflag, int yflag, int zflag)
40
40
  }
41
41
 
42
42
  reduceBlock(sharedmem);
43
 
  ENERGY_FLOAT* buffer = (ENERGY_FLOAT*) _buffer;
 
43
  ENERGY_CFLOAT* buffer = (ENERGY_CFLOAT*) _buffer;
44
44
 
45
45
  if(threadIdx.x == 0) {
46
46
    buffer[blockIdx.x * gridDim.y + blockIdx.y] = sharedmem[0];
59
59
 
60
60
  if(i < _nlocal)
61
61
    if(_mask[i] & groupbit) {
62
 
      V_FLOAT massone;
 
62
      V_CFLOAT massone;
63
63
 
64
64
      if(_rmass_flag) massone = _rmass[i];
65
65
      else massone = _mass[_type[i]];
78
78
  reduceBlock(&sharedmem[3 * blockDim.x]);
79
79
  reduceBlock(&sharedmem[4 * blockDim.x]);
80
80
  reduceBlock(&sharedmem[5 * blockDim.x]);
81
 
  ENERGY_FLOAT* buffer = (ENERGY_FLOAT*) _buffer;
 
81
  ENERGY_CFLOAT* buffer = (ENERGY_CFLOAT*) _buffer;
82
82
 
83
83
  if(threadIdx.x == 0) {
84
84
    buffer[blockIdx.x * gridDim.y + blockIdx.y] = sharedmem[0];
91
91
}
92
92
 
93
93
 
94
 
__global__ void Cuda_ComputeTempPartialCuda_Reduce_Kernel(int n, ENERGY_FLOAT* t)
 
94
__global__ void Cuda_ComputeTempPartialCuda_Reduce_Kernel(int n, ENERGY_CFLOAT* t)
95
95
{
96
96
  int i = 0;
97
97
  sharedmem[threadIdx.x] = 0;
98
 
  ENERGY_FLOAT myforig = 0.0;
99
 
  ENERGY_FLOAT* buf = (ENERGY_FLOAT*) _buffer;
 
98
  ENERGY_CFLOAT myforig = 0.0;
 
99
  ENERGY_CFLOAT* buf = (ENERGY_CFLOAT*) _buffer;
100
100
  buf = &buf[blockIdx.x * n];
101
101
 
102
102
  while(i < n) {
117
117
    t[blockIdx.x] = myforig;
118
118
}
119
119
 
120
 
__global__ void Cuda_ComputeTempPartialCuda_RemoveBiasAll_Kernel(int groupbit, int xflag, int yflag, int zflag, V_FLOAT* vbiasall)
 
120
__global__ void Cuda_ComputeTempPartialCuda_RemoveBiasAll_Kernel(int groupbit, int xflag, int yflag, int zflag, V_CFLOAT* vbiasall)
121
121
{
122
122
  int i = (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x;
123
123
 
140
140
    }
141
141
}
142
142
 
143
 
__global__ void Cuda_ComputeTempPartialCuda_RestoreBiasAll_Kernel(int groupbit, int xflag, int yflag, int zflag, V_FLOAT* vbiasall)
 
143
__global__ void Cuda_ComputeTempPartialCuda_RestoreBiasAll_Kernel(int groupbit, int xflag, int yflag, int zflag, V_CFLOAT* vbiasall)
144
144
{
145
145
  int i = (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x;
146
146