Skip to content

[SYCL] Sync SYCL with latest CUDA backend #44

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 14 commits into from
Sep 30, 2025
Merged

[SYCL] Sync SYCL with latest CUDA backend #44

Show file tree
Hide file tree
Changes from 7 commits
Commits
Show all changes
14 commits
Select commit Hold shift + click to select a range
8bac863
[SYCL] refactor SYCL
abagusetty Sep 23, 2025
0d13be1
[SYCL] replace sycl::printf with printf
abagusetty Sep 23, 2025
a296a72
Revert "[SYCL] replace sycl::printf with printf"
abagusetty Sep 23, 2025
71ecc13
[SYCL] add AoT build options for SYCL backend
abagusetty Sep 24, 2025
4184d5d
[SYCL] fix sycl kernel names
abagusetty Sep 25, 2025
460737d
[SYCL] AoT compilation options
abagusetty Sep 25, 2025
4102e6e
[SYCL] remove fsycl-device-only flag
abagusetty Sep 26, 2025
2105ab7
[SYCL] fix PR comments
abagusetty Sep 29, 2025
0dc84b0
[SYCL] restore white-spaces
abagusetty Sep 29, 2025
6d008d7
[SYCL] AoT builds
abagusetty Sep 29, 2025
ecb0556
[SYCL] fix cmake changes
abagusetty Sep 30, 2025
089c587
[SYCL] address white-space
abagusetty Sep 30, 2025
b38d8f0
restore the entire headers
abagusetty Sep 30, 2025
85fd075
redo util.hpp
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
16 changes: 11 additions & 5 deletions CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,6 +30,12 @@ if( EXCHCXX_ENABLE_SYCL AND EXCHCXX_ENABLE_HIP )
endif()


if(EXCHCXX_ENABLE_SYCL)
# e.g. intel_gpu_pvc | nvidia_gpu_sm_80 | nvidia_gpu_sm_90 | amd_gpu_gfx90a | amd_gpu_gfx942
set(EXCHCXX_SYCL_TARGET "" CACHE STRING "Alias for -fsycl-targets (see Users Manual)")
endif()


# Append local cmake directory to find CMAKE Modules
if( CMAKE_MODULE_PATH )
list( APPEND CMAKE_MODULE_PATH "${PROJECT_SOURCE_DIR}/cmake")
Expand Down Expand Up @@ -69,7 +75,7 @@ if( ${Libxc_FOUND} )
else()

option( FETCHCONTENT_LIBXC_GIT_SHALLOW "Whether to use GIT_SHALLOW for FetchContent'ing libxc" ON )

