Skip to content

Commit 14746d3

Browse files
dzzz2001claude
dzzz2001
and
claude
authored
Refactor: Device module modernization and GPU initialization consolidation (Useful information for GPU coding) (deepmodeling#6936)
* Refactor: introduce DeviceContext singleton for unified GPU initialization - Add DeviceContext singleton class (device_context.h/cpp) to manage GPU device binding with thread-safe initialization using std::mutex - Move GPU initialization from get_device_kpar() side-effect to explicit DeviceContext::init() call in read_input.cpp after INPUT parsing - Use MPI_COMM_TYPE_SHARED for modern node-local rank detection - Update callers to use DeviceContext::instance().get_device_id(): - hsolver_lcao.cpp: parakSolve_cusolver() - diag_cusolvermp.cu: constructor - gint_gpu_vars.cpp: constructor - td_nonlocal_lcao.cpp: remove redundant set_device_by_rank() call This is Phase 1 of GPU resource initialization refactoring, establishing a single entry point for GPU device binding instead of scattered calls. Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com> * Refactor: cleanup device.cpp and merge DeviceContext Phase 2 of GPU device module refactoring: - Remove deprecated functions: stringCmp, get_node_rank(), set_device_by_rank() - Merge device_context.h/cpp into device.h/cpp - Update include statements in dependent files - Remove device_context.cpp from CMakeLists.txt This reduces ~70 lines of dead code and consolidates the device module into fewer files with a unified interface. Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com> * Refactor: modularize device module with single responsibility per file - Create kernel_compat.h: move atomicAdd polyfill for pre-Pascal GPUs - Create device_helpers.h/cpp: extract get_device_type and get_current_precision templates - Create gpu_runtime.h: unified CUDA/ROCm API macros for portable GPU code - Refactor output_device.cpp: unify duplicated CUDA/ROCm implementations (~150 lines reduced) - Update device.h: clean interface with only DeviceContext and information namespace - Update CMakeLists.txt: add device_helpers.cpp to build Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com> * Refactor: simplify parakSolve_cusolver using DeviceContext - Utilize DeviceContext for node-local rank and device count - Remove redundant MPI_Comm_split_type and manual CUDA calls - Cleanup unused variables and communicator management * Refactor: remove redundant cudaGetDeviceCount in snap_psibeta_gpu - Remove manual device count check in initialize_gpu_resources - Rely on DeviceContext to handle device availability and binding - Simplify GPU initialization logic in module_rt * Refactor: rename initialize_gpu_resources to init_snap_psibeta_gpu - Rename function to reflect its module-specific scope - Update comments to clarify that general GPU setup is handled by DeviceContext - Remove unused finalize_gpu_resources declaration - Update caller in td_nonlocal_lcao.cpp * Refactor: separate kernel_compat.h from device.h Remove unconditional include of kernel_compat.h from device.h to properly separate Host/Device code. The kernel_compat.h header contains CUDA-specific code (__device__ keyword) and should only be included by .cu files that actually use atomicAdd. Add explicit includes to the 4 CUDA files that need it: - stress_op.cu - force_op.cu - exx_cal_energy_op.cu - phi_operator_kernel.cu Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com> * Refactor: move get_device_kpar logic to read_input_item_system Move GPU kpar validation from device module to kpar's reset_value in read_input_item_system.cpp. This keeps parameter validation logic with its related input item rather than in the device module. Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com> * Refactor: consolidate CUDA error checking macros into device_check.h - Update device_check.h with proper error handling (exit on error, stderr output) - Add cuSOLVER error string function with full IRS status codes - Add CHECK_LAST_CUDA_ERROR, CHECK_CUDA_SYNC, CHECK_CAL macros - Add ROCm CHECK_CUSOLVER support with hipsolver - Remove error checking code from module_container/base/macros/cuda.h - Remove error checking macros from helper_cuda.h - Delete helper_cusolver.h (functionality merged into device_check.h) - Update diag_cusolver.cu and diag_cusolvermp.cu to use new macros Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com> * Refactor: replace old CUDA error macros with CHECK_* macros - Replace cudaErrcheck -> CHECK_CUDA - Replace cublasErrcheck -> CHECK_CUBLAS - Replace cusolverErrcheck -> CHECK_CUSOLVER - Replace checkCudaErrors -> CHECK_CUDA - Replace CUSOLVER_CHECK -> CHECK_CUSOLVER - Replace CAL_CHECK -> CHECK_CAL - Replace cudaCheckOnDebug -> CHECK_CUDA_SYNC - Replace getLastCudaError -> CHECK_LAST_CUDA_ERROR - Update gpuErrcheck alias in gpu_runtime.h - Remove deprecated compatibility aliases from device_check.h Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com> * Refactor: remove CAL_CHECK compatibility alias from device_check.h Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com> * Refactor: unify cufftGetErrorString into device_check.h - Add _cufftGetErrorString static function to device_check.h - Remove dependency on cuda_compat.h for cufft error strings - Update CHECK_CUFFT macro to use the local function Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com> * Refactor: remove cufftGetErrorStringCompat from cuda_compat The function is now unified into device_check.h as _cufftGetErrorString. Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com> * Refactor: remove helper_cuda.h and helper_string.h - Replaced helper_cuda.h with device_check.h in hegvd_op.cu and diag_cusolvermp.cu - Removed helper_cuda.h and helper_string.h as they are no longer used Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com> * Refactor: unify DeviceContext::init() interface for MPI and non-MPI builds Remove conditional compilation from DeviceContext::init() signature. Now both MPI and non-MPI builds use the same void init() interface, with MPI_COMM_WORLD used internally in MPI builds. Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com> * Build: Fix CI failures for PyTest and CUDA test 1. PyTest: Add device_helpers.cpp to pyabacus ModuleBase build - Fixes undefined symbol: get_device_type<DEVICE_CPU> 2. CUDA test: Add __CUDA definition for cusolver test target - The test includes diag_cusolver.cu which needs CHECK_CUDA macros - remove_definitions(-D__CUDA) at top of file removed the macro - Re-add it specifically for this CUDA test target Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com> * Test: Fix DiagoCusolver test to match new diag() interface Update the test to use the new 4-parameter diag() interface: - Old: diag(&hmtest, psi, eigenvalue) - New: diag(h_mat, s_mat, psi, eigenvalue) The test now properly extracts MatrixBlock from HamiltTEST before calling diag(). Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com> * Refactor: replace custom CUDA error checks with unified device_check.h macros in module_gint Replace checkCuda/checkCudaLastError with CHECK_CUDA/CHECK_LAST_CUDA_ERROR from device_check.h to unify CUDA error handling across the codebase. Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com> * Refactor: replace remaining checkCuda with CHECK_CUDA in module_gint GPU files Continue refactoring to use unified device_check.h macros in all module_gint GPU implementation files. Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com> * Refactor: replace CUDA_CHECK with CHECK_CUDA in module_rt/kernels/cuda Remove custom CUDA_CHECK macro from snap_psibeta_kernel.cuh and use unified CHECK_CUDA from device_check.h for consistent CUDA error handling. Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com> * Fix CUDA pre-6.0 atomicAdd compat in vbatched GEMM * Add device_helpers to Makefile objects --------- Co-authored-by: Claude Opus 4.5 <noreply@anthropic.com>
1 parent 515e323 commit 14746d3

76 files changed

Lines changed: 1636 additions & 2991 deletions

File tree

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

‎python/pyabacus/src/ModuleBase/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5,6 +5,7 @@ list(APPEND pymodule_base
55
${BASE_PATH}/kernels/math_kernel_op_vec.cpp
66
${BASE_PATH}/module_device/memory_op.cpp
77
${BASE_PATH}/module_device/device.cpp
8+
${BASE_PATH}/module_device/device_helpers.cpp
89
)
910

1011
pybind11_add_module(_base_pack MODULE ${pymodule_base})

‎python/pyabacus/src/ModuleNAO/CMakeLists.txt‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
# add nao shared library
1+
# add nao shared library
22
list(APPEND _naos
33
${NAO_PATH}/atomic_radials.cpp
44
${NAO_PATH}/beta_radials.cpp
@@ -19,6 +19,7 @@ list(APPEND _naos
1919
# ${ABACUS_SOURCE_DIR}/source_psi/kernels/psi_memory_op.cpp
2020
${ABACUS_SOURCE_DIR}/source_base/module_device/memory_op.cpp
2121
${ABACUS_SOURCE_DIR}/source_base/module_device/device.cpp
22+
${ABACUS_SOURCE_DIR}/source_base/module_device/device_helpers.cpp
2223
)
2324
add_library(naopack SHARED
2425
${_naos}

‎python/pyabacus/src/hsolver/CMakeLists.txt‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -16,8 +16,9 @@ list(APPEND _diago
1616
${BASE_PATH}/kernels/math_kernel_op_vec.cpp
1717
${BASE_PATH}/kernels/math_ylm_op.cpp
1818
${BASE_PATH}/module_device/device.cpp
19+
${BASE_PATH}/module_device/device_helpers.cpp
1920
${BASE_PATH}/module_device/memory_op.cpp
20-
21+
2122
${PSI_PATH}/psi.cpp
2223
)
2324
add_library(diagopack SHARED

‎source/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -43,6 +43,7 @@ list(APPEND device_srcs
4343
# source_psi/kernels/device.cpp
4444

4545
source_base/module_device/device.cpp
46+
source_base/module_device/device_helpers.cpp
4647
source_base/module_device/output_device.cpp
4748
source_base/module_device/memory_op.cpp
4849
source_base/kernels/math_kernel_op.cpp

‎source/Makefile.Objects‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -176,6 +176,7 @@ OBJS_BASE=abfs-vector3_order.o\
176176
broyden_mixing.o\
177177
memory_op.o\
178178
device.o\
179+
device_helpers.o\
179180
output_device.o\
180181
parallel_2d.o\
181182

‎source/source_base/kernels/cuda/math_kernel_op.cu‎

Lines changed: 29 additions & 29 deletions
Original file line numberDiff line numberDiff line change
@@ -27,22 +27,22 @@ struct GetTypeThrust<std::complex<double>> {
2727
static cublasHandle_t cublas_handle = nullptr;
2828

2929
void xdot_wrapper(const int &n, const float * x, const int &incx, const float * y, const int &incy, float &result) {
30-
cublasErrcheck(cublasSdot(cublas_handle, n, x, incx, y, incy, &result));
30+
CHECK_CUBLAS(cublasSdot(cublas_handle, n, x, incx, y, incy, &result));
3131
}
3232

3333
void xdot_wrapper(const int &n, const double * x, const int &incx, const double * y, const int &incy, double &result) {
34-
cublasErrcheck(cublasDdot(cublas_handle, n, x, incx, y, incy, &result));
34+
CHECK_CUBLAS(cublasDdot(cublas_handle, n, x, incx, y, incy, &result));
3535
}
3636

3737
void createGpuBlasHandle(){
3838
if (cublas_handle == nullptr) {
39-
cublasErrcheck(cublasCreate(&cublas_handle));
39+
CHECK_CUBLAS(cublasCreate(&cublas_handle));
4040
}
4141
}
4242

4343
void destoryBLAShandle(){
4444
if (cublas_handle != nullptr) {
45-
cublasErrcheck(cublasDestroy(cublas_handle));
45+
CHECK_CUBLAS(cublasDestroy(cublas_handle));
4646
cublas_handle = nullptr;
4747
}
4848
}
@@ -58,7 +58,7 @@ void scal_op<float, base_device::DEVICE_GPU>::operator()(const int& N,
5858
std::complex<float>* X,
5959
const int& incx)
6060
{
61-
cublasErrcheck(cublasCscal(cublas_handle, N, (float2*)alpha, (float2*)X, incx));
61+
CHECK_CUBLAS(cublasCscal(cublas_handle, N, (float2*)alpha, (float2*)X, incx));
6262
}
6363

6464
template <>
@@ -67,7 +67,7 @@ void scal_op<double, base_device::DEVICE_GPU>::operator()(const int& N,
6767
std::complex<double>* X,
6868
const int& incx)
6969
{
70-
cublasErrcheck(cublasZscal(cublas_handle, N, (double2*)alpha, (double2*)X, incx));
70+
CHECK_CUBLAS(cublasZscal(cublas_handle, N, (double2*)alpha, (double2*)X, incx));
7171
}
7272

7373
template <>
@@ -78,7 +78,7 @@ void axpy_op<double, base_device::DEVICE_GPU>::operator()(const int& N,
7878
double* Y,
7979
const int& incY)
8080
{
81-
cublasErrcheck(cublasDaxpy(cublas_handle, N, alpha, X, incX, Y, incY));
81+
CHECK_CUBLAS(cublasDaxpy(cublas_handle, N, alpha, X, incX, Y, incY));
8282
}
8383

8484
template <>
@@ -89,7 +89,7 @@ void axpy_op<std::complex<float>, base_device::DEVICE_GPU>::operator()(const int
8989
std::complex<float>* Y,
9090
const int& incY)
9191
{
92-
cublasErrcheck(cublasCaxpy(cublas_handle, N, (float2*)alpha, (float2*)X, incX, (float2*)Y, incY));
92+
CHECK_CUBLAS(cublasCaxpy(cublas_handle, N, (float2*)alpha, (float2*)X, incX, (float2*)Y, incY));
9393
}
9494

9595
template <>
@@ -100,7 +100,7 @@ void axpy_op<std::complex<double>, base_device::DEVICE_GPU>::operator()(const in
100100
std::complex<double>* Y,
101101
const int& incY)
102102
{
103-
cublasErrcheck(cublasZaxpy(cublas_handle, N, (double2*)alpha, (double2*)X, incX, (double2*)Y, incY));
103+
CHECK_CUBLAS(cublasZaxpy(cublas_handle, N, (double2*)alpha, (double2*)X, incX, (double2*)Y, incY));
104104
}
105105

106106

@@ -175,7 +175,7 @@ void gemv_op<double, base_device::DEVICE_GPU>::operator()(const char& trans,
175175
const int& incy)
176176
{
177177
cublasOperation_t cutrans = judge_trans_op(false, trans, "gemv_op");
178-
cublasErrcheck(cublasDgemv(cublas_handle, cutrans, m, n, alpha, A, lda, X, incx, beta, Y, incx));
178+
CHECK_CUBLAS(cublasDgemv(cublas_handle, cutrans, m, n, alpha, A, lda, X, incx, beta, Y, incx));
179179
}
180180

181181
template <>
@@ -194,7 +194,7 @@ void gemv_op<std::complex<float>, base_device::DEVICE_GPU>::operator()(const cha
194194
cublasOperation_t cutrans = judge_trans_op(true, trans, "gemv_op");
195195
cuFloatComplex alpha = make_cuFloatComplex(alpha_in->real(), alpha_in->imag());
196196
cuFloatComplex beta = make_cuFloatComplex(beta_in->real(), beta_in->imag());
197-
cublasErrcheck(cublasCgemv(cublas_handle, cutrans, m, n, &alpha, (cuFloatComplex*)A, lda, (cuFloatComplex*)X, incx, &beta, (cuFloatComplex*)Y, incx));
197+
CHECK_CUBLAS(cublasCgemv(cublas_handle, cutrans, m, n, &alpha, (cuFloatComplex*)A, lda, (cuFloatComplex*)X, incx, &beta, (cuFloatComplex*)Y, incx));
198198
}
199199

200200
template <>
@@ -215,7 +215,7 @@ void gemv_op<std::complex<double>, base_device::DEVICE_GPU>::operator()(const ch
215215
cuDoubleComplex beta = make_cuDoubleComplex(beta_in->real(), beta_in->imag());
216216
// icpc and nvcc have some compatible problems
217217
// We must use cuDoubleComplex instead of converting std::complex<double>* to cuDoubleComplex*
218-
cublasErrcheck(cublasZgemv(cublas_handle, cutrans, m, n, &alpha, (cuDoubleComplex*)A, lda, (cuDoubleComplex*)X, incx, &beta, (cuDoubleComplex*)Y, incx));
218+
CHECK_CUBLAS(cublasZgemv(cublas_handle, cutrans, m, n, &alpha, (cuDoubleComplex*)A, lda, (cuDoubleComplex*)X, incx, &beta, (cuDoubleComplex*)Y, incx));
219219
}
220220

221221
template <>
@@ -235,7 +235,7 @@ void gemm_op<float, base_device::DEVICE_GPU>::operator()(const char& transa,
235235
{
236236
cublasOperation_t cutransA = judge_trans_op(false, transa, "gemm_op");
237237
cublasOperation_t cutransB = judge_trans_op(false, transb, "gemm_op");
238-
cublasErrcheck(cublasSgemm(cublas_handle, cutransA, cutransB, m, n, k, alpha, a, lda, b, ldb, beta, c, ldc));
238+
CHECK_CUBLAS(cublasSgemm(cublas_handle, cutransA, cutransB, m, n, k, alpha, a, lda, b, ldb, beta, c, ldc));
239239
}
240240

241241
template <>
@@ -255,7 +255,7 @@ void gemm_op<double, base_device::DEVICE_GPU>::operator()(const char& transa,
255255
{
256256
cublasOperation_t cutransA = judge_trans_op(false, transa, "gemm_op");
257257
cublasOperation_t cutransB = judge_trans_op(false, transb, "gemm_op");
258-
cublasErrcheck(cublasDgemm(cublas_handle, cutransA, cutransB, m, n, k, alpha, a, lda, b, ldb, beta, c, ldc));
258+
CHECK_CUBLAS(cublasDgemm(cublas_handle, cutransA, cutransB, m, n, k, alpha, a, lda, b, ldb, beta, c, ldc));
259259
}
260260
template <>
261261
void gemm_op<std::complex<float>, base_device::DEVICE_GPU>::operator()(const char& transa,
@@ -274,7 +274,7 @@ void gemm_op<std::complex<float>, base_device::DEVICE_GPU>::operator()(const cha
274274
{
275275
cublasOperation_t cutransA = judge_trans_op(true, transa, "gemm_op");
276276
cublasOperation_t cutransB = judge_trans_op(true, transb, "gemm_op");
277-
cublasErrcheck(cublasCgemm(cublas_handle, cutransA, cutransB, m, n ,k, (float2*)alpha, (float2*)a , lda, (float2*)b, ldb, (float2*)beta, (float2*)c, ldc));
277+
CHECK_CUBLAS(cublasCgemm(cublas_handle, cutransA, cutransB, m, n ,k, (float2*)alpha, (float2*)a , lda, (float2*)b, ldb, (float2*)beta, (float2*)c, ldc));
278278
}
279279

280280
template <>
@@ -294,7 +294,7 @@ void gemm_op<std::complex<double>, base_device::DEVICE_GPU>::operator()(const ch
294294
{
295295
cublasOperation_t cutransA = judge_trans_op(true, transa, "gemm_op");
296296
cublasOperation_t cutransB = judge_trans_op(true, transb, "gemm_op");
297-
cublasErrcheck(cublasZgemm(cublas_handle, cutransA, cutransB, m, n ,k, (double2*)alpha, (double2*)a , lda, (double2*)b, ldb, (double2*)beta, (double2*)c, ldc));
297+
CHECK_CUBLAS(cublasZgemm(cublas_handle, cutransA, cutransB, m, n ,k, (double2*)alpha, (double2*)a , lda, (double2*)b, ldb, (double2*)beta, (double2*)c, ldc));
298298
}
299299

300300
template <>
@@ -311,15 +311,15 @@ void matrixTranspose_op<double, base_device::DEVICE_GPU>::operator()(const int&
311311
double ONE = 1.0, ZERO = 0.0;
312312

313313
// use 'geam' API todo transpose.
314-
cublasErrcheck(cublasDgeam(cublas_handle, CUBLAS_OP_T, CUBLAS_OP_N, col, row, &ONE, input_matrix, col, &ZERO, input_matrix, col, device_temp, col));
314+
CHECK_CUBLAS(cublasDgeam(cublas_handle, CUBLAS_OP_T, CUBLAS_OP_N, col, row, &ONE, input_matrix, col, &ZERO, input_matrix, col, device_temp, col));
315315
}
316316
else
317317
{
318318
int thread = 1024;
319319
int block = (row + col + thread - 1) / thread;
320320
matrix_transpose_kernel<double> <<<block, thread >>> (row, col, input_matrix, device_temp);
321321

322-
cudaCheckOnDebug();
322+
CHECK_CUDA_SYNC();
323323
}
324324

325325
base_device::memory::synchronize_memory_op<double, base_device::DEVICE_GPU, base_device::DEVICE_GPU>()(
@@ -348,7 +348,7 @@ void matrixTranspose_op<std::complex<float>, base_device::DEVICE_GPU>::operator(
348348
ZERO.x = ZERO.y = 0.0;
349349

350350
// use 'geam' API todo transpose.
351-
cublasErrcheck(cublasCgeam(cublas_handle, CUBLAS_OP_T, CUBLAS_OP_N, col, row,
351+
CHECK_CUBLAS(cublasCgeam(cublas_handle, CUBLAS_OP_T, CUBLAS_OP_N, col, row,
352352
reinterpret_cast<const float2 *>(&ONE), (float2*)input_matrix, col,
353353
reinterpret_cast<const float2 *>(&ZERO), (float2*)input_matrix, col, (float2*)device_temp, col));
354354
} else
@@ -357,7 +357,7 @@ void matrixTranspose_op<std::complex<float>, base_device::DEVICE_GPU>::operator(
357357
int block = (row + col + thread - 1) / thread;
358358
matrix_transpose_kernel<thrust::complex<float>> <<<block, thread >>> (row, col, (thrust::complex<float>*)input_matrix, (thrust::complex<float>*)device_temp);
359359

360-
cudaCheckOnDebug();
360+
CHECK_CUDA_SYNC();
361361
}
362362

363363
base_device::memory::synchronize_memory_op<std::complex<float>, base_device::DEVICE_GPU, base_device::DEVICE_GPU>()(
@@ -367,7 +367,7 @@ void matrixTranspose_op<std::complex<float>, base_device::DEVICE_GPU>::operator(
367367

368368
base_device::memory::delete_memory_op<std::complex<float>, base_device::DEVICE_GPU>()(device_temp);
369369

370-
cudaCheckOnDebug();
370+
CHECK_CUDA_SYNC();
371371

372372
}
373373

@@ -389,13 +389,13 @@ void matrixTranspose_op<std::complex<double>, base_device::DEVICE_GPU>::operator
389389
ZERO.x = ZERO.y = 0.0;
390390

391391
// use 'geam' API todo transpose.
392-
cublasErrcheck(cublasZgeam(cublas_handle, CUBLAS_OP_T, CUBLAS_OP_N, col, row, &ONE, (double2*)input_matrix, col, &ZERO, (double2*)input_matrix, col, (double2*)device_temp, col));
392+
CHECK_CUBLAS(cublasZgeam(cublas_handle, CUBLAS_OP_T, CUBLAS_OP_N, col, row, &ONE, (double2*)input_matrix, col, &ZERO, (double2*)input_matrix, col, (double2*)device_temp, col));
393393
} else
394394
{
395395
int thread = 1024;
396396
int block = (row + col + thread - 1) / thread;
397397
matrix_transpose_kernel<thrust::complex<double>> <<<block, thread >>> (row, col, (thrust::complex<double>*)input_matrix, (thrust::complex<double>*)device_temp);
398-
cudaCheckOnDebug();
398+
CHECK_CUDA_SYNC();
399399
}
400400

401401
base_device::memory::synchronize_memory_op<std::complex<double>,
@@ -416,7 +416,7 @@ void matrixCopy<double, base_device::DEVICE_GPU>::operator()(const int& n1,
416416
const dim3 blockSize(16, 16);
417417
const dim3 gridSize((n1 + blockSize.x - 1) / blockSize.x, (n2 + blockSize.y - 1) / blockSize.y);
418418
matrix_copy_kernel<double> <<<gridSize, blockSize >>> (n1, n2, A, LDA, B, LDB);
419-
cudaCheckOnDebug();
419+
CHECK_CUDA_SYNC();
420420
}
421421
template <>
422422
void matrixCopy<std::complex<float>, base_device::DEVICE_GPU>::operator()(const int& n1,
@@ -429,7 +429,7 @@ void matrixCopy<std::complex<float>, base_device::DEVICE_GPU>::operator()(const
429429
const dim3 blockSize(16, 16);
430430
const dim3 gridSize((n1 + blockSize.x - 1) / blockSize.x, (n2 + blockSize.y - 1) / blockSize.y);
431431
matrix_copy_kernel<thrust::complex<float>> <<<gridSize, blockSize >>> (n1, n2, reinterpret_cast<const thrust::complex<float>*>(A), LDA, reinterpret_cast<thrust::complex<float>*>(B), LDB);
432-
cudaCheckOnDebug();
432+
CHECK_CUDA_SYNC();
433433

434434
}
435435
template <>
@@ -443,7 +443,7 @@ void matrixCopy<std::complex<double>, base_device::DEVICE_GPU>::operator()(const
443443
const dim3 blockSize(16, 16);
444444
const dim3 gridSize((n1 + blockSize.x - 1) / blockSize.x, (n2 + blockSize.y - 1) / blockSize.y);
445445
matrix_copy_kernel<thrust::complex<double>> <<<gridSize, blockSize >>> (n1, n2, reinterpret_cast<const thrust::complex<double>*>(A), LDA, reinterpret_cast<thrust::complex<double>*>(B), LDB);
446-
cudaCheckOnDebug();
446+
CHECK_CUDA_SYNC();
447447
}
448448

449449
template <>
@@ -453,7 +453,7 @@ void matrix_mul_vector_op<double, base_device::DEVICE_GPU>::operator()(const int
453453
dim3 block((m + thread.x - 1) / thread.x, (n + thread.y - 1) / thread.y, 1);
454454
matrix_multiply_vector_kernel<double, double> <<<block, thread >>>(m, n, a, lda,
455455
b, alpha, c, ldc);
456-
cudaCheckOnDebug();
456+
CHECK_CUDA_SYNC();
457457
}
458458

459459
template <>
@@ -463,7 +463,7 @@ void matrix_mul_vector_op<std::complex<float>, base_device::DEVICE_GPU>::operato
463463
dim3 block((m + thread.x - 1) / thread.x, (n + thread.y - 1) / thread.y, 1);
464464
matrix_multiply_vector_kernel<thrust::complex<float>, float> <<<block, thread >>>(m, n, reinterpret_cast<thrust::complex<float>*>(a), lda,
465465
b, alpha, reinterpret_cast<thrust::complex<float>*>(c), ldc);
466-
cudaCheckOnDebug();
466+
CHECK_CUDA_SYNC();
467467
}
468468

469469
template <>
@@ -474,7 +474,7 @@ void matrix_mul_vector_op<std::complex<double>, base_device::DEVICE_GPU>::operat
474474
dim3 block((m + thread.x - 1) / thread.x, (n + thread.y - 1) / thread.y, 1);
475475
matrix_multiply_vector_kernel<thrust::complex<double>, double> <<<block, thread >>>(m, n, reinterpret_cast<thrust::complex<double>*>(a), lda,
476476
b, alpha, reinterpret_cast<thrust::complex<double>*>(c), ldc);
477-
cudaCheckOnDebug();
477+
CHECK_CUDA_SYNC();
478478
}
479479

480480
// Explicitly instantiate functors for the types of functor registered.

‎source/source_base/kernels/cuda/math_kernel_op_vec.cu‎

Lines changed: 9 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -105,7 +105,7 @@ void vector_mul_real_op<double, base_device::DEVICE_GPU>::operator()(const int d
105105
int block = (dim + thread - 1) / thread;
106106
vector_mul_real_kernel<double><<<block, thread>>>(dim, result, vector, constant);
107107

108-
cudaCheckOnDebug();
108+
CHECK_CUDA_SYNC();
109109
}
110110

111111
template <typename FPTYPE>
@@ -121,7 +121,7 @@ inline void vector_mul_real_wrapper(const int dim,
121121
int block = (dim + thread - 1) / thread;
122122
vector_mul_real_kernel<thrust::complex<FPTYPE>><<<block, thread>>>(dim, result_tmp, vector_tmp, constant);
123123

124-
cudaCheckOnDebug();
124+
CHECK_CUDA_SYNC();
125125
}
126126
template <>
127127
void vector_mul_real_op<std::complex<float>, base_device::DEVICE_GPU>::operator()(const int dim,
@@ -152,7 +152,7 @@ void vector_div_constant_op<double, base_device::DEVICE_GPU>::operator()(const i
152152
int block = (dim + thread - 1) / thread;
153153
vector_div_constant_kernel<double><<<block, thread>>>(dim, result, vector, constant);
154154

155-
cudaCheckOnDebug();
155+
CHECK_CUDA_SYNC();
156156
}
157157

158158
template <typename FPTYPE>
@@ -168,7 +168,7 @@ inline void vector_div_constant_wrapper(const int& dim,
168168
int block = (dim + thread - 1) / thread;
169169
vector_div_constant_kernel<thrust::complex<FPTYPE>><<<block, thread>>>(dim, result_tmp, vector_tmp, constant);
170170

171-
cudaCheckOnDebug();
171+
CHECK_CUDA_SYNC();
172172
}
173173

174174
template <>
@@ -201,7 +201,7 @@ void vector_mul_vector_op<double, base_device::DEVICE_GPU>::operator()(const int
201201
int block = (dim + thread - 1) / thread;
202202
vector_mul_vector_kernel<double><<<block, thread>>>(dim, result, vector1, vector2, add);
203203

204-
cudaCheckOnDebug();
204+
CHECK_CUDA_SYNC();
205205
}
206206
// vector operator: result[i] = vector1[i](complex) * vector2[i](not complex)
207207
template <typename FPTYPE>
@@ -217,7 +217,7 @@ inline void vector_mul_vector_complex_wrapper(const int& dim,
217217
int block = (dim + thread - 1) / thread;
218218
vector_mul_vector_kernel<thrust::complex<FPTYPE>><<<block, thread>>>(dim, result_tmp, vector1_tmp, vector2, add);
219219

220-
cudaCheckOnDebug();
220+
CHECK_CUDA_SYNC();
221221
}
222222
template <>
223223
void vector_mul_vector_op<std::complex<float>, base_device::DEVICE_GPU>::operator()(const int& dim,
@@ -250,7 +250,7 @@ void vector_div_vector_op<double, base_device::DEVICE_GPU>::operator()(const int
250250
int block = (dim + thread - 1) / thread;
251251
vector_div_vector_kernel<double><<<block, thread>>>(dim, result, vector1, vector2);
252252

253-
cudaCheckOnDebug();
253+
CHECK_CUDA_SYNC();
254254
}
255255
// vector operator: result[i] = vector1[i](complex) / vector2[i](not complex)
256256
template <typename FPTYPE>
@@ -265,7 +265,7 @@ inline void vector_div_vector_complex_wrapper(const int& dim,
265265
int block = (dim + thread - 1) / thread;
266266
vector_div_vector_kernel<thrust::complex<FPTYPE>><<<block, thread>>>(dim, result_tmp, vector1_tmp, vector2);
267267

268-
cudaCheckOnDebug();
268+
CHECK_CUDA_SYNC();
269269
}
270270
template <>
271271
void vector_div_vector_op<std::complex<float>, base_device::DEVICE_GPU>::operator()(const int& dim,
@@ -306,7 +306,7 @@ void vector_add_vector_op<T, base_device::DEVICE_GPU>::operator()(const int& dim
306306
constantvector_addORsub_constantVector_kernel<Type, Real>
307307
<<<block, thread>>>(dim, result_tmp, vector1_tmp, constant1, vector2_tmp, constant2);
308308

309-
cudaCheckOnDebug();
309+
CHECK_CUDA_SYNC();
310310
}
311311

312312
template <>

‎source/source_base/kernels/cuda/math_ylm_op.cu‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -148,7 +148,7 @@ void cal_ylm_real_op<FPTYPE, base_device::DEVICE_GPU>::operator()(const base_dev
148148
p,
149149
ylm);
150150

151-
cudaCheckOnDebug();
151+
CHECK_CUDA_SYNC();
152152
}
153153

154154
template struct cal_ylm_real_op<float, base_device::DEVICE_GPU>;

0 commit comments

Comments
 (0)