Merge pull request #1589 from IntelPython/improve-coverage-sycl-inter…

…face

Improve coverage sycl interface

libsyclinterface/dbg_build.sh

@@ -7,6 +7,11 @@ pushd build || exit 1
 INSTALL_PREFIX=$(pwd)/../install
 rm -rf ${INSTALL_PREFIX}
 
+# With DPC++ 2024.0 adn newer set these to ensure that
+# cmake can find llvm-cov and other utilities
+LLVM_TOOLS_HOME=${CMPLR_ROOT}/bin/compiler
+PATH=$PATH:${CMPLR_ROOT}/bin/compiler
+
 cmake                                                       \
     -DCMAKE_BUILD_TYPE=Debug                                \
     -DCMAKE_C_COMPILER=icx                                  \
@@ -16,13 +21,19 @@ cmake                                                       \
     -DCMAKE_PREFIX_PATH=${INSTALL_PREFIX}                   \
     -DDPCTL_ENABLE_L0_PROGRAM_CREATION=ON                   \
     -DDPCTL_BUILD_CAPI_TESTS=ON                             \
+    -DDPCTL_GENERATE_COVERAGE=OFF                           \
     ..
 
-make V=1 -n -j 4 && make check && make install
+# build
+make V=1 -n -j 4
+# run ctest
+make check
+# install
+make install
 
 # Turn on to generate coverage report html files reconfigure with
 # -DDPCTL_GENERATE_COVERAGE=ON and then
-# make lcov-genhtml
+# make llvm-cov-report
 
 # For more verbose tests use:
 # cd tests

libsyclinterface/include/dpctl_sycl_queue_interface.h

@@ -171,6 +171,18 @@ DPCTL_API
 __dpctl_give DPCTLSyclDeviceRef
 DPCTLQueue_GetDevice(__dpctl_keep const DPCTLSyclQueueRef QRef);
 
