From 106ceeb26889fb58d03bcbf79923791b504fbb8a Mon Sep 17 00:00:00 2001 From: Qianruipku Date: Thu, 14 Nov 2024 17:17:06 +0800 Subject: [PATCH] Fix: compile cuda without openmp --- source/module_hamilt_lcao/module_gint/gint_force_gpu.cu | 9 ++++++++- source/module_hamilt_lcao/module_gint/gint_rho_gpu.cu | 8 ++++++++ source/module_hamilt_lcao/module_gint/gint_vl_gpu.cu | 8 ++++++++ 3 files changed, 24 insertions(+), 1 deletion(-) diff --git a/source/module_hamilt_lcao/module_gint/gint_force_gpu.cu b/source/module_hamilt_lcao/module_gint/gint_force_gpu.cu index ac7f9e89c8..d15dc8416a 100644 --- a/source/module_hamilt_lcao/module_gint/gint_force_gpu.cu +++ b/source/module_hamilt_lcao/module_gint/gint_force_gpu.cu @@ -1,4 +1,6 @@ +#ifdef _OPENMP #include +#endif #include "gint_force_gpu.h" #include "kernels/cuda/cuda_tools.cuh" @@ -87,8 +89,9 @@ void gint_fvl_gpu(const hamilt::HContainer* dm, dm->get_wrapper(), dm->get_nnr() * sizeof(double), cudaMemcpyHostToDevice)); - +#ifdef _OPENMP #pragma omp parallel for num_threads(num_streams) collapse(2) +#endif for (int i = 0; i < gridt.nbx; i++) { for (int j = 0; j < gridt.nby; j++) @@ -96,7 +99,11 @@ void gint_fvl_gpu(const hamilt::HContainer* dm, // 20240620 Note that it must be set again here because // cuda's device is not safe in a multi-threaded environment. checkCuda(cudaSetDevice(gridt.dev_id)); +#ifdef _OPENMP const int sid = omp_get_thread_num(); +#else + const int sid = 0; +#endif int max_m = 0; int max_n = 0; diff --git a/source/module_hamilt_lcao/module_gint/gint_rho_gpu.cu b/source/module_hamilt_lcao/module_gint/gint_rho_gpu.cu index 4b18d50438..0fb6accad4 100644 --- a/source/module_hamilt_lcao/module_gint/gint_rho_gpu.cu +++ b/source/module_hamilt_lcao/module_gint/gint_rho_gpu.cu @@ -4,7 +4,9 @@ #include "gint_tools.h" #include "kernels/cuda/gint_rho.cuh" +#ifdef _OPENMP #include +#endif namespace GintKernel { @@ -68,7 +70,9 @@ void gint_rho_gpu(const hamilt::HContainer* dm, cudaMemcpyHostToDevice)); // calculate the rho for every nbzp bigcells +#ifdef _OPENMP #pragma omp parallel for num_threads(num_streams) collapse(2) +#endif for (int i = 0; i < gridt.nbx; i++) { for (int j = 0; j < gridt.nby; j++) @@ -78,7 +82,11 @@ void gint_rho_gpu(const hamilt::HContainer* dm, checkCuda(cudaSetDevice(gridt.dev_id)); // get stream id +#ifdef _OPENMP const int sid = omp_get_thread_num(); +#else + const int sid = 0; +#endif int max_m = 0; int max_n = 0; diff --git a/source/module_hamilt_lcao/module_gint/gint_vl_gpu.cu b/source/module_hamilt_lcao/module_gint/gint_vl_gpu.cu index 145f37417d..e9b5806a67 100644 --- a/source/module_hamilt_lcao/module_gint/gint_vl_gpu.cu +++ b/source/module_hamilt_lcao/module_gint/gint_vl_gpu.cu @@ -1,4 +1,6 @@ +#ifdef _OPENMP #include +#endif #include "kernels/cuda/cuda_tools.cuh" #include "module_base/ylm.h" @@ -71,7 +73,9 @@ void gint_vl_gpu(hamilt::HContainer* hRGint, Cuda_Mem_Wrapper gemm_B(max_atompair_per_z, num_streams, true); Cuda_Mem_Wrapper gemm_C(max_atompair_per_z, num_streams, true); +#ifdef _OPENMP #pragma omp parallel for num_threads(num_streams) collapse(2) +#endif for (int i = 0; i < gridt.nbx; i++) { for (int j = 0; j < gridt.nby; j++) @@ -79,7 +83,11 @@ void gint_vl_gpu(hamilt::HContainer* hRGint, // 20240620 Note that it must be set again here because // cuda's device is not safe in a multi-threaded environment. checkCuda(cudaSetDevice(gridt.dev_id)); +#ifdef _OPENMP const int sid = omp_get_thread_num(); +#else + const int sid = 0; +#endif int max_m = 0; int max_n = 0;