Skip to content

Commit b7940fc

Browse files
LuLu
authored andcommitted
fix bug of 0 output when sec_a.back() is lager than 1024
1 parent 715c7a0 commit b7940fc

File tree

1 file changed

+10
-5
lines changed

1 file changed

+10
-5
lines changed

source/op/cuda/descrpt_se_a.cu

Lines changed: 10 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -194,13 +194,15 @@ __global__ void compute_descriptor_se_a (VALUETYPE* descript,
194194
const VALUETYPE* coord,
195195
const VALUETYPE rmin,
196196
const VALUETYPE rmax,
197-
compute_t* sel_a_diff_dev)
197+
compute_t* sel_a_diff_dev,
198+
const int sec_a_size)
198199
{
199200
// <<<nloc, sec_a.back()>>>
200-
const unsigned int idx = blockIdx.x;
201-
const unsigned int idy = threadIdx.x;
201+
const unsigned int idx = blockIdx.y;
202+
const unsigned int idy = blockIdx.x * blockDim.x + threadIdx.x;
202203
const int idx_deriv = idy * 4 * 3; // 4 components time 3 directions
203204
const int idx_value = idy * 4; // 4 components
205+
if (idy >= sec_a_size) {return;}
204206

205207
// else {return;}
206208
VALUETYPE * row_descript = descript + idx * ndescrpt;
@@ -341,7 +343,9 @@ void DescrptSeALauncher(const VALUETYPE* coord,
341343
);
342344
}
343345

344-
compute_descriptor_se_a<<<nloc, sec_a.back()>>> (
346+
const int nblock_ = (sec_a.back() + LEN -1) / LEN;
347+
dim3 block_grid(nblock_, nloc);
348+
compute_descriptor_se_a<<<block_grid, LEN>>> (
345349
descript,
346350
ndescrpt,
347351
descript_deriv,
@@ -356,7 +360,8 @@ void DescrptSeALauncher(const VALUETYPE* coord,
356360
coord,
357361
rcut_r_smth,
358362
rcut_r,
359-
sel_a_diff
363+
sel_a_diff,
364+
sec_a.back()
360365
);
361366
////
362367
// res = cudaFree(sec_a_dev); cudaErrcheck(res);

0 commit comments

Comments
 (0)