diff --git a/attic/src/integrator/cuda/cuda_eval_denvars.cu b/attic/src/integrator/cuda/cuda_eval_denvars.cu index 0f1cd2f0c..c5f49ffa6 100644 --- a/attic/src/integrator/cuda/cuda_eval_denvars.cu +++ b/attic/src/integrator/cuda/cuda_eval_denvars.cu @@ -31,7 +31,7 @@ __global__ void eval_uvars_lda_kernel( size_t ntasks, const int tid_x = blockIdx.x * blockDim.x + threadIdx.x; const int tid_y = blockIdx.y * blockDim.y + threadIdx.y; - register double den_reg = 0.; + double den_reg = 0.; if( tid_x < nbf and tid_y < npts ) { @@ -115,10 +115,10 @@ __global__ void eval_uvars_gga_kernel( size_t ntasks, for (int sm_y = threadIdx.y; sm_y < GGA_KERNEL_SM_BLOCK_Y; sm_y += blockDim.y) { const int tid_y = bid_y + sm_y; - register double den_reg = den_shared[0][sm_y][threadIdx.x]; - register double dx_reg = den_shared[1][sm_y][threadIdx.x]; - register double dy_reg = den_shared[2][sm_y][threadIdx.x]; - register double dz_reg = den_shared[3][sm_y][threadIdx.x]; + double den_reg = den_shared[0][sm_y][threadIdx.x]; + double dx_reg = den_shared[1][sm_y][threadIdx.x]; + double dy_reg = den_shared[2][sm_y][threadIdx.x]; + double dz_reg = den_shared[3][sm_y][threadIdx.x]; // Warp blocks are stored col major den_reg = 2 * warpReduceSum( den_reg ); diff --git a/attic/src/integrator/cuda/cuda_weights.cu b/attic/src/integrator/cuda/cuda_weights.cu index 3013324b0..94ab86e9c 100644 --- a/attic/src/integrator/cuda/cuda_weights.cu +++ b/attic/src/integrator/cuda/cuda_weights.cu @@ -30,7 +30,7 @@ __global__ void compute_point_center_dist( ) { __shared__ double3 point_buffer[warp_size]; - register double3 coord_reg; + double3 coord_reg; const int natoms_block = (natoms + warp_size-1) / warp_size; const int coords_block = (npts + warp_size-1) / warp_size; diff --git a/attic/src/integrator/cuda/gauxc-cuda_integrator.cmake b/attic/src/integrator/cuda/gauxc-cuda_integrator.cmake index 864739b9b..20d957ce7 100644 --- a/attic/src/integrator/cuda/gauxc-cuda_integrator.cmake +++ b/attic/src/integrator/cuda/gauxc-cuda_integrator.cmake @@ -32,7 +32,7 @@ target_sources( gauxc PRIVATE cuda/collocation_device.cu cuda/cuda_device_properties.cxx ) -target_compile_features( gauxc PRIVATE cuda_std_14 ) +target_compile_features( gauxc PRIVATE cuda_std_17 ) #target_include_directories( gauxc # PRIVATE # $ diff --git a/attic/src/new_integrator/device/cuda/cuda_eval_denvars.cu b/attic/src/new_integrator/device/cuda/cuda_eval_denvars.cu index 813634370..54c17bb25 100644 --- a/attic/src/new_integrator/device/cuda/cuda_eval_denvars.cu +++ b/attic/src/new_integrator/device/cuda/cuda_eval_denvars.cu @@ -31,7 +31,7 @@ __global__ void eval_uvars_lda_kernel( size_t ntasks, const int tid_x = blockIdx.x * blockDim.x + threadIdx.x; const int tid_y = blockIdx.y * blockDim.y + threadIdx.y; - register double den_reg = 0.; + double den_reg = 0.; if( tid_x < nbf and tid_y < npts ) { @@ -115,10 +115,10 @@ __global__ void eval_uvars_gga_kernel( size_t ntasks, for (int sm_y = threadIdx.y; sm_y < GGA_KERNEL_SM_BLOCK_Y; sm_y += blockDim.y) { const int tid_y = bid_y + sm_y; - register double den_reg = den_shared[0][sm_y][threadIdx.x]; - register double dx_reg = den_shared[1][sm_y][threadIdx.x]; - register double dy_reg = den_shared[2][sm_y][threadIdx.x]; - register double dz_reg = den_shared[3][sm_y][threadIdx.x]; + double den_reg = den_shared[0][sm_y][threadIdx.x]; + double dx_reg = den_shared[1][sm_y][threadIdx.x]; + double dy_reg = den_shared[2][sm_y][threadIdx.x]; + double dz_reg = den_shared[3][sm_y][threadIdx.x]; // Warp blocks are stored col major den_reg = 2 * warpReduceSum( den_reg ); diff --git a/attic/src/new_integrator/device/cuda/cuda_weights.cu b/attic/src/new_integrator/device/cuda/cuda_weights.cu index f8da49bfc..98bae8268 100644 --- a/attic/src/new_integrator/device/cuda/cuda_weights.cu +++ b/attic/src/new_integrator/device/cuda/cuda_weights.cu @@ -30,7 +30,7 @@ __global__ void compute_point_center_dist( ) { __shared__ double3 point_buffer[warp_size]; - register double3 coord_reg; + double3 coord_reg; const int natoms_block = (natoms + warp_size-1) / warp_size; const int coords_block = (npts + warp_size-1) / warp_size; diff --git a/attic/src/new_integrator/device/cuda/gauxc-cuda.cmake b/attic/src/new_integrator/device/cuda/gauxc-cuda.cmake index e95000a37..e2bb0e328 100644 --- a/attic/src/new_integrator/device/cuda/gauxc-cuda.cmake +++ b/attic/src/new_integrator/device/cuda/gauxc-cuda.cmake @@ -22,7 +22,7 @@ target_sources( gauxc PRIVATE ) -target_compile_features( gauxc PRIVATE cuda_std_14 ) +target_compile_features( gauxc PRIVATE cuda_std_17 ) target_compile_options( gauxc PRIVATE $<$: -Xcudafe --diag_suppress=partial_override -Xptxas -v > diff --git a/src/runtime_environment/device/cuda/CMakeLists.txt b/src/runtime_environment/device/cuda/CMakeLists.txt index 50ea945c4..18c3039b5 100644 --- a/src/runtime_environment/device/cuda/CMakeLists.txt +++ b/src/runtime_environment/device/cuda/CMakeLists.txt @@ -13,7 +13,7 @@ if( NOT TARGET CUDA::cublas ) find_package( CUDAToolkit REQUIRED ) endif() -target_compile_features( gauxc PRIVATE cuda_std_14 ) +target_compile_features( gauxc PRIVATE cuda_std_17 ) target_sources( gauxc PRIVATE cuda_backend.cxx ) if(NOT GAUXC_LINK_CUDA_STATIC) diff --git a/src/xc_integrator/local_work_driver/device/cuda/kernels/grid_to_center.cu b/src/xc_integrator/local_work_driver/device/cuda/kernels/grid_to_center.cu index 03600010c..f90e2056d 100644 --- a/src/xc_integrator/local_work_driver/device/cuda/kernels/grid_to_center.cu +++ b/src/xc_integrator/local_work_driver/device/cuda/kernels/grid_to_center.cu @@ -28,7 +28,7 @@ __global__ void compute_grid_to_center_dist( ) { __shared__ double3 point_buffer[cuda::warp_size]; - register double3 coord_reg; + double3 coord_reg; const int natoms_block = (natoms + cuda::warp_size-1) / cuda::warp_size; const int coords_block = (npts + cuda::warp_size-1) / cuda::warp_size; diff --git a/src/xc_integrator/local_work_driver/device/cuda/kernels/uvvars_gga.hpp b/src/xc_integrator/local_work_driver/device/cuda/kernels/uvvars_gga.hpp index 9b466e24a..be703841c 100644 --- a/src/xc_integrator/local_work_driver/device/cuda/kernels/uvvars_gga.hpp +++ b/src/xc_integrator/local_work_driver/device/cuda/kernels/uvvars_gga.hpp @@ -131,10 +131,10 @@ __global__ void eval_vvar_gga_kern( size_t ntasks, for (int sm_y = threadIdx.y; sm_y < VVAR_KERNEL_SM_BLOCK; sm_y += blockDim.y) { const int tid_y = bid_y + sm_y; - register double den_reg = den_shared[0][sm_y][threadIdx.x]; - register double dx_reg = den_shared[1][sm_y][threadIdx.x]; - register double dy_reg = den_shared[2][sm_y][threadIdx.x]; - register double dz_reg = den_shared[3][sm_y][threadIdx.x]; + double den_reg = den_shared[0][sm_y][threadIdx.x]; + double dx_reg = den_shared[1][sm_y][threadIdx.x]; + double dy_reg = den_shared[2][sm_y][threadIdx.x]; + double dz_reg = den_shared[3][sm_y][threadIdx.x]; // Warp blocks are stored col major den_reg = cuda::warp_reduce_sum( den_reg ); diff --git a/src/xc_integrator/local_work_driver/device/cuda/kernels/uvvars_lda.hpp b/src/xc_integrator/local_work_driver/device/cuda/kernels/uvvars_lda.hpp index 54dc5043c..344acb6c6 100644 --- a/src/xc_integrator/local_work_driver/device/cuda/kernels/uvvars_lda.hpp +++ b/src/xc_integrator/local_work_driver/device/cuda/kernels/uvvars_lda.hpp @@ -49,7 +49,7 @@ __global__ void eval_vvar_lda_kern( size_t ntasks, const int tid_x = blockIdx.x * blockDim.x + threadIdx.x; const int tid_y = blockIdx.y * blockDim.y + threadIdx.y; - register double den_reg = 0.; + double den_reg = 0.; if( tid_x < nbf and tid_y < npts ) { diff --git a/src/xc_integrator/local_work_driver/device/cuda/kernels/uvvars_mgga.hpp b/src/xc_integrator/local_work_driver/device/cuda/kernels/uvvars_mgga.hpp index 82b5207e1..427a5923a 100644 --- a/src/xc_integrator/local_work_driver/device/cuda/kernels/uvvars_mgga.hpp +++ b/src/xc_integrator/local_work_driver/device/cuda/kernels/uvvars_mgga.hpp @@ -128,16 +128,16 @@ __global__ void eval_vvar_mgga_kern( size_t ntasks, for (int sm_y = threadIdx.y; sm_y < MGGA_KERNEL_SM_BLOCK; sm_y += blockDim.y) { const int tid_y = bid_y + sm_y; - register double tx_reg = den_shared[0][sm_y][threadIdx.x]; - register double ty_reg = den_shared[1][sm_y][threadIdx.x]; - register double tz_reg = den_shared[2][sm_y][threadIdx.x]; + double tx_reg = den_shared[0][sm_y][threadIdx.x]; + double ty_reg = den_shared[1][sm_y][threadIdx.x]; + double tz_reg = den_shared[2][sm_y][threadIdx.x]; // Warp blocks are stored col major - register double tau_reg = 0.0; + double tau_reg = 0.0; tau_reg = 0.5 * cuda::warp_reduce_sum( tx_reg ); tau_reg += 0.5 * cuda::warp_reduce_sum( ty_reg ); tau_reg += 0.5 * cuda::warp_reduce_sum( tz_reg ); - register double lapl_reg = 0.0; + double lapl_reg = 0.0; if constexpr (need_lapl) { lapl_reg = den_shared[3][sm_y][threadIdx.x]; lapl_reg = cuda::warp_reduce_sum(lapl_reg);