reduce header dependency

Author: dzzz2001

source/module_hamilt_lcao/module_gint/kernels/cuda/cuda_tools.cu

@@ -4,26 +4,6 @@
 
 #include "cuda_tools.cuh"
 
-cudaError_t check(cudaError_t result, const char *const func, const char *const file, const int line)
-{
-    if (result != cudaSuccess)
-    {
-        fprintf(stderr, "CUDA Runtime Error at %s:%d code=%s \"%s\" \n", file, line, cudaGetErrorString(result), func);
-        exit(EXIT_FAILURE);
-    }
-    return result;
-}
-cudaError_t __checkCudaLastError(const char *file, const int line)
-{
-    cudaError_t result = cudaGetLastError();
-    if (result != cudaSuccess)
-    {
-        fprintf(stderr, "%s(%i) : getLastCudaError():%s\n", file, line, cudaGetErrorString(result));
-        assert(result == cudaSuccess);
-    }
-    return result;
-}
-
 void dump_cuda_array_to_file(const double* cuda_array,
                              int width,
                              int hight,

source/module_hamilt_lcao/module_gint/kernels/cuda/cuda_tools.cuh

@@ -9,11 +9,32 @@
 #include <iostream>
 #include <sstream>
 
-#define checkCuda(val) check(val, #val, __FILE__, __LINE__)
-#define checkCudaLastError() __checkCudaLastError(__FILE__, __LINE__)
+#define checkCuda(val) check((val), #val, __FILE__, __LINE__)
+#define checkCudaLastError() __getLastCudaError(__FILE__, __LINE__)
 
-cudaError_t check(cudaError_t result, const char *const func, const char *const file, const int line);
-cudaError_t __checkCudaLastError(const char *file, const int line);
+inline void check(cudaError_t result, char const *const func, const char *const file,
+           int const line) {
+  if (result) {
+    fprintf(stderr, "CUDA error at %s:%d code=%d \"%s\" \n", file, line,
+            static_cast<unsigned int>(result), func);
+    exit(EXIT_FAILURE);
+  }
+}
+
+inline void __getLastCudaError(const char *file,
+                               const int line) 
+{
+  cudaError_t err = cudaGetLastError();
+
+  if (cudaSuccess != err) {
+    fprintf(stderr,
+            "%s(%i) : getLastCudaError() CUDA error :"
+            " (%d) %s.\n",
+            file, line, static_cast<int>(err),
+            cudaGetErrorString(err));
+    exit(EXIT_FAILURE);
+  }
+}
 
 static inline int ceildiv(int x, int y)
 {

source/module_hamilt_lcao/module_gint/temp_gint/gint_helper.h

@@ -7,18 +7,11 @@
 
 namespace ModuleGint
 {
-
-template <typename T>
-std::shared_ptr<const T> toConstSharedPtr(std::shared_ptr<T> ptr) {
-    return std::static_pointer_cast<const T>(ptr);
-}
-
-
 inline int index3Dto1D(const int id_x, const int id_y, const int id_z,
                         const int dim_x, const int dim_y, const int dim_z)
 {
     return id_z + id_y * dim_z + id_x * dim_y * dim_z;
-};
+}
 
 inline Vec3i index1Dto3D(const int index_1d,
                             const int dim_x, const int dim_y, const int dim_z)
@@ -27,7 +20,7 @@ inline Vec3i index1Dto3D(const int index_1d,
     int id_y = (index_1d - id_x * dim_y * dim_z) / dim_z;
     int id_z = index_1d % dim_z;
     return Vec3i(id_x, id_y, id_z);
-};
+}
 
 // if exponent is an integer between 0 and 5 (the most common cases in gint) and
 // and exp is a variable that cannot be determined at compile time (which means the compiler cannot optimize the code),
@@ -52,17 +45,17 @@ inline double pow_int(const double base, const int exp)
         double result = std::pow(base, exp);
         return result;
     }
-};
+}
 
 inline int floor_div(const int a, const int b)
 {
     // a ^ b < 0 means a and b have different signs
     return a / b - (a % b != 0 && (a ^ b) < 0);
-};
+}
 
 inline int ceil_div(const int a, const int b)
 {
     return a / b + (a % b != 0 && (a ^ b) > 0); 
-};
+}
 
 }

source/module_hamilt_lcao/module_gint/temp_gint/kernel/cuda_mem_wrapper.h

@@ -1,7 +1,7 @@
 #pragma once
 #include <cuda_runtime.h>
 #include "module_base/tool_quit.h"
-#include "module_hamilt_lcao/module_gint/kernels/cuda/cuda_tools.cuh"
+#include "gint_helper.cuh"
 
 template <typename T>
 class CudaMemWrapper

source/module_hamilt_lcao/module_gint/temp_gint/kernel/gemm_nn_vbatch.cuh

@@ -6,7 +6,7 @@
 #include <cuda_runtime.h>
 #include <stdio.h> // for fprintf and stderr
 
-#include "module_hamilt_lcao/module_gint/kernels/cuda/cuda_tools.cuh"
+#include "gint_helper.cuh"
 #include <functional>
 
 
