Merge Fix new algorithm script and DPC++ flags.

Some DPC++ related improvements.

+ Add DPC++ to the new algorithm script.
+ Extract DPC++ version in CMake
+ Add the proper flag for GPU complex upport.
+ Remove one of the issues from the customized DPC++
   `abs` implementation.

Related PR: #661

cmake/GinkgoConfig.cmake.in

@@ -67,10 +67,14 @@ set(GINKGO_CUDA_HOST_COMPILER @CMAKE_CUDA_HOST_COMPILER@)
 set(GINKGO_HIP_COMPILER_FLAGS @GINKGO_HIP_COMPILER_FLAGS@)
 set(GINKGO_HIP_HCC_COMPILER_FLAGS @GINKGO_HIP_HCC_COMPILER_FLAGS@)
 set(GINKGO_HIP_NVCC_COMPILER_FLAGS @GINKGO_HIP_NVCC_COMPILER_FLAGS@)
+set(GINKGO_HIP_CLANG_COMPILER_FLAGS @GINKGO_HIP_CLANG_COMPILER_FLAGS@)
 set(GINKGO_HIP_PLATFORM @GINKGO_HIP_PLATFORM@)
 set(GINKGO_HIP_AMDGPU @GINKGO_HIP_AMDGPU@)
 set(GINKGO_HIP_VERSION @GINKGO_HIP_VERSION@)
 
+set(GINKGO_DPCPP_VERSION @GINKGO_DPCPP_VERSION@)
+set(GINKGO_DPCPP_FLAGS @GINKGO_DPCPP_FLAGS@)
+
 set(GINKGO_HAVE_PAPI_SDE @GINKGO_HAVE_PAPI_SDE@)
 
 # Ginkgo external package variables

cmake/build_helpers.cmake

@@ -156,3 +156,23 @@ function(ginkgo_extract_clang_version CLANG_COMPILER GINKGO_CLANG_VERSION)
     file(REMOVE ${CMAKE_CURRENT_BINARY_DIR}/extract_clang_ver.cpp)
     file(REMOVE ${CMAKE_CURRENT_BINARY_DIR}/extract_clang_ver)
 endfunction()
+
+# Extract the DPC++ version
+function(ginkgo_extract_dpcpp_version DPCPP_COMPILER GINKGO_DPCPP_VERSION)
+    set(DPCPP_VERSION_PROG "#include <CL/sycl.hpp>\n#include <iostream>\n"
+        "int main() {std::cout << __SYCL_COMPILER_VERSION << '\\n'\;"
+        "return 0\;}")
+    file(WRITE "${CMAKE_CURRENT_BINARY_DIR}/extract_dpcpp_ver.cpp" ${DPCPP_VERSION_PROG})
+    execute_process(COMMAND ${DPCPP_COMPILER} ${CMAKE_CURRENT_BINARY_DIR}/extract_dpcpp_ver.cpp
+        -o ${CMAKE_CURRENT_BINARY_DIR}/extract_dpcpp_ver
+        ERROR_VARIABLE DPCPP_EXTRACT_VER_ERROR)
+    execute_process(COMMAND ${CMAKE_CURRENT_BINARY_DIR}/extract_dpcpp_ver
+        OUTPUT_VARIABLE FOUND_DPCPP_VERSION
+        OUTPUT_STRIP_TRAILING_WHITESPACE
+        ERROR_STRIP_TRAILING_WHITESPACE
+        )
+
+    set (${GINKGO_DPCPP_VERSION} "${FOUND_DPCPP_VERSION}" PARENT_SCOPE)
+    file(REMOVE ${CMAKE_CURRENT_BINARY_DIR}/extract_dpcpp_ver.cpp)
+    file(REMOVE ${CMAKE_CURRENT_BINARY_DIR}/extract_dpcpp_ver)
+endfunction()

dev_tools/scripts/create_new_algorithm.sh

