From 616615f334c56be1195c1fb480d90f8fc37bf1bf Mon Sep 17 00:00:00 2001 From: Ryan Stocks Date: Thu, 18 Jul 2024 15:56:46 +0800 Subject: [PATCH 1/6] HIP compilation --- src/hip/builtin.hip | 315 ++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 315 insertions(+) diff --git a/src/hip/builtin.hip b/src/hip/builtin.hip index 0b71482..8071265 100644 --- a/src/hip/builtin.hip +++ b/src/hip/builtin.hip @@ -399,6 +399,209 @@ __global__ GGA_EXC_VXC_INC_GENERATOR( device_eval_exc_vxc_inc_helper_polar_kerne } +template +__global__ MGGA_EXC_GENERATOR( device_eval_exc_helper_unpolar_kernel ) { + + using traits = kernel_traits; + int tid = threadIdx.x + blockIdx.x * blockDim.x; + + if( tid < N ) { + + const double lapl_use = traits::needs_laplacian ? lapl[tid] : 0.0; + traits::eval_exc_unpolar( rho[tid], sigma[tid], lapl_use, tau[tid], eps[tid] ); + + } + +} + + +template +__global__ MGGA_EXC_GENERATOR( device_eval_exc_helper_polar_kernel ) { + + using traits = kernel_traits; + int tid = threadIdx.x + blockIdx.x * blockDim.x; + + if( tid < N ) { + + auto* rho_i = rho + 2*tid; + auto* sigma_i = sigma + 3*tid; + auto* lapl_i = traits::needs_laplacian ? (lapl + 2*tid) : nullptr; + auto* tau_i = tau + 2*tid; + + const double lapl_a_use = traits::needs_laplacian ? lapl_i[0] : 0.0; + const double lapl_b_use = traits::needs_laplacian ? lapl_i[1] : 0.0; + + traits::eval_exc_polar( rho_i[0], rho_i[1], sigma_i[0], + sigma_i[1], sigma_i[2], lapl_a_use, lapl_b_use, tau_i[0], + tau_i[1], eps[tid] ); + + } + +} + +template +__global__ MGGA_EXC_VXC_GENERATOR( device_eval_exc_vxc_helper_unpolar_kernel ) { + + using traits = kernel_traits; + int tid = threadIdx.x + blockIdx.x * blockDim.x; + + if( tid < N ) { + + const double lapl_use = traits::needs_laplacian ? lapl[tid] : 0.0; + + double dummy; + auto& vlapl_return = traits::needs_laplacian ? vlapl[tid] : dummy; + traits::eval_exc_vxc_unpolar( rho[tid], sigma[tid], lapl_use, tau[tid], + eps[tid], vrho[tid], vsigma[tid], vlapl_return, vtau[tid] ); + + } + +} + +template +__global__ MGGA_EXC_VXC_GENERATOR( device_eval_exc_vxc_helper_polar_kernel ) { + + using traits = kernel_traits; + int tid = threadIdx.x + blockIdx.x * blockDim.x; + + double dummy_vlapl[2]; + + if( tid < N ) { + + auto* rho_i = rho + 2*tid; + auto* sigma_i = sigma + 3*tid; + auto* lapl_i = traits::needs_laplacian ? (lapl + 2*tid) : lapl; + auto* tau_i = tau + 2*tid; + + auto* vrho_i = vrho + 2*tid; + auto* vsigma_i = vsigma + 3*tid; + auto* vlapl_i = traits::needs_laplacian ? vlapl + 2*tid : dummy_vlapl; + auto* vtau_i = vtau + 2*tid; + const double lapl_a_use = traits::needs_laplacian ? lapl_i[0] : 0.0; + const double lapl_b_use = traits::needs_laplacian ? lapl_i[1] : 0.0; + + traits::eval_exc_vxc_polar( rho_i[0], rho_i[1], sigma_i[0], + sigma_i[1], sigma_i[2], lapl_a_use, lapl_b_use, tau_i[0], + tau_i[1], eps[tid], vrho_i[0], vrho_i[1], vsigma_i[0], vsigma_i[1], + vsigma_i[2], vlapl_i[0], vlapl_i[1], vtau_i[0], vtau_i[1] ); + + } + +} + + +template +__global__ MGGA_EXC_INC_GENERATOR( device_eval_exc_inc_helper_unpolar_kernel ) { + + using traits = kernel_traits; + int tid = threadIdx.x + blockIdx.x * blockDim.x; + + double e; + if( tid < N ) { + + const double lapl_use = traits::needs_laplacian ? lapl[tid] : 0.0; + traits::eval_exc_unpolar( rho[tid], sigma[tid], lapl_use, tau[tid], e ); + eps[tid] += scal_fact * e; + + + } + +} + +template +__global__ MGGA_EXC_INC_GENERATOR( device_eval_exc_inc_helper_polar_kernel ) { + + using traits = kernel_traits; + int tid = threadIdx.x + blockIdx.x * blockDim.x; + + if( tid < N ) { + + auto* rho_i = rho + 2*tid; + auto* sigma_i = sigma + 3*tid; + auto* lapl_i = traits::needs_laplacian ? (lapl + 2*tid) : lapl; + auto* tau_i = tau + 2*tid; + + const double lapl_a_use = traits::needs_laplacian ? lapl_i[0] : 0.0; + const double lapl_b_use = traits::needs_laplacian ? lapl_i[1] : 0.0; + + double e; + traits::eval_exc_polar( rho_i[0], rho_i[1], sigma_i[0], + sigma_i[1], sigma_i[2], lapl_a_use, lapl_b_use, tau_i[0], + tau_i[1], e ); + eps[tid] += scal_fact * e; + + + } + +} + +template +__global__ MGGA_EXC_VXC_INC_GENERATOR( device_eval_exc_vxc_inc_helper_unpolar_kernel ) { + + using traits = kernel_traits; + int tid = threadIdx.x + blockIdx.x * blockDim.x; + + double e, vr, vs, vl, vt; + if( tid < N ) { + + const double lapl_use = traits::needs_laplacian ? lapl[tid] : 0.0; + + traits::eval_exc_vxc_unpolar( rho[tid], sigma[tid], lapl_use, tau[tid], + e, vr, vs, vl, vt ); + eps[tid] += scal_fact * e; + vrho[tid] += scal_fact * vr; + vsigma[tid] += scal_fact * vs; + vtau[tid] += scal_fact * vt; + if(traits::needs_laplacian) vlapl[tid] += scal_fact * vl; + + } + +} + +template +__global__ MGGA_EXC_VXC_INC_GENERATOR( device_eval_exc_vxc_inc_helper_polar_kernel ) { + + using traits = kernel_traits; + int tid = threadIdx.x + blockIdx.x * blockDim.x; + + double dummy_vlapl[2]; + if( tid < N ) { + + auto* rho_i = rho + 2*tid; + auto* sigma_i = sigma + 3*tid; + auto* lapl_i = traits::needs_laplacian ? (lapl + 2*tid) : lapl; + auto* tau_i = tau + 2*tid; + + auto* vrho_i = vrho + 2*tid; + auto* vsigma_i = vsigma + 3*tid; + auto* vlapl_i = traits::needs_laplacian ? vlapl + 2*tid : dummy_vlapl; + auto* vtau_i = vtau + 2*tid; + + const double lapl_a_use = traits::needs_laplacian ? lapl_i[0] : 0.0; + const double lapl_b_use = traits::needs_laplacian ? lapl_i[1] : 0.0; + + + double e, vra, vrb, vsaa,vsab,vsbb, vla, vlb, vta, vtb; + traits::eval_exc_vxc_polar( rho_i[0], rho_i[1], sigma_i[0], + sigma_i[1], sigma_i[2], lapl_a_use, lapl_b_use, tau_i[0], + tau_i[1], e, vra, vrb, vsaa, vsab, vsbb, vla, vlb, vta, vtb ); + + eps[tid] += scal_fact * e; + vrho_i[0] += scal_fact * vra; + vrho_i[1] += scal_fact * vrb; + vsigma_i[0] += scal_fact * vsaa; + vsigma_i[1] += scal_fact * vsab; + vsigma_i[2] += scal_fact * vsbb; + vtau_i[0] += scal_fact * vta; + vtau_i[1] += scal_fact * vtb; + if(traits::needs_laplacian) { + vlapl_i[0] += scal_fact * vla; + vlapl_i[1] += scal_fact * vlb; + } + + } + +} template LDA_EXC_GENERATOR_DEVICE( device_eval_exc_helper_unpolar ) { @@ -582,6 +785,99 @@ GGA_EXC_VXC_INC_GENERATOR_DEVICE( device_eval_exc_vxc_inc_helper_polar ) { } +template +MGGA_EXC_GENERATOR_DEVICE( device_eval_exc_helper_unpolar ) { + + dim3 threads(32); + dim3 blocks( util::div_ceil( N, threads.x) ); + device_eval_exc_helper_unpolar_kernel<<>>( + N, rho, sigma, lapl, tau, eps + ); + +} + +template +MGGA_EXC_GENERATOR_DEVICE( device_eval_exc_helper_polar ) { + + dim3 threads(32); + dim3 blocks( util::div_ceil( N, threads.x) ); + device_eval_exc_helper_polar_kernel<<>>( + N, rho, sigma, lapl, tau, eps + ); + +} + +template +MGGA_EXC_VXC_GENERATOR_DEVICE( device_eval_exc_vxc_helper_unpolar ) { + + dim3 threads(32); + dim3 blocks( util::div_ceil( N, threads.x) ); + + device_eval_exc_vxc_helper_unpolar_kernel<<>>( + N, rho, sigma, lapl, tau, eps, vrho, vsigma, vlapl, vtau + ); + +} + +template +MGGA_EXC_VXC_GENERATOR_DEVICE( device_eval_exc_vxc_helper_polar ) { + + dim3 threads(32); + dim3 blocks( util::div_ceil( N, threads.x) ); + + device_eval_exc_vxc_helper_polar_kernel<<>>( + N, rho, sigma, lapl, tau, eps, vrho, vsigma, vlapl, vtau + ); + +} + + +template +MGGA_EXC_INC_GENERATOR_DEVICE( device_eval_exc_inc_helper_unpolar ) { + + dim3 threads(32); + dim3 blocks( util::div_ceil( N, threads.x) ); + device_eval_exc_inc_helper_unpolar_kernel<<>>( + scal_fact, N, rho, sigma, lapl, tau, eps + ); + +} + +template +MGGA_EXC_INC_GENERATOR_DEVICE( device_eval_exc_inc_helper_polar ) { + + dim3 threads(32); + dim3 blocks( util::div_ceil( N, threads.x) ); + device_eval_exc_inc_helper_polar_kernel<<>>( + scal_fact, N, rho, sigma, lapl, tau, eps + ); + +} + +template +MGGA_EXC_VXC_INC_GENERATOR_DEVICE( device_eval_exc_vxc_inc_helper_unpolar ) { + + dim3 threads(32); + dim3 blocks( util::div_ceil( N, threads.x) ); + + device_eval_exc_vxc_inc_helper_unpolar_kernel<<>>( + scal_fact, N, rho, sigma, lapl, tau, eps, vrho, vsigma, vlapl, vtau + ); + +} + +template +MGGA_EXC_VXC_INC_GENERATOR_DEVICE( device_eval_exc_vxc_inc_helper_polar ) { + + dim3 threads(32); + dim3 blocks( util::div_ceil( N, threads.x) ); + + device_eval_exc_vxc_inc_helper_polar_kernel<<>>( + scal_fact, N, rho, sigma, lapl, tau, eps, vrho, vsigma, vlapl, vtau + ); + +} + #define LDA_GENERATE_DEVICE_HELPERS(KERN) \ template LDA_EXC_GENERATOR_DEVICE( device_eval_exc_helper_unpolar ); \ template LDA_EXC_VXC_GENERATOR_DEVICE( device_eval_exc_vxc_helper_unpolar ); \ @@ -602,6 +898,16 @@ GGA_EXC_VXC_INC_GENERATOR_DEVICE( device_eval_exc_vxc_inc_helper_polar ) { template GGA_EXC_INC_GENERATOR_DEVICE( device_eval_exc_inc_helper_polar ); \ template GGA_EXC_VXC_INC_GENERATOR_DEVICE( device_eval_exc_vxc_inc_helper_polar ); +#define MGGA_GENERATE_DEVICE_HELPERS(KERN) \ + template MGGA_EXC_GENERATOR_DEVICE( device_eval_exc_helper_unpolar ); \ + template MGGA_EXC_VXC_GENERATOR_DEVICE( device_eval_exc_vxc_helper_unpolar ); \ + template MGGA_EXC_INC_GENERATOR_DEVICE( device_eval_exc_inc_helper_unpolar ); \ + template MGGA_EXC_VXC_INC_GENERATOR_DEVICE( device_eval_exc_vxc_inc_helper_unpolar );\ + template MGGA_EXC_GENERATOR_DEVICE( device_eval_exc_helper_polar ); \ + template MGGA_EXC_VXC_GENERATOR_DEVICE( device_eval_exc_vxc_helper_polar ); \ + template MGGA_EXC_INC_GENERATOR_DEVICE( device_eval_exc_inc_helper_polar ); \ + template MGGA_EXC_VXC_INC_GENERATOR_DEVICE( device_eval_exc_vxc_inc_helper_polar ); + LDA_GENERATE_DEVICE_HELPERS( BuiltinSlaterExchange ); LDA_GENERATE_DEVICE_HELPERS( BuiltinVWN3 ); LDA_GENERATE_DEVICE_HELPERS( BuiltinVWN_RPA ); @@ -624,6 +930,15 @@ MGGA_GENERATE_DEVICE_HELPERS( BuiltinSCAN_X ); MGGA_GENERATE_DEVICE_HELPERS( BuiltinSCAN_C ); MGGA_GENERATE_DEVICE_HELPERS( BuiltinR2SCAN_X ); MGGA_GENERATE_DEVICE_HELPERS( BuiltinR2SCAN_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinFT98_X ); + +MGGA_GENERATE_DEVICE_HELPERS( BuiltinPC07_K ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinPC07OPT_K ); + +MGGA_GENERATE_DEVICE_HELPERS( BuiltinSCANL_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinSCANL_X ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinR2SCANL_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinR2SCANL_X ); LDA_GENERATE_DEVICE_HELPERS( BuiltinEPC17_1 ) LDA_GENERATE_DEVICE_HELPERS( BuiltinEPC17_2 ) From cd04c1b724ddb5826542296a4adb46bb11c080f2 Mon Sep 17 00:00:00 2001 From: Ryan Stocks Date: Sat, 20 Jul 2024 15:22:52 +1000 Subject: [PATCH 2/6] Fixed CUDA tests failing due to unsupported unpolarized kernels --- .gitignore | 3 + include/exchcxx/enums/kernels.hpp | 13 +++++ src/builtin_interface.cxx | 10 ++-- test/xc_kernel_test.cxx | 92 +++++++++++++++++-------------- 4 files changed, 73 insertions(+), 45 deletions(-) diff --git a/.gitignore b/.gitignore index 0a37f33..9d25f2f 100644 --- a/.gitignore +++ b/.gitignore @@ -1 +1,4 @@ *.*.swp +.cache/ +build/ +compile_commands.json diff --git a/include/exchcxx/enums/kernels.hpp b/include/exchcxx/enums/kernels.hpp index d7eeefa..babb883 100644 --- a/include/exchcxx/enums/kernels.hpp +++ b/include/exchcxx/enums/kernels.hpp @@ -109,6 +109,19 @@ enum class Kernel { EPC18_2, }; +inline static bool supports_unpolarized(ExchCXX::Kernel kern) { + switch (kern) { + case ExchCXX::Kernel::EPC17_1: + case ExchCXX::Kernel::EPC17_2: + case ExchCXX::Kernel::EPC18_1: + case ExchCXX::Kernel::EPC18_2: + return false; + default: + return true; + } +} + + extern BidirectionalMap kernel_map; std::ostream& operator<<( std::ostream& out, Kernel kern ); diff --git a/src/builtin_interface.cxx b/src/builtin_interface.cxx index a470cf0..105dc7e 100644 --- a/src/builtin_interface.cxx +++ b/src/builtin_interface.cxx @@ -59,6 +59,11 @@ namespace detail { std::unique_ptr gen_from_kern( Kernel kern, Spin polar ) { + if (!supports_unpolarized(kern)) { + EXCHCXX_BOOL_CHECK(kernel_map.key(kern) + " Needs to be Spin-Polarized!", + polar == Spin::Polarized); + } + if( kern == Kernel::SlaterExchange ) return std::make_unique( polar ); else if( kern == Kernel::VWN3 ) @@ -119,22 +124,17 @@ std::unique_ptr return std::make_unique( polar ); else if( kern == Kernel::EPC17_1) { - EXCHCXX_BOOL_CHECK("EPC17_1 Needs to be Spin-Polarized!",polar==Spin::Polarized); return std::make_unique( polar ); } else if( kern == Kernel::EPC17_2) { - EXCHCXX_BOOL_CHECK("EPC17_2 Needs to be Spin-Polarized!",polar==Spin::Polarized); return std::make_unique( polar ); } else if( kern == Kernel::EPC18_1) { - EXCHCXX_BOOL_CHECK("EPC18_1 Needs to be Spin-Polarized!",polar==Spin::Polarized); return std::make_unique( polar ); } else if( kern == Kernel::EPC18_2) { - EXCHCXX_BOOL_CHECK("EPC18_2 Needs to be Spin-Polarized!",polar==Spin::Polarized); return std::make_unique( polar ); } else throw std::runtime_error("Specified kernel does not have a builtin implementation"); - } BuiltinKernelInterface::~BuiltinKernelInterface() noexcept = default; diff --git a/test/xc_kernel_test.cxx b/test/xc_kernel_test.cxx index 29d7d24..cca4782 100644 --- a/test/xc_kernel_test.cxx +++ b/test/xc_kernel_test.cxx @@ -1613,83 +1613,95 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { SECTION( "Builtin Functionals" ) { SECTION("EXC Regular: Unpolarized") { - for( auto kern : builtin_supported_kernels ) - test_cuda_interface( TestInterface::EXC, EvalType::Regular, - Backend::builtin, kern, Spin::Unpolarized ); + for (auto kern : builtin_supported_kernels) + if (supports_unpolarized(kern)) + test_cuda_interface(TestInterface::EXC, EvalType::Regular, + Backend::builtin, kern, Spin::Unpolarized); } SECTION("EXC + VXC Regular: Unpolarized") { - for( auto kern : builtin_supported_kernels ) - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Regular, - Backend::builtin, kern, Spin::Unpolarized ); + for (auto kern : builtin_supported_kernels) + if (supports_unpolarized(kern)) + test_cuda_interface(TestInterface::EXC_VXC, EvalType::Regular, + Backend::builtin, kern, Spin::Unpolarized); } SECTION("EXC + INC Regular: Unpolarized") { - for( auto kern : builtin_supported_kernels ) - test_cuda_interface( TestInterface::EXC_INC, EvalType::Regular, - Backend::builtin, kern, Spin::Unpolarized ); + for (auto kern : builtin_supported_kernels) + if (supports_unpolarized(kern)) + test_cuda_interface(TestInterface::EXC_INC, EvalType::Regular, + Backend::builtin, kern, Spin::Unpolarized); } SECTION("EXC + VXC + INC Regular: Unpolarized") { - for( auto kern : builtin_supported_kernels ) - test_cuda_interface( TestInterface::EXC_VXC_INC, EvalType::Regular, - Backend::builtin, kern, Spin::Unpolarized ); + for (auto kern : builtin_supported_kernels) + if (supports_unpolarized(kern)) + test_cuda_interface(TestInterface::EXC_VXC_INC, EvalType::Regular, + Backend::builtin, kern, Spin::Unpolarized); } SECTION("EXC Small: Unpolarized") { - for( auto kern : builtin_supported_kernels ) { - if(is_unstable_small(kern)) continue; - test_cuda_interface( TestInterface::EXC, EvalType::Small, - Backend::builtin, kern, Spin::Unpolarized ); + for (auto kern : builtin_supported_kernels) { + if (is_unstable_small(kern) || !supports_unpolarized(kern)) + continue; + test_cuda_interface(TestInterface::EXC, EvalType::Small, + Backend::builtin, kern, Spin::Unpolarized); } } SECTION("EXC + VXC Small: Unpolarized") { - for( auto kern : builtin_supported_kernels ) { - if(is_unstable_small(kern)) continue; - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Small, - Backend::builtin, kern, Spin::Unpolarized ); + for (auto kern : builtin_supported_kernels) { + if (is_unstable_small(kern) || !supports_unpolarized(kern)) + continue; + test_cuda_interface(TestInterface::EXC_VXC, EvalType::Small, + Backend::builtin, kern, Spin::Unpolarized); } } SECTION("EXC + INC Small: Unpolarized") { - for( auto kern : builtin_supported_kernels ) { - if(is_unstable_small(kern)) continue; - test_cuda_interface( TestInterface::EXC_INC, EvalType::Small, - Backend::builtin, kern, Spin::Unpolarized ); + for (auto kern : builtin_supported_kernels) { + if (is_unstable_small(kern) || !supports_unpolarized(kern)) + continue; + test_cuda_interface(TestInterface::EXC_INC, EvalType::Small, + Backend::builtin, kern, Spin::Unpolarized); } } SECTION("EXC + VXC + INC Small: Unpolarized") { - for( auto kern : builtin_supported_kernels ) { - if(is_unstable_small(kern)) continue; - test_cuda_interface( TestInterface::EXC_VXC_INC, EvalType::Small, - Backend::builtin, kern, Spin::Unpolarized ); + for (auto kern : builtin_supported_kernels) { + if (is_unstable_small(kern) || !supports_unpolarized(kern)) + continue; + test_cuda_interface(TestInterface::EXC_VXC_INC, EvalType::Small, + Backend::builtin, kern, Spin::Unpolarized); } } SECTION("EXC Zero: Unpolarized") { - for( auto kern : builtin_supported_kernels ) - test_cuda_interface( TestInterface::EXC, EvalType::Zero, - Backend::builtin, kern, Spin::Unpolarized ); + for (auto kern : builtin_supported_kernels) + if (supports_unpolarized(kern)) + test_cuda_interface(TestInterface::EXC, EvalType::Zero, + Backend::builtin, kern, Spin::Unpolarized); } SECTION("EXC + VXC Zero: Unpolarized") { - for( auto kern : builtin_supported_kernels ) - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Zero, - Backend::builtin, kern, Spin::Unpolarized ); + for (auto kern : builtin_supported_kernels) + if (supports_unpolarized(kern)) + test_cuda_interface(TestInterface::EXC_VXC, EvalType::Zero, + Backend::builtin, kern, Spin::Unpolarized); } SECTION("EXC + INC Zero: Unpolarized") { - for( auto kern : builtin_supported_kernels ) - test_cuda_interface( TestInterface::EXC_INC, EvalType::Zero, - Backend::builtin, kern, Spin::Unpolarized ); + for (auto kern : builtin_supported_kernels) + if (supports_unpolarized(kern)) + test_cuda_interface(TestInterface::EXC_INC, EvalType::Zero, + Backend::builtin, kern, Spin::Unpolarized); } SECTION("EXC + VXC + INC Zero: Unpolarized") { - for( auto kern : builtin_supported_kernels ) - test_cuda_interface( TestInterface::EXC_VXC_INC, EvalType::Zero, - Backend::builtin, kern, Spin::Unpolarized ); + for (auto kern : builtin_supported_kernels) + if (supports_unpolarized(kern)) + test_cuda_interface(TestInterface::EXC_VXC_INC, EvalType::Zero, + Backend::builtin, kern, Spin::Unpolarized); } SECTION("EXC Regular: Polarized") { From 07908a00b99a4924b704c422b1ee70a20bd841db Mon Sep 17 00:00:00 2001 From: Ryan Stocks Date: Sat, 20 Jul 2024 16:16:46 +1000 Subject: [PATCH 3/6] Cleaner checking unpolarized is supported --- test/xc_kernel_test.cxx | 58 ++++++++++++++++++----------------------- 1 file changed, 26 insertions(+), 32 deletions(-) diff --git a/test/xc_kernel_test.cxx b/test/xc_kernel_test.cxx index cca4782..d8ff648 100644 --- a/test/xc_kernel_test.cxx +++ b/test/xc_kernel_test.cxx @@ -43,6 +43,7 @@ * in binary and source code form. */ +#include "exchcxx/enums/kernels.hpp" #include "ut_common.hpp" using namespace ExchCXX; @@ -1121,6 +1122,10 @@ void test_cuda_interface( TestInterface interface, EvalType evaltype, const int npts = npts_lda; + if (polar == Spin::Unpolarized && !supports_unpolarized(kern)){ + CHECK_THROWS( XCKernel( backend, kern, polar ) ); + return; + } XCKernel func( backend, kern, polar ); size_t len_rho_buffer = func.rho_buffer_len(npts); @@ -1610,39 +1615,35 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { } - SECTION( "Builtin Functionals" ) { + SECTION("Builtin Functionals") { SECTION("EXC Regular: Unpolarized") { for (auto kern : builtin_supported_kernels) - if (supports_unpolarized(kern)) - test_cuda_interface(TestInterface::EXC, EvalType::Regular, - Backend::builtin, kern, Spin::Unpolarized); + test_cuda_interface(TestInterface::EXC, EvalType::Regular, + Backend::builtin, kern, Spin::Unpolarized); } SECTION("EXC + VXC Regular: Unpolarized") { for (auto kern : builtin_supported_kernels) - if (supports_unpolarized(kern)) - test_cuda_interface(TestInterface::EXC_VXC, EvalType::Regular, - Backend::builtin, kern, Spin::Unpolarized); + test_cuda_interface(TestInterface::EXC_VXC, EvalType::Regular, + Backend::builtin, kern, Spin::Unpolarized); } SECTION("EXC + INC Regular: Unpolarized") { for (auto kern : builtin_supported_kernels) - if (supports_unpolarized(kern)) - test_cuda_interface(TestInterface::EXC_INC, EvalType::Regular, - Backend::builtin, kern, Spin::Unpolarized); + test_cuda_interface(TestInterface::EXC_INC, EvalType::Regular, + Backend::builtin, kern, Spin::Unpolarized); } SECTION("EXC + VXC + INC Regular: Unpolarized") { for (auto kern : builtin_supported_kernels) - if (supports_unpolarized(kern)) - test_cuda_interface(TestInterface::EXC_VXC_INC, EvalType::Regular, - Backend::builtin, kern, Spin::Unpolarized); + test_cuda_interface(TestInterface::EXC_VXC_INC, EvalType::Regular, + Backend::builtin, kern, Spin::Unpolarized); } SECTION("EXC Small: Unpolarized") { for (auto kern : builtin_supported_kernels) { - if (is_unstable_small(kern) || !supports_unpolarized(kern)) + if (is_unstable_small(kern)) continue; test_cuda_interface(TestInterface::EXC, EvalType::Small, Backend::builtin, kern, Spin::Unpolarized); @@ -1651,7 +1652,7 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { SECTION("EXC + VXC Small: Unpolarized") { for (auto kern : builtin_supported_kernels) { - if (is_unstable_small(kern) || !supports_unpolarized(kern)) + if (is_unstable_small(kern)) continue; test_cuda_interface(TestInterface::EXC_VXC, EvalType::Small, Backend::builtin, kern, Spin::Unpolarized); @@ -1660,7 +1661,7 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { SECTION("EXC + INC Small: Unpolarized") { for (auto kern : builtin_supported_kernels) { - if (is_unstable_small(kern) || !supports_unpolarized(kern)) + if (is_unstable_small(kern)) continue; test_cuda_interface(TestInterface::EXC_INC, EvalType::Small, Backend::builtin, kern, Spin::Unpolarized); @@ -1669,7 +1670,7 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { SECTION("EXC + VXC + INC Small: Unpolarized") { for (auto kern : builtin_supported_kernels) { - if (is_unstable_small(kern) || !supports_unpolarized(kern)) + if (is_unstable_small(kern)) continue; test_cuda_interface(TestInterface::EXC_VXC_INC, EvalType::Small, Backend::builtin, kern, Spin::Unpolarized); @@ -1678,30 +1679,26 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { SECTION("EXC Zero: Unpolarized") { for (auto kern : builtin_supported_kernels) - if (supports_unpolarized(kern)) - test_cuda_interface(TestInterface::EXC, EvalType::Zero, - Backend::builtin, kern, Spin::Unpolarized); + test_cuda_interface(TestInterface::EXC, EvalType::Zero, + Backend::builtin, kern, Spin::Unpolarized); } SECTION("EXC + VXC Zero: Unpolarized") { for (auto kern : builtin_supported_kernels) - if (supports_unpolarized(kern)) - test_cuda_interface(TestInterface::EXC_VXC, EvalType::Zero, - Backend::builtin, kern, Spin::Unpolarized); + test_cuda_interface(TestInterface::EXC_VXC, EvalType::Zero, + Backend::builtin, kern, Spin::Unpolarized); } SECTION("EXC + INC Zero: Unpolarized") { for (auto kern : builtin_supported_kernels) - if (supports_unpolarized(kern)) - test_cuda_interface(TestInterface::EXC_INC, EvalType::Zero, - Backend::builtin, kern, Spin::Unpolarized); + test_cuda_interface(TestInterface::EXC_INC, EvalType::Zero, + Backend::builtin, kern, Spin::Unpolarized); } SECTION("EXC + VXC + INC Zero: Unpolarized") { for (auto kern : builtin_supported_kernels) - if (supports_unpolarized(kern)) - test_cuda_interface(TestInterface::EXC_VXC_INC, EvalType::Zero, - Backend::builtin, kern, Spin::Unpolarized); + test_cuda_interface(TestInterface::EXC_VXC_INC, EvalType::Zero, + Backend::builtin, kern, Spin::Unpolarized); } SECTION("EXC Regular: Polarized") { @@ -1783,10 +1780,7 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { test_cuda_interface( TestInterface::EXC_VXC_INC, EvalType::Zero, Backend::builtin, kern, Spin::Polarized ); } - } - - } #endif From 2b146b5b1297afd1b2403303f9ba95a94632ab07 Mon Sep 17 00:00:00 2001 From: Ryan Stocks Date: Sat, 20 Jul 2024 21:14:29 +0800 Subject: [PATCH 4/6] Combined cuda and hip kernel tests --- test/xc_kernel_test.cxx | 559 +--------------------------------------- 1 file changed, 12 insertions(+), 547 deletions(-) diff --git a/test/xc_kernel_test.cxx b/test/xc_kernel_test.cxx index d8ff648..d34acd7 100644 --- a/test/xc_kernel_test.cxx +++ b/test/xc_kernel_test.cxx @@ -1060,7 +1060,18 @@ TEST_CASE( "kernel_map Test", "[xc-kernel-map]") { } -#ifdef EXCHCXX_ENABLE_CUDA +#if defined(EXCHCXX_ENABLE_CUDA) || defined(EXCHCXX_ENABLE_HIP) + +#ifdef EXCHCXX_ENABLE_HIP +#define cudaFree hipFree +#define cudaGetErrorString hipGetErrorString +#define cudaMalloc hipMalloc +#define cudaStream_t hipStream_t +#define cudaSuccess hipSuccess +#define cudaMemcpy hipMemcpy +#define cudaMemcpyDefault hipMemcpyDefault +#define cudaDeviceSynchronize hipDeviceSynchronize +#endif template T* safe_cuda_malloc( size_t n ) { @@ -1787,552 +1798,6 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { -#ifdef EXCHCXX_ENABLE_HIP - -template -T* safe_hip_malloc( size_t n ) { - - T* ptr = nullptr;; - if( n ) { - auto stat = hipMalloc( (void**)&ptr, n*sizeof(T) ); - if( stat != hipSuccess ) - throw std::runtime_error(hipGetErrorString( stat )); - } - return ptr; - -} - -template -void safe_hip_cpy( T* dest, const T* src, size_t len ) { - - auto stat = hipMemcpy( dest, src, len*sizeof(T), hipMemcpyDefault ); - if( stat != hipSuccess ) - throw std::runtime_error(hipGetErrorString( stat )); - -} - -void hip_free_all(){ } -template -void hip_free_all( T* ptr, Args&&... args ) { - - if( ptr ) { - auto stat = hipFree( ptr ); - if( stat != hipSuccess ) - throw std::runtime_error(hipGetErrorString( stat )); - } - - hip_free_all( std::forward(args)... ); - - -} - -void device_synchronize() { - auto stat = hipDeviceSynchronize(); - if( stat != hipSuccess ) - throw std::runtime_error(hipGetErrorString( stat )); -} - - -void test_hip_interface( TestInterface interface, EvalType evaltype, - Backend backend, Kernel kern, Spin polar ) { - - size_t npts_lda, npts_gga; - std::vector ref_rho, ref_sigma; - std::tie(npts_lda, ref_rho ) = load_reference_density( polar ); - std::tie(npts_gga, ref_sigma) = load_reference_sigma ( polar ); - - REQUIRE( npts_lda == npts_gga ); - - const int npts = npts_lda; - - XCKernel func( backend, kern, polar ); - - size_t len_rho_buffer = func.rho_buffer_len(npts); - size_t len_sigma_buffer = func.sigma_buffer_len(npts); - size_t len_exc_buffer = func.exc_buffer_len(npts); - size_t len_vrho_buffer = func.vrho_buffer_len(npts); - size_t len_vsigma_buffer = func.vsigma_buffer_len(npts); - - - std::vector rho_small(len_rho_buffer, 1e-13); - std::vector sigma_small(len_sigma_buffer, 1e-14); - - std::vector rho_zero(len_rho_buffer, 0.); - std::vector sigma_zero(len_sigma_buffer, 0.); - - std::vector rho, sigma; - - if( evaltype == EvalType::Regular ) { - rho = ref_rho; - sigma = ref_sigma; - } - - if( evaltype == EvalType::Small ) { - rho = rho_small; - sigma = sigma_small; - } - - if( evaltype == EvalType::Zero ) { - rho = rho_zero; - sigma = sigma_zero; - } - - // Get Reference Values - std::vector - exc_ref( len_exc_buffer ), - vrho_ref( len_vrho_buffer ), - vsigma_ref( len_vsigma_buffer ); - - if( interface == TestInterface::EXC or interface == TestInterface::EXC_INC ) { - - if( func.is_lda() ) - func.eval_exc( npts, rho.data(), exc_ref.data() ); - else if( func.is_gga() ) - func.eval_exc( npts, rho.data(), sigma.data(), exc_ref.data() ); - - } else if( interface == TestInterface::EXC_VXC or interface == TestInterface::EXC_VXC_INC ) { - - if( func.is_lda() ) - func.eval_exc_vxc( npts, rho.data(), exc_ref.data(), vrho_ref.data() ); - else if( func.is_gga() ) - func.eval_exc_vxc( npts, rho.data(), sigma.data(), exc_ref.data(), - vrho_ref.data(), vsigma_ref.data() ); - - } - - - - - - - // Allocate device memory - double* rho_device = safe_hip_malloc( len_rho_buffer ); - double* sigma_device = safe_hip_malloc( len_sigma_buffer ); - double* exc_device = safe_hip_malloc( len_exc_buffer ); - double* vrho_device = safe_hip_malloc( len_vrho_buffer ); - double* vsigma_device = safe_hip_malloc( len_vsigma_buffer ); - - // H2D Copy of rho / sigma - safe_hip_cpy( rho_device, rho.data(), len_rho_buffer ); - if( func.is_gga() ) - safe_hip_cpy( sigma_device, sigma.data(), len_sigma_buffer ); - - const double alpha = 3.14; - const double fill_val_e = 2.; - const double fill_val_vr = 10.; - const double fill_val_vs = 50.; - - std::vector - exc( len_exc_buffer, fill_val_e ), vrho( len_vrho_buffer, fill_val_vr ), - vsigma( len_vsigma_buffer, fill_val_vs ); - - // H2D copy of initial values, tests clobber / increment - safe_hip_cpy( exc_device, exc.data(), len_exc_buffer ); - safe_hip_cpy( vrho_device, vrho.data(), len_vrho_buffer ); - if( func.is_gga() ) - safe_hip_cpy( vsigma_device, vsigma.data(), len_vsigma_buffer ); - - // Evaluate functional on device - hipStream_t stream = 0; - if( interface == TestInterface::EXC ) { - - if( func.is_lda() ) - func.eval_exc_device( npts, rho_device, exc_device, stream ); - else if( func.is_gga() ) - func.eval_exc_device( npts, rho_device, sigma_device, exc_device, - stream ); - - } else if( interface == TestInterface::EXC_INC ) { - - if( func.is_lda() ) - func.eval_exc_inc_device( alpha, npts, rho_device, exc_device, stream ); - else if( func.is_gga() ) - func.eval_exc_inc_device( alpha, npts, rho_device, sigma_device, exc_device, - stream ); - - } else if( interface == TestInterface::EXC_VXC ) { - - if( func.is_lda() ) - func.eval_exc_vxc_device( npts, rho_device, exc_device, vrho_device, stream ); - else if( func.is_gga() ) - func.eval_exc_vxc_device( npts, rho_device, sigma_device, exc_device, - vrho_device, vsigma_device, stream ); - - } else if( interface == TestInterface::EXC_VXC_INC ) { - - if( func.is_lda() ) - func.eval_exc_vxc_inc_device( alpha, npts, rho_device, exc_device, - vrho_device, stream ); - else if( func.is_gga() ) - func.eval_exc_vxc_inc_device( alpha, npts, rho_device, sigma_device, - exc_device, vrho_device, vsigma_device, stream ); - - } - - device_synchronize(); - - // D2H of results - safe_hip_cpy( exc.data(), exc_device, len_exc_buffer ); - safe_hip_cpy( vrho.data(), vrho_device, len_vrho_buffer ); - if(func.is_gga()) - safe_hip_cpy( vsigma.data(), vsigma_device, len_vsigma_buffer ); - - // Check correctness - if( interface == TestInterface::EXC_INC or interface == TestInterface::EXC_VXC_INC ) { - for( auto i = 0ul; i < len_exc_buffer; ++i ) - CHECK( exc[i] == Approx(fill_val_e + alpha * exc_ref[i]) ); - } else { - for( auto i = 0ul; i < len_exc_buffer; ++i ) - CHECK( exc[i] == Approx(exc_ref[i]) ); - } - - if( interface == TestInterface::EXC_VXC_INC ) { - - for( auto i = 0ul; i < len_vrho_buffer; ++i ) - CHECK( vrho[i] == Approx(fill_val_vr + alpha * vrho_ref[i]) ); - for( auto i = 0ul; i < len_vsigma_buffer; ++i ) - CHECK( vsigma[i] == Approx(fill_val_vs + alpha * vsigma_ref[i]) ); - - } else if(interface == TestInterface::EXC_VXC) { - - for( auto i = 0ul; i < len_vrho_buffer; ++i ) - CHECK( vrho[i] == Approx(vrho_ref[i]) ); - for( auto i = 0ul; i < len_vsigma_buffer; ++i ) { - INFO( "Kernel is " << kern ); - CHECK( vsigma[i] == Approx(vsigma_ref[i]) ); - } - - } - - hip_free_all( rho_device, sigma_device, exc_device, vrho_device, vsigma_device ); -} - - - -TEST_CASE( "HIP Interfaces", "[xc-device]" ) { - - SECTION( "Libxc Functionals" ) { - - SECTION( "LDA Functionals: EXC Regular Eval Unpolarized" ) { - for( auto kern : lda_kernels ) - test_hip_interface( TestInterface::EXC, EvalType::Regular, - Backend::libxc, kern, Spin::Unpolarized ); - } - - - SECTION( "LDA Functionals: EXC + VXC Regular Eval Unpolarized" ) { - for( auto kern : lda_kernels ) - test_hip_interface( TestInterface::EXC_VXC, EvalType::Regular, - Backend::libxc, kern, Spin::Unpolarized ); - } - - SECTION( "GGA Functionals: EXC Regular Eval Unpolarized" ) { - for( auto kern : gga_kernels ) - test_hip_interface( TestInterface::EXC, EvalType::Regular, - Backend::libxc, kern, Spin::Unpolarized ); - } - - SECTION( "GGA Functionals: EXC + VXC Regular Eval Unpolarized" ) { - for( auto kern : gga_kernels ) - test_hip_interface( TestInterface::EXC_VXC, EvalType::Regular, - Backend::libxc, kern, Spin::Unpolarized ); - } - - SECTION( "LDA Functionals: EXC Small Eval Unpolarized" ) { - for( auto kern : lda_kernels ) - test_hip_interface( TestInterface::EXC, EvalType::Small, - Backend::libxc, kern, Spin::Unpolarized ); - } - - - SECTION( "LDA Functionals: EXC + VXC Small Eval Unpolarized" ) { - for( auto kern : lda_kernels ) - test_hip_interface( TestInterface::EXC_VXC, EvalType::Small, - Backend::libxc, kern, Spin::Unpolarized ); - } - - SECTION( "GGA Functionals: EXC Small Eval Unpolarized" ) { - for( auto kern : gga_kernels ) - test_hip_interface( TestInterface::EXC, EvalType::Small, - Backend::libxc, kern, Spin::Unpolarized ); - } - - SECTION( "GGA Functionals: EXC + VXC Small Eval Unpolarized" ) { - for( auto kern : gga_kernels ) - test_hip_interface( TestInterface::EXC_VXC, EvalType::Small, - Backend::libxc, kern, Spin::Unpolarized ); - } - - SECTION( "LDA Functionals: EXC Zero Eval Unpolarized" ) { - for( auto kern : lda_kernels ) - test_hip_interface( TestInterface::EXC, EvalType::Zero, - Backend::libxc, kern, Spin::Unpolarized ); - } - - - SECTION( "LDA Functionals: EXC + VXC Zero Eval Unpolarized" ) { - for( auto kern : lda_kernels ) - test_hip_interface( TestInterface::EXC_VXC, EvalType::Zero, - Backend::libxc, kern, Spin::Unpolarized ); - } - - SECTION( "GGA Functionals: EXC Zero Eval Unpolarized" ) { - for( auto kern : gga_kernels ) - test_hip_interface( TestInterface::EXC, EvalType::Zero, - Backend::libxc, kern, Spin::Unpolarized ); - } - - SECTION( "GGA Functionals: EXC + VXC Zero Eval Unpolarized" ) { - for( auto kern : gga_kernels ) - test_hip_interface( TestInterface::EXC_VXC, EvalType::Zero, - Backend::libxc, kern, Spin::Unpolarized ); - } - - - - - - - - - SECTION( "LDA Functionals: EXC Regular Eval Polarized" ) { - for( auto kern : lda_kernels ) - test_hip_interface( TestInterface::EXC, EvalType::Regular, - Backend::libxc, kern, Spin::Polarized ); - } - - - SECTION( "LDA Functionals: EXC + VXC Regular Eval Polarized" ) { - for( auto kern : lda_kernels ) - test_hip_interface( TestInterface::EXC_VXC, EvalType::Regular, - Backend::libxc, kern, Spin::Polarized ); - } - - SECTION( "GGA Functionals: EXC Regular Eval Polarized" ) { - for( auto kern : gga_kernels ) - test_hip_interface( TestInterface::EXC, EvalType::Regular, - Backend::libxc, kern, Spin::Polarized ); - } - - SECTION( "GGA Functionals: EXC + VXC Regular Eval Polarized" ) { - for( auto kern : gga_kernels ) - test_hip_interface( TestInterface::EXC_VXC, EvalType::Regular, - Backend::libxc, kern, Spin::Polarized ); - } - - SECTION( "LDA Functionals: EXC Small Eval Polarized" ) { - for( auto kern : lda_kernels ) - test_hip_interface( TestInterface::EXC, EvalType::Small, - Backend::libxc, kern, Spin::Polarized ); - } - - - SECTION( "LDA Functionals: EXC + VXC Small Eval Polarized" ) { - for( auto kern : lda_kernels ) - test_hip_interface( TestInterface::EXC_VXC, EvalType::Small, - Backend::libxc, kern, Spin::Polarized ); - } - - SECTION( "GGA Functionals: EXC Small Eval Polarized" ) { - for( auto kern : gga_kernels ) - test_hip_interface( TestInterface::EXC, EvalType::Small, - Backend::libxc, kern, Spin::Polarized ); - } - - SECTION( "GGA Functionals: EXC + VXC Small Eval Polarized" ) { - for( auto kern : gga_kernels ) - test_hip_interface( TestInterface::EXC_VXC, EvalType::Small, - Backend::libxc, kern, Spin::Polarized ); - } - - SECTION( "LDA Functionals: EXC Zero Eval Polarized" ) { - for( auto kern : lda_kernels ) - test_hip_interface( TestInterface::EXC, EvalType::Zero, - Backend::libxc, kern, Spin::Polarized ); - } - - - SECTION( "LDA Functionals: EXC + VXC Zero Eval Polarized" ) { - for( auto kern : lda_kernels ) - test_hip_interface( TestInterface::EXC_VXC, EvalType::Zero, - Backend::libxc, kern, Spin::Polarized ); - } - - SECTION( "GGA Functionals: EXC Zero Eval Polarized" ) { - for( auto kern : gga_kernels ) - test_hip_interface( TestInterface::EXC, EvalType::Zero, - Backend::libxc, kern, Spin::Polarized ); - } - - SECTION( "GGA Functionals: EXC + VXC Zero Eval Polarized" ) { - for( auto kern : gga_kernels ) - test_hip_interface( TestInterface::EXC_VXC, EvalType::Zero, - Backend::libxc, kern, Spin::Polarized ); - } - - } - - SECTION( "Builtin Functionals" ) { - - SECTION("EXC Regular: Unpolarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC, EvalType::Regular, - Backend::builtin, kern, Spin::Unpolarized ); - } - - SECTION("EXC + VXC Regular: Unpolarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC_VXC, EvalType::Regular, - Backend::builtin, kern, Spin::Unpolarized ); - } - - SECTION("EXC + INC Regular: Unpolarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC_INC, EvalType::Regular, - Backend::builtin, kern, Spin::Unpolarized ); - } - - SECTION("EXC + VXC + INC Regular: Unpolarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC_VXC_INC, EvalType::Regular, - Backend::builtin, kern, Spin::Unpolarized ); - } - - SECTION("EXC Small: Unpolarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC, EvalType::Small, - Backend::builtin, kern, Spin::Unpolarized ); - } - - SECTION("EXC + VXC Small: Unpolarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC_VXC, EvalType::Small, - Backend::builtin, kern, Spin::Unpolarized ); - } - - SECTION("EXC + INC Small: Unpolarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC_INC, EvalType::Small, - Backend::builtin, kern, Spin::Unpolarized ); - } - - SECTION("EXC + VXC + INC Small: Unpolarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC_VXC_INC, EvalType::Small, - Backend::builtin, kern, Spin::Unpolarized ); - } - - SECTION("EXC Zero: Unpolarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC, EvalType::Zero, - Backend::builtin, kern, Spin::Unpolarized ); - } - - SECTION("EXC + VXC Zero: Unpolarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC_VXC, EvalType::Zero, - Backend::builtin, kern, Spin::Unpolarized ); - } - - SECTION("EXC + INC Zero: Unpolarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC_INC, EvalType::Zero, - Backend::builtin, kern, Spin::Unpolarized ); - } - - SECTION("EXC + VXC + INC Zero: Unpolarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC_VXC_INC, EvalType::Zero, - Backend::builtin, kern, Spin::Unpolarized ); - } - - SECTION("EXC Regular: Polarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC, EvalType::Regular, - Backend::builtin, kern, Spin::Polarized ); - } - - SECTION("EXC + VXC Regular: Polarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC_VXC, EvalType::Regular, - Backend::builtin, kern, Spin::Polarized ); - } - - SECTION("EXC + INC Regular: Polarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC_INC, EvalType::Regular, - Backend::builtin, kern, Spin::Polarized ); - } - - SECTION("EXC + VXC + INC Regular: Polarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC_VXC_INC, EvalType::Regular, - Backend::builtin, kern, Spin::Polarized ); - } - - SECTION("EXC Small: Polarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC, EvalType::Small, - Backend::builtin, kern, Spin::Polarized ); - } - - SECTION("EXC + VXC Small: Polarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC_VXC, EvalType::Small, - Backend::builtin, kern, Spin::Polarized ); - } - - SECTION("EXC + INC Small: Polarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC_INC, EvalType::Small, - Backend::builtin, kern, Spin::Polarized ); - } - - SECTION("EXC + VXC + INC Small: Polarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC_VXC_INC, EvalType::Small, - Backend::builtin, kern, Spin::Polarized ); - } - - SECTION("EXC Zero: Polarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC, EvalType::Zero, - Backend::builtin, kern, Spin::Polarized ); - } - - SECTION("EXC + VXC Zero: Polarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC_VXC, EvalType::Zero, - Backend::builtin, kern, Spin::Polarized ); - } - - SECTION("EXC + INC Zero: Polarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC_INC, EvalType::Zero, - Backend::builtin, kern, Spin::Polarized ); - } - - SECTION("EXC + VXC + INC Zero: Polarized") { - for( auto kern : builtin_supported_kernels ) - test_hip_interface( TestInterface::EXC_VXC_INC, EvalType::Zero, - Backend::builtin, kern, Spin::Polarized ); - } - - - } - - -} - -#endif - - - - - - - - #ifdef EXCHCXX_ENABLE_SYCL template From 63f2682144aa59b5df0eb6762152b062cbdda670 Mon Sep 17 00:00:00 2001 From: Ryan Stocks <42725471+ryanstocks00@users.noreply.github.com> Date: Wed, 24 Jul 2024 12:36:47 +1000 Subject: [PATCH 5/6] Update src/builtin_interface.cxx Co-authored-by: David Williams-Young --- src/builtin_interface.cxx | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/src/builtin_interface.cxx b/src/builtin_interface.cxx index 105dc7e..7c5f91e 100644 --- a/src/builtin_interface.cxx +++ b/src/builtin_interface.cxx @@ -59,10 +59,9 @@ namespace detail { std::unique_ptr gen_from_kern( Kernel kern, Spin polar ) { - if (!supports_unpolarized(kern)) { - EXCHCXX_BOOL_CHECK(kernel_map.key(kern) + " Needs to be Spin-Polarized!", - polar == Spin::Polarized); - } + // Bail if polarized eval is requested and not supported + EXCHCXX_BOOL_CHECK(kernel_map.key(kern) + " Needs to be Spin-Polarized!", + (not supports_unpolarized(kern)) and polar == Spin::Polarized); if( kern == Kernel::SlaterExchange ) return std::make_unique( polar ); From 3c28073ed085b41e2370001c0a76972c54fbee1f Mon Sep 17 00:00:00 2001 From: Ryan Stocks Date: Wed, 24 Jul 2024 12:55:07 +1000 Subject: [PATCH 6/6] PR review comments (undo formatting, clean gitignore) --- .gitignore | 3 - src/builtin_interface.cxx | 2 +- test/xc_kernel_test.cxx | 187 +++++++++++++++++++------------------- 3 files changed, 94 insertions(+), 98 deletions(-) diff --git a/.gitignore b/.gitignore index 9d25f2f..0a37f33 100644 --- a/.gitignore +++ b/.gitignore @@ -1,4 +1 @@ *.*.swp -.cache/ -build/ -compile_commands.json diff --git a/src/builtin_interface.cxx b/src/builtin_interface.cxx index 7c5f91e..275cfea 100644 --- a/src/builtin_interface.cxx +++ b/src/builtin_interface.cxx @@ -61,7 +61,7 @@ std::unique_ptr // Bail if polarized eval is requested and not supported EXCHCXX_BOOL_CHECK(kernel_map.key(kern) + " Needs to be Spin-Polarized!", - (not supports_unpolarized(kern)) and polar == Spin::Polarized); + supports_unpolarized(kern) or polar == Spin::Polarized); if( kern == Kernel::SlaterExchange ) return std::make_unique( polar ); diff --git a/test/xc_kernel_test.cxx b/test/xc_kernel_test.cxx index d34acd7..77ffa8d 100644 --- a/test/xc_kernel_test.cxx +++ b/test/xc_kernel_test.cxx @@ -1117,7 +1117,7 @@ void device_synchronize() { } -void test_cuda_interface( TestInterface interface, EvalType evaltype, +void test_cuda_hip_interface( TestInterface interface, EvalType evaltype, Backend backend, Kernel kern, Spin polar ) { size_t npts_lda, npts_gga, npts_mgga, npts_lapl; @@ -1373,45 +1373,45 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { SECTION( "LDA Functionals: EXC Regular Eval Unpolarized" ) { for( auto kern : lda_kernels ) - test_cuda_interface( TestInterface::EXC, EvalType::Regular, + test_cuda_hip_interface( TestInterface::EXC, EvalType::Regular, Backend::libxc, kern, Spin::Unpolarized ); } SECTION( "LDA Functionals: EXC + VXC Regular Eval Unpolarized" ) { for( auto kern : lda_kernels ) - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Regular, + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Regular, Backend::libxc, kern, Spin::Unpolarized ); } SECTION( "GGA Functionals: EXC Regular Eval Unpolarized" ) { for( auto kern : gga_kernels ) - test_cuda_interface( TestInterface::EXC, EvalType::Regular, + test_cuda_hip_interface( TestInterface::EXC, EvalType::Regular, Backend::libxc, kern, Spin::Unpolarized ); } SECTION( "GGA Functionals: EXC + VXC Regular Eval Unpolarized" ) { for( auto kern : gga_kernels ) - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Regular, + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Regular, Backend::libxc, kern, Spin::Unpolarized ); } SECTION( "MGGA Functionals: EXC Regular Eval Unpolarized" ) { for( auto kern : mgga_kernels ) - test_cuda_interface( TestInterface::EXC, EvalType::Regular, + test_cuda_hip_interface( TestInterface::EXC, EvalType::Regular, Backend::libxc, kern, Spin::Unpolarized ); } SECTION( "MGGA Functionals: EXC + VXC Regular Eval Unpolarized" ) { for( auto kern : mgga_kernels ) - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Regular, + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Regular, Backend::libxc, kern, Spin::Unpolarized ); } SECTION( "LDA Functionals: EXC Small Eval Unpolarized" ) { for( auto kern : lda_kernels ) { if(is_unstable_small(kern)) continue; - test_cuda_interface( TestInterface::EXC, EvalType::Small, + test_cuda_hip_interface( TestInterface::EXC, EvalType::Small, Backend::libxc, kern, Spin::Unpolarized ); } } @@ -1420,7 +1420,7 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { SECTION( "LDA Functionals: EXC + VXC Small Eval Unpolarized" ) { for( auto kern : lda_kernels ){ if(is_unstable_small(kern)) continue; - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Small, + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Small, Backend::libxc, kern, Spin::Unpolarized ); } } @@ -1428,7 +1428,7 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { SECTION( "GGA Functionals: EXC Small Eval Unpolarized" ) { for( auto kern : gga_kernels ){ if(is_unstable_small(kern)) continue; - test_cuda_interface( TestInterface::EXC, EvalType::Small, + test_cuda_hip_interface( TestInterface::EXC, EvalType::Small, Backend::libxc, kern, Spin::Unpolarized ); } } @@ -1436,7 +1436,7 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { SECTION( "GGA Functionals: EXC + VXC Small Eval Unpolarized" ) { for( auto kern : gga_kernels ){ if(is_unstable_small(kern)) continue; - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Small, + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Small, Backend::libxc, kern, Spin::Unpolarized ); } } @@ -1444,7 +1444,7 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { SECTION( "MGGA Functionals: EXC Small Eval Unpolarized" ) { for( auto kern : mgga_kernels ){ if(is_unstable_small(kern)) continue; - test_cuda_interface( TestInterface::EXC, EvalType::Small, + test_cuda_hip_interface( TestInterface::EXC, EvalType::Small, Backend::libxc, kern, Spin::Unpolarized ); } } @@ -1452,45 +1452,45 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { SECTION( "MGGA Functionals: EXC + VXC Small Eval Unpolarized" ) { for( auto kern : mgga_kernels ) { if(is_unstable_small(kern)) continue; - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Small, + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Small, Backend::libxc, kern, Spin::Unpolarized ); } } SECTION( "LDA Functionals: EXC Zero Eval Unpolarized" ) { for( auto kern : lda_kernels ) - test_cuda_interface( TestInterface::EXC, EvalType::Zero, + test_cuda_hip_interface( TestInterface::EXC, EvalType::Zero, Backend::libxc, kern, Spin::Unpolarized ); } SECTION( "LDA Functionals: EXC + VXC Zero Eval Unpolarized" ) { for( auto kern : lda_kernels ) - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Zero, + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Zero, Backend::libxc, kern, Spin::Unpolarized ); } SECTION( "GGA Functionals: EXC Zero Eval Unpolarized" ) { for( auto kern : gga_kernels ) - test_cuda_interface( TestInterface::EXC, EvalType::Zero, + test_cuda_hip_interface( TestInterface::EXC, EvalType::Zero, Backend::libxc, kern, Spin::Unpolarized ); } SECTION( "GGA Functionals: EXC + VXC Zero Eval Unpolarized" ) { for( auto kern : gga_kernels ) - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Zero, + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Zero, Backend::libxc, kern, Spin::Unpolarized ); } SECTION( "MGGA Functionals: EXC Zero Eval Unpolarized" ) { for( auto kern : mgga_kernels ) - test_cuda_interface( TestInterface::EXC, EvalType::Zero, + test_cuda_hip_interface( TestInterface::EXC, EvalType::Zero, Backend::libxc, kern, Spin::Unpolarized ); } SECTION( "MGGA Functionals: EXC + VXC Zero Eval Unpolarized" ) { for( auto kern : mgga_kernels ) - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Zero, + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Zero, Backend::libxc, kern, Spin::Unpolarized ); } @@ -1503,45 +1503,45 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { SECTION( "LDA Functionals: EXC Regular Eval Polarized" ) { for( auto kern : lda_kernels ) - test_cuda_interface( TestInterface::EXC, EvalType::Regular, + test_cuda_hip_interface( TestInterface::EXC, EvalType::Regular, Backend::libxc, kern, Spin::Polarized ); } SECTION( "LDA Functionals: EXC + VXC Regular Eval Polarized" ) { for( auto kern : lda_kernels ) - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Regular, + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Regular, Backend::libxc, kern, Spin::Polarized ); } SECTION( "GGA Functionals: EXC Regular Eval Polarized" ) { for( auto kern : gga_kernels ) - test_cuda_interface( TestInterface::EXC, EvalType::Regular, + test_cuda_hip_interface( TestInterface::EXC, EvalType::Regular, Backend::libxc, kern, Spin::Polarized ); } SECTION( "GGA Functionals: EXC + VXC Regular Eval Polarized" ) { for( auto kern : gga_kernels ) - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Regular, + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Regular, Backend::libxc, kern, Spin::Polarized ); } SECTION( "MGGA Functionals: EXC Regular Eval Polarized" ) { for( auto kern : mgga_kernels ) - test_cuda_interface( TestInterface::EXC, EvalType::Regular, + test_cuda_hip_interface( TestInterface::EXC, EvalType::Regular, Backend::libxc, kern, Spin::Polarized ); } SECTION( "MGGA Functionals: EXC + VXC Regular Eval Polarized" ) { for( auto kern : mgga_kernels ) - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Regular, + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Regular, Backend::libxc, kern, Spin::Polarized ); } SECTION( "LDA Functionals: EXC Small Eval Polarized" ) { for( auto kern : lda_kernels ){ if(is_unstable_small(kern)) continue; - test_cuda_interface( TestInterface::EXC, EvalType::Small, + test_cuda_hip_interface( TestInterface::EXC, EvalType::Small, Backend::libxc, kern, Spin::Polarized ); } } @@ -1550,7 +1550,7 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { SECTION( "LDA Functionals: EXC + VXC Small Eval Polarized" ) { for( auto kern : lda_kernels ) { if(is_unstable_small(kern)) continue; - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Small, + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Small, Backend::libxc, kern, Spin::Polarized ); } } @@ -1558,7 +1558,7 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { SECTION( "GGA Functionals: EXC Small Eval Polarized" ) { for( auto kern : gga_kernels ) { if(is_unstable_small(kern)) continue; - test_cuda_interface( TestInterface::EXC, EvalType::Small, + test_cuda_hip_interface( TestInterface::EXC, EvalType::Small, Backend::libxc, kern, Spin::Polarized ); } } @@ -1566,7 +1566,7 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { SECTION( "GGA Functionals: EXC + VXC Small Eval Polarized" ) { for( auto kern : gga_kernels ) { if(is_unstable_small(kern)) continue; - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Small, + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Small, Backend::libxc, kern, Spin::Polarized ); } } @@ -1574,7 +1574,7 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { SECTION( "MGGA Functionals: EXC Small Eval Polarized" ) { for( auto kern : mgga_kernels ) { if(is_unstable_small(kern)) continue; - test_cuda_interface( TestInterface::EXC, EvalType::Small, + test_cuda_hip_interface( TestInterface::EXC, EvalType::Small, Backend::libxc, kern, Spin::Polarized ); } } @@ -1582,164 +1582,160 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { SECTION( "MGGA Functionals: EXC + VXC Small Eval Polarized" ) { for( auto kern : mgga_kernels ) { if(is_unstable_small(kern)) continue; - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Small, + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Small, Backend::libxc, kern, Spin::Polarized ); } } SECTION( "LDA Functionals: EXC Zero Eval Polarized" ) { for( auto kern : lda_kernels ) - test_cuda_interface( TestInterface::EXC, EvalType::Zero, + test_cuda_hip_interface( TestInterface::EXC, EvalType::Zero, Backend::libxc, kern, Spin::Polarized ); } SECTION( "LDA Functionals: EXC + VXC Zero Eval Polarized" ) { for( auto kern : lda_kernels ) - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Zero, + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Zero, Backend::libxc, kern, Spin::Polarized ); } SECTION( "GGA Functionals: EXC Zero Eval Polarized" ) { for( auto kern : gga_kernels ) - test_cuda_interface( TestInterface::EXC, EvalType::Zero, + test_cuda_hip_interface( TestInterface::EXC, EvalType::Zero, Backend::libxc, kern, Spin::Polarized ); } SECTION( "GGA Functionals: EXC + VXC Zero Eval Polarized" ) { for( auto kern : gga_kernels ) - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Zero, + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Zero, Backend::libxc, kern, Spin::Polarized ); } SECTION( "MGGA Functionals: EXC Zero Eval Polarized" ) { for( auto kern : mgga_kernels ) - test_cuda_interface( TestInterface::EXC, EvalType::Zero, + test_cuda_hip_interface( TestInterface::EXC, EvalType::Zero, Backend::libxc, kern, Spin::Polarized ); } SECTION( "MGGA Functionals: EXC + VXC Zero Eval Polarized" ) { for( auto kern : mgga_kernels ) - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Zero, + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Zero, Backend::libxc, kern, Spin::Polarized ); } } - SECTION("Builtin Functionals") { + SECTION( "Builtin Functionals" ) { SECTION("EXC Regular: Unpolarized") { - for (auto kern : builtin_supported_kernels) - test_cuda_interface(TestInterface::EXC, EvalType::Regular, - Backend::builtin, kern, Spin::Unpolarized); + for( auto kern : builtin_supported_kernels ) + test_cuda_hip_interface( TestInterface::EXC, EvalType::Regular, + Backend::builtin, kern, Spin::Unpolarized ); } SECTION("EXC + VXC Regular: Unpolarized") { - for (auto kern : builtin_supported_kernels) - test_cuda_interface(TestInterface::EXC_VXC, EvalType::Regular, - Backend::builtin, kern, Spin::Unpolarized); + for( auto kern : builtin_supported_kernels ) + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Regular, + Backend::builtin, kern, Spin::Unpolarized ); } SECTION("EXC + INC Regular: Unpolarized") { - for (auto kern : builtin_supported_kernels) - test_cuda_interface(TestInterface::EXC_INC, EvalType::Regular, - Backend::builtin, kern, Spin::Unpolarized); + for( auto kern : builtin_supported_kernels ) + test_cuda_hip_interface( TestInterface::EXC_INC, EvalType::Regular, + Backend::builtin, kern, Spin::Unpolarized ); } SECTION("EXC + VXC + INC Regular: Unpolarized") { - for (auto kern : builtin_supported_kernels) - test_cuda_interface(TestInterface::EXC_VXC_INC, EvalType::Regular, - Backend::builtin, kern, Spin::Unpolarized); + for( auto kern : builtin_supported_kernels ) + test_cuda_hip_interface( TestInterface::EXC_VXC_INC, EvalType::Regular, + Backend::builtin, kern, Spin::Unpolarized ); } SECTION("EXC Small: Unpolarized") { - for (auto kern : builtin_supported_kernels) { - if (is_unstable_small(kern)) - continue; - test_cuda_interface(TestInterface::EXC, EvalType::Small, - Backend::builtin, kern, Spin::Unpolarized); + for( auto kern : builtin_supported_kernels ) { + if(is_unstable_small(kern)) continue; + test_cuda_hip_interface( TestInterface::EXC, EvalType::Small, + Backend::builtin, kern, Spin::Unpolarized ); } } SECTION("EXC + VXC Small: Unpolarized") { - for (auto kern : builtin_supported_kernels) { - if (is_unstable_small(kern)) - continue; - test_cuda_interface(TestInterface::EXC_VXC, EvalType::Small, - Backend::builtin, kern, Spin::Unpolarized); + for( auto kern : builtin_supported_kernels ) { + if(is_unstable_small(kern)) continue; + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Small, + Backend::builtin, kern, Spin::Unpolarized ); } } SECTION("EXC + INC Small: Unpolarized") { - for (auto kern : builtin_supported_kernels) { - if (is_unstable_small(kern)) - continue; - test_cuda_interface(TestInterface::EXC_INC, EvalType::Small, - Backend::builtin, kern, Spin::Unpolarized); + for( auto kern : builtin_supported_kernels ) { + if(is_unstable_small(kern)) continue; + test_cuda_hip_interface( TestInterface::EXC_INC, EvalType::Small, + Backend::builtin, kern, Spin::Unpolarized ); } } SECTION("EXC + VXC + INC Small: Unpolarized") { - for (auto kern : builtin_supported_kernels) { - if (is_unstable_small(kern)) - continue; - test_cuda_interface(TestInterface::EXC_VXC_INC, EvalType::Small, - Backend::builtin, kern, Spin::Unpolarized); + for( auto kern : builtin_supported_kernels ) { + if(is_unstable_small(kern)) continue; + test_cuda_hip_interface( TestInterface::EXC_VXC_INC, EvalType::Small, + Backend::builtin, kern, Spin::Unpolarized ); } } SECTION("EXC Zero: Unpolarized") { - for (auto kern : builtin_supported_kernels) - test_cuda_interface(TestInterface::EXC, EvalType::Zero, - Backend::builtin, kern, Spin::Unpolarized); + for( auto kern : builtin_supported_kernels ) + test_cuda_hip_interface( TestInterface::EXC, EvalType::Zero, + Backend::builtin, kern, Spin::Unpolarized ); } SECTION("EXC + VXC Zero: Unpolarized") { - for (auto kern : builtin_supported_kernels) - test_cuda_interface(TestInterface::EXC_VXC, EvalType::Zero, - Backend::builtin, kern, Spin::Unpolarized); + for( auto kern : builtin_supported_kernels ) + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Zero, + Backend::builtin, kern, Spin::Unpolarized ); } SECTION("EXC + INC Zero: Unpolarized") { - for (auto kern : builtin_supported_kernels) - test_cuda_interface(TestInterface::EXC_INC, EvalType::Zero, - Backend::builtin, kern, Spin::Unpolarized); + for( auto kern : builtin_supported_kernels ) + test_cuda_hip_interface( TestInterface::EXC_INC, EvalType::Zero, + Backend::builtin, kern, Spin::Unpolarized ); } SECTION("EXC + VXC + INC Zero: Unpolarized") { - for (auto kern : builtin_supported_kernels) - test_cuda_interface(TestInterface::EXC_VXC_INC, EvalType::Zero, - Backend::builtin, kern, Spin::Unpolarized); + for( auto kern : builtin_supported_kernels ) + test_cuda_hip_interface( TestInterface::EXC_VXC_INC, EvalType::Zero, + Backend::builtin, kern, Spin::Unpolarized ); } SECTION("EXC Regular: Polarized") { for( auto kern : builtin_supported_kernels ) - test_cuda_interface( TestInterface::EXC, EvalType::Regular, + test_cuda_hip_interface( TestInterface::EXC, EvalType::Regular, Backend::builtin, kern, Spin::Polarized ); } SECTION("EXC + VXC Regular: Polarized") { for( auto kern : builtin_supported_kernels ) - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Regular, + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Regular, Backend::builtin, kern, Spin::Polarized ); } SECTION("EXC + INC Regular: Polarized") { for( auto kern : builtin_supported_kernels ) - test_cuda_interface( TestInterface::EXC_INC, EvalType::Regular, + test_cuda_hip_interface( TestInterface::EXC_INC, EvalType::Regular, Backend::builtin, kern, Spin::Polarized ); } SECTION("EXC + VXC + INC Regular: Polarized") { for( auto kern : builtin_supported_kernels ) - test_cuda_interface( TestInterface::EXC_VXC_INC, EvalType::Regular, + test_cuda_hip_interface( TestInterface::EXC_VXC_INC, EvalType::Regular, Backend::builtin, kern, Spin::Polarized ); } SECTION("EXC Small: Polarized") { for( auto kern : builtin_supported_kernels ) { if(is_unstable_small(kern)) continue; - test_cuda_interface( TestInterface::EXC, EvalType::Small, + test_cuda_hip_interface( TestInterface::EXC, EvalType::Small, Backend::builtin, kern, Spin::Polarized ); } } @@ -1747,7 +1743,7 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { SECTION("EXC + VXC Small: Polarized") { for( auto kern : builtin_supported_kernels ) { if(is_unstable_small(kern)) continue; - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Small, + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Small, Backend::builtin, kern, Spin::Polarized ); } } @@ -1755,7 +1751,7 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { SECTION("EXC + INC Small: Polarized") { for( auto kern : builtin_supported_kernels ) { if(is_unstable_small(kern)) continue; - test_cuda_interface( TestInterface::EXC_INC, EvalType::Small, + test_cuda_hip_interface( TestInterface::EXC_INC, EvalType::Small, Backend::builtin, kern, Spin::Polarized ); } } @@ -1763,35 +1759,38 @@ TEST_CASE( "CUDA Interfaces", "[xc-device]" ) { SECTION("EXC + VXC + INC Small: Polarized") { for( auto kern : builtin_supported_kernels ) { if(is_unstable_small(kern)) continue; - test_cuda_interface( TestInterface::EXC_VXC_INC, EvalType::Small, + test_cuda_hip_interface( TestInterface::EXC_VXC_INC, EvalType::Small, Backend::builtin, kern, Spin::Polarized ); } } SECTION("EXC Zero: Polarized") { for( auto kern : builtin_supported_kernels ) - test_cuda_interface( TestInterface::EXC, EvalType::Zero, + test_cuda_hip_interface( TestInterface::EXC, EvalType::Zero, Backend::builtin, kern, Spin::Polarized ); } SECTION("EXC + VXC Zero: Polarized") { for( auto kern : builtin_supported_kernels ) - test_cuda_interface( TestInterface::EXC_VXC, EvalType::Zero, + test_cuda_hip_interface( TestInterface::EXC_VXC, EvalType::Zero, Backend::builtin, kern, Spin::Polarized ); } SECTION("EXC + INC Zero: Polarized") { for( auto kern : builtin_supported_kernels ) - test_cuda_interface( TestInterface::EXC_INC, EvalType::Zero, + test_cuda_hip_interface( TestInterface::EXC_INC, EvalType::Zero, Backend::builtin, kern, Spin::Polarized ); } SECTION("EXC + VXC + INC Zero: Polarized") { for( auto kern : builtin_supported_kernels ) - test_cuda_interface( TestInterface::EXC_VXC_INC, EvalType::Zero, + test_cuda_hip_interface( TestInterface::EXC_VXC_INC, EvalType::Zero, Backend::builtin, kern, Spin::Polarized ); } + } + + } #endif