-
Notifications
You must be signed in to change notification settings - Fork 34
Expand file tree
/
Copy path_program.pyx
More file actions
339 lines (280 loc) · 11.4 KB
/
_program.pyx
File metadata and controls
339 lines (280 loc) · 11.4 KB
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
# Data Parallel Control (dpctl)
#
# Copyright 2020-2025 Intel Corporation
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
# distutils: language = c++
# cython: language_level=3
# cython: linetrace=True
"""Implements a Python interface for SYCL's program and kernel runtime classes.
The module also provides functions to create a SYCL program from either
a OpenCL source string or a SPIR-V binary file.
"""
from libc.stdint cimport uint32_t
from dpctl._backend cimport ( # noqa: E211, E402;
DPCTLKernel_Copy,
DPCTLKernel_Delete,
DPCTLKernel_GetCompileNumSubGroups,
DPCTLKernel_GetCompileSubGroupSize,
DPCTLKernel_GetMaxNumSubGroups,
DPCTLKernel_GetMaxSubGroupSize,
DPCTLKernel_GetNumArgs,
DPCTLKernel_GetPreferredWorkGroupSizeMultiple,
DPCTLKernel_GetPrivateMemSize,
DPCTLKernel_GetWorkGroupSize,
DPCTLKernelBundle_Copy,
DPCTLKernelBundle_CreateFromOCLSource,
DPCTLKernelBundle_CreateFromSpirv,
DPCTLKernelBundle_Delete,
DPCTLKernelBundle_GetKernel,
DPCTLKernelBundle_HasKernel,
DPCTLSyclContextRef,
DPCTLSyclDeviceRef,
DPCTLSyclKernelBundleRef,
DPCTLSyclKernelRef,
)
__all__ = [
"create_program_from_source",
"create_program_from_spirv",
"SyclKernel",
"SyclProgram",
"SyclProgramCompilationError",
]
cdef class SyclProgramCompilationError(Exception):
"""This exception is raised when a ``sycl::kernel_bundle`` could not be
built from either a SPIR-V binary file or a string source.
"""
pass
cdef class SyclKernel:
"""
"""
@staticmethod
cdef SyclKernel _create(DPCTLSyclKernelRef kref, str name):
cdef SyclKernel ret = SyclKernel.__new__(SyclKernel)
ret._kernel_ref = kref
ret._function_name = name
return ret
def __dealloc__(self):
DPCTLKernel_Delete(self._kernel_ref)
def get_function_name(self):
""" Returns the name of the ``sycl::kernel`` function.
"""
return self._function_name
def get_num_args(self):
""" Returns the number of arguments for this kernel function.
"""
return DPCTLKernel_GetNumArgs(self._kernel_ref)
cdef DPCTLSyclKernelRef get_kernel_ref(self):
""" Returns the ``DPCTLSyclKernelRef`` pointer for this SyclKernel.
"""
return self._kernel_ref
def addressof_ref(self):
""" Returns the address of the C API ``DPCTLSyclKernelRef`` pointer
as a ``size_t``.
Returns:
The address of the ``DPCTLSyclKernelRef`` pointer used to create
this :class:`dpctl.SyclKernel` object cast to a ``size_t``.
"""
return int(<size_t>self._kernel_ref)
@property
def num_args(self):
""" Property equivalent to method call `SyclKernel.get_num_args()`
"""
return self.get_num_args()
@property
def work_group_size(self):
""" Returns the maximum number of work-items in a work-group that can
be used to execute the kernel on device it was built for.
"""
cdef size_t v = DPCTLKernel_GetWorkGroupSize(self._kernel_ref)
return v
@property
def preferred_work_group_size_multiple(self):
""" Returns a value, of which work-group size is preferred to be
a multiple, for executing the kernel on the device it was built for.
"""
cdef size_t v = DPCTLKernel_GetPreferredWorkGroupSizeMultiple(
self._kernel_ref
)
return v
@property
def private_mem_size(self):
""" Returns the minimum amount of private memory, in bytes, used by each
work-item in the kernel.
"""
cdef size_t v = DPCTLKernel_GetPrivateMemSize(self._kernel_ref)
return v
@property
def max_num_sub_groups(self):
""" Returns the maximum number of sub-groups for this kernel.
"""
cdef uint32_t n = DPCTLKernel_GetMaxNumSubGroups(self._kernel_ref)
return n
@property
def max_sub_group_size(self):
""" Returns the maximum sub-groups size for this kernel.
"""
cdef uint32_t sz = DPCTLKernel_GetMaxSubGroupSize(self._kernel_ref)
return sz
@property
def compile_num_sub_groups(self):
""" Returns the number of sub-groups specified by this kernel,
or 0 (if not specified).
"""
cdef size_t n = DPCTLKernel_GetCompileNumSubGroups(self._kernel_ref)
return n
@property
def compile_sub_group_size(self):
""" Returns the required sub-group size specified by this kernel,
or 0 (if not specified).
"""
cdef size_t n = DPCTLKernel_GetCompileSubGroupSize(self._kernel_ref)
return n
cdef api DPCTLSyclKernelRef SyclKernel_GetKernelRef(SyclKernel ker):
""" C-API function to access opaque kernel reference from
Python object of type :class:`dpctl.program.SyclKernel`.
"""
return ker.get_kernel_ref()
cdef api SyclKernel SyclKernel_Make(DPCTLSyclKernelRef KRef, const char *name):
"""
C-API function to create :class:`dpctl.program.SyclKernel`
instance from opaque sycl kernel reference.
"""
cdef DPCTLSyclKernelRef copied_KRef = DPCTLKernel_Copy(KRef)
if (name is NULL):
return SyclKernel._create(copied_KRef, "default_name")
else:
return SyclKernel._create(copied_KRef, name.decode("utf-8"))
cdef class SyclProgram:
""" Wraps a ``sycl::kernel_bundle<sycl::bundle_state::executable>`` object
created using SYCL interoperability layer with underlying backends. Only the
OpenCL and Level-Zero backends are currently supported.
SyclProgram exposes the C API from ``dpctl_sycl_kernel_bundle_interface.h``.
A SyclProgram can be created from either a source string or a SPIR-V
binary file.
"""
@staticmethod
cdef SyclProgram _create(DPCTLSyclKernelBundleRef KBRef):
cdef SyclProgram ret = SyclProgram.__new__(SyclProgram)
ret._program_ref = KBRef
return ret
def __dealloc__(self):
DPCTLKernelBundle_Delete(self._program_ref)
cdef DPCTLSyclKernelBundleRef get_program_ref(self):
return self._program_ref
cpdef SyclKernel get_sycl_kernel(self, str kernel_name):
name = kernel_name.encode("utf8")
return SyclKernel._create(
DPCTLKernelBundle_GetKernel(self._program_ref, name),
kernel_name
)
def has_sycl_kernel(self, str kernel_name):
name = kernel_name.encode("utf8")
return DPCTLKernelBundle_HasKernel(self._program_ref, name)
def addressof_ref(self):
"""Returns the address of the C API DPCTLSyclKernelBundleRef pointer
as a long.
Returns:
The address of the ``DPCTLSyclKernelBundleRef`` pointer used to
create this :class:`dpctl.SyclProgram` object cast to a ``size_t``.
"""
return int(<size_t>self._program_ref)
cpdef create_program_from_source(SyclQueue q, str src, str copts=""):
"""
Creates a Sycl interoperability program from an OpenCL source string.
We use the ``DPCTLKernelBundle_CreateFromOCLSource()`` C API function
to create a ``sycl::kernel_bundle<sycl::bundle_state::executable>``
from an OpenCL source program that can contain multiple kernels.
Note: This function is currently only supported for the OpenCL backend.
Parameters:
q (:class:`dpctl.SyclQueue`)
The :class:`dpctl.SyclQueue` for which the
:class:`.SyclProgram` is going to be built.
src (str)
Source string for an OpenCL program.
copts (str, optional)
Optional compilation flags that will be used
when compiling the program. Default: ``""``.
Returns:
program (:class:`.SyclProgram`)
A :class:`.SyclProgram` object wrapping the
``sycl::kernel_bundle<sycl::bundle_state::executable>``
returned by the C API.
Raises:
SyclProgramCompilationError
If a SYCL kernel bundle could not be created.
"""
cdef DPCTLSyclKernelBundleRef KBref
cdef bytes bSrc = src.encode("utf8")
cdef bytes bCOpts = copts.encode("utf8")
cdef const char *Src = <const char*>bSrc
cdef const char *COpts = <const char*>bCOpts
cdef DPCTLSyclContextRef CRef = q.get_sycl_context().get_context_ref()
cdef DPCTLSyclDeviceRef DRef = q.get_sycl_device().get_device_ref()
KBref = DPCTLKernelBundle_CreateFromOCLSource(CRef, DRef, Src, COpts)
if KBref is NULL:
raise SyclProgramCompilationError()
return SyclProgram._create(KBref)
cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL,
str copts=""):
"""
Creates a Sycl interoperability program from an SPIR-V binary.
We use the :c:func:`DPCTLKernelBundle_CreateFromOCLSpirv` C API function
to create a ``sycl::kernel_bundle<sycl::bundle_state::executable>``
object from an compiled SPIR-V binary file.
Parameters:
q (:class:`dpctl.SyclQueue`)
The :class:`dpctl.SyclQueue` for which the
:class:`.SyclProgram` is going to be built.
IL (bytes)
SPIR-V binary IL file for an OpenCL program.
copts (str, optional)
Optional compilation flags that will be used
when compiling the program. Default: ``""``.
Returns:
program (:class:`.SyclProgram`)
A :class:`.SyclProgram` object wrapping the
``sycl::kernel_bundle<sycl::bundle_state::executable>``
returned by the C API.
Raises:
SyclProgramCompilationError
If a SYCL kernel bundle could not be created.
"""
cdef DPCTLSyclKernelBundleRef KBref
cdef const unsigned char *dIL = &IL[0]
cdef DPCTLSyclContextRef CRef = q.get_sycl_context().get_context_ref()
cdef DPCTLSyclDeviceRef DRef = q.get_sycl_device().get_device_ref()
cdef size_t length = IL.shape[0]
cdef bytes bCOpts = copts.encode("utf8")
cdef const char *COpts = <const char*>bCOpts
KBref = DPCTLKernelBundle_CreateFromSpirv(
CRef, DRef, <const void*>dIL, length, COpts
)
if KBref is NULL:
raise SyclProgramCompilationError()
return SyclProgram._create(KBref)
cdef api DPCTLSyclKernelBundleRef SyclProgram_GetKernelBundleRef(
SyclProgram pro
):
""" C-API function to access opaque kernel bundle reference from
Python object of type :class:`dpctl.program.SyclKernel`.
"""
return pro.get_program_ref()
cdef api SyclProgram SyclProgram_Make(DPCTLSyclKernelBundleRef KBRef):
"""
C-API function to create :class:`dpctl.program.SyclProgram`
instance from opaque ``sycl::kernel_bundle<sycl::bundle_state::executable>``
reference.
"""
cdef DPCTLSyclKernelBundleRef copied_KBRef = DPCTLKernelBundle_Copy(KBRef)
return SyclProgram._create(copied_KBRef)