Skip to content

Commit 5f08828

Browse files
author
Diptorup Deb
authored
Merge pull request #1558 from IntelPython/feature/local_accessor_kernel_arg
Allow setting local_accessors as kernel arguments in DPCTLQueue_Submit
2 parents d20e08d + 0da6390 commit 5f08828

18 files changed

+622
-33
lines changed

‎dpctl/_backend.pxd

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -67,7 +67,8 @@ cdef extern from "syclinterface/dpctl_sycl_enum_types.h":
6767
_UINT64_T 'DPCTL_UINT64_T',
6868
_FLOAT 'DPCTL_FLOAT32_T',
6969
_DOUBLE 'DPCTL_FLOAT64_T',
70-
_VOID_PTR 'DPCTL_VOID_PTR'
70+
_VOID_PTR 'DPCTL_VOID_PTR',
71+
_LOCAL_ACCESSOR 'DPCTL_LOCAL_ACCESSOR'
7172

7273
ctypedef enum _queue_property_type 'DPCTLQueuePropertyType':
7374
_DEFAULT_PROPERTY 'DPCTL_DEFAULT_PROPERTY'

‎dpctl/_sycl_queue.pyx

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -233,6 +233,15 @@ cdef class _kernel_arg_type:
233233
_arg_data_type._VOID_PTR
234234
)
235235

236+
@property
237+
def dpctl_local_accessor(self):
238+
cdef str p_name = "dpctl_local_accessor"
239+
return kernel_arg_type_attribute(
240+
self._name,
241+
p_name,
242+
_arg_data_type._LOCAL_ACCESSOR
243+
)
244+
236245

237246
kernel_arg_type = _kernel_arg_type()
238247

‎dpctl/enum_types.py