@@ -99,11 +99,13 @@ TEMPLATE_FILES=(
     "${name}_kernels.cpp"
     "${name}_*.[ch]*"
     "${name}_kernels.hip.cpp"
+    "${name}_kernels.dp.cpp"
     "${name}.cpp"
     "${name}_kernels.cpp"
     "${name}_kernels.cpp"
     "${name}_kernels.cpp"
     "${name}_kernels.*"
+    "${name}_kernels.*"
 )
 CMAKE_FILES=(
     "core/CMakeLists.txt"
@@ -113,11 +115,13 @@ CMAKE_FILES=(
     "omp/CMakeLists.txt"
     "cuda/CMakeLists.txt"
     "hip/CMakeLists.txt"
+    "dpcpp/CMakeLists.txt"
     "core/test/$source_type/CMakeLists.txt"
     "reference/test/$source_type/CMakeLists.txt"
     "omp/test/$source_type/CMakeLists.txt"
     "cuda/test/$source_type/CMakeLists.txt"
     "hip/test/$source_type/CMakeLists.txt"
+    "dpcpp/test/$source_type/CMakeLists.txt"
 )
 TEMPLATE_FILES_LOCATIONS=(
     "core/$source_type"
@@ -127,11 +131,13 @@ TEMPLATE_FILES_LOCATIONS=(
     "omp/$source_type"
     "cuda/$source_type"
     "hip/$source_type"
+    "dpcpp/$source_type"
     "core/test/$source_type"
     "reference/test/$source_type"
     "omp/test/$source_type"
     "cuda/test/$source_type"
     "hip/test/$source_type"
+    "dpcpp/test/$source_type"
 )
 TEMPLATE_FILES_TYPES=(
     "$source_type file"
@@ -141,11 +147,13 @@ TEMPLATE_FILES_TYPES=(
     "OpenMP kernel file"
     "CUDA kernel file"
     "HIP kernel file"
+    "DPC++ kernel file"
     "unit tests for ${name} $source_type"
     "unit tests for ${name} reference kernels"
     "unit tests for ${name} OMP kernels"
     "unit tests for ${name} CUDA kernels"
     "unit tests for ${name} HIP kernels"
+    "unit tests for ${name} DPC++ kernels"
 )
 TEMPLATE_FILES_DESCRIPTIONS=(
     "This is where the ${name} algorithm needs to be implemented."
@@ -155,11 +163,13 @@ TEMPLATE_FILES_DESCRIPTIONS=(
     "OMP kernels for ${name} need to be implemented here."
     "CUDA kernels for ${name} need to be implemented here."
     "HIP kernels for ${name} need to be implemented here."
+    "DPC++ kernels for ${name} need to be implemented here."
     "This is where core related unit tests should be implemented, i.e. relating to the interface without executor usage."
     "This is where tests with the Reference executor should be implemented. Usually, this means comparing against previously known values."
     "This is where tests with the OpenMP executor should be implemented. Usually, this means comparing against a Reference execution."
     "This is where tests with the CUDA executor should be implemented. Usually, this means comparing against a Reference execution."
     "This is where tests with the HIP executor should be implemented. Usually, this means comparing against a Reference execution."
+    "This is where tests with the DPC++ executor should be implemented. Usually, this means comparing against a Reference execution."
 )
 
 mkdir ${TMPDIR}
@@ -389,6 +399,10 @@ then
     echo "hip/test/${source_type}/CMakeLists.txt"             | tee -a todo_${name}.txt
     echo ""                                                   | tee -a todo_${name}.txt
     echo ""                                                   | tee -a todo_${name}.txt
+    echo "dpcpp/CMakeLists.txt"                               | tee -a todo_${name}.txt
+    echo "dpcpp/test/${source_type}/CMakeLists.txt"           | tee -a todo_${name}.txt
+    echo ""                                                   | tee -a todo_${name}.txt
+    echo ""                                                   | tee -a todo_${name}.txt
     echo "The following header file has to be modified:"      | tee -a todo_${name}.txt
     echo "core/device_hooks/common_kernels.inc.cpp"           | tee -a todo_${name}.txt
     echo "Equivalent to the other solvers, the following part has to be appended:"  | tee -a todo_${name}.txt

dpcpp/CMakeLists.txt

@@ -3,6 +3,9 @@ if (NOT GKO_CAN_COMPILE_DPCPP)
         "${CMAKE_CXX_COMPILER} cannot compile DPC++ code!")
 endif()
 
+ginkgo_extract_dpcpp_version(${CMAKE_CXX_COMPILER} GINKGO_DPCPP_VERSION)
+set(GINKGO_DPCPP_VERSION ${GINKGO_DPCPP_VERSION} PARENT_SCOPE)
+
 add_library(ginkgo_dpcpp $<TARGET_OBJECTS:ginkgo_dpcpp_device> "")
 target_sources(ginkgo_dpcpp
     PRIVATE
@@ -42,7 +45,8 @@ target_sources(ginkgo_dpcpp
 
 ginkgo_compile_features(ginkgo_dpcpp)
 
-set(GINKGO_DPCPP_FLAGS ${GINKGO_COMPILER_FLAGS} -fsycl)
+set(GINKGO_DPCPP_FLAGS ${GINKGO_COMPILER_FLAGS} -fsycl -fsycl-device-lib=libc)
+set(GINKGO_DPCPP_FLAGS ${GINKGO_DPCPP_FLAGS} PARENT_SCOPE)
 target_compile_options(ginkgo_dpcpp PRIVATE "${GINKGO_DPCPP_FLAGS}")
 target_compile_features(ginkgo_dpcpp PRIVATE cxx_std_17)
 

include/ginkgo/core/base/math.hpp

@@ -914,12 +914,9 @@ GKO_INLINE GKO_ATTRIBUTES constexpr xstd::enable_if_t<is_complex_s<T>::value,
 abs(const T &x)
 {
 #ifdef CL_SYCL_LANGUAGE_VERSION
-    // FIXME: This implementation is due to two DPC++ issues
-    // 1) plain `sqrt` call evaluates to `std::sqrt` which fails on GPUs
-    // 2) DPC++ GPUs require the OpenCL backend (`SYCL_BE=OPENCL`) to support
-    //    complex multiplication, on top of extra binary objects. For now, we
-    //    avoid this issue with the current implementation.
-    return cl::sycl::sqrt(real(x) * real(x) + imag(x) * imag(x));
+    // FIXME: This implementation is due to a DPC++ issue:
+    // plain `sqrt` call evaluates to `std::sqrt` which fails on GPUs
+    return cl::sycl::sqrt(squared_norm(x));
 #else
     return sqrt(squared_norm(x));
 #endif