Merge pull request #499 from IntelPython/use-dpcpp-built-kernels

Added a test that retrieved DPCPP compiled kernel and submits them via DPCTLQueue_SubmitRange

Author: Diptorup Deb

dpctl-capi/CMakeLists.txt

@@ -86,6 +86,7 @@ elseif(UNIX)
         "-Wuninitialized "
         "-Wmissing-declarations "
         "-fdiagnostics-color=auto "
+        "-Wno-deprecated-declarations "
     )
     string(CONCAT SDL_FLAGS
         "-fstack-protector "

dpctl-capi/dbg_build.sh

@@ -13,7 +13,7 @@ cmake                                                       \
     -DCMAKE_CXX_COMPILER=dpcpp                              \
     -DCMAKE_INSTALL_PREFIX=${INSTALL_PREFIX}                \
     -DCMAKE_PREFIX_PATH=${INSTALL_PREFIX}                   \
-    -DDPCTL_ENABLE_LO_PROGRAM_CREATION=${USE_LO_HEADERS}    \
+    -DDPCTL_ENABLE_LO_PROGRAM_CREATION=ON                   \
     -DDPCTL_BUILD_CAPI_TESTS=ON                             \
     -DDPCTL_GENERATE_COVERAGE=ON                            \
     ..

dpctl-capi/tests/CMakeLists.txt

@@ -35,11 +35,12 @@ if(DPCTL_GENERATE_COVERAGE)
     list(REMOVE_ITEM
         dpctl_sources
         "${CMAKE_CURRENT_SOURCE_DIR}/../source/dpctl_vector_templ.cpp"
+        "${CMAKE_CURRENT_SOURCE_DIR}/../source/dpcpp_kernels.cpp"
     )
 
     # Add profiling flags
     set(CMAKE_CXX_FLAGS
-        "${CMAKE_CXX_FLAGS} -fprofile-instr-generate -fcoverage-mapping"
+        "${CMAKE_CXX_FLAGS} -fprofile-instr-generate -fcoverage-mapping -DDPCTL_COVERAGE"
     )
 
     # Add all dpctl sources into a single executable so that we can run coverage
@@ -89,13 +90,18 @@ if(DPCTL_GENERATE_COVERAGE)
         WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}
     )
 else()
+    add_library(dpcpp_kernels
+        STATIC
+        ${CMAKE_CURRENT_SOURCE_DIR}/dpcpp_kernels.cpp
+    )
     file(GLOB_RECURSE sources ${CMAKE_CURRENT_SOURCE_DIR}/*.cpp)
     add_executable(dpctl_c_api_tests EXCLUDE_FROM_ALL ${sources})
     target_link_libraries(dpctl_c_api_tests
         ${CMAKE_THREAD_LIBS_INIT}
         GTest::GTest
         DPCTLSyclInterface
         ${LEVEL_ZERO_LIBRARY}
+        dpcpp_kernels
     )
 endif()
 

dpctl-capi/tests/dpcpp_kernels.cpp

@@ -0,0 +1,89 @@
+#include "dpcpp_kernels.hpp"
+#include <CL/sycl.hpp>
+#include <cstddef>
+
+template sycl::kernel
+dpcpp_kernels::get_fill_kernel<int>(sycl::queue &, size_t, int *, int);
+
+template sycl::kernel
+dpcpp_kernels::get_fill_kernel<unsigned int>(sycl::queue &,
+                                             size_t,
+                                             unsigned int *,
+                                             unsigned int);
+
+template sycl::kernel
+dpcpp_kernels::get_fill_kernel<double>(sycl::queue &, size_t, double *, double);
+
+template sycl::kernel
+dpcpp_kernels::get_fill_kernel<float>(sycl::queue &, size_t, float *, float);
+
+template sycl::kernel
+dpcpp_kernels::get_range_kernel<int>(sycl::queue &, size_t, int *);
+
+template sycl::kernel
+dpcpp_kernels::get_range_kernel<unsigned int>(sycl::queue &,
+                                              size_t,
+                                              unsigned int *);
+
+template sycl::kernel
+dpcpp_kernels::get_range_kernel<float>(sycl::queue &, size_t, float *);
+
+template sycl::kernel
+dpcpp_kernels::get_range_kernel<double>(sycl::queue &, size_t, double *);
+
+template sycl::kernel dpcpp_kernels::get_mad_kernel<int, int>(sycl::queue &,
+                                                              size_t,
+                                                              int *,
+                                                              int *,
+                                                              int *,
+                                                              int);
+
+template sycl::kernel
+dpcpp_kernels::get_mad_kernel<unsigned int, unsigned int>(sycl::queue &,
+                                                          size_t,
+                                                          unsigned int *,
+                                                          unsigned int *,
+                                                          unsigned int *,
+                                                          unsigned int);
+
+template sycl::kernel dpcpp_kernels::get_local_sort_kernel<int>(sycl::queue &,
+                                                                size_t,
+                                                                size_t,
+                                                                int *,
+                                                                size_t);
+
+template sycl::kernel
+dpcpp_kernels::get_local_count_exceedance_kernel<int>(sycl::queue &,
+                                                      size_t,
+                                                      size_t,
+                                                      int *,
+                                                      size_t,
+                                                      int,
+                                                      int *);
+
+template sycl::kernel
+dpcpp_kernels::get_local_count_exceedance_kernel<unsigned int>(sycl::queue &,
+                                                               size_t,
+                                                               size_t,
+                                                               unsigned int *,
+                                                               size_t,
+                                                               unsigned int,
+                                                               int *);
+
+template sycl::kernel
+dpcpp_kernels::get_local_count_exceedance_kernel<float>(sycl::queue &,
+                                                        size_t,
+                                                        size_t,
+                                                        float *,
+                                                        size_t,
+                                                        float,
+                                                        int *);
+
+template sycl::kernel
+dpcpp_kernels::get_local_count_exceedance_kernel<double>(sycl::queue &,
+                                                         size_t,
+                                                         size_t,
+                                                         double *,
+                                                         size_t,
+                                                         double,
+                                                         int *);