Skip to content

Commit aa5bc87

Browse files
abagusettyabagusetty
authored
[SYCL] Replace printf with syclex::printf for device-side (#45)
* [SYCL] fix printf for SYCL * fix readme * Update deorbitalized.hpp * fix build issue
1 parent 74a914b commit aa5bc87

File tree

3 files changed

+14
-9
lines changed

3 files changed

+14
-9
lines changed

‎README.md‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -144,7 +144,7 @@ b3lyp.eval_exc_vxc_device( npts, rho_device, gamma_device, exc_device, vrho_devi
144144
| Vosko-Wilk-Nusair V | `XC_LDA_C_VWN_RPA` | `Kernel::VWN5` | Y |
145145
| Perdew-Burke-Ernzerhof (Exchange) | `XC_GGA_X_PBE` | `Kernel::PBE_X` | Y |
146146
| Perdew-Burke-Ernzerhof (Correlation) | `XC_GGA_C_PBE` | `Kernel::PBE_C` | Y |
147-
| Revised PBE from Zhang & Yang | `XC_GGA_X_PBE_R` | `Kernel::revPBE_X | Y |
147+
| Revised PBE from Zhang & Yang | `XC_GGA_X_PBE_R` | `Kernel::revPBE_X` | Y |
148148
| Perdew-Wang 91 (LDA) | `XC_LDA_C_PW` | `Kernel::PW91_LDA` | Y |
149149
| Perdew-Wang 91 (LDA) Modified | `XC_LDA_C_PW_MOD` | `Kernel::PW91_LDA_MOD` | Y |
150150
| Perdew-Wang 91 (LDA) RPA | `XC_LDA_C_PW_RPA` | `Kernel::PW91_LDA_RPA` | Y |

‎include/exchcxx/impl/builtin/kernels/deorbitalized.hpp‎

Lines changed: 12 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -141,8 +141,10 @@ struct kernel_traits<Deorbitalized<XCEF,KEDF>> {
141141
double& v2rho2, double& v2rhosigma, double& v2rholapl, double& v2rhotau,
142142
double& v2sigma2, double& v2sigmalapl, double& v2sigmatau,
143143
double& v2lapl2, double& v2lapltau, double& v2tau2 ) {
144-
#if defined(__CUDACC__) || defined(__HIPCC__) || defined(__SYCL_DEVICE_ONLY__) || defined(EXCHCXX_ENABLE_SYCL)
144+
#if defined(__CUDACC__) || defined(__HIPCC__)
145145
printf("eval_vxc_fxc_unpolar not implemented for deorbitalized kernels\n");
146+
#elif defined(__SYCL_DEVICE_ONLY__) || defined(EXCHCXX_ENABLE_SYCL)
147+
sycl::ext::oneapi::experimental::printf("eval_vxc_fxc_unpolar not implemented for deorbitalized kernels\n");
146148
#else
147149
unused(rho, sigma, lapl, tau, vrho, vsigma, vlapl, vtau, v2rho2, v2rhosigma, v2rholapl, v2rhotau, v2sigma2, v2sigmalapl, v2sigmatau, v2lapl2, v2lapltau, v2tau2);
148150
throw std::runtime_error("eval_vxc_fxc_unpolar not implemented for deorbitalized kernels");
@@ -170,8 +172,10 @@ struct kernel_traits<Deorbitalized<XCEF,KEDF>> {
170172
double& v2lapl2_aa, double& v2lapl2_ab, double& v2lapl2_bb,
171173
double& v2lapltau_a_a, double& v2lapltau_a_b, double& v2lapltau_b_a, double& v2lapltau_b_b,
172174
double& v2tau2_aa, double& v2tau2_ab, double& v2tau2_bb ) {
173-
#if defined(__CUDACC__) || defined(__HIPCC__) || defined(__SYCL_DEVICE_ONLY__) || defined(EXCHCXX_ENABLE_SYCL)
175+
#if defined(__CUDACC__) || defined(__HIPCC__)
174176
printf("eval_vxc_fxc_polar not implemented for deorbitalized kernels\n");
177+
#elif defined(__SYCL_DEVICE_ONLY__) || defined(EXCHCXX_ENABLE_SYCL)
178+
sycl::ext::oneapi::experimental::printf("eval_vxc_fxc_polar not implemented for deorbitalized kernels\n");
175179
#else
176180
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);
177181
throw std::runtime_error("eval_vxc_fxc_polar not implemented for deorbitalized kernels");
@@ -184,8 +188,10 @@ struct kernel_traits<Deorbitalized<XCEF,KEDF>> {
184188
double& v2rho2, double& v2rhosigma, double& v2rholapl, double& v2rhotau,
185189
double& v2sigma2, double& v2sigmalapl, double& v2sigmatau,
186190
double& v2lapl2, double& v2lapltau, double& v2tau2 ) {
187-
#if defined(__CUDACC__) || defined(__HIPCC__) || defined(__SYCL_DEVICE_ONLY__) || defined(EXCHCXX_ENABLE_SYCL)
191+
#if defined(__CUDACC__) || defined(__HIPCC__)
188192
printf("eval_fxc_unpolar not implemented for deorbitalized kernels\n");
193+
#elif defined(__SYCL_DEVICE_ONLY__) || defined(EXCHCXX_ENABLE_SYCL)
194+
sycl::ext::oneapi::experimental::printf("eval_fxc_unpolar not implemented for deorbitalized kernels\n");
189195
#else
190196
unused(rho, sigma, lapl, tau, v2rho2, v2rhosigma, v2rholapl, v2rhotau, v2sigma2, v2sigmalapl, v2sigmatau, v2lapl2, v2lapltau, v2tau2);
191197
throw std::runtime_error("eval_fxc_unpolar not implemented for deorbitalized kernels");
@@ -210,8 +216,10 @@ struct kernel_traits<Deorbitalized<XCEF,KEDF>> {
210216
double& v2lapl2_aa, double& v2lapl2_ab, double& v2lapl2_bb,
211217
double& v2lapltau_a_a, double& v2lapltau_a_b, double& v2lapltau_b_a, double& v2lapltau_b_b,
212218
double& v2tau2_aa, double& v2tau2_ab, double& v2tau2_bb ) {
213-
#if defined(__CUDACC__) || defined(__HIPCC__) || defined(__SYCL_DEVICE_ONLY__) || defined(EXCHCXX_ENABLE_SYCL)
219+
#if defined(__CUDACC__) || defined(__HIPCC__)
214220
printf("eval_fxc_polar not implemented for deorbitalized kernels\n");
221+
#elif defined(__SYCL_DEVICE_ONLY__) || defined(EXCHCXX_ENABLE_SYCL)
222+
sycl::ext::oneapi::experimental::printf("eval_fxc_polar not implemented for deorbitalized kernels\n");
215223
#else
216224
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);
217225
throw std::runtime_error("eval_fxc_polar not implemented for deorbitalized kernels");

‎include/exchcxx/impl/builtin/util.hpp‎

Lines changed: 1 addition & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -147,10 +147,7 @@ SAFE_INLINE(auto) xc_cheb_eval(const double x, const double *cs, const int N)
147147
return 0.5*(b0 - b2);
148148
}
149149
// The following data is taken from libxc
150-
#if defined(__CUDACC__) || defined(__HIPCC__)
151-
__device__
152-
#endif
153-
static double AE11_data[39] = {
150+
EXCHCXX_READONLY_TABLE double AE11_data[39] = {
154151
0.121503239716065790, -0.065088778513550150, 0.004897651357459670, -0.000649237843027216, 0.000093840434587471,
155152
0.000000420236380882, -0.000008113374735904, 0.000002804247688663, 0.000000056487164441, -0.000000344809174450,
156153
0.000000058209273578, 0.000000038711426349, -0.000000012453235014, -0.000000005118504888, 0.000000002148771527,

0 commit comments

Comments
 (0)