From 1e8410ad562e00aab4963710ea404e68059786c1 Mon Sep 17 00:00:00 2001 From: Abhishek Bagusetty Date: Tue, 30 Sep 2025 20:06:40 +0000 Subject: [PATCH 1/4] [SYCL] fix printf for SYCL --- .../impl/builtin/kernels/deorbitalized.hpp | 17 ++++++++++++----- 1 file changed, 12 insertions(+), 5 deletions(-) diff --git a/include/exchcxx/impl/builtin/kernels/deorbitalized.hpp b/include/exchcxx/impl/builtin/kernels/deorbitalized.hpp index 474f04c..dbdce6f 100644 --- a/include/exchcxx/impl/builtin/kernels/deorbitalized.hpp +++ b/include/exchcxx/impl/builtin/kernels/deorbitalized.hpp @@ -57,7 +57,6 @@ namespace ExchCXX { - template struct kernel_traits> { @@ -141,8 +140,10 @@ struct kernel_traits> { double& v2rho2, double& v2rhosigma, double& v2rholapl, double& v2rhotau, double& v2sigma2, double& v2sigmalapl, double& v2sigmatau, double& v2lapl2, double& v2lapltau, double& v2tau2 ) { - #if defined(__CUDACC__) || defined(__HIPCC__) || defined(__SYCL_DEVICE_ONLY__) || defined(EXCHCXX_ENABLE_SYCL) + #if defined(__CUDACC__) || defined(__HIPCC__) printf("eval_vxc_fxc_unpolar not implemented for deorbitalized kernels\n"); + #elif defined(__SYCL_DEVICE_ONLY__) || defined(EXCHCXX_ENABLE_SYCL) + sycl::ext::oneapi::experimental::printf("eval_vxc_fxc_unpolar not implemented for deorbitalized kernels\n"); #else unused(rho, sigma, lapl, tau, vrho, vsigma, vlapl, vtau, v2rho2, v2rhosigma, v2rholapl, v2rhotau, v2sigma2, v2sigmalapl, v2sigmatau, v2lapl2, v2lapltau, v2tau2); throw std::runtime_error("eval_vxc_fxc_unpolar not implemented for deorbitalized kernels"); @@ -170,8 +171,10 @@ struct kernel_traits> { double& v2lapl2_aa, double& v2lapl2_ab, double& v2lapl2_bb, double& v2lapltau_a_a, double& v2lapltau_a_b, double& v2lapltau_b_a, double& v2lapltau_b_b, double& v2tau2_aa, double& v2tau2_ab, double& v2tau2_bb ) { - #if defined(__CUDACC__) || defined(__HIPCC__) || defined(__SYCL_DEVICE_ONLY__) || defined(EXCHCXX_ENABLE_SYCL) + #if defined(__CUDACC__) || defined(__HIPCC__) printf("eval_vxc_fxc_polar not implemented for deorbitalized kernels\n"); + #elif defined(__SYCL_DEVICE_ONLY__) || defined(EXCHCXX_ENABLE_SYCL) + sycl::ext::oneapi::experimental::printf("eval_vxc_fxc_polar not implemented for deorbitalized kernels\n"); #else unused(rho_a, rho_b, sigma_aa, sigma_ab, sigma_bb, lapl_a, lapl_b, tau_a, tau_b, vrho_a, vrho_b, vsigma_aa, vsigma_ab, vsigma_bb, vlapl_a, vlapl_b, vtau_a, vtau_b, v2rho2_aa, v2rho2_ab, v2rho2_bb, v2rhosigma_a_aa, v2rhosigma_a_ab, v2rhosigma_a_bb, v2rhosigma_b_aa, v2rhosigma_b_ab, v2rhosigma_b_bb, v2rholapl_a_a, v2rholapl_a_b, v2rholapl_b_a, v2rholapl_b_b, v2rhotau_a_a, v2rhotau_a_b, v2rhotau_b_a, v2rhotau_b_b, v2sigma2_aa_aa, v2sigma2_aa_ab, v2sigma2_aa_bb, v2sigma2_ab_ab, v2sigma2_ab_bb, v2sigma2_bb_bb, v2sigmalapl_aa_a, v2sigmalapl_aa_b, v2sigmalapl_ab_a, v2sigmalapl_ab_b, v2sigmalapl_bb_a, v2sigmalapl_bb_b, v2sigmatau_aa_a, v2sigmatau_aa_b, v2sigmatau_ab_a, v2sigmatau_ab_b, v2sigmatau_bb_a, v2sigmatau_bb_b, v2lapl2_aa, v2lapl2_ab, v2lapl2_bb, v2lapltau_a_a, v2lapltau_a_b, v2lapltau_b_a, v2lapltau_b_b, v2tau2_aa, v2tau2_ab, v2tau2_bb); throw std::runtime_error("eval_vxc_fxc_polar not implemented for deorbitalized kernels"); @@ -184,8 +187,10 @@ struct kernel_traits> { double& v2rho2, double& v2rhosigma, double& v2rholapl, double& v2rhotau, double& v2sigma2, double& v2sigmalapl, double& v2sigmatau, double& v2lapl2, double& v2lapltau, double& v2tau2 ) { - #if defined(__CUDACC__) || defined(__HIPCC__) || defined(__SYCL_DEVICE_ONLY__) || defined(EXCHCXX_ENABLE_SYCL) + #if defined(__CUDACC__) || defined(__HIPCC__) printf("eval_fxc_unpolar not implemented for deorbitalized kernels\n"); + #elif defined(__SYCL_DEVICE_ONLY__) || defined(EXCHCXX_ENABLE_SYCL) + sycl::ext::oneapi::experimental::printf("eval_fxc_unpolar not implemented for deorbitalized kernels\n"); #else unused(rho, sigma, lapl, tau, v2rho2, v2rhosigma, v2rholapl, v2rhotau, v2sigma2, v2sigmalapl, v2sigmatau, v2lapl2, v2lapltau, v2tau2); throw std::runtime_error("eval_fxc_unpolar not implemented for deorbitalized kernels"); @@ -210,8 +215,10 @@ struct kernel_traits> { double& v2lapl2_aa, double& v2lapl2_ab, double& v2lapl2_bb, double& v2lapltau_a_a, double& v2lapltau_a_b, double& v2lapltau_b_a, double& v2lapltau_b_b, double& v2tau2_aa, double& v2tau2_ab, double& v2tau2_bb ) { - #if defined(__CUDACC__) || defined(__HIPCC__) || defined(__SYCL_DEVICE_ONLY__) || defined(EXCHCXX_ENABLE_SYCL) + #if defined(__CUDACC__) || defined(__HIPCC__) printf("eval_fxc_polar not implemented for deorbitalized kernels\n"); + #elif defined(__SYCL_DEVICE_ONLY__) || defined(EXCHCXX_ENABLE_SYCL) + sycl::ext::oneapi::experimental::printf("eval_fxc_polar not implemented for deorbitalized kernels\n"); #else unused(rho_a, rho_b, sigma_aa, sigma_ab, sigma_bb, lapl_a, lapl_b, tau_a, tau_b, v2rho2_aa, v2rho2_ab, v2rho2_bb, v2rhosigma_a_aa, v2rhosigma_a_ab, v2rhosigma_a_bb, v2rhosigma_b_aa, v2rhosigma_b_ab, v2rhosigma_b_bb, v2rholapl_a_a, v2rholapl_a_b, v2rholapl_b_a, v2rholapl_b_b, v2rhotau_a_a, v2rhotau_a_b, v2rhotau_b_a, v2rhotau_b_b, v2sigma2_aa_aa, v2sigma2_aa_ab, v2sigma2_aa_bb, v2sigma2_ab_ab, v2sigma2_ab_bb, v2sigma2_bb_bb, v2sigmalapl_aa_a, v2sigmalapl_aa_b, v2sigmalapl_ab_a, v2sigmalapl_ab_b, v2sigmalapl_bb_a, v2sigmalapl_bb_b, v2sigmatau_aa_a, v2sigmatau_aa_b, v2sigmatau_ab_a, v2sigmatau_ab_b, v2sigmatau_bb_a, v2sigmatau_bb_b, v2lapl2_aa, v2lapl2_ab, v2lapl2_bb, v2lapltau_a_a, v2lapltau_a_b, v2lapltau_b_a, v2lapltau_b_b, v2tau2_aa, v2tau2_ab, v2tau2_bb); throw std::runtime_error("eval_fxc_polar not implemented for deorbitalized kernels"); From 0d5dc21cb5271b193bde34b0fad17bc97ddb0a45 Mon Sep 17 00:00:00 2001 From: Abhishek Bagusetty Date: Tue, 30 Sep 2025 20:07:20 +0000 Subject: [PATCH 2/4] fix readme --- README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/README.md b/README.md index 34617a3..9b9c79e 100644 --- a/README.md +++ b/README.md @@ -144,7 +144,7 @@ b3lyp.eval_exc_vxc_device( npts, rho_device, gamma_device, exc_device, vrho_devi | Vosko-Wilk-Nusair V | `XC_LDA_C_VWN_RPA` | `Kernel::VWN5` | Y | | Perdew-Burke-Ernzerhof (Exchange) | `XC_GGA_X_PBE` | `Kernel::PBE_X` | Y | | Perdew-Burke-Ernzerhof (Correlation) | `XC_GGA_C_PBE` | `Kernel::PBE_C` | Y | -| Revised PBE from Zhang & Yang | `XC_GGA_X_PBE_R` | `Kernel::revPBE_X | Y | +| Revised PBE from Zhang & Yang | `XC_GGA_X_PBE_R` | `Kernel::revPBE_X` | Y | | Perdew-Wang 91 (LDA) | `XC_LDA_C_PW` | `Kernel::PW91_LDA` | Y | | Perdew-Wang 91 (LDA) Modified | `XC_LDA_C_PW_MOD` | `Kernel::PW91_LDA_MOD` | Y | | Perdew-Wang 91 (LDA) RPA | `XC_LDA_C_PW_RPA` | `Kernel::PW91_LDA_RPA` | Y | From 2740380a35e489de5eb7b10ddfce579a8644e296 Mon Sep 17 00:00:00 2001 From: Abhishek Bagusetty <59661409+abagusetty@users.noreply.github.com> Date: Tue, 30 Sep 2025 15:09:09 -0500 Subject: [PATCH 3/4] Update deorbitalized.hpp --- include/exchcxx/impl/builtin/kernels/deorbitalized.hpp | 1 + 1 file changed, 1 insertion(+) diff --git a/include/exchcxx/impl/builtin/kernels/deorbitalized.hpp b/include/exchcxx/impl/builtin/kernels/deorbitalized.hpp index dbdce6f..e00bc23 100644 --- a/include/exchcxx/impl/builtin/kernels/deorbitalized.hpp +++ b/include/exchcxx/impl/builtin/kernels/deorbitalized.hpp @@ -57,6 +57,7 @@ namespace ExchCXX { + template struct kernel_traits> { From 210b503bce561c462b203523a6748ffe22350520 Mon Sep 17 00:00:00 2001 From: Abhishek Bagusetty Date: Tue, 30 Sep 2025 20:25:10 +0000 Subject: [PATCH 4/4] fix build issue --- include/exchcxx/impl/builtin/util.hpp | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/include/exchcxx/impl/builtin/util.hpp b/include/exchcxx/impl/builtin/util.hpp index f7dbd8a..c5bb7bf 100644 --- a/include/exchcxx/impl/builtin/util.hpp +++ b/include/exchcxx/impl/builtin/util.hpp @@ -147,10 +147,7 @@ SAFE_INLINE(auto) xc_cheb_eval(const double x, const double *cs, const int N) return 0.5*(b0 - b2); } // The following data is taken from libxc -#if defined(__CUDACC__) || defined(__HIPCC__) -__device__ -#endif -static double AE11_data[39] = { +EXCHCXX_READONLY_TABLE double AE11_data[39] = { 0.121503239716065790, -0.065088778513550150, 0.004897651357459670, -0.000649237843027216, 0.000093840434587471, 0.000000420236380882, -0.000008113374735904, 0.000002804247688663, 0.000000056487164441, -0.000000344809174450, 0.000000058209273578, 0.000000038711426349, -0.000000012453235014, -0.000000005118504888, 0.000000002148771527,