← Back to branch summary

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

« back to all changes in this revision

Viewing changes to lib/cuda/fix_nh_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_nh_cuda_kernel.cu
21
21
   This software is distributed under the GNU General Public License.
22
22
------------------------------------------------------------------------- */
23
23
 
24
 
static inline __device__ void check_distance(X_FLOAT &xtmp, X_FLOAT &ytmp, X_FLOAT &ztmp, int &i, int groupbit)
 
24
static inline __device__ void check_distance(X_CFLOAT &xtmp, X_CFLOAT &ytmp, X_CFLOAT &ztmp, int &i, int groupbit)
25
25
{
26
26
  if(_dist_check) {
27
27
 
28
 
    X_FLOAT d = X_F(0.0);
 
28
    X_CFLOAT d = X_F(0.0);
29
29
 
30
30
    if(i < _nlocal) {
31
 
      X_FLOAT tmp = xtmp - _xhold[i];
 
31
      X_CFLOAT tmp = xtmp - _xhold[i];
32
32
      d = tmp * tmp;
33
33
      tmp = ytmp - _xhold[i + _maxhold];
34
34
      d += tmp * tmp;
43
43
  }
44
44
}
45
45
 
46
 
__global__ void FixNHCuda_nh_v_press_Kernel(int groupbit, F_FLOAT3 factor, int p_triclinic, F_FLOAT3 factor2)
 
46
__global__ void FixNHCuda_nh_v_press_Kernel(int groupbit, F_CFLOAT3 factor, int p_triclinic, F_CFLOAT3 factor2)
47
47
{
48
48
  int i = (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x;
49
49
 
50
50
  if(i < _nlocal && _mask[i] & groupbit) {
51
 
    V_FLOAT* my_v = _v + i;
52
 
    V_FLOAT vx = my_v[0];
53
 
    V_FLOAT vy = my_v[_nmax];
54
 
    V_FLOAT vz = my_v[2 * _nmax];
 
51
    V_CFLOAT* my_v = _v + i;
 
52
    V_CFLOAT vx = my_v[0];
 
53
    V_CFLOAT vy = my_v[_nmax];
 
54
    V_CFLOAT vz = my_v[2 * _nmax];
55
55
    vx *= factor.x;
56
56
    vy *= factor.y;
57
57
    vz *= factor.z;
71
71
 
72
72
}
73
73
 
74
 
__global__ void FixNHCuda_nh_v_temp_Kernel(int groupbit, F_FLOAT factor_eta)
 
74
__global__ void FixNHCuda_nh_v_temp_Kernel(int groupbit, F_CFLOAT factor_eta)
75
75
{
76
76
  int i = (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x;
77
77
 
78
78
  if(i < _nlocal && _mask[i] & groupbit) {
79
 
    V_FLOAT* my_v = _v + i;
 
79
    V_CFLOAT* my_v = _v + i;
80
80
    my_v[0] *= factor_eta;
81
81
    my_v[_nmax] *= factor_eta;
82
82
    my_v[2 * _nmax] *= factor_eta;
84
84
 
85
85
}
86
86
 
87
 
__global__ void FixNHCuda_nh_v_press_and_nve_v_NoBias_Kernel(int groupbit, F_FLOAT3 factor, int p_triclinic, F_FLOAT3 factor2)
 
87
__global__ void FixNHCuda_nh_v_press_and_nve_v_NoBias_Kernel(int groupbit, F_CFLOAT3 factor, int p_triclinic, F_CFLOAT3 factor2)
88
88
{
89
89
  int i = (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x;
90
90
 
91
91
  if(i < _nlocal && _mask[i] & groupbit) {
92
 
    F_FLOAT* my_f = _f + i;
93
 
    V_FLOAT* my_v = _v + i;
 
92
    F_CFLOAT* my_f = _f + i;
 
93
    V_CFLOAT* my_v = _v + i;
94
94
 
95
 
    V_FLOAT             dtfm = _dtf;
 
95
    V_CFLOAT            dtfm = _dtf;
96
96
 
97
97
    if(_rmass_flag) dtfm *= V_F(1.0) / _rmass[i];
98
98
    else                        dtfm *= V_F(1.0) / _mass[_type[i]];
99
99
 
100
 
    V_FLOAT vx = my_v[0];
101
 
    V_FLOAT vy = my_v[_nmax];
102
 
    V_FLOAT vz = my_v[2 * _nmax];
 
100
    V_CFLOAT vx = my_v[0];
 
101
    V_CFLOAT vy = my_v[_nmax];
 
102
    V_CFLOAT vz = my_v[2 * _nmax];
103
103
    vx *= factor.x;
104
104
    vy *= factor.y;
105
105
    vz *= factor.z;
125
125
  int i = (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x;
126
126
 
127
127
  if(i < _nlocal && _mask[i] & groupbit) {
128
 
    F_FLOAT* my_f = _f + i;
129
 
    V_FLOAT* my_v = _v + i;
 
128
    F_CFLOAT* my_f = _f + i;
 
129
    V_CFLOAT* my_v = _v + i;
130
130
 
131
 
    V_FLOAT             dtfm = _dtf;
 
131
    V_CFLOAT            dtfm = _dtf;
132
132
 
133
133
    if(_rmass_flag) dtfm *= V_F(1.0) / _rmass[i];
134
134
    else                        dtfm *= V_F(1.0) / _mass[_type[i]];
145
145
 
146
146
__global__ void FixNHCuda_nve_x_Kernel(int groupbit)
147
147
{
148
 
  X_FLOAT xtmp, ytmp, ztmp;
 
148
  X_CFLOAT xtmp, ytmp, ztmp;
149
149
 
150
150
  int i = (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x;
151
151
 
152
152
  if(i < _nlocal && _mask[i] & groupbit) {
153
 
    V_FLOAT* my_v = _v + i;
154
 
    X_FLOAT* my_x = _x + i;
 
153
    V_CFLOAT* my_v = _v + i;
 
154
    X_CFLOAT* my_x = _x + i;
155
155
 
156
156
    xtmp = *my_x += _dtv * *my_v;
157
157
    my_v += _nmax;
166
166
}
167
167
 
168
168
 
169
 
__global__ void FixNHCuda_nve_v_and_nh_v_press_NoBias_Kernel(int groupbit, F_FLOAT3 factor, int p_triclinic, F_FLOAT3 factor2)
 
169
__global__ void FixNHCuda_nve_v_and_nh_v_press_NoBias_Kernel(int groupbit, F_CFLOAT3 factor, int p_triclinic, F_CFLOAT3 factor2)
170
170
{
171
171
 
172
172
  int i = (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x;
173
173
 
174
174
  if(i < _nlocal && _mask[i] & groupbit) {
175
 
    F_FLOAT* my_f = _f + i;
176
 
    V_FLOAT* my_v = _v + i;
 
175
    F_CFLOAT* my_f = _f + i;
 
176
    V_CFLOAT* my_v = _v + i;
177
177
 
178
 
    V_FLOAT             dtfm = _dtf;
 
178
    V_CFLOAT            dtfm = _dtf;
179
179
 
180
180
    if(_rmass_flag) dtfm *= V_F(1.0) / _rmass[i];
181
181
    else                        dtfm *= V_F(1.0) / _mass[_type[i]];
182
182
 
183
 
    V_FLOAT vx = my_v[0] + dtfm * my_f[0];
184
 
    V_FLOAT vy = my_v[_nmax] + dtfm * my_f[_nmax];
185
 
    V_FLOAT vz = my_v[2 * _nmax] + dtfm * my_f[2 * _nmax];
 
183
    V_CFLOAT vx = my_v[0] + dtfm * my_f[0];
 
184
    V_CFLOAT vy = my_v[_nmax] + dtfm * my_f[_nmax];
 
185
    V_CFLOAT vz = my_v[2 * _nmax] + dtfm * my_f[2 * _nmax];
186
186
 
187
187
    vx *= factor.x;
188
188
    vy *= factor.y;