Skip to content

Commit 4023e87

Browse files
author
Diptorup Deb
committed
Support local_accessor kernel arguments.
 - Adds support in libsyclinterface:: dpctl_sycl_queue_interface for sycl::local_accessor as kernel arguments. - Refactoring to get rid of compiler warnings.
1 parent 545dff2 commit 4023e87

File tree

5 files changed

+119
-8
lines changed

5 files changed

+119
-8
lines changed

‎dpctl/enum_types.py

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -26,6 +26,7 @@
2626
"device_type",
2727
"backend_type",
2828
"event_status_type",
29+
"kernel_arg_type",
2930
]
3031

3132

@@ -132,3 +133,4 @@ class kernel_arg_type(Enum):
132133
dpctl_float32 = auto()
133134
dpctl_float64 = auto()
134135
dpctl_void_ptr = auto()
136+
dpctl_local_accessor = auto()

‎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/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: 112 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,
@@ -51,6 +121,15 @@ typedef struct complex
51121
uint64_t imag;
52122
} complexNumber;
53123

124+
typedef struct MDLocalAccessorTy
125+
{
126+
size_t ndim;
127+
DPCTLKernelArgType dpctl_type_id;
128+
size_t dim0;
129+
size_t dim1;
130+
size_t dim2;
131+
} MDLocalAccessor;
132+
54133
void set_dependent_events(handler &cgh,
55134
__dpctl_keep const DPCTLSyclEventRef *DepEvents,
56135
size_t NDepEvents)
@@ -62,11 +141,39 @@ void set_dependent_events(handler &cgh,
62141
}
63142
}
64143

144+
bool set_local_accessor_arg(handler &cgh,
145+
size_t idx,
146+
const MDLocalAccessor *mdstruct)
147+
{
148+
switch (mdstruct->ndim) {
149+
case 1:
150+
{
151+
auto r = range<1>(mdstruct->dim0);
152+
SET_LOCAL_ACCESSOR_ARG(cgh, 1, mdstruct->dpctl_type_id, r, idx);
153+
}
154+
case 2:
155+
{
156+
auto r = range<2>(mdstruct->dim0, mdstruct->dim1);
157+
SET_LOCAL_ACCESSOR_ARG(cgh, 2, mdstruct->dpctl_type_id, r, idx);
158+
}
159+
case 3:
160+
{
161+
auto r = range<3>(mdstruct->dim0, mdstruct->dim1, mdstruct->dim2);
162+
SET_LOCAL_ACCESSOR_ARG(cgh, 3, mdstruct->dpctl_type_id, r, idx);
163+
}
164+
default:
165+
return false;
166+
}
167+
}
65168
/*!
66169
* @brief Set the kernel arg object
67170
*
68-
* @param cgh My Param doc
69-
* @param Arg My Param doc
171+
* @param cgh SYCL command group handler using which a kernel is going to
172+
* be submitted.
173+
* @param idx The position of the argument in the list of arguments passed
174+
* to a kernel.
175+
* @param Arg A void* representing a kernel argument.
176+
* @param Argty A typeid specifying the C++ type of the Arg parameter.
70177
*/
71178
bool set_kernel_arg(handler &cgh,
72179
size_t idx,
@@ -109,10 +216,11 @@ bool set_kernel_arg(handler &cgh,
109216
case DPCTL_VOID_PTR:
110217
cgh.set_arg(idx, Arg);
111218
break;
219+
case DPCTL_LOCAL_ACCESSOR:
220+
arg_set = set_local_accessor_arg(cgh, idx, (MDLocalAccessor *)Arg);
221+
break;
112222
default:
113223
arg_set = false;
114-
error_handler("Kernel argument could not be created.", __FILE__,
115-
__func__, __LINE__);
116224
break;
117225
}
118226
return arg_set;

0 commit comments

Comments
 (0)