+/*! @brief Structure to be used to specify dimensionality and type of
+ * local_accessor kernel type argument.
+ */
+typedef struct MDLocalAccessorTy
+{
+    size_t ndim;
+    DPCTLKernelArgType dpctl_type_id;
+    size_t dim0;
+    size_t dim1;
+    size_t dim2;
+} MDLocalAccessor;
+
 /*!
  * @brief Submits the kernel to the specified queue with the provided range
  * argument.

libsyclinterface/source/dpctl_sycl_queue_interface.cpp

@@ -121,15 +121,6 @@ typedef struct complex
     uint64_t imag;
 } complexNumber;
 
-typedef struct MDLocalAccessorTy
-{
-    size_t ndim;
-    DPCTLKernelArgType dpctl_type_id;
-    size_t dim0;
-    size_t dim1;
-    size_t dim2;
-} MDLocalAccessor;
-
 void set_dependent_events(handler &cgh,
                           __dpctl_keep const DPCTLSyclEventRef *DepEvents,
                           size_t NDepEvents)

libsyclinterface/tests/CMakeLists.txt

@@ -89,8 +89,35 @@ if(DPCTL_GENERATE_COVERAGE)
         ${CMAKE_DL_LIBS}
     )
     set(object_arg "-object;")
-    add_custom_target(llvm-cov
+    add_custom_target(run-c-api-tests
         COMMAND ${CMAKE_COMMAND} -E env DPCTL_VERBOSITY=warning ${CMAKE_CURRENT_BINARY_DIR}/dpctl_c_api_tests
+        WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}
+        COMMAND_EXPAND_LISTS
+        DEPENDS dpctl_c_api_tests
+    )
+    add_custom_target(llvm-cov-show
+        COMMAND ${LLVMProfdata_EXE}
+        merge
+        -sparse default.profraw
+        -o
+        dpctl.profdata
+        COMMAND ${LLVMCov_EXE}
+        export
+        -format=lcov
+        -ignore-filename-regex=/tmp/icpx*
+        -instr-profile=dpctl.profdata
+        "${object_arg}$<JOIN:$<TARGET_OBJECTS:DPCTLSyclInterface>,;${object_arg}>"
+        > dpctl.lcov
+        COMMAND ${LLVMCov_EXE}
+        show
+        -instr-profile=dpctl.profdata
+        "${object_arg}$<JOIN:$<TARGET_OBJECTS:DPCTLSyclInterface>,;${object_arg}>"
+        WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}
+        COMMAND_EXPAND_LISTS
+        DEPENDS run-c-api-tests
+    )
+
+    add_custom_target(llvm-cov-report
         COMMAND ${LLVMProfdata_EXE}
         merge
         -sparse default.profraw
@@ -109,11 +136,10 @@ if(DPCTL_GENERATE_COVERAGE)
         "${object_arg}$<JOIN:$<TARGET_OBJECTS:DPCTLSyclInterface>,;${object_arg}>"
         WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}
         COMMAND_EXPAND_LISTS
-        DEPENDS dpctl_c_api_tests
+        DEPENDS run-c-api-tests
     )
 
     add_custom_target(lcov-genhtml
-        COMMAND ${CMAKE_COMMAND} -E env DPCTL_VERBOSITY=warning ${CMAKE_CURRENT_BINARY_DIR}/dpctl_c_api_tests
         COMMAND ${LLVMProfdata_EXE}
         merge
         -sparse default.profraw
@@ -132,7 +158,7 @@ if(DPCTL_GENERATE_COVERAGE)
         ${COVERAGE_OUTPUT_DIR}/dpctl-c-api-coverage
         WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}
         COMMAND_EXPAND_LISTS
-        DEPENDS dpctl_c_api_tests
+        DEPENDS run-c-api-tests
     )
 else()
     target_link_libraries(dpctl_c_api_tests

libsyclinterface/tests/test_sycl_device_aspects.cpp

@@ -97,7 +97,7 @@ auto build_gtest_values(const std::array<std::pair<T1, T2>, N> &params)
 auto build_params()
 {
     constexpr auto param_1 = get_param_list<const char *>(
-        "opencl:gpu", "opencl:cpu", "level_zero:gpu", "host");
+        "opencl:gpu", "opencl:cpu", "level_zero:gpu");
 
     constexpr auto param_2 =
         get_param_list<std::pair<const char *, sycl::aspect>>(

libsyclinterface/tests/test_sycl_queue_interface.cpp

@@ -90,6 +90,34 @@ struct TestDPCTLQueueMemberFunctions
 
 } /* End of anonymous namespace */
 
