Skip to content

[SYCL] Replace printf with syclex::printf for device-side #45

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
wavefunction91 merged 4 commits into from
Sep 30, 2025
Merged

[SYCL] Replace printf with syclex::printf for device-side #45

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
1e8410a
[SYCL] fix printf for SYCL
abagusetty Sep 30, 2025
0d5dc21
fix readme
abagusetty Sep 30, 2025
2740380
Update deorbitalized.hpp
abagusetty Sep 30, 2025
210b503
fix build issue
abagusetty Sep 30, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion README.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 |
Expand Down
16 changes: 12 additions & 4 deletions include/exchcxx/impl/builtin/kernels/deorbitalized.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -141,8 +141,10 @@ struct kernel_traits<Deorbitalized<XCEF,KEDF>> {
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");
Expand Down Expand Up @@ -170,8 +172,10 @@ struct kernel_traits<Deorbitalized<XCEF,KEDF>> {
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");
Expand All @@ -184,8 +188,10 @@ struct kernel_traits<Deorbitalized<XCEF,KEDF>> {
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");
Expand All @@ -210,8 +216,10 @@ struct kernel_traits<Deorbitalized<XCEF,KEDF>> {
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");
Expand Down
5 changes: 1 addition & 4 deletions include/exchcxx/impl/builtin/util.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand Down