[SYCL] Sync SYCL with latest CUDA backend (#44)

* [SYCL] refactor SYCL

* [SYCL] replace sycl::printf with printf

* Revert "[SYCL] replace sycl::printf with printf"

This reverts commit 0d13be1.

* [SYCL] add AoT build options for SYCL backend

* [SYCL] fix sycl kernel names

* [SYCL] AoT compilation options

* [SYCL] remove fsycl-device-only flag

* [SYCL] fix PR comments

* [SYCL] restore white-spaces

* [SYCL] AoT builds

* [SYCL] fix cmake changes

* [SYCL] address white-space

* restore the entire headers

* redo util.hpp

Author: abagusetty

CMakeLists.txt

@@ -12,7 +12,6 @@ option( EXCHCXX_ENABLE_SYCL      "Enable Device Code (SYCL)"    OFF )
 option( EXCHCXX_ENABLE_LIBXC     "Enable Libxc Backend"         ON  )
 option( BUILD_SHARED_LIBS        "Build Shared Libs"            OFF )
 
-
 # Decided if we're compiling device bindings
 if( EXCHCXX_ENABLE_CUDA OR EXCHCXX_ENABLE_SYCL OR EXCHCXX_ENABLE_HIP )
   set( EXCHCXX_ENABLE_DEVICE TRUE CACHE BOOL "Enable Device Code" )
@@ -30,6 +29,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")

README.md

@@ -22,8 +22,8 @@ small subset of XC functionals which may be evaluated either on the host (CPU)
 or device (GPU, FPGA, etc). Currently GPU support is provided through the
 [CUDA](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html) for NVIDIA
 GPUs, [HIP](https://rocmdocs.amd.com/en/latest/Programming_Guides/HIP-GUIDE.html) for
-AMD GPUs and [SYCL](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html) (experimental)
-for generic accelerator backends (including Intel GPUs).
+AMD GPUs and [SYCL](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html) (experimental,
+supports only oneAPI implementaion) for generic accelerator backends (including Intel GPUs).
 
 
 ExchCXX is a work in progress. Its development has been funded by the U.S.

include/exchcxx/impl/builtin/util.hpp

@@ -58,6 +58,13 @@
 
 namespace ExchCXX {
 
+#if defined(__CUDACC__) || defined(__HIPCC__)
+#define EXCHCXX_READONLY_TABLE  static __device__
+#elif defined(__SYCL_DEVICE_ONLY__)
+#define EXCHCXX_READONLY_TABLE  inline constexpr
+#else
+#define EXCHCXX_READONLY_TABLE  static
+#endif
 
 #if defined(__CUDACC__) || defined(__HIPCC__)
 
@@ -154,52 +161,37 @@ static double AE11_data[39] = {
   -0.000000000000000024, -0.000000000000000201, -0.000000000000000082,  0.000000000000000017
 };
 
-#if defined(__CUDACC__) || defined(__HIPCC__)
-__device__
-#endif
-static double AE12_data[25] = {
+EXCHCXX_READONLY_TABLE 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,
   -0.000000000015830222,  0.000000000002905732,  0.000000000001769356,  0.000000000000492735,  0.000000000000093709,
    0.000000000000010707, -0.000000000000000537, -0.000000000000000716, -0.000000000000000244, -0.000000000000000058
 };
 
-#if defined(__CUDACC__) || defined(__HIPCC__)
-__device__
-#endif
-static double E11_data[19] = {
+EXCHCXX_READONLY_TABLE 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__
-#endif
-static double E12_data[16] = {
+EXCHCXX_READONLY_TABLE 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__
-#endif
-static double AE13_data[25] = {
+EXCHCXX_READONLY_TABLE 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,
   -0.000000000008733026,  0.000000000002023672, -0.000000000000474132,  0.000000000000112211, -0.000000000000026804,
    0.000000000000006457, -0.000000000000001568,  0.000000000000000383, -0.000000000000000094,  0.000000000000000023
 };
 
-#if defined(__CUDACC__) || defined(__HIPCC__)
-__device__
-#endif
-static double AE14_data[26] = {
+EXCHCXX_READONLY_TABLE 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,

include/exchcxx/util/exchcxx_macros.hpp

@@ -79,7 +79,7 @@
   #define DEVICE_PARAMS sycl::queue* queue
   #define DEVICE_PARAMS_NOTYPE queue
 
-  #define SYCL_KERNEL_PARAMS sycl::id<1> idx
+  #define SYCL_KERNEL_PARAMS sycl::id<1> tid
 
 #endif