Allow setting local_accessors as kernel arguments in DPCTLQueue_Submit

dpctl/_backend.pxd

@@ -67,7 +67,8 @@ cdef extern from "syclinterface/dpctl_sycl_enum_types.h":
         _UINT64_T           'DPCTL_UINT64_T',
         _FLOAT              'DPCTL_FLOAT32_T',
         _DOUBLE             'DPCTL_FLOAT64_T',
-        _VOID_PTR           'DPCTL_VOID_PTR'
+        _VOID_PTR           'DPCTL_VOID_PTR',
+        _LOCAL_ACCESSOR     'DPCTL_LOCAL_ACCESSOR'
 
     ctypedef enum _queue_property_type 'DPCTLQueuePropertyType':
         _DEFAULT_PROPERTY   'DPCTL_DEFAULT_PROPERTY'

dpctl/_sycl_queue.pyx

@@ -233,6 +233,15 @@ cdef class _kernel_arg_type:
             _arg_data_type._VOID_PTR
         )
 
+    @property
+    def dpctl_local_accessor(self):
+        cdef str p_name = "dpctl_local_accessor"
+        return kernel_arg_type_attribute(
+            self._name,
+            p_name,
+            _arg_data_type._LOCAL_ACCESSOR
+        )
+
 
 kernel_arg_type = _kernel_arg_type()
 

dpctl/enum_types.py

@@ -22,11 +22,7 @@
 """
 from enum import Enum, auto
 
-__all__ = [
-    "device_type",
-    "backend_type",
-    "event_status_type",
-]
+__all__ = ["device_type", "backend_type", "event_status_type"]
 
 
 class device_type(Enum):

dpctl/tests/test_sycl_kernel_submit.py

@@ -274,3 +274,4 @@ def test_kernel_arg_type():
     _check_kernel_arg_type_instance(kernel_arg_type.dpctl_float32)
     _check_kernel_arg_type_instance(kernel_arg_type.dpctl_float64)
     _check_kernel_arg_type_instance(kernel_arg_type.dpctl_void_ptr)
+    _check_kernel_arg_type_instance(kernel_arg_type.dpctl_local_accessor)

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/helper/include/dpctl_error_handlers.h

@@ -20,7 +20,7 @@
 ///
 /// \file
 /// A functor to use for passing an error handler callback function to sycl
-/// context and queue contructors.
+/// context and queue constructors.
 //===----------------------------------------------------------------------===//
 
 #pragma once

libsyclinterface/include/dpctl_sycl_enum_types.h

@@ -98,6 +98,7 @@ typedef enum
     DPCTL_FLOAT32_T,
     DPCTL_FLOAT64_T,
     DPCTL_VOID_PTR,
+    DPCTL_LOCAL_ACCESSOR,
     DPCTL_UNSUPPORTED_KERNEL_ARG
 } DPCTLKernelArgType;
 

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_kernel_bundle_interface.cpp

@@ -530,7 +530,7 @@ _GetKernel_ze_impl(const kernel_bundle<bundle_state::executable> &kb,
     else {
         error_handler("Kernel named " + std::string(kernel_name) +
                           " could not be found.",
-                      __FILE__, __func__, __LINE__);
+                      __FILE__, __func__, __LINE__, error_level::error);
         return nullptr;
     }
 }
@@ -541,7 +541,7 @@ bool _HasKernel_ze_impl(const kernel_bundle<bundle_state::executable> &kb,
     auto zeKernelCreateFn = get_zeKernelCreate();
     if (zeKernelCreateFn == nullptr) {
         error_handler("Could not load zeKernelCreate function.", __FILE__,
-                      __func__, __LINE__);
+                      __func__, __LINE__, error_level::error);
         return false;
     }
 
@@ -564,7 +564,7 @@ bool _HasKernel_ze_impl(const kernel_bundle<bundle_state::executable> &kb,
             if (ze_status != ZE_RESULT_ERROR_INVALID_KERNEL_NAME) {
                 error_handler("zeKernelCreate failed: " +
                                   _GetErrorCode_ze_impl(ze_status),
-                              __FILE__, __func__, __LINE__);
+                              __FILE__, __func__, __LINE__, error_level::error);
                 return false;
             }
         }

libsyclinterface/source/dpctl_sycl_queue_interface.cpp

@@ -38,6 +38,76 @@
 
 using namespace sycl;
 
+#define SET_LOCAL_ACCESSOR_ARG(CGH, NDIM, ARGTY, R, IDX)                       \
+    do {                                                                       \
+        switch ((ARGTY)) {                                                     \
+        case DPCTL_INT8_T:                                                     \
+        {                                                                      \
+            auto la = local_accessor<int8_t, NDIM>(R, CGH);                    \
+            CGH.set_arg(IDX, la);                                              \
+            return true;                                                       \
+        }                                                                      \
+        case DPCTL_UINT8_T:                                                    \
+        {                                                                      \
+            auto la = local_accessor<uint8_t, NDIM>(R, CGH);                   \
+            CGH.set_arg(IDX, la);                                              \
+            return true;                                                       \
+        }                                                                      \
+        case DPCTL_INT16_T:                                                    \
+        {                                                                      \
+            auto la = local_accessor<int16_t, NDIM>(R, CGH);                   \
+            CGH.set_arg(IDX, la);                                              \
+            return true;                                                       \
+        }                                                                      \
+        case DPCTL_UINT16_T:                                                   \
+        {                                                                      \
+            auto la = local_accessor<uint16_t, NDIM>(R, CGH);                  \
+            CGH.set_arg(IDX, la);                                              \
+            return true;                                                       \
+        }                                                                      \
+        case DPCTL_INT32_T:                                                    \
+        {                                                                      \
+            auto la = local_accessor<int32_t, NDIM>(R, CGH);                   \
+            CGH.set_arg(IDX, la);                                              \
+            return true;                                                       \
+        }                                                                      \
+        case DPCTL_UINT32_T:                                                   \
+        {                                                                      \
+            auto la = local_accessor<uint32_t, NDIM>(R, CGH);                  \
+            CGH.set_arg(IDX, la);                                              \
+            return true;                                                       \
+        }                                                                      \
+        case DPCTL_INT64_T:                                                    \
+        {                                                                      \
+            auto la = local_accessor<int64_t, NDIM>(R, CGH);                   \
+            CGH.set_arg(IDX, la);                                              \
+            return true;                                                       \
+        }                                                                      \
+        case DPCTL_UINT64_T:                                                   \
+        {                                                                      \
+            auto la = local_accessor<uint64_t, NDIM>(R, CGH);                  \
+            CGH.set_arg(IDX, la);                                              \
+            return true;                                                       \
+        }                                                                      \
+        case DPCTL_FLOAT32_T:                                                  \
+        {                                                                      \
+            auto la = local_accessor<float, NDIM>(R, CGH);                     \
+            CGH.set_arg(IDX, la);                                              \
+            return true;                                                       \
+        }                                                                      \
+        case DPCTL_FLOAT64_T:                                                  \
+        {                                                                      \
+            auto la = local_accessor<double, NDIM>(R, CGH);                    \
+            CGH.set_arg(IDX, la);                                              \
+            return true;                                                       \
+        }                                                                      \
+        default:                                                               \
+            error_handler("Kernel argument could not be created.", __FILE__,   \
+                          __func__, __LINE__, error_level::error);             \
+            return false;                                                      \
+        }                                                                      \
+    } while (0);
+
 namespace
 {
 static_assert(__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_VERSION_REQUIRED,
@@ -62,11 +132,39 @@ void set_dependent_events(handler &cgh,
     }
 }
 
+bool set_local_accessor_arg(handler &cgh,
+                            size_t idx,
+                            const MDLocalAccessor *mdstruct)
+{
+    switch (mdstruct->ndim) {
+    case 1:
+    {
+        auto r = range<1>(mdstruct->dim0);
+        SET_LOCAL_ACCESSOR_ARG(cgh, 1, mdstruct->dpctl_type_id, r, idx);
+    }
+    case 2:
+    {
+        auto r = range<2>(mdstruct->dim0, mdstruct->dim1);
+        SET_LOCAL_ACCESSOR_ARG(cgh, 2, mdstruct->dpctl_type_id, r, idx);
+    }
+    case 3:
+    {
+        auto r = range<3>(mdstruct->dim0, mdstruct->dim1, mdstruct->dim2);
+        SET_LOCAL_ACCESSOR_ARG(cgh, 3, mdstruct->dpctl_type_id, r, idx);
+    }
+    default:
+        return false;
+    }
+}
 /*!
  * @brief Set the kernel arg object
  *
- * @param    cgh            My Param doc
- * @param    Arg            My Param doc
+ * @param cgh   SYCL command group handler using which a kernel is going to
+ *              be submitted.
+ * @param idx   The position of the argument in the list of arguments passed
+ * to a kernel.
+ * @param Arg   A void* representing a kernel argument.
+ * @param Argty A typeid specifying the C++ type of the Arg parameter.
  */
 bool set_kernel_arg(handler &cgh,
                     size_t idx,
@@ -109,10 +207,11 @@ bool set_kernel_arg(handler &cgh,
     case DPCTL_VOID_PTR:
         cgh.set_arg(idx, Arg);
         break;
+    case DPCTL_LOCAL_ACCESSOR:
+        arg_set = set_local_accessor_arg(cgh, idx, (MDLocalAccessor *)Arg);
+        break;
     default:
         arg_set = false;
-        error_handler("Kernel argument could not be created.", __FILE__,
-                      __func__, __LINE__);
         break;
     }
     return arg_set;

libsyclinterface/tests/CMakeLists.txt

@@ -21,6 +21,8 @@ set(spirv-test-files
     multi_kernel.spv
     oneD_range_kernel_inttys_fp32.spv
     oneD_range_kernel_fp64.spv
+    local_accessor_kernel_inttys_fp32.spv
+    local_accessor_kernel_fp64.spv
 )
 
 foreach(tf ${spirv-test-files})
@@ -55,6 +57,7 @@ add_sycl_to_target(
     ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_platform_invalid_filters.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_queue_manager.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_queue_submit.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_queue_submit_local_accessor_arg.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_queue_interface.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_usm_interface.cpp
 )
@@ -86,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
@@ -106,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
@@ -129,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. */