Skip to content

fix bug 824 and Synchronize updates to CUDA code #828

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
amcadmus merged 4 commits into from
Jul 6, 2021
Merged

fix bug 824 and Synchronize updates to CUDA code #828

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
70cb226
Synchronize CUDA _r modifications to ROCM
galeselee Jun 29, 2021
d6e9d13
Merge pull request #1 from deepmodeling/devel
galeselee Jul 5, 2021
8588bbb
fix bug 824 and Synchronize updates to CUDA code
galeselee Jul 5, 2021
1f0cb9a
Update prod_env_mat.hip.cu
galeselee Jul 5, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions source/lib/src/rocm/coord.hip.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -66,8 +66,8 @@ __global__ void normalize_one(
FPTYPE inter[3];
phys2Inter(inter,out_c+idy*3,rec_boxt);
for (int dd = 0; dd < 3; ++dd) {
while(inter[dd] >= 1.) inter[dd] -= 1.;
while(inter[dd] < 0.) inter[dd] += 1.;
inter[dd]=(FPTYPE)fmod((double)inter[dd], 1.);
if (inter[dd] < 0.) inter[dd] += 1.;
}
inter2Phys(out_c+idy*3,inter,boxt);
}
Expand Down
12 changes: 12 additions & 0 deletions source/lib/src/rocm/gelu.hip.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -56,6 +56,10 @@ namespace deepmd {
const FPTYPE * xx,
const int size)
{
if(size <= 0)
{
return ;
}
const int THREAD_ITEMS = 1024;
const int BLOCK_NUMS = (size + THREAD_ITEMS - 1) / THREAD_ITEMS;

Expand All @@ -69,6 +73,10 @@ namespace deepmd {
const FPTYPE * dy,
const int size)
{
if(size <= 0)
{
return;
}
const int THREAD_ITEMS = 1024;
const int BLOCK_NUMS = (size + THREAD_ITEMS - 1) / THREAD_ITEMS;

Expand All @@ -83,6 +91,10 @@ namespace deepmd {
const FPTYPE * dy_2,
const int size)
{
if(size <= 0)
{
return;
}
const int THREAD_ITEMS = 1024;
const int BLOCK_NUMS = (size + THREAD_ITEMS - 1) / THREAD_ITEMS;

Expand Down
4 changes: 2 additions & 2 deletions source/lib/src/rocm/prod_env_mat.hip.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -504,7 +504,7 @@ void prod_env_mat_a_gpu_rocm(
const int ndescrpt = nnei * 4;
hipErrcheck(hipMemset(em, 0.0, sizeof(FPTYPE) * nloc * ndescrpt));
hipErrcheck(hipMemset(em_deriv, 0.0, sizeof(FPTYPE) * nloc * ndescrpt * 3));
hipErrcheck(hipMemset(rij, 0., sizeof(FPTYPE) * nloc * nnei * 3));
hipErrcheck(hipMemset(rij, 0.0, sizeof(FPTYPE) * nloc * nnei * 3));

format_nbor_list_gpu_rocm(
nlist,
Expand Down Expand Up @@ -541,7 +541,7 @@ void prod_env_mat_r_gpu_rocm(
const int ndescrpt = nnei * 1;
hipErrcheck(hipMemset(em, 0.0, sizeof(FPTYPE) * nloc * ndescrpt));
hipErrcheck(hipMemset(em_deriv, 0.0, sizeof(FPTYPE) * nloc * ndescrpt * 3));
hipErrcheck(hipMemset(rij, 0., sizeof(FPTYPE) * nloc * nnei * 3));
hipErrcheck(hipMemset(rij, 0.0, sizeof(FPTYPE) * nloc * nnei * 3));

format_nbor_list_gpu_rocm(
nlist,
Expand Down
2 changes: 2 additions & 0 deletions source/lib/src/rocm/tabulate.hip.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -209,6 +209,7 @@ template<typename FPTYPE>
const int nnei,
const int last_layer_size)
{
if(nloc <= 0){return;}
hipLaunchKernelGGL(HIP_KERNEL_NAME(tabulate_fusion_fifth_order_polynomial<FPTYPE, MM, KK>), nloc, last_layer_size, sizeof(FPTYPE) * MM * last_layer_size, 0,
out,
table, em_x, em, table_info[0], table_info[1], table_info[2], table_info[3], table_info[4], nnei, last_layer_size);
Expand All @@ -227,6 +228,7 @@ template<typename FPTYPE>
const int nnei,
const int last_layer_size)
{
if( nloc<=0 ) { return;}
hipErrcheck(hipMemset(
dy_dem_x,
0.0, sizeof(FPTYPE) * nloc * nnei));
Expand Down