Skip to content

Commit ac68881

Browse files
LuHan Wang
authored andcommitted
fix bug of 0 output when sec_a.back() is lager than 1024
1 parent 0f9edb4 commit ac68881

File tree

1 file changed

+11
-6
lines changed

1 file changed

+11
-6
lines changed

source/op/cuda/descrpt_se_r.cu

Lines changed: 11 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -209,14 +209,16 @@ __global__ void compute_descriptor_se_r (VALUETYPE* descript,
209209
const VALUETYPE* coord,
210210
const VALUETYPE rmin,
211211
const VALUETYPE rmax,
212-
compute_t* sel_diff_dev)
212+
compute_t* sel_diff_dev,
213+
const int sec_size)
213214
{
214215
// <<<nloc, sec.back()>>>
215-
const unsigned int idx = blockIdx.x;
216-
const unsigned int idy = threadIdx.x;
216+
const unsigned int idx = blockIdx.y;
217+
const unsigned int idy = blockIdx.x * blockDim.x + threadIdx.x;
217218
const int idx_deriv = idy * 3; // 1 components time 3 directions
218219
const int idx_value = idy; // 1 components
219-
220+
if (idy >= sec_size) {return;}
221+
220222
// else {return;}
221223
VALUETYPE * row_descript = descript + idx * ndescrpt;
222224
VALUETYPE * row_descript_deriv = descript_deriv + idx * descript_deriv_size;
@@ -324,7 +326,9 @@ void DescrptSeRLauncher(const VALUETYPE* coord,
324326
nei_iter
325327
);
326328
}
327-
compute_descriptor_se_r<<<nloc, sec.back()>>> (
329+
const int nblock_ = (sec.back() + LEN -1) / LEN;
330+
dim3 block_grid(nblock_, nloc);
331+
compute_descriptor_se_r<<<block_grid, LEN>>> (
328332
descript,
329333
ndescrpt,
330334
descript_deriv,
@@ -339,6 +343,7 @@ void DescrptSeRLauncher(const VALUETYPE* coord,
339343
coord,
340344
rcut_smth,
341345
rcut,
342-
sel_diff
346+
sel_diff,
347+
sec.back()
343348
);
344349
}

0 commit comments

Comments
 (0)