FetchContent_Declare(
libxc
GIT_REPOSITORY https://gitlab.com/libxc/libxc.git
Expand All @@ -85,8 +91,8 @@ else()

FetchContent_MakeAvailable( libxc )
add_library( Libxc::xc ALIAS xc )
target_include_directories( xc
PUBLIC
target_include_directories( xc
PUBLIC
$<BUILD_INTERFACE:${libxc_SOURCE_DIR}/src>
$<BUILD_INTERFACE:${libxc_BINARY_DIR}/src>
$<BUILD_INTERFACE:${libxc_BINARY_DIR}>
Expand All @@ -98,15 +104,15 @@ else()
set_target_properties(xc PROPERTIES UNITY_BUILD OFF)
message(STATUS "Will disable unity-build for Libxc::xc")
endif()

set( BUILD_TESTING ${OLD_BUILD_TESTING} CACHE BOOL "" FORCE )

endif()

else( EXCHCXX_ENABLE_LIBXC )
set( Libxc_FOUND FALSE )
endif( EXCHCXX_ENABLE_LIBXC )

add_subdirectory( src )

# Testing
Expand Down
52 changes: 30 additions & 22 deletions include/exchcxx/impl/builtin/kernels/deorbitalized.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,30 +1,30 @@
/**
* ExchCXX
* ExchCXX
*
* Copyright (c) 2020-2024, The Regents of the University of California,
* through Lawrence Berkeley National Laboratory (subject to receipt of
* any required approvals from the U.S. Dept. of Energy).
* any required approvals from the U.S. Dept. of Energy).
*
* Portions Copyright (c) Microsoft Corporation.
*
* All rights reserved.
*
*
* Redistribution and use in source and binary forms, with or without
* modification, are permitted provided that the following conditions are met:
*
*
* (1) Redistributions of source code must retain the above copyright notice,
* this list of conditions and the following disclaimer.
*
*
* (2) Redistributions in binary form must reproduce the above copyright
* notice, this list of conditions and the following disclaimer in the
* documentation and/or other materials provided with the distribution.
*
*
* (3) Neither the name of the University of California, Lawrence Berkeley
* National Laboratory, U.S. Dept. of Energy nor the names of its contributors
* may be used to endorse or promote products derived from this software
* without specific prior written permission.
*
*
*
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
* AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
* IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
Expand All @@ -36,7 +36,7 @@
* CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
* ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
* POSSIBILITY OF SUCH DAMAGE.
*
*
* You are under no obligation whatsoever to provide any bug fixes, patches,
* or upgrades to the features, functionality or performance of the source
* code ("Enhancements") to anyone; however, if you choose to make your
Expand All @@ -49,7 +49,7 @@
* in binary and source code form.
*/

#pragma once
#pragma once
#include <exchcxx/impl/builtin/fwd.hpp>
#include <exchcxx/impl/builtin/kernel_type.hpp>
#include <exchcxx/impl/builtin/kernels/screening_interface.hpp>
Expand Down 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 All @@ -163,15 +165,17 @@ struct kernel_traits<Deorbitalized<XCEF,KEDF>> {
double& v2rhotau_a_a, double& v2rhotau_a_b, double& v2rhotau_b_a, double& v2rhotau_b_b,
double& v2sigma2_aa_aa, double& v2sigma2_aa_ab, double& v2sigma2_aa_bb,
double& v2sigma2_ab_ab, double& v2sigma2_ab_bb, double& v2sigma2_bb_bb,
double& v2sigmalapl_aa_a, double& v2sigmalapl_aa_b, double& v2sigmalapl_ab_a,
double& v2sigmalapl_aa_a, double& v2sigmalapl_aa_b, double& v2sigmalapl_ab_a,
double& v2sigmalapl_ab_b, double& v2sigmalapl_bb_a, double& v2sigmalapl_bb_b,
double& v2sigmatau_aa_a, double& v2sigmatau_aa_b, double& v2sigmatau_ab_a,
double& v2sigmatau_aa_a, double& v2sigmatau_aa_b, double& v2sigmatau_ab_a,
double& v2sigmatau_ab_b, double& v2sigmatau_bb_a, double& v2sigmatau_bb_b,
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 @@ -181,20 +185,22 @@ struct kernel_traits<Deorbitalized<XCEF,KEDF>> {

BUILTIN_KERNEL_EVAL_RETURN
eval_fxc_unpolar( double rho, double sigma, double lapl, double tau,
double& v2rho2, double& v2rhosigma, double& v2rholapl, double& v2rhotau,
double& v2sigma2, double& v2sigmalapl, double& v2sigmatau,
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");
#endif
}

BUILTIN_KERNEL_EVAL_RETURN
eval_fxc_polar( double rho_a, double rho_b,
double sigma_aa, double sigma_ab, double sigma_bb,
eval_fxc_polar( double rho_a, double rho_b,
double sigma_aa, double sigma_ab, double sigma_bb,
double lapl_a, double lapl_b, double tau_a, double tau_b,
double& v2rho2_aa, double& v2rho2_ab, double& v2rho2_bb,
double& v2rhosigma_a_aa, double& v2rhosigma_a_ab, double& v2rhosigma_a_bb,
Expand All @@ -203,15 +209,17 @@ struct kernel_traits<Deorbitalized<XCEF,KEDF>> {
double& v2rhotau_a_a, double& v2rhotau_a_b, double& v2rhotau_b_a, double& v2rhotau_b_b,
double& v2sigma2_aa_aa, double& v2sigma2_aa_ab, double& v2sigma2_aa_bb,
double& v2sigma2_ab_ab, double& v2sigma2_ab_bb, double& v2sigma2_bb_bb,
double& v2sigmalapl_aa_a, double& v2sigmalapl_aa_b, double& v2sigmalapl_ab_a,
double& v2sigmalapl_aa_a, double& v2sigmalapl_aa_b, double& v2sigmalapl_ab_a,
double& v2sigmalapl_ab_b, double& v2sigmalapl_bb_a, double& v2sigmalapl_bb_b,
double& v2sigmatau_aa_a, double& v2sigmatau_aa_b, double& v2sigmatau_ab_a,
double& v2sigmatau_aa_a, double& v2sigmatau_aa_b, double& v2sigmatau_ab_a,
double& v2sigmatau_ab_b, double& v2sigmatau_bb_a, double& v2sigmatau_bb_b,
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
69 changes: 47 additions & 22 deletions include/exchcxx/impl/builtin/util.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,30 +1,30 @@
/**
* ExchCXX
* ExchCXX
*
* Copyright (c) 2020-2024, The Regents of the University of California,
* through Lawrence Berkeley National Laboratory (subject to receipt of
* any required approvals from the U.S. Dept. of Energy).
* any required approvals from the U.S. Dept. of Energy).
*
* Portions Copyright (c) Microsoft Corporation.
*
* All rights reserved.
*
*
* Redistribution and use in source and binary forms, with or without
* modification, are permitted provided that the following conditions are met:
*
*
* (1) Redistributions of source code must retain the above copyright notice,
* this list of conditions and the following disclaimer.
*
*
* (2) Redistributions in binary form must reproduce the above copyright
* notice, this list of conditions and the following disclaimer in the
* documentation and/or other materials provided with the distribution.
*
*
* (3) Neither the name of the University of California, Lawrence Berkeley
* National Laboratory, U.S. Dept. of Energy nor the names of its contributors
* may be used to endorse or promote products derived from this software
* without specific prior written permission.
*
*
*
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
* AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
* IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
Expand All @@ -36,7 +36,7 @@
* CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
* ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
* POSSIBILITY OF SUCH DAMAGE.
*
*
* You are under no obligation whatsoever to provide any bug fixes, patches,
* or upgrades to the features, functionality or performance of the source
* code ("Enhancements") to anyone; however, if you choose to make your
Expand Down Expand Up @@ -118,7 +118,7 @@ SAFE_INLINE(auto) erf( T x ) { return sm::erf(x); }
template <typename T, typename U>
SAFE_INLINE(auto) pow( T x, U e ) { return sm::pow(x,e); }
template <typename T>
SAFE_INLINE(auto) xc_erfcx( T x ) { return sm::exp(x*x)*sm::erfc(x); }
SAFE_INLINE(auto) xc_erfcx( T x ) { return sm::exp(x*x)*sm::erfc(x); }



Expand All @@ -139,11 +139,16 @@ 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__
__device__ static
#elif defined(__SYCL_DEVICE_ONLY__)
inline constexpr
#else
static
#endif
static double AE11_data[39] = {
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 All @@ -155,9 +160,13 @@ static double AE11_data[39] = {
};

#if defined(__CUDACC__) || defined(__HIPCC__)
__device__
__device__ static
#elif defined(__SYCL_DEVICE_ONLY__)
inline constexpr
#else
static
#endif
static double AE12_data[25] = {
double AE12_data[25] = {
0.582417495134726740, -0.158348850905782750, -0.006764275590323141, 0.005125843950185725, 0.000435232492169391,
-0.000143613366305483, -0.000041801320556301, -0.000002713395758640, 0.000001151381913647, 0.000000420650022012,
0.000000066581901391, 0.000000000662143777, -0.000000002844104870, -0.000000000940724197, -0.000000000177476602,
Expand All @@ -166,29 +175,41 @@ static double AE12_data[25] = {
};

#if defined(__CUDACC__) || defined(__HIPCC__)
__device__
__device__ static
#elif defined(__SYCL_DEVICE_ONLY__)
inline constexpr
#else
static
#endif
static double E11_data[19] = {
double E11_data[19] = {
-16.11346165557149402600, 7.79407277874268027690, -1.95540581886314195070, 0.37337293866277945612, -0.05692503191092901938,
0.00721107776966009185, -0.00078104901449841593, 0.00007388093356262168, -0.00000620286187580820, 0.00000046816002303176,
-0.00000003209288853329, 0.00000000201519974874, -0.00000000011673686816, 0.00000000000627627066, -0.00000000000031481541,
0.00000000000001479904, -0.00000000000000065457, 0.00000000000000002733, -0.00000000000000000108
};

#if defined(__CUDACC__) || defined(__HIPCC__)
__device__
__device__ static
#elif defined(__SYCL_DEVICE_ONLY__)
inline constexpr
#else
static
#endif
static double E12_data[16] = {
double E12_data[16] = {
-0.03739021479220279500, 0.04272398606220957700, -0.13031820798497005440, 0.01441912402469889073, -0.00134617078051068022,
0.00010731029253063780, -0.00000742999951611943, 0.00000045377325690753, -0.00000002476417211390, 0.00000000122076581374,
-0.00000000005485141480, 0.00000000000226362142, -0.00000000000008635897, 0.00000000000000306291, -0.00000000000000010148,
0.00000000000000000315
};

#if defined(__CUDACC__) || defined(__HIPCC__)
__device__
__device__ static
#elif defined(__SYCL_DEVICE_ONLY__)
inline constexpr
#else
static
#endif
static double AE13_data[25] = {
double AE13_data[25] = {
-0.605773246640603460, -0.112535243483660900, 0.013432266247902779, -0.001926845187381145, 0.000309118337720603,
-0.000053564132129618, 0.000009827812880247, -0.000001885368984916, 0.000000374943193568, -0.000000076823455870,
0.000000016143270567, -0.000000003466802211, 0.000000000758754209, -0.000000000168864333, 0.000000000038145706,
Expand All @@ -197,9 +218,13 @@ static double AE13_data[25] = {
};

#if defined(__CUDACC__) || defined(__HIPCC__)
__device__
__device__ static
#elif defined(__SYCL_DEVICE_ONLY__)
inline constexpr
#else
static
#endif
static double AE14_data[26] = {
double AE14_data[26] = {
-0.18929180007530170, -0.08648117855259871, 0.00722410154374659, -0.00080975594575573, 0.00010999134432661,
-0.00001717332998937, 0.00000298562751447, -0.00000056596491457, 0.00000011526808397, -0.00000002495030440,
0.00000000569232420, -0.00000000135995766, 0.00000000033846628, -0.00000000008737853, 0.00000000002331588,
Expand Down
Loading