+TEST(TestDPCTLSyclQueueInterface, CheckCreate)
+{
+    /* We are testing that we do not crash even when input is NULL. */
+    DPCTLSyclQueueRef QRef = nullptr;
+
+    EXPECT_NO_FATAL_FAILURE(
+        QRef = DPCTLQueue_Create(nullptr, nullptr, nullptr, 0));
+    ASSERT_TRUE(QRef == nullptr);
+}
+
+TEST(TestDPCTLSyclQueueInterface, CheckCreate2)
+{
+    /* We are testing that we do not crash even when input is NULL. */
+    DPCTLSyclQueueRef QRef = nullptr;
+    DPCTLSyclDeviceSelectorRef DSRef = nullptr;
+    DPCTLSyclDeviceRef DRef = nullptr;
+
+    EXPECT_NO_FATAL_FAILURE(DSRef = DPCTLDefaultSelector_Create());
+    EXPECT_NO_FATAL_FAILURE(DRef = DPCTLDevice_CreateFromSelector(DSRef));
+    EXPECT_NO_FATAL_FAILURE(DPCTLDeviceSelector_Delete(DSRef));
+
+    EXPECT_NO_FATAL_FAILURE(QRef =
+                                DPCTLQueue_Create(nullptr, DRef, nullptr, 0));
+    ASSERT_TRUE(QRef == nullptr);
+
+    EXPECT_NO_FATAL_FAILURE(DPCTLDevice_Delete(DRef));
+}
+
 TEST(TestDPCTLSyclQueueInterface, CheckCreateForDevice)
 {
     /* We are testing that we do not crash even when input is NULL. */

libsyclinterface/tests/test_sycl_queue_submit.cpp

@@ -56,7 +56,9 @@ void submit_kernel(DPCTLSyclQueueRef QRef,
 {
     T scalarVal = 3;
     constexpr size_t NARGS = 4;
-    constexpr size_t RANGE_NDIMS = 1;
+    constexpr size_t RANGE_NDIMS_1 = 1;
+    constexpr size_t RANGE_NDIMS_2 = 2;
+    constexpr size_t RANGE_NDIMS_3 = 3;
 
     ASSERT_TRUE(DPCTLKernelBundle_HasKernel(KBRef, kernelName.c_str()));
     auto kernel = DPCTLKernelBundle_GetKernel(KBRef, kernelName.c_str());
@@ -75,13 +77,33 @@ void submit_kernel(DPCTLSyclQueueRef QRef,
                          (void *)&scalarVal};
     DPCTLKernelArgType addKernelArgTypes[] = {DPCTL_VOID_PTR, DPCTL_VOID_PTR,
                                               DPCTL_VOID_PTR, kernelArgTy};
-    auto ERef = DPCTLQueue_SubmitRange(kernel, QRef, args, addKernelArgTypes,
-                                       NARGS, Range, RANGE_NDIMS, nullptr, 0);
-    ASSERT_TRUE(ERef != nullptr);
-    DPCTLQueue_Wait(QRef);
+    auto E1Ref =
+        DPCTLQueue_SubmitRange(kernel, QRef, args, addKernelArgTypes, NARGS,
+                               Range, RANGE_NDIMS_1, nullptr, 0);
+    ASSERT_TRUE(E1Ref != nullptr);
+
+    // Create kernel args for vector_add
+    size_t Range2D[] = {SIZE, 1};
+    DPCTLSyclEventRef DepEvs[] = {E1Ref};
+    auto E2Ref =
+        DPCTLQueue_SubmitRange(kernel, QRef, args, addKernelArgTypes, NARGS,
+                               Range2D, RANGE_NDIMS_2, DepEvs, 1);
+    ASSERT_TRUE(E2Ref != nullptr);
+
+    // Create kernel args for vector_add
+    size_t Range3D[] = {SIZE, 1, 1};
+    DPCTLSyclEventRef DepEvs2[] = {E1Ref, E2Ref};
+    auto E3Ref =
+        DPCTLQueue_SubmitRange(kernel, QRef, args, addKernelArgTypes, NARGS,
+                               Range3D, RANGE_NDIMS_3, DepEvs2, 2);
+    ASSERT_TRUE(E3Ref != nullptr);
+
+    DPCTLEvent_Wait(E3Ref);
 
     // clean ups
-    DPCTLEvent_Delete(ERef);
+    DPCTLEvent_Delete(E1Ref);
+    DPCTLEvent_Delete(E2Ref);
+    DPCTLEvent_Delete(E3Ref);
     DPCTLKernel_Delete(kernel);
     DPCTLfree_with_queue((DPCTLSyclUSMRef)a, QRef);
     DPCTLfree_with_queue((DPCTLSyclUSMRef)b, QRef);
@@ -234,13 +256,13 @@ struct TestQueueSubmitFP64 : public ::testing::Test
     std::ifstream spirvFile;
     size_t spirvFileSize_;
     std::vector<char> spirvBuffer_;
+    DPCTLSyclDeviceRef DRef = nullptr;
     DPCTLSyclQueueRef QRef = nullptr;
     DPCTLSyclKernelBundleRef KBRef = nullptr;
 
     TestQueueSubmitFP64()
     {
         DPCTLSyclDeviceSelectorRef DSRef = nullptr;
-        DPCTLSyclDeviceRef DRef = nullptr;
 
         spirvFile.open("./oneD_range_kernel_fp64.spv",
                        std::ios::binary | std::ios::ate);
@@ -257,13 +279,13 @@ struct TestQueueSubmitFP64 : public ::testing::Test
 
         KBRef = DPCTLKernelBundle_CreateFromSpirv(
             CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr);
-        DPCTLDevice_Delete(DRef);
         DPCTLDeviceSelector_Delete(DSRef);
     }
 
     ~TestQueueSubmitFP64()
     {
         spirvFile.close();
+        DPCTLDevice_Delete(DRef);
         DPCTLQueue_Delete(QRef);
         DPCTLKernelBundle_Delete(KBRef);
     }
@@ -334,9 +356,11 @@ TEST_F(TestQueueSubmit, CheckForFloat)
 
 TEST_F(TestQueueSubmitFP64, CheckForDouble)
 {
-    submit_kernel<double>(QRef, KBRef, spirvBuffer_, spirvFileSize_,
-                          DPCTLKernelArgType::DPCTL_FLOAT64_T,
-                          "_ZTS11RangeKernelIdE");
+    if (DPCTLDevice_HasAspect(DRef, DPCTLSyclAspectType::fp64)) {
+        submit_kernel<double>(QRef, KBRef, spirvBuffer_, spirvFileSize_,
+                              DPCTLKernelArgType::DPCTL_FLOAT64_T,
+                              "_ZTS11RangeKernelIdE");
+    }
 }
 
 TEST_F(TestQueueSubmit, CheckForUnsupportedArgTy)

libsyclinterface/tests/test_sycl_queue_submit_local_accessor_arg.cpp

@@ -44,15 +44,6 @@ constexpr size_t SIZE = 100;
 
 using namespace dpctl::syclinterface;
 
-typedef struct MDLocalAccessorTy
-{
-    size_t ndim;
-    DPCTLKernelArgType dpctl_type_id;
-    size_t dim0;
-    size_t dim1;
-    size_t dim2;
-} MDLocalAccessor;
-
 template <typename T>
 void submit_kernel(DPCTLSyclQueueRef QRef,
                    DPCTLSyclKernelBundleRef KBRef,
@@ -75,28 +66,49 @@ void submit_kernel(DPCTLSyclQueueRef QRef,
         a_ptr[i] = 0;
     }
 
-    auto la = MDLocalAccessor{1, kernelArgTy, SIZE / 10, 1, 1};
+    auto la1 = MDLocalAccessor{1, kernelArgTy, SIZE / 10, 1, 1};
 
     // Create kernel args for vector_add
     size_t gRange[] = {SIZE};
     size_t lRange[] = {SIZE / 10};
-    void *args[NARGS] = {unwrap<void>(a), (void *)&la};
+    void *args_1d[NARGS] = {unwrap<void>(a), (void *)&la1};
     DPCTLKernelArgType addKernelArgTypes[] = {DPCTL_VOID_PTR,
                                               DPCTL_LOCAL_ACCESSOR};
 
-    auto ERef =
-        DPCTLQueue_SubmitNDRange(kernel, QRef, args, addKernelArgTypes, NARGS,
-                                 gRange, lRange, RANGE_NDIMS, nullptr, 0);
-    ASSERT_TRUE(ERef != nullptr);
-    DPCTLQueue_Wait(QRef);
+    DPCTLSyclEventRef E1Ref = DPCTLQueue_SubmitNDRange(
+        kernel, QRef, args_1d, addKernelArgTypes, NARGS, gRange, lRange,
+        RANGE_NDIMS, nullptr, 0);
+    ASSERT_TRUE(E1Ref != nullptr);
+
+    DPCTLSyclEventRef DepEv1[] = {E1Ref};
+    auto la2 = MDLocalAccessor{2, kernelArgTy, SIZE / 10, 1, 1};
+    void *args_2d[NARGS] = {unwrap<void>(a), (void *)&la2};
+
+    DPCTLSyclEventRef E2Ref =
+        DPCTLQueue_SubmitNDRange(kernel, QRef, args_2d, addKernelArgTypes,
+                                 NARGS, gRange, lRange, RANGE_NDIMS, DepEv1, 1);
+    ASSERT_TRUE(E2Ref != nullptr);
+
+    DPCTLSyclEventRef DepEv2[] = {E1Ref, E2Ref};
+    auto la3 = MDLocalAccessor{3, kernelArgTy, SIZE / 10, 1, 1};
+    void *args_3d[NARGS] = {unwrap<void>(a), (void *)&la3};
+
+    DPCTLSyclEventRef E3Ref =
+        DPCTLQueue_SubmitNDRange(kernel, QRef, args_3d, addKernelArgTypes,
+                                 NARGS, gRange, lRange, RANGE_NDIMS, DepEv2, 2);
+    ASSERT_TRUE(E3Ref != nullptr);
+
+    DPCTLEvent_Wait(E3Ref);
 
     if (kernelArgTy != DPCTL_FLOAT32_T && kernelArgTy != DPCTL_FLOAT64_T)
         ASSERT_TRUE(a_ptr[0] == 20);
     else
         ASSERT_TRUE(a_ptr[0] == 20.0);
 
     // clean ups
-    DPCTLEvent_Delete(ERef);
+    DPCTLEvent_Delete(E1Ref);
+    DPCTLEvent_Delete(E2Ref);
+    DPCTLEvent_Delete(E3Ref);
     DPCTLKernel_Delete(kernel);
     DPCTLfree_with_queue((DPCTLSyclUSMRef)a, QRef);
 }
@@ -239,13 +251,13 @@ struct TestQueueSubmitWithLocalAccessorFP64 : public ::testing::Test
     std::ifstream spirvFile;
     size_t spirvFileSize_;
     std::vector<char> spirvBuffer_;
+    DPCTLSyclDeviceRef DRef = nullptr;
     DPCTLSyclQueueRef QRef = nullptr;
     DPCTLSyclKernelBundleRef KBRef = nullptr;
 
     TestQueueSubmitWithLocalAccessorFP64()
     {
         DPCTLSyclDeviceSelectorRef DSRef = nullptr;
-        DPCTLSyclDeviceRef DRef = nullptr;
 
         spirvFile.open("./local_accessor_kernel_fp64.spv",
                        std::ios::binary | std::ios::ate);
@@ -262,13 +274,13 @@ struct TestQueueSubmitWithLocalAccessorFP64 : public ::testing::Test
 
         KBRef = DPCTLKernelBundle_CreateFromSpirv(
             CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr);
-        DPCTLDevice_Delete(DRef);
         DPCTLDeviceSelector_Delete(DSRef);
     }
 
     ~TestQueueSubmitWithLocalAccessorFP64()
     {
         spirvFile.close();
+        DPCTLDevice_Delete(DRef);
         DPCTLQueue_Delete(QRef);
         DPCTLKernelBundle_Delete(KBRef);
     }
@@ -339,9 +351,11 @@ TEST_F(TestQueueSubmitWithLocalAccessor, CheckForFloat)
 
 TEST_F(TestQueueSubmitWithLocalAccessorFP64, CheckForDouble)
 {
-    submit_kernel<double>(QRef, KBRef, spirvBuffer_, spirvFileSize_,
-                          DPCTLKernelArgType::DPCTL_FLOAT64_T,
-                          "_ZTS14SyclKernel_SLMIdE");
+    if (DPCTLDevice_HasAspect(DRef, DPCTLSyclAspectType::fp64)) {
+        submit_kernel<double>(QRef, KBRef, spirvBuffer_, spirvFileSize_,
+                              DPCTLKernelArgType::DPCTL_FLOAT64_T,
+                              "_ZTS14SyclKernel_SLMIdE");
+    }
 }
 
 TEST_F(TestQueueSubmitWithLocalAccessor, CheckForUnsupportedArgTy)

scripts/gen_coverage.py

@@ -82,7 +82,7 @@ def run(
         .strip("\n")
     )
     subprocess.check_call(
-        ["cmake", "--build", ".", "--target", "llvm-cov"],
+        ["cmake", "--build", ".", "--target", "llvm-cov-report"],
         cwd=cmake_build_dir,
     )
     env["LLVM_PROFILE_FILE"] = "dpctl_pytest.profraw"