@@ -395,8 +395,8 @@ void vbatched_gemm_nn_impl(int max_m,
     for (int i = 0; i < batchCount; i += max_batch_count)
     {
         const int ibatch = min(max_batch_count, batchCount - i);
-        dim3 dimGrid(ceildiv(max_n, BLK_M),
-                     ceildiv(max_m, BLK_N),
+        dim3 dimGrid(ceil_div(max_n, BLK_M),
+                     ceil_div(max_m, BLK_N),
                      ibatch);
         const T* alpha_tmp = nullptr;
         if (alpha != nullptr)

source/module_hamilt_lcao/module_gint/temp_gint/kernel/gemm_tn_vbatch.cuh

@@ -6,7 +6,7 @@
 #include <cuda_runtime.h>
 #include <stdio.h> // for fprintf and stderr
 
-#include "module_hamilt_lcao/module_gint/kernels/cuda/cuda_tools.cuh"
+#include "gint_helper.cuh"
 #include <functional>
 
 
@@ -420,8 +420,8 @@ void vbatched_gemm_tn_impl(int max_m,
     for (int i = 0; i < batchCount; i += max_batch_count)
     {
         const int ibatch = min(max_batch_count, batchCount - i);
-        dim3 dimGrid(ceildiv(max_n, BLK_M),
-                     ceildiv(max_m, BLK_N),
+        dim3 dimGrid(ceil_div(max_n, BLK_M),
+                     ceil_div(max_m, BLK_N),
                      ibatch);
         const T* alpha_tmp = nullptr;
         if (alpha != nullptr)

source/module_hamilt_lcao/module_gint/temp_gint/kernel/gint_gpu_vars.h

@@ -6,7 +6,7 @@
 #include "module_cell/unitcell.h"
 #include "module_cell/atom_spec.h"
 #include "module_hamilt_lcao/module_gint/temp_gint/biggrid_info.h"
-#include "module_hamilt_lcao/module_gint/kernels/cuda/cuda_tools.cuh"
+#include "gint_helper.cuh"
 #include "module_hamilt_lcao/module_gint/kernels/cuda/gemm_selector.cuh"
 
 namespace ModuleGint

source/module_hamilt_lcao/module_gint/temp_gint/kernel/gint_helper.cuh

@@ -1,4 +1,5 @@
 #pragma once
+#include <cstdio>
 
 // if exponent is an integer between 0 and 5 (the most common cases in gint) and
 // and exp is a variable that cannot be determined at compile time (which means the compiler cannot optimize the code),
@@ -24,7 +25,7 @@ __forceinline__ __device__ T pow_int(const T base, const int exp)
         double result = std::pow(base, exp);
         return result;
     }
-};
+}
 
 template<typename T>
 __forceinline__ __device__ T warpReduceSum(T val)
@@ -35,4 +36,40 @@ __forceinline__ __device__ T warpReduceSum(T val)
     val += __shfl_xor_sync(0xffffffff, val, 2, 32);
     val += __shfl_xor_sync(0xffffffff, val, 1, 32);
     return val;
-}
+}
+
+inline int ceil_div(const int a, const int b)
+{
+    return a / b + (a % b != 0 && (a ^ b) > 0); 
+}
+
+inline void check(cudaError_t result, char const *const func, const char *const file,
+           int const line) {
+  if (result) {
+    fprintf(stderr, "CUDA error at %s:%d code=%d \"%s\" \n", file, line,
+            static_cast<unsigned int>(result), func);
+    exit(EXIT_FAILURE);
+  }
+}
+
+inline void __getLastCudaError(const char *file,
+                               const int line) 
+{
+  cudaError_t err = cudaGetLastError();
+
+  if (cudaSuccess != err) {
+    fprintf(stderr,
+            "%s(%i) : getLastCudaError() CUDA error :"
+            " (%d) %s.\n",
+            file, line, static_cast<int>(err),
+            cudaGetErrorString(err));
+    exit(EXIT_FAILURE);
+  }
+}
+
+// This will output the proper CUDA error strings in the event
+// that a CUDA host call returns an error
+#define checkCuda(val) check((val), #val, __FILE__, __LINE__)
+
+// This will output the proper error string when calling cudaGetLastError
+#define checkCudaLastError() __getLastCudaError(__FILE__, __LINE__)

source/module_hamilt_lcao/module_gint/temp_gint/kernel/phi_operator_gpu.h

@@ -3,7 +3,7 @@
 #include <cuda_runtime.h>
 
 #include "module_hamilt_lcao/module_gint/temp_gint/batch_biggrid.h"
-#include "module_hamilt_lcao/module_gint/kernels/cuda/cuda_tools.cuh"
+#include "gint_helper.cuh"
 #include "gint_gpu_vars.h"
 #include "cuda_mem_wrapper.h"
 

source/module_hamilt_lcao/module_gint/temp_gint/kernel/set_const_mem.cu

@@ -1,5 +1,5 @@
 #include "set_const_mem.cuh"
-#include "module_hamilt_lcao/module_gint/kernels/cuda/cuda_tools.cuh"
+#include "gint_helper.cuh"
 
 __constant__ double ylmcoe_d[100];