Fix: Compiling error for ROCM (#5464)

dyzheng · dyzheng · web-flow · commit 9718c41c5a6c · 2024-11-13T14:00:31.000+08:00
* Fix: compiling error for ROCM

* Fix: compiling error of DCU

---------

Co-authored-by: dyzheng &lt;zhengdy@bjaisi.com&gt;
diff --git a/source/module_hamilt_pw/hamilt_pwdft/kernels/rocm/stress_op.hip.cu b/source/module_hamilt_pw/hamilt_pwdft/kernels/rocm/stress_op.hip.cu
@@ -291,7 +291,6 @@ __global__ void cal_stress_nl(
         ++iat;
         sum+=nproj;
     }//ia
-    __syncwarp();
     warp_reduce(stress_var);
     if (threadIdx.x % WARP_SIZE == 0) {
         atomicAdd(stress + ipol * 3 + jpol, stress_var);
diff --git a/source/module_hamilt_pw/hamilt_stodft/kernels/rocm/hpsi_norm_op.hip.cu b/source/module_hamilt_pw/hamilt_stodft/kernels/rocm/hpsi_norm_op.hip.cu
@@ -45,7 +45,7 @@ void hamilt::hpsi_norm_op<FPTYPE, base_device::DEVICE_GPU>::operator()(const bas
       npwk_max, npwk, Ebar, DeltaE,
       reinterpret_cast<thrust::complex<FPTYPE>*>(hpsi),
       reinterpret_cast<const thrust::complex<FPTYPE>*>(psi_in));
-    cudaCheckOnDebug();
+    hipCheckOnDebug();
 }
 
 template struct hpsi_norm_op<float, base_device::DEVICE_GPU>;

Original file line number	Diff line number	Diff line change
`@@ -45,7 +45,7 @@ void hamilt::hpsi_norm_op<FPTYPE, base_device::DEVICE_GPU>::operator()(const bas`
`45`	`45`	`npwk_max, npwk, Ebar, DeltaE,`
`46`	`46`	`reinterpret_cast<thrust::complex<FPTYPE>*>(hpsi),`
`47`	`47`	`reinterpret_cast<const thrust::complex<FPTYPE>*>(psi_in));`
`48`		`- cudaCheckOnDebug();`
	`48`	`+ hipCheckOnDebug();`
`49`	`49`	`}`
`50`	`50`
`51`	`51`	`template struct hpsi_norm_op<float, base_device::DEVICE_GPU>;`