Lines changed: 1 addition & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -22,11 +22,7 @@
2222
"""
2323
from enum import Enum, auto
2424

25-
__all__ = [
26-
"device_type",
27-
"backend_type",
28-
"event_status_type",
29-
]
25+
__all__ = ["device_type", "backend_type", "event_status_type"]
3026

3127

3228
class device_type(Enum):

‎dpctl/tests/test_sycl_kernel_submit.py

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -274,3 +274,4 @@ def test_kernel_arg_type():
274274
_check_kernel_arg_type_instance(kernel_arg_type.dpctl_float32)
275275
_check_kernel_arg_type_instance(kernel_arg_type.dpctl_float64)
276276
_check_kernel_arg_type_instance(kernel_arg_type.dpctl_void_ptr)
277+
_check_kernel_arg_type_instance(kernel_arg_type.dpctl_local_accessor)

‎libsyclinterface/dbg_build.sh

Lines changed: 13 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -7,6 +7,11 @@ pushd build || exit 1
77
INSTALL_PREFIX=$(pwd)/../install
88
rm -rf ${INSTALL_PREFIX}
99

10+
# With DPC++ 2024.0 adn newer set these to ensure that
11+
# cmake can find llvm-cov and other utilities
12+
LLVM_TOOLS_HOME=${CMPLR_ROOT}/bin/compiler
13+
PATH=$PATH:${CMPLR_ROOT}/bin/compiler
14+
1015
cmake \
1116
-DCMAKE_BUILD_TYPE=Debug \
1217
-DCMAKE_C_COMPILER=icx \
@@ -16,13 +21,19 @@ cmake \
1621
-DCMAKE_PREFIX_PATH=${INSTALL_PREFIX} \
1722
-DDPCTL_ENABLE_L0_PROGRAM_CREATION=ON \
1823
-DDPCTL_BUILD_CAPI_TESTS=ON \
24+
-DDPCTL_GENERATE_COVERAGE=OFF \
1925
..
2026

21-
make V=1 -n -j 4 && make check && make install
27+
# build
28+
make V=1 -n -j 4
29+
# run ctest
30+
make check
31+
# install
32+
make install
2233

2334
# Turn on to generate coverage report html files reconfigure with
2435
# -DDPCTL_GENERATE_COVERAGE=ON and then
25-
# make lcov-genhtml
36+
# make llvm-cov-report
2637

2738
# For more verbose tests use:
2839
# cd tests

‎libsyclinterface/helper/include/dpctl_error_handlers.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -20,7 +20,7 @@
2020
///
2121
/// \file
2222
/// A functor to use for passing an error handler callback function to sycl
23-
/// context and queue contructors.
23+
/// context and queue constructors.
2424
//===----------------------------------------------------------------------===//
2525

2626
#pragma once

‎libsyclinterface/include/dpctl_sycl_enum_types.h

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -98,6 +98,7 @@ typedef enum
9898
DPCTL_FLOAT32_T,
9999
DPCTL_FLOAT64_T,
100100
DPCTL_VOID_PTR,
101+
DPCTL_LOCAL_ACCESSOR,
101102
DPCTL_UNSUPPORTED_KERNEL_ARG
102103
} DPCTLKernelArgType;
103104

‎libsyclinterface/include/dpctl_sycl_queue_interface.h

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -171,6 +171,18 @@ DPCTL_API
171171
__dpctl_give DPCTLSyclDeviceRef
172172
DPCTLQueue_GetDevice(__dpctl_keep const DPCTLSyclQueueRef QRef);
173173

174+
/*! @brief Structure to be used to specify dimensionality and type of
175+
* local_accessor kernel type argument.
176+
*/
177+
typedef struct MDLocalAccessorTy
178+
{
179+
size_t ndim;
180+
DPCTLKernelArgType dpctl_type_id;
181+
size_t dim0;
182+
size_t dim1;
183+
size_t dim2;
184+
} MDLocalAccessor;
185+
174186
/*!
175187
* @brief Submits the kernel to the specified queue with the provided range
176188
* argument.

‎libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -530,7 +530,7 @@ _GetKernel_ze_impl(const kernel_bundle<bundle_state::executable> &kb,
530530
else {
531531
error_handler("Kernel named " + std::string(kernel_name) +
532532
" could not be found.",
533-
__FILE__, __func__, __LINE__);
533+
__FILE__, __func__, __LINE__, error_level::error);
534534
return nullptr;
535535
}
536536
}
@@ -541,7 +541,7 @@ bool _HasKernel_ze_impl(const kernel_bundle<bundle_state::executable> &kb,
541541
auto zeKernelCreateFn = get_zeKernelCreate();
542542
if (zeKernelCreateFn == nullptr) {
543543
error_handler("Could not load zeKernelCreate function.", __FILE__,
544-
__func__, __LINE__);
544+
__func__, __LINE__, error_level::error);
545545
return false;
546546
}
547547

@@ -564,7 +564,7 @@ bool _HasKernel_ze_impl(const kernel_bundle<bundle_state::executable> &kb,
564564
if (ze_status != ZE_RESULT_ERROR_INVALID_KERNEL_NAME) {
565565
error_handler("zeKernelCreate failed: " +
566566
_GetErrorCode_ze_impl(ze_status),
567-
__FILE__, __func__, __LINE__);
567+
__FILE__, __func__, __LINE__, error_level::error);
568568
return false;
569569
}
570570
}

‎libsyclinterface/source/dpctl_sycl_queue_interface.cpp

Lines changed: 103 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -38,6 +38,76 @@
3838

3939
using namespace sycl;
4040

41+
#define SET_LOCAL_ACCESSOR_ARG(CGH, NDIM, ARGTY, R, IDX) \
42+
do { \
43+
switch ((ARGTY)) { \
44+
case DPCTL_INT8_T: \
45+
{ \
46+
auto la = local_accessor<int8_t, NDIM>(R, CGH); \
47+
CGH.set_arg(IDX, la); \
48+
return true; \
49+
} \
50+
case DPCTL_UINT8_T: \
51+
{ \
52+
auto la = local_accessor<uint8_t, NDIM>(R, CGH); \
53+
CGH.set_arg(IDX, la); \
54+
return true; \
55+
} \
56+
case DPCTL_INT16_T: \
57+
{ \
58+
auto la = local_accessor<int16_t, NDIM>(R, CGH); \
59+
CGH.set_arg(IDX, la); \
60+
return true; \
61+
} \
62+
case DPCTL_UINT16_T: \
63+
{ \
64+
auto la = local_accessor<uint16_t, NDIM>(R, CGH); \
65+
CGH.set_arg(IDX, la); \
66+
return true; \
67+
} \
68+
case DPCTL_INT32_T: \
69+
{ \
70+
auto la = local_accessor<int32_t, NDIM>(R, CGH); \
71+
CGH.set_arg(IDX, la); \
72+
return true; \
73+
} \
74+
case DPCTL_UINT32_T: \
75+
{ \
76+
auto la = local_accessor<uint32_t, NDIM>(R, CGH); \
77+
CGH.set_arg(IDX, la); \
78+
return true; \
79+
} \
80+
case DPCTL_INT64_T: \
81+
{ \
82+
auto la = local_accessor<int64_t, NDIM>(R, CGH); \
83+
CGH.set_arg(IDX, la); \
84+
return true; \
85+
} \
86+
case DPCTL_UINT64_T: \
87+
{ \
88+
auto la = local_accessor<uint64_t, NDIM>(R, CGH); \
89+
CGH.set_arg(IDX, la); \
90+
return true; \
91+
} \
92+
case DPCTL_FLOAT32_T: \
93+
{ \
94+
auto la = local_accessor<float, NDIM>(R, CGH); \
95+
CGH.set_arg(IDX, la); \
96+
return true; \
97+
} \
98+
case DPCTL_FLOAT64_T: \
99+
{ \
100+
auto la = local_accessor<double, NDIM>(R, CGH); \
101+
CGH.set_arg(IDX, la); \
102+
return true; \
103+
} \
104+
default: \
105+
error_handler("Kernel argument could not be created.", __FILE__, \
106+
__func__, __LINE__, error_level::error); \
107+
return false; \
108+
} \
109+
} while (0);
110+
41111
namespace
42112
{
43113
static_assert(__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_VERSION_REQUIRED,
@@ -62,11 +132,39 @@ void set_dependent_events(handler &cgh,
62132
}
63133
}
64134

135+
bool set_local_accessor_arg(handler &cgh,
136+
size_t idx,
137+
const MDLocalAccessor *mdstruct)
138+
{
139+
switch (mdstruct->ndim) {
140+
case 1:
141+
{
142+
auto r = range<1>(mdstruct->dim0);
143+
SET_LOCAL_ACCESSOR_ARG(cgh, 1, mdstruct->dpctl_type_id, r, idx);
144+
}
145+
case 2:
146+
{
147+
auto r = range<2>(mdstruct->dim0, mdstruct->dim1);
148+
SET_LOCAL_ACCESSOR_ARG(cgh, 2, mdstruct->dpctl_type_id, r, idx);
149+
}
150+
case 3:
151+
{
152+
auto r = range<3>(mdstruct->dim0, mdstruct->dim1, mdstruct->dim2);
153+
SET_LOCAL_ACCESSOR_ARG(cgh, 3, mdstruct->dpctl_type_id, r, idx);
154+
}
155+
default:
156+
return false;
157+
}
158+
}
65159
/*!
66160
* @brief Set the kernel arg object
67161
*
68-
* @param cgh My Param doc
69-
* @param Arg My Param doc
162+
* @param cgh SYCL command group handler using which a kernel is going to
163+
* be submitted.
164+
* @param idx The position of the argument in the list of arguments passed
165+
* to a kernel.
166+
* @param Arg A void* representing a kernel argument.
167+
* @param Argty A typeid specifying the C++ type of the Arg parameter.
70168
*/
71169
bool set_kernel_arg(handler &cgh,
72170
size_t idx,
@@ -109,10 +207,11 @@ bool set_kernel_arg(handler &cgh,
109207
case DPCTL_VOID_PTR:
110208
cgh.set_arg(idx, Arg);
111209
break;
210+
case DPCTL_LOCAL_ACCESSOR:
211+
arg_set = set_local_accessor_arg(cgh, idx, (MDLocalAccessor *)Arg);
212+
break;
112213
default:
113214
arg_set = false;
114-
error_handler("Kernel argument could not be created.", __FILE__,
115-
__func__, __LINE__);
116215
break;
117216
}
118217
return arg_set;

‎libsyclinterface/tests/CMakeLists.txt

Lines changed: 33 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -21,6 +21,8 @@ set(spirv-test-files
2121
multi_kernel.spv
2222
oneD_range_kernel_inttys_fp32.spv
2323
oneD_range_kernel_fp64.spv
24+
local_accessor_kernel_inttys_fp32.spv
25+
local_accessor_kernel_fp64.spv
2426
)
2527

2628
foreach(tf ${spirv-test-files})
@@ -55,6 +57,7 @@ add_sycl_to_target(
5557
${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_platform_invalid_filters.cpp
5658
${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_queue_manager.cpp
5759
${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_queue_submit.cpp
60+
${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_queue_submit_local_accessor_arg.cpp
5861
${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_queue_interface.cpp
5962
${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_usm_interface.cpp
6063
)
@@ -86,8 +89,35 @@ if(DPCTL_GENERATE_COVERAGE)
8689
${CMAKE_DL_LIBS}
8790
)
8891
set(object_arg "-object;")
89-
add_custom_target(llvm-cov
92+
add_custom_target(run-c-api-tests
9093
COMMAND ${CMAKE_COMMAND} -E env DPCTL_VERBOSITY=warning ${CMAKE_CURRENT_BINARY_DIR}/dpctl_c_api_tests
94+
WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}
95+
COMMAND_EXPAND_LISTS
96+
DEPENDS dpctl_c_api_tests
97+
)
98+
add_custom_target(llvm-cov-show
99+
COMMAND ${LLVMProfdata_EXE}
100+
merge
101+
-sparse default.profraw
102+
-o
103+
dpctl.profdata
104+
COMMAND ${LLVMCov_EXE}
105+
export
106+
-format=lcov
107+
-ignore-filename-regex=/tmp/icpx*
108+
-instr-profile=dpctl.profdata
109+
"${object_arg}$<JOIN:$<TARGET_OBJECTS:DPCTLSyclInterface>,;${object_arg}>"
110+
> dpctl.lcov
111+
COMMAND ${LLVMCov_EXE}
112+
show
113+
-instr-profile=dpctl.profdata
114+
"${object_arg}$<JOIN:$<TARGET_OBJECTS:DPCTLSyclInterface>,;${object_arg}>"
115+
WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}
116+
COMMAND_EXPAND_LISTS
117+
DEPENDS run-c-api-tests
118+
)
119+
120+
add_custom_target(llvm-cov-report
91121
COMMAND ${LLVMProfdata_EXE}
92122
merge
93123
-sparse default.profraw
@@ -106,11 +136,10 @@ if(DPCTL_GENERATE_COVERAGE)
106136
"${object_arg}$<JOIN:$<TARGET_OBJECTS:DPCTLSyclInterface>,;${object_arg}>"
107137
WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}
108138
COMMAND_EXPAND_LISTS
109-
DEPENDS dpctl_c_api_tests
139+
DEPENDS run-c-api-tests
110140
)
111141

112142
add_custom_target(lcov-genhtml
113-
COMMAND ${CMAKE_COMMAND} -E env DPCTL_VERBOSITY=warning ${CMAKE_CURRENT_BINARY_DIR}/dpctl_c_api_tests
114143
COMMAND ${LLVMProfdata_EXE}
115144
merge
116145
-sparse default.profraw
@@ -129,7 +158,7 @@ if(DPCTL_GENERATE_COVERAGE)
129158
${COVERAGE_OUTPUT_DIR}/dpctl-c-api-coverage
130159
WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}
131160
COMMAND_EXPAND_LISTS
132-
DEPENDS dpctl_c_api_tests
161+
DEPENDS run-c-api-tests
133162
)
134163
else()
135164
target_link_libraries(dpctl_c_api_tests

‎libsyclinterface/tests/local_accessor_kernel_fp64.spv

5.56 KB
Binary file not shown.

‎libsyclinterface/tests/local_accessor_kernel_inttys_fp32.spv

12.6 KB
Binary file not shown.

‎libsyclinterface/tests/test_sycl_device_aspects.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -97,7 +97,7 @@ auto build_gtest_values(const std::array<std::pair<T1, T2>, N> &params)
9797
auto build_params()
9898
{
9999
constexpr auto param_1 = get_param_list<const char *>(
100-
"opencl:gpu", "opencl:cpu", "level_zero:gpu", "host");
100+
"opencl:gpu", "opencl:cpu", "level_zero:gpu");
101101

102102
constexpr auto param_2 =
103103
get_param_list<std::pair<const char *, sycl::aspect>>(

‎libsyclinterface/tests/test_sycl_queue_interface.cpp

Lines changed: 28 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -90,6 +90,34 @@ struct TestDPCTLQueueMemberFunctions
9090

9191
} /* End of anonymous namespace */
9292

93+
TEST(TestDPCTLSyclQueueInterface, CheckCreate)
94+
{
95+
/* We are testing that we do not crash even when input is NULL. */
96+
DPCTLSyclQueueRef QRef = nullptr;
97+
98+
EXPECT_NO_FATAL_FAILURE(
99+
QRef = DPCTLQueue_Create(nullptr, nullptr, nullptr, 0));
100+
ASSERT_TRUE(QRef == nullptr);
101+
}
102+
103+
TEST(TestDPCTLSyclQueueInterface, CheckCreate2)
104+
{
105+
/* We are testing that we do not crash even when input is NULL. */
106+
DPCTLSyclQueueRef QRef = nullptr;
107+
DPCTLSyclDeviceSelectorRef DSRef = nullptr;
108+
DPCTLSyclDeviceRef DRef = nullptr;
109+
110+
EXPECT_NO_FATAL_FAILURE(DSRef = DPCTLDefaultSelector_Create());
111+
EXPECT_NO_FATAL_FAILURE(DRef = DPCTLDevice_CreateFromSelector(DSRef));
112+
EXPECT_NO_FATAL_FAILURE(DPCTLDeviceSelector_Delete(DSRef));
113+
114+
EXPECT_NO_FATAL_FAILURE(QRef =
115+
DPCTLQueue_Create(nullptr, DRef, nullptr, 0));
116+
ASSERT_TRUE(QRef == nullptr);
117+
118+
EXPECT_NO_FATAL_FAILURE(DPCTLDevice_Delete(DRef));
119+
}
120+
93121
TEST(TestDPCTLSyclQueueInterface, CheckCreateForDevice)
94122
{
95123
/* We are testing that we do not crash even when input is NULL. */

0 commit comments

Comments
 (0)