From 3c294f3288a862bb6b805a75d8d86e746d11f4ad Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Thu, 10 Apr 2025 10:39:23 +0100 Subject: [PATCH 1/5] Support compilation from SYCL source code --- dpctl/_backend.pxd | 40 +++++ dpctl/_sycl_device.pxd | 1 + dpctl/_sycl_device.pyx | 32 ++++ dpctl/program/__init__.py | 2 + dpctl/program/_program.pxd | 6 +- dpctl/program/_program.pyx | 130 ++++++++++++++- dpctl/tests/test_sycl_program.py | 105 +++++++++++- .../dpctl_sycl_device_interface.h | 36 ++++ .../dpctl_sycl_kernel_bundle_interface.h | 157 ++++++++++++++++++ .../source/dpctl_sycl_device_interface.cpp | 25 +++ .../dpctl_sycl_kernel_bundle_interface.cpp | 140 ++++++++++++++++ .../test_sycl_kernel_bundle_interface.cpp | 119 +++++++++++++ 12 files changed, 785 insertions(+), 8 deletions(-) diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index ca9e9ccb9f..677e3699ad 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -278,6 +278,9 @@ cdef extern from "syclinterface/dpctl_sycl_device_interface.h": cdef DPCTLDeviceVectorRef DPCTLDevice_GetComponentDevices( const DPCTLSyclDeviceRef DRef ) + cdef bool DPCTLDevice_CanCompileSPIRV(const DPCTLSyclDeviceRef DRef) + cdef bool DPCTLDevice_CanCompileOpenCL(const DPCTLSyclDeviceRef DRef) + cdef bool DPCTLDevice_CanCompileSYCL(const DPCTLSyclDeviceRef DRef) cdef extern from "syclinterface/dpctl_sycl_device_manager.h": @@ -441,6 +444,43 @@ cdef extern from "syclinterface/dpctl_sycl_kernel_bundle_interface.h": cdef DPCTLSyclKernelBundleRef DPCTLKernelBundle_Copy( const DPCTLSyclKernelBundleRef KBRef) + cdef struct DPCTLBuildOptionList + cdef struct DPCTLKernelNameList + cdef struct DPCTLVirtualHeaderList + ctypedef DPCTLBuildOptionList* DPCTLBuildOptionListRef + ctypedef DPCTLKernelNameList* DPCTLKernelNameListRef + ctypedef DPCTLVirtualHeaderList* DPCTLVirtualHeaderListRef + + cdef DPCTLBuildOptionListRef DPCTLBuildOptionList_Create() + cdef void DPCTLBuildOptionList_Delete(DPCTLBuildOptionListRef Ref) + cdef void DPCTLBuildOptionList_Append(DPCTLBuildOptionListRef Ref, + const char *Option) + + cdef DPCTLKernelNameListRef DPCTLKernelNameList_Create() + cdef void DPCTLKernelNameList_Delete(DPCTLKernelNameListRef Ref) + cdef void DPCTLKernelNameList_Append(DPCTLKernelNameListRef Ref, + const char *Option) + + cdef DPCTLVirtualHeaderListRef DPCTLVirtualHeaderList_Create() + cdef void DPCTLVirtualHeaderList_Delete(DPCTLVirtualHeaderListRef Ref) + cdef void DPCTLVirtualHeaderList_Append(DPCTLVirtualHeaderListRef Ref, + const char *Name, + const char *Content) + + cdef DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource( + const DPCTLSyclContextRef Ctx, + const DPCTLSyclDeviceRef Dev, + const char *Source, + DPCTLVirtualHeaderListRef Headers, + DPCTLKernelNameListRef Names, + DPCTLBuildOptionListRef BuildOptions) + + cdef DPCTLSyclKernelRef DPCTLKernelBundle_GetSyclKernel(DPCTLSyclKernelBundleRef KBRef, + const char *KernelName) + + cdef bool DPCTLKernelBundle_HasSyclKernel(DPCTLSyclKernelBundleRef KBRef, + const char *KernelName); + cdef extern from "syclinterface/dpctl_sycl_queue_interface.h": ctypedef struct _md_local_accessor "MDLocalAccessor": diff --git a/dpctl/_sycl_device.pxd b/dpctl/_sycl_device.pxd index 190d981cd0..d9378f0897 100644 --- a/dpctl/_sycl_device.pxd +++ b/dpctl/_sycl_device.pxd @@ -61,3 +61,4 @@ cdef public api class SyclDevice(_SyclDevice) [ cdef int get_overall_ordinal(self) cdef int get_backend_ordinal(self) cdef int get_backend_and_device_type_ordinal(self) + cpdef bint can_compile(self, str language) diff --git a/dpctl/_sycl_device.pyx b/dpctl/_sycl_device.pyx index 5b43ffed1a..bda700cf11 100644 --- a/dpctl/_sycl_device.pyx +++ b/dpctl/_sycl_device.pyx @@ -25,6 +25,9 @@ from ._backend cimport ( # noqa: E211 DPCTLCString_Delete, DPCTLDefaultSelector_Create, DPCTLDevice_AreEq, + DPCTLDevice_CanCompileOpenCL, + DPCTLDevice_CanCompileSPIRV, + DPCTLDevice_CanCompileSYCL, DPCTLDevice_Copy, DPCTLDevice_CreateFromSelector, DPCTLDevice_CreateSubDevicesByAffinity, @@ -2160,6 +2163,35 @@ cdef class SyclDevice(_SyclDevice): raise ValueError("device could not be found") return dev_id + cpdef bint can_compile(self, str language): + """ + Check whether it is possible to create an executable kernel_bundle + for this device from the given source language. + + Parameters: + language + Input language. Possible values are "spirv" for SPIR-V binary + files, "opencl" for OpenCL C device code and "sycl" for SYCL + device code. + + Returns: + bool: + True if compilation is supported, False otherwise. + + Raises: + ValueError: + If an unknown source language is used. + """ + if language == "spirv" or language == "spv": + return DPCTLDevice_CanCompileSYCL(self._device_ref) + if language == "opencl" or language == "ocl": + return DPCTLDevice_CanCompileOpenCL(self._device_ref) + if language == "sycl": + return DPCTLDevice_CanCompileSYCL(self._device_ref) + + raise ValueError(f"Unknown source language {language}") + + cdef api DPCTLSyclDeviceRef SyclDevice_GetDeviceRef(SyclDevice dev): """ diff --git a/dpctl/program/__init__.py b/dpctl/program/__init__.py index a96d33f04a..e209b68b40 100644 --- a/dpctl/program/__init__.py +++ b/dpctl/program/__init__.py @@ -26,11 +26,13 @@ SyclProgramCompilationError, create_program_from_source, create_program_from_spirv, + create_program_from_sycl_source, ) __all__ = [ "create_program_from_source", "create_program_from_spirv", + "create_program_from_sycl_source", "SyclKernel", "SyclProgram", "SyclProgramCompilationError", diff --git a/dpctl/program/_program.pxd b/dpctl/program/_program.pxd index dc4208a29b..37d7fbaac5 100644 --- a/dpctl/program/_program.pxd +++ b/dpctl/program/_program.pxd @@ -49,9 +49,10 @@ cdef api class SyclProgram [object PySyclProgramObject, type PySyclProgramType]: binary file. """ cdef DPCTLSyclKernelBundleRef _program_ref + cdef bint _is_sycl_source @staticmethod - cdef SyclProgram _create (DPCTLSyclKernelBundleRef pref) + cdef SyclProgram _create (DPCTLSyclKernelBundleRef pref, bint _is_sycl_source) cdef DPCTLSyclKernelBundleRef get_program_ref (self) cpdef SyclKernel get_sycl_kernel(self, str kernel_name) @@ -59,3 +60,6 @@ cdef api class SyclProgram [object PySyclProgramObject, type PySyclProgramType]: cpdef create_program_from_source (SyclQueue q, unicode source, unicode copts=*) cpdef create_program_from_spirv (SyclQueue q, const unsigned char[:] IL, unicode copts=*) +cpdef create_program_from_sycl_source(SyclQueue q, unicode source, + list headers=*, list registered_names=*, + list copts=*) diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index 3859314505..f175cf23a7 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -28,6 +28,10 @@ a OpenCL source string or a SPIR-V binary file. from libc.stdint cimport uint32_t from dpctl._backend cimport ( # noqa: E211, E402; + DPCTLBuildOptionList_Append, + DPCTLBuildOptionList_Create, + DPCTLBuildOptionList_Delete, + DPCTLBuildOptionListRef, DPCTLKernel_Copy, DPCTLKernel_Delete, DPCTLKernel_GetCompileNumSubGroups, @@ -41,13 +45,24 @@ from dpctl._backend cimport ( # noqa: E211, E402; DPCTLKernelBundle_Copy, DPCTLKernelBundle_CreateFromOCLSource, DPCTLKernelBundle_CreateFromSpirv, + DPCTLKernelBundle_CreateFromSYCLSource, DPCTLKernelBundle_Delete, DPCTLKernelBundle_GetKernel, + DPCTLKernelBundle_GetSyclKernel, DPCTLKernelBundle_HasKernel, + DPCTLKernelBundle_HasSyclKernel, + DPCTLKernelNameList_Append, + DPCTLKernelNameList_Create, + DPCTLKernelNameList_Delete, + DPCTLKernelNameListRef, DPCTLSyclContextRef, DPCTLSyclDeviceRef, DPCTLSyclKernelBundleRef, DPCTLSyclKernelRef, + DPCTLVirtualHeaderList_Append, + DPCTLVirtualHeaderList_Create, + DPCTLVirtualHeaderList_Delete, + DPCTLVirtualHeaderListRef, ) __all__ = [ @@ -196,9 +211,10 @@ cdef class SyclProgram: """ @staticmethod - cdef SyclProgram _create(DPCTLSyclKernelBundleRef KBRef): + cdef SyclProgram _create(DPCTLSyclKernelBundleRef KBRef, bint is_sycl_source): cdef SyclProgram ret = SyclProgram.__new__(SyclProgram) ret._program_ref = KBRef + ret._is_sycl_source = is_sycl_source return ret def __dealloc__(self): @@ -209,6 +225,10 @@ cdef class SyclProgram: cpdef SyclKernel get_sycl_kernel(self, str kernel_name): name = kernel_name.encode("utf8") + if self._is_sycl_source: + return SyclKernel._create( + DPCTLKernelBundle_GetSyclKernel(self._program_ref, name), + kernel_name) return SyclKernel._create( DPCTLKernelBundle_GetKernel(self._program_ref, name), kernel_name @@ -216,6 +236,8 @@ cdef class SyclProgram: def has_sycl_kernel(self, str kernel_name): name = kernel_name.encode("utf8") + if self._is_sycl_source: + return DPCTLKernelBundle_HasSyclKernel(self._program_ref, name) return DPCTLKernelBundle_HasKernel(self._program_ref, name) def addressof_ref(self): @@ -271,7 +293,7 @@ cpdef create_program_from_source(SyclQueue q, str src, str copts=""): if KBref is NULL: raise SyclProgramCompilationError() - return SyclProgram._create(KBref) + return SyclProgram._create(KBref, False) cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL, @@ -317,7 +339,107 @@ cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL, if KBref is NULL: raise SyclProgramCompilationError() - return SyclProgram._create(KBref) + return SyclProgram._create(KBref, False) + + +cpdef create_program_from_sycl_source(SyclQueue q, unicode source, list headers=[], list registered_names=[], list copts=[]): + """ + Creates an executable SYCL kernel_bundle from SYCL source code. + + This uses the DPC++ ``kernel_compiler`` extension to create a + ``sycl::kernel_bundle`` object from + SYCL source code. + + Parameters: + q (:class:`dpctl.SyclQueue`) + The :class:`dpctl.SyclQueue` for which the + :class:`.SyclProgram` is going to be built. + source (unicode) + SYCL source code string. + headers (list) + Optional list of virtual headers, where each entry in the list + needs to be a tuple of header name and header content. See the + documentation of the ``include_files`` property in the DPC++ + ``kernel_compiler`` extension for more information. + Default: [] + registered_names (list, optional) + Optional list of kernel names to register. See the + documentation of the ``registered_names`` property in the DPC++ + ``kernel_compiler`` extension for more information. + Default: [] + copts (list) + Optional list of compilation flags that will be used + when compiling the program. Default: ``""``. + + Returns: + program (:class:`.SyclProgram`) + A :class:`.SyclProgram` object wrapping the + ``sycl::kernel_bundle`` + returned by the C API. + + Raises: + SyclProgramCompilationError + If a SYCL kernel bundle could not be created. + """ + cdef DPCTLSyclKernelBundleRef KBref + cdef DPCTLSyclContextRef CRef = q.get_sycl_context().get_context_ref() + cdef DPCTLSyclDeviceRef DRef = q.get_sycl_device().get_device_ref() + cdef bytes bSrc = source.encode('utf8') + cdef const char *Src = bSrc + cdef DPCTLBuildOptionListRef BuildOpts = DPCTLBuildOptionList_Create() + cdef bytes bOpt + cdef const char* sOpt + cdef bytes bName + cdef const char* sName + cdef bytes bContent + cdef const char* sContent + for opt in copts: + if not isinstance(opt, unicode): + DPCTLBuildOptionList_Delete(BuildOpts) + raise SyclProgramCompilationError() + bOpt = opt.encode('utf8') + sOpt = bOpt + DPCTLBuildOptionList_Append(BuildOpts, sOpt) + + cdef DPCTLKernelNameListRef KernelNames = DPCTLKernelNameList_Create() + for name in registered_names: + if not isinstance(name, unicode): + DPCTLBuildOptionList_Delete(BuildOpts) + DPCTLKernelNameList_Delete(KernelNames) + raise SyclProgramCompilationError() + bName = name.encode('utf8') + sName = bName + DPCTLKernelNameList_Append(KernelNames, sName) + + + cdef DPCTLVirtualHeaderListRef VirtualHeaders = DPCTLVirtualHeaderList_Create() + for name, content in headers: + if not isinstance(name, unicode) or not isinstance(content, unicode): + DPCTLBuildOptionList_Delete(BuildOpts) + DPCTLKernelNameList_Delete(KernelNames) + DPCTLVirtualHeaderList_Delete(VirtualHeaders) + raise SyclProgramCompilationError() + bName = name.encode('utf8') + sName = bName + bContent = content.encode('utf8') + sContent = bContent + DPCTLVirtualHeaderList_Append(VirtualHeaders, sName, sContent) + + KBref = DPCTLKernelBundle_CreateFromSYCLSource(CRef, DRef, Src, + VirtualHeaders, KernelNames, + BuildOpts) + + if KBref is NULL: + DPCTLBuildOptionList_Delete(BuildOpts) + DPCTLKernelNameList_Delete(KernelNames) + DPCTLVirtualHeaderList_Delete(VirtualHeaders) + raise SyclProgramCompilationError() + + DPCTLBuildOptionList_Delete(BuildOpts) + DPCTLKernelNameList_Delete(KernelNames) + DPCTLVirtualHeaderList_Delete(VirtualHeaders) + + return SyclProgram._create(KBref, True) cdef api DPCTLSyclKernelBundleRef SyclProgram_GetKernelBundleRef( @@ -336,4 +458,4 @@ cdef api SyclProgram SyclProgram_Make(DPCTLSyclKernelBundleRef KBRef): reference. """ cdef DPCTLSyclKernelBundleRef copied_KBRef = DPCTLKernelBundle_Copy(KBRef) - return SyclProgram._create(copied_KBRef) + return SyclProgram._create(copied_KBRef, False) diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index a791a59daa..8cf8f06215 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -81,8 +81,7 @@ def _check_cpython_api_SyclProgram_Make(sycl_prog): make_prog_fn = callable_maker(make_prog_fn_ptr) p2 = make_prog_fn(sycl_prog.addressof_ref()) - assert p2.has_sycl_kernel("add") - assert p2.has_sycl_kernel("axpy") + return p2 def _check_cpython_api_SyclKernel_GetKernelRef(krn): @@ -187,7 +186,9 @@ def _check_multi_kernel_program(prog): assert type(cmsgsz) is int _check_cpython_api_SyclProgram_GetKernelBundleRef(prog) - _check_cpython_api_SyclProgram_Make(prog) + p2 = _check_cpython_api_SyclProgram_Make(prog) + assert p2.has_sycl_kernel("add") + assert p2.has_sycl_kernel("axpy") def test_create_program_from_source_ocl(): @@ -263,3 +264,101 @@ def test_create_program_from_invalid_src_ocl(): }" with pytest.raises(dpctl_prog.SyclProgramCompilationError): dpctl_prog.create_program_from_source(q, invalid_oclSrc) + + +def test_create_program_from_sycl_source(): + try: + q = dpctl.SyclQueue("level_zero") + except dpctl.SyclQueueCreationError: + pytest.skip("No Level-zero queue is available") + + if not q.get_sycl_device().can_compile("sycl"): + pytest.skip("SYCL source compilation not supported") + + sycl_source = """ + #include + #include "math_ops.hpp" + #include "math_template_ops.hpp" + + namespace syclext = sycl::ext::oneapi::experimental; + + extern "C" SYCL_EXTERNAL + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclext::nd_range_kernel<1>)) + void vector_add(int* in1, int* in2, int* out){ + sycl::nd_item<1> item = + sycl::ext::oneapi::this_work_item::get_nd_item<1>(); + size_t globalID = item.get_global_linear_id(); + out[globalID] = math_op(in1[globalID],in2[globalID]); + } + + template + SYCL_EXTERNAL + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclext::nd_range_kernel<1>)) + void vector_add_template(T* in1, T* in2, T* out){ + sycl::nd_item<1> item = + sycl::ext::oneapi::this_work_item::get_nd_item<1>(); + size_t globalID = item.get_global_linear_id(); + out[globalID] = math_op_template(in1[globalID], in2[globalID]); + } + """ + + header_content = """ + int math_op(int a, int b){ + return a + b; + } + """ + + header2_content = """ + template + T math_op_template(T a, T b){ + return a + b; + } + """ + + prog = dpctl.program.create_program_from_sycl_source( + q, + sycl_source, + headers=[ + ("math_ops.hpp", header_content), + ("math_template_ops.hpp", header2_content), + ], + registered_names=["vector_add_template"], + copts=["-fno-fast-math"], + ) + + assert type(prog) is dpctl_prog.SyclProgram + + assert type(prog.addressof_ref()) is int + assert prog.has_sycl_kernel("vector_add") + assert prog.has_sycl_kernel("vector_add_template") + + regularKernel = prog.get_sycl_kernel("vector_add") + templateKernel = prog.get_sycl_kernel("vector_add_template") + + assert "vector_add" == regularKernel.get_function_name() + assert type(regularKernel.addressof_ref()) is int + assert type(templateKernel.addressof_ref()) is int + + for krn in [regularKernel, templateKernel]: + _check_cpython_api_SyclKernel_GetKernelRef(krn) + _check_cpython_api_SyclKernel_Make(krn) + + assert 3 == krn.get_num_args() + na = krn.num_args + assert na == krn.get_num_args() + wgsz = krn.work_group_size + assert type(wgsz) is int + pwgszm = krn.preferred_work_group_size_multiple + assert type(pwgszm) is int + pmsz = krn.private_mem_size + assert type(pmsz) is int + vmnsg = krn.max_num_sub_groups + assert type(vmnsg) is int + v = krn.max_sub_group_size + assert type(v) is int + cmnsg = krn.compile_num_sub_groups + assert type(cmnsg) is int + cmsgsz = krn.compile_sub_group_size + assert type(cmsgsz) is int + + _check_cpython_api_SyclProgram_GetKernelBundleRef(prog) diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_device_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_device_interface.h index 6fddb2967f..b712e7cb41 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_device_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_device_interface.h @@ -792,4 +792,40 @@ DPCTL_API __dpctl_give DPCTLDeviceVectorRef DPCTLDevice_GetComponentDevices(__dpctl_keep const DPCTLSyclDeviceRef DRef); +/*! + * @brief Checks whether it is possible to create executables kernel bundles + * from SPIR-V binaries on this device. + * + * @param DRef Opaque pointer to a ``sycl::device``. + * @return True if creation is supported. + * #DPCTLSyclDeviceRef objects + * @ingroup DeviceInterface + */ +DPCTL_API +bool DPCTLDevice_CanCompileSPIRV(__dpctl_keep const DPCTLSyclDeviceRef DRef); + +/*! + * @brief Checks whether it is possible to create executables kernel bundles + * from OpenCL source code on this device. + * + * @param DRef Opaque pointer to a ``sycl::device``. + * @return True if creation is supported. + * #DPCTLSyclDeviceRef objects + * @ingroup DeviceInterface + */ +DPCTL_API +bool DPCTLDevice_CanCompileOpenCL(__dpctl_keep const DPCTLSyclDeviceRef DRef); + +/*! + * @brief Checks whether it is possible to create executables kernel bundles + * from SYCL source code on this device. + * + * @param DRef Opaque pointer to a ``sycl::device``. + * @return True if creation is supported. + * #DPCTLSyclDeviceRef objects + * @ingroup DeviceInterface + */ +DPCTL_API +bool DPCTLDevice_CanCompileSYCL(__dpctl_keep const DPCTLSyclDeviceRef DRef); + DPCTL_C_EXTERN_C_END diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h index 529bc3cca1..979ea4c7af 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h @@ -32,6 +32,8 @@ #include "Support/MemOwnershipAttrs.h" #include "dpctl_data_types.h" #include "dpctl_sycl_types.h" +#include +#include DPCTL_C_EXTERN_C_BEGIN @@ -129,4 +131,159 @@ DPCTL_API __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_Copy(__dpctl_keep const DPCTLSyclKernelBundleRef KBRef); +struct DPCTLBuildOptionList +{ + std::vector options; +}; + +struct DPCTLKernelNameList +{ + std::vector names; +}; + +struct DPCTLVirtualHeaderList +{ + std::vector> headers; +}; + +using DPCTLBuildOptionListRef = DPCTLBuildOptionList *; +using DPCTLKernelNameListRef = DPCTLKernelNameList *; +using DPCTLVirtualHeaderListRef = DPCTLVirtualHeaderList *; + +/*! + * @brief Create an empty list of build options. + * + * @return Opaque pointer to the build option file list. + * @ingroup KernelBundleInterface + */ +DPCTL_API +__dpctl_give DPCTLBuildOptionListRef DPCTLBuildOptionList_Create(); + +/*! + * @brief Frees the DPCTLBuildOptionListRef pointer. + * + * @param KBRef Opaque pointer to a list of build options + * @ingroup KernelBundleInterface + */ +DPCTL_API void +DPCTLBuildOptionList_Delete(__dpctl_take DPCTLBuildOptionListRef Ref); + +/*! + * @brief Append a build option to the list of build options + * + * @param Ref Opaque pointer to the list of build options + * @param Option Option to append + */ +DPCTL_API +void DPCTLBuildOptionList_Append(__dpctl_keep DPCTLBuildOptionListRef Ref, + __dpctl_keep const char *Option); + +/*! + * @brief Create an empty list of kernel names to register. + * + * @return Opaque pointer to the list of kernel names to register. + * @ingroup KernelBundleInterface + */ +DPCTL_API +__dpctl_give DPCTLKernelNameListRef DPCTLKernelNameList_Create(); + +/*! + * @brief Frees the DPCTLKernelNameListRef pointer. + * + * @param KBRef Opaque pointer to a list of kernels to register + * @ingroup KernelBundleInterface + */ +DPCTL_API void +DPCTLKernelNameList_Delete(__dpctl_take DPCTLKernelNameListRef Ref); + +/*! + * @brief Append a kernel name to register to the list of build options + * + * @param Ref Opaque pointer to the list of kernel names + * @param Option Kernel name to append + */ +DPCTL_API +void DPCTLKernelNameList_Append(__dpctl_keep DPCTLKernelNameListRef Ref, + __dpctl_keep const char *Option); +/*! + * @brief Create an empty list of virtual header files. + * + * @return Opaque pointer to the virtual header file list. + * @ingroup KernelBundleInterface + */ +DPCTL_API +__dpctl_give DPCTLVirtualHeaderListRef DPCTLVirtualHeaderList_Create(); + +/*! + * @brief Frees the DPCTLVirtualHeaderListRef pointer. + * + * @param KBRef Opaque pointer to a list of virtual headers + * @ingroup KernelBundleInterface + */ +DPCTL_API void +DPCTLVirtualHeaderList_Delete(__dpctl_take DPCTLVirtualHeaderListRef Ref); + +/*! + * @brief Append a kernel name to register to the list of virtual header files + * + * @param Ref Opaque pointer to the list of header files + * @param Name Name of the virtual header file + * @param Content Content of the virtual header + */ +DPCTL_API +void DPCTLVirtualHeaderList_Append(__dpctl_keep DPCTLVirtualHeaderListRef Ref, + __dpctl_keep const char *Name, + __dpctl_keep const char *Content); + +/*! + * @brief Create a SYCL kernel bundle from an SYCL kernel source string. + * + * @param Ctx An opaque pointer to a sycl::context + * @param Dev An opaque pointer to a sycl::device + * @param Source SYCL source string + * @param Headers List of virtual headers + * @param Names List of kernel names to register + * @param CompileOpts List of extra compiler flags (refer Sycl spec.) + * @return A new SyclKernelBundleRef pointer if the program creation + * succeeded, else returns NULL. + * @ingroup KernelBundleInterface + */ +DPCTL_API +__dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource( + __dpctl_keep const DPCTLSyclContextRef Ctx, + __dpctl_keep const DPCTLSyclDeviceRef Dev, + __dpctl_keep const char *Source, + __dpctl_keep DPCTLVirtualHeaderListRef Headers, + __dpctl_keep DPCTLKernelNameListRef Names, + __dpctl_keep DPCTLBuildOptionListRef BuildOptions); + +/*! + * @brief Returns the SyclKernel with given name from the program compiled from + * SYCL source code, if not found then return NULL. + * + * @param KBRef Opaque pointer to a sycl::kernel_bundle + * @param KernelName Name of kernel + * @return A SyclKernel reference if the kernel exists, else NULL + * @ingroup KernelBundleInterface + */ +DPCTL_API +__dpctl_give DPCTLSyclKernelRef +DPCTLKernelBundle_GetSyclKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, + __dpctl_keep const char *KernelName); + +/*! + * @brief Return True if a SyclKernel with given name exists in the program + * compiled from SYCL source code, if not found then returns False. + * + * @param KBRef Opaque pointer to a sycl::kernel_bundle + * @param KernelName Name of kernel + * @return True if the kernel exists, else False + * @ingroup KernelBundleInterface + */ + +DPCTL_API +bool DPCTLKernelBundle_HasSyclKernel(__dpctl_keep DPCTLSyclKernelBundleRef + KBRef, + __dpctl_keep const char *KernelName); + DPCTL_C_EXTERN_C_END diff --git a/libsyclinterface/source/dpctl_sycl_device_interface.cpp b/libsyclinterface/source/dpctl_sycl_device_interface.cpp index 7b1e900b58..7eff0b242d 100644 --- a/libsyclinterface/source/dpctl_sycl_device_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_device_interface.cpp @@ -903,3 +903,28 @@ DPCTLDevice_GetCompositeDevice(__dpctl_keep const DPCTLSyclDeviceRef DRef) else return nullptr; } + +bool DPCTLDevice_CanCompileSPIRV(__dpctl_keep const DPCTLSyclDeviceRef DRef) +{ + auto Dev = unwrap(DRef); + auto Backend = Dev->get_platform().get_backend(); + return Backend == backend::opencl || + Backend == backend::ext_oneapi_level_zero; +} + +bool DPCTLDevice_CanCompileOpenCL(__dpctl_keep const DPCTLSyclDeviceRef DRef) +{ + auto Dev = unwrap(DRef); + return Dev->get_platform().get_backend() == backend::opencl; +} + +bool DPCTLDevice_CanCompileSYCL(__dpctl_keep const DPCTLSyclDeviceRef DRef) +{ +#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER + auto Dev = unwrap(DRef); + return Dev->ext_oneapi_can_compile( + ext::oneapi::experimental::source_language::sycl); +#else + return false; +#endif +} diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp index c702018687..55bca0ab45 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -761,3 +761,143 @@ DPCTLKernelBundle_Copy(__dpctl_keep const DPCTLSyclKernelBundleRef KBRef) return nullptr; } } + +__dpctl_give DPCTLBuildOptionListRef DPCTLBuildOptionList_Create() +{ + return new DPCTLBuildOptionList; +} + +void DPCTLBuildOptionList_Delete(__dpctl_take DPCTLBuildOptionListRef Ref) +{ + delete Ref; +} + +void DPCTLBuildOptionList_Append(__dpctl_keep DPCTLBuildOptionListRef Ref, + __dpctl_keep const char *Option) +{ + Ref->options.emplace_back(Option); +} + +__dpctl_give DPCTLKernelNameListRef DPCTLKernelNameList_Create() +{ + return new DPCTLKernelNameList; +} + +void DPCTLKernelNameList_Delete(__dpctl_take DPCTLKernelNameListRef Ref) +{ + delete Ref; +} + +void DPCTLKernelNameList_Append(__dpctl_keep DPCTLKernelNameListRef Ref, + __dpctl_keep const char *Option) +{ + Ref->names.emplace_back(Option); +} + +__dpctl_give DPCTLVirtualHeaderListRef DPCTLVirtualHeaderList_Create() +{ + return new DPCTLVirtualHeaderList; +} + +void DPCTLVirtualHeaderList_Delete(__dpctl_take DPCTLVirtualHeaderListRef Ref) +{ + delete Ref; +} + +void DPCTLVirtualHeaderList_Append(__dpctl_keep DPCTLVirtualHeaderListRef Ref, + __dpctl_keep const char *Name, + __dpctl_keep const char *Content) +{ + auto Header = std::make_pair(Name, Content); + Ref->headers.push_back(Header); +} + +namespace syclex = sycl::ext::oneapi::experimental; + +__dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource( + __dpctl_keep const DPCTLSyclContextRef Ctx, + __dpctl_keep const DPCTLSyclDeviceRef Dev, + __dpctl_keep const char *Source, + __dpctl_keep DPCTLVirtualHeaderListRef Headers, + __dpctl_keep DPCTLKernelNameListRef Names, + __dpctl_keep DPCTLBuildOptionListRef BuildOptions) +{ +#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER + context *SyclCtx = unwrap(Ctx); + device *SyclDev = unwrap(Dev); + if (!SyclDev->ext_oneapi_can_compile(syclex::source_language::sycl)) { + return nullptr; + } + try { + syclex::include_files IncludeFiles; + for (auto &Include : Headers->headers) { + const auto &[Name, Content] = Include; + IncludeFiles.add(Name, Content); + } + + std::string Src(Source); + auto SrcBundle = syclex::create_kernel_bundle_from_source( + *SyclCtx, syclex::source_language::sycl, Src, + syclex::properties{IncludeFiles}); + + syclex::registered_names RegisteredNames; + for (const std::string &Name : Names->names) { + RegisteredNames.add(Name); + } + + syclex::build_options Opts{BuildOptions->options}; + + std::vector Devices({*SyclDev}); + + auto ExeBundle = syclex::build( + SrcBundle, Devices, syclex::properties{RegisteredNames, Opts}); + auto ResultBundle = + std::make_unique>( + ExeBundle); + return wrap>( + ResultBundle.release()); + } catch (const std::exception &e) { + error_handler(e, __FILE__, __func__, __LINE__); + return nullptr; + } +#else + return nullptr; +#endif +} + +__dpctl_give DPCTLSyclKernelRef +DPCTLKernelBundle_GetSyclKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, + __dpctl_keep const char *KernelName) +{ +#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER + try { + auto KernelBundle = + unwrap>(KBRef); + auto Kernel = KernelBundle->ext_oneapi_get_kernel(KernelName); + return wrap(new sycl::kernel(Kernel)); + } catch (const std::exception &e) { + error_handler(e, __FILE__, __func__, __LINE__); + return nullptr; + } +#else + return nullptr; +#endif +} + +bool DPCTLKernelBundle_HasSyclKernel(__dpctl_keep DPCTLSyclKernelBundleRef + KBRef, + __dpctl_keep const char *KernelName) +{ +#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER + try { + auto KernelBundle = + unwrap>(KBRef); + return KernelBundle->ext_oneapi_has_kernel(KernelName); + } catch (const std::exception &e) { + error_handler(e, __FILE__, __func__, __LINE__); + return false; + } +#else + return false; +#endif +} diff --git a/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp index d136c700b6..252b3260ed 100644 --- a/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp @@ -273,6 +273,119 @@ TEST_P(TestOCLKernelBundleFromSource, CheckGetKernelOCLSource) DPCTLKernel_Delete(AxpyKernel); } +struct TestSYCLKernelBundleFromSource + : public ::testing::TestWithParam +{ + const char *sycl_source = R"===( + #include + #include "math_ops.hpp" + #include "math_template_ops.hpp" + + namespace syclext = sycl::ext::oneapi::experimental; + + extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclext::nd_range_kernel<1>)) + void vector_add(int* in1, int* in2, int* out){ + sycl::nd_item<1> item = sycl::ext::oneapi::this_work_item::get_nd_item<1>(); + size_t globalID = item.get_global_linear_id(); + out[globalID] = math_op(in1[globalID],in2[globalID]); + } + + template + SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclext::nd_range_kernel<1>)) + void vector_add_template(T* in1, T* in2, T* out){ + sycl::nd_item<1> item = sycl::ext::oneapi::this_work_item::get_nd_item<1>(); + size_t globalID = item.get_global_linear_id(); + out[globalID] = math_op_template(in1[globalID], in2[globalID]); + } + )==="; + + const char *header1_content = R"===( + int math_op(int a, int b){ + return a + b; + } + )==="; + + const char *header2_content = R"===( + template + T math_op_template(T a, T b){ + return a + b; + } + )==="; + + const char *CompileOpt = "-fno-fast-math"; + const char *KernelName = "vector_add_template"; + const char *Header1Name = "math_ops.hpp"; + const char *Header2Name = "math_template_ops.hpp"; + DPCTLSyclDeviceRef DRef = nullptr; + DPCTLSyclContextRef CRef = nullptr; + DPCTLSyclKernelBundleRef KBRef = nullptr; + + TestSYCLKernelBundleFromSource() + { + auto DS = DPCTLFilterSelector_Create(GetParam()); + DRef = DPCTLDevice_CreateFromSelector(DS); + DPCTLDeviceSelector_Delete(DS); + CRef = DPCTLDeviceMgr_GetCachedContext(DRef); + + if (DRef) { + DPCTLBuildOptionListRef BORef = DPCTLBuildOptionList_Create(); + DPCTLBuildOptionList_Append(BORef, CompileOpt); + DPCTLKernelNameListRef KNRef = DPCTLKernelNameList_Create(); + DPCTLKernelNameList_Append(KNRef, KernelName); + DPCTLVirtualHeaderListRef VHRef = DPCTLVirtualHeaderList_Create(); + DPCTLVirtualHeaderList_Append(VHRef, Header1Name, header1_content); + DPCTLVirtualHeaderList_Append(VHRef, Header2Name, header2_content); + KBRef = DPCTLKernelBundle_CreateFromSYCLSource( + CRef, DRef, sycl_source, VHRef, KNRef, BORef); + DPCTLVirtualHeaderList_Delete(VHRef); + DPCTLKernelNameList_Delete(KNRef); + DPCTLBuildOptionList_Delete(BORef); + } + } + + void SetUp() + { + if (!DRef) { + auto message = "Skipping as no device of type " + + std::string(GetParam()) + "."; + GTEST_SKIP_(message.c_str()); + } + if (!DPCTLDevice_CanCompileSYCL(DRef)) { + const char *message = "Skipping as SYCL compilation not supported"; + GTEST_SKIP_(message); + } + } + + ~TestSYCLKernelBundleFromSource() + { + if (DRef) + DPCTLDevice_Delete(DRef); + if (CRef) + DPCTLContext_Delete(CRef); + if (KBRef) + DPCTLKernelBundle_Delete(KBRef); + } +}; + +TEST_P(TestSYCLKernelBundleFromSource, CheckCreateFromSYCLSource) +{ + ASSERT_TRUE(KBRef != nullptr); + ASSERT_TRUE(DPCTLKernelBundle_HasSyclKernel(KBRef, "vector_add")); + ASSERT_TRUE( + DPCTLKernelBundle_HasSyclKernel(KBRef, "vector_add_template")); +} + +TEST_P(TestSYCLKernelBundleFromSource, CheckGetKernelSYCLSource) +{ + auto AddKernel = DPCTLKernelBundle_GetSyclKernel(KBRef, "vector_add"); + auto AxpyKernel = + DPCTLKernelBundle_GetSyclKernel(KBRef, "vector_add_template"); + ASSERT_TRUE(AddKernel != nullptr); + ASSERT_TRUE(AxpyKernel != nullptr); + DPCTLKernel_Delete(AddKernel); + DPCTLKernel_Delete(AxpyKernel); +} + INSTANTIATE_TEST_SUITE_P(KernelBundleCreationFromSpirv, TestDPCTLSyclKernelBundleInterface, ::testing::Values("opencl", @@ -289,6 +402,12 @@ INSTANTIATE_TEST_SUITE_P(KernelBundleCreationFromSource, TestOCLKernelBundleFromSource, ::testing::Values("opencl:gpu", "opencl:cpu")); +INSTANTIATE_TEST_SUITE_P(KernelBundleCreationFromSYCL, + TestSYCLKernelBundleFromSource, + ::testing::Values("opencl:gpu", + "opencl:cpu", + "level_zero:gpu")); + struct TestKernelBundleUnsupportedBackend : public ::testing::Test { DPCTLSyclDeviceRef DRef = nullptr; From 01fc7f997ea1edc034cde6dc32cbe0cdcc60b9a8 Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Wed, 23 Apr 2025 17:07:22 +0100 Subject: [PATCH 2/5] Enable conditionally and make API C-compatible Enable SYCL source compilation only for DPC++ versions that actually support the compilation, based on the __SYCL_COMPILER_VERSION reported. Use the correct naming for the property based on DPC++ version. Remove all mentions of `std::vector` and other STL types from the header and use opaque pointers instead. Signed-off-by: Lukas Sommer --- dpctl/_backend.pxd | 15 ++-- dpctl/_sycl_device.pyx | 2 - dpctl/program/_program.pxd | 3 +- dpctl/program/_program.pyx | 21 +++-- .../dpctl_sycl_kernel_bundle_interface.h | 23 +----- .../dpctl_sycl_kernel_bundle_interface.cpp | 81 +++++++++++++++---- 6 files changed, 91 insertions(+), 54 deletions(-) diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index 677e3699ad..351c9bf8eb 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -454,18 +454,18 @@ cdef extern from "syclinterface/dpctl_sycl_kernel_bundle_interface.h": cdef DPCTLBuildOptionListRef DPCTLBuildOptionList_Create() cdef void DPCTLBuildOptionList_Delete(DPCTLBuildOptionListRef Ref) cdef void DPCTLBuildOptionList_Append(DPCTLBuildOptionListRef Ref, - const char *Option) + const char *Option) cdef DPCTLKernelNameListRef DPCTLKernelNameList_Create() cdef void DPCTLKernelNameList_Delete(DPCTLKernelNameListRef Ref) cdef void DPCTLKernelNameList_Append(DPCTLKernelNameListRef Ref, - const char *Option) + const char *Option) cdef DPCTLVirtualHeaderListRef DPCTLVirtualHeaderList_Create() cdef void DPCTLVirtualHeaderList_Delete(DPCTLVirtualHeaderListRef Ref) cdef void DPCTLVirtualHeaderList_Append(DPCTLVirtualHeaderListRef Ref, - const char *Name, - const char *Content) + const char *Name, + const char *Content) cdef DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource( const DPCTLSyclContextRef Ctx, @@ -475,11 +475,12 @@ cdef extern from "syclinterface/dpctl_sycl_kernel_bundle_interface.h": DPCTLKernelNameListRef Names, DPCTLBuildOptionListRef BuildOptions) - cdef DPCTLSyclKernelRef DPCTLKernelBundle_GetSyclKernel(DPCTLSyclKernelBundleRef KBRef, - const char *KernelName) + cdef DPCTLSyclKernelRef DPCTLKernelBundle_GetSyclKernel( + DPCTLSyclKernelBundleRef KBRef, + const char *KernelName) cdef bool DPCTLKernelBundle_HasSyclKernel(DPCTLSyclKernelBundleRef KBRef, - const char *KernelName); + const char *KernelName) cdef extern from "syclinterface/dpctl_sycl_queue_interface.h": diff --git a/dpctl/_sycl_device.pyx b/dpctl/_sycl_device.pyx index bda700cf11..b5be0af1ee 100644 --- a/dpctl/_sycl_device.pyx +++ b/dpctl/_sycl_device.pyx @@ -26,7 +26,6 @@ from ._backend cimport ( # noqa: E211 DPCTLDefaultSelector_Create, DPCTLDevice_AreEq, DPCTLDevice_CanCompileOpenCL, - DPCTLDevice_CanCompileSPIRV, DPCTLDevice_CanCompileSYCL, DPCTLDevice_Copy, DPCTLDevice_CreateFromSelector, @@ -2192,7 +2191,6 @@ cdef class SyclDevice(_SyclDevice): raise ValueError(f"Unknown source language {language}") - cdef api DPCTLSyclDeviceRef SyclDevice_GetDeviceRef(SyclDevice dev): """ C-API function to get opaque device reference from diff --git a/dpctl/program/_program.pxd b/dpctl/program/_program.pxd index 37d7fbaac5..880843c27f 100644 --- a/dpctl/program/_program.pxd +++ b/dpctl/program/_program.pxd @@ -52,7 +52,8 @@ cdef api class SyclProgram [object PySyclProgramObject, type PySyclProgramType]: cdef bint _is_sycl_source @staticmethod - cdef SyclProgram _create (DPCTLSyclKernelBundleRef pref, bint _is_sycl_source) + cdef SyclProgram _create (DPCTLSyclKernelBundleRef pref, + bint _is_sycl_source) cdef DPCTLSyclKernelBundleRef get_program_ref (self) cpdef SyclKernel get_sycl_kernel(self, str kernel_name) diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index f175cf23a7..f371149bfd 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -211,7 +211,8 @@ cdef class SyclProgram: """ @staticmethod - cdef SyclProgram _create(DPCTLSyclKernelBundleRef KBRef, bint is_sycl_source): + cdef SyclProgram _create(DPCTLSyclKernelBundleRef KBRef, + bint is_sycl_source): cdef SyclProgram ret = SyclProgram.__new__(SyclProgram) ret._program_ref = KBRef ret._is_sycl_source = is_sycl_source @@ -342,7 +343,10 @@ cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL, return SyclProgram._create(KBref, False) -cpdef create_program_from_sycl_source(SyclQueue q, unicode source, list headers=[], list registered_names=[], list copts=[]): +cpdef create_program_from_sycl_source(SyclQueue q, unicode source, + list headers=None, + list registered_names=None, + list copts=None): """ Creates an executable SYCL kernel_bundle from SYCL source code. @@ -384,7 +388,7 @@ cpdef create_program_from_sycl_source(SyclQueue q, unicode source, list headers= cdef DPCTLSyclKernelBundleRef KBref cdef DPCTLSyclContextRef CRef = q.get_sycl_context().get_context_ref() cdef DPCTLSyclDeviceRef DRef = q.get_sycl_device().get_device_ref() - cdef bytes bSrc = source.encode('utf8') + cdef bytes bSrc = source.encode("utf8") cdef const char *Src = bSrc cdef DPCTLBuildOptionListRef BuildOpts = DPCTLBuildOptionList_Create() cdef bytes bOpt @@ -397,7 +401,7 @@ cpdef create_program_from_sycl_source(SyclQueue q, unicode source, list headers= if not isinstance(opt, unicode): DPCTLBuildOptionList_Delete(BuildOpts) raise SyclProgramCompilationError() - bOpt = opt.encode('utf8') + bOpt = opt.encode("utf8") sOpt = bOpt DPCTLBuildOptionList_Append(BuildOpts, sOpt) @@ -407,21 +411,22 @@ cpdef create_program_from_sycl_source(SyclQueue q, unicode source, list headers= DPCTLBuildOptionList_Delete(BuildOpts) DPCTLKernelNameList_Delete(KernelNames) raise SyclProgramCompilationError() - bName = name.encode('utf8') + bName = name.encode("utf8") sName = bName DPCTLKernelNameList_Append(KernelNames, sName) + cdef DPCTLVirtualHeaderListRef VirtualHeaders + VirtualHeaders = DPCTLVirtualHeaderList_Create() - cdef DPCTLVirtualHeaderListRef VirtualHeaders = DPCTLVirtualHeaderList_Create() for name, content in headers: if not isinstance(name, unicode) or not isinstance(content, unicode): DPCTLBuildOptionList_Delete(BuildOpts) DPCTLKernelNameList_Delete(KernelNames) DPCTLVirtualHeaderList_Delete(VirtualHeaders) raise SyclProgramCompilationError() - bName = name.encode('utf8') + bName = name.encode("utf8") sName = bName - bContent = content.encode('utf8') + bContent = content.encode("utf8") sContent = bContent DPCTLVirtualHeaderList_Append(VirtualHeaders, sName, sContent) diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h index 979ea4c7af..32cd289f20 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h @@ -32,8 +32,6 @@ #include "Support/MemOwnershipAttrs.h" #include "dpctl_data_types.h" #include "dpctl_sycl_types.h" -#include -#include DPCTL_C_EXTERN_C_BEGIN @@ -131,24 +129,9 @@ DPCTL_API __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_Copy(__dpctl_keep const DPCTLSyclKernelBundleRef KBRef); -struct DPCTLBuildOptionList -{ - std::vector options; -}; - -struct DPCTLKernelNameList -{ - std::vector names; -}; - -struct DPCTLVirtualHeaderList -{ - std::vector> headers; -}; - -using DPCTLBuildOptionListRef = DPCTLBuildOptionList *; -using DPCTLKernelNameListRef = DPCTLKernelNameList *; -using DPCTLVirtualHeaderListRef = DPCTLVirtualHeaderList *; +typedef struct DPCTLBuildOptionList *DPCTLBuildOptionListRef; +typedef struct DPCTLKernelNameList *DPCTLKernelNameListRef; +typedef struct DPCTLVirtualHeaderList *DPCTLVirtualHeaderListRef; /*! * @brief Create an empty list of build options. diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp index 55bca0ab45..0a6afacf70 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -762,46 +762,67 @@ DPCTLKernelBundle_Copy(__dpctl_keep const DPCTLSyclKernelBundleRef KBRef) } } +using build_option_list_t = std::vector; + __dpctl_give DPCTLBuildOptionListRef DPCTLBuildOptionList_Create() { - return new DPCTLBuildOptionList; + auto BuildOptionList = + std::unique_ptr(new build_option_list_t()); + auto *RetVal = + reinterpret_cast(BuildOptionList.get()); + BuildOptionList.release(); + return RetVal; } void DPCTLBuildOptionList_Delete(__dpctl_take DPCTLBuildOptionListRef Ref) { - delete Ref; + delete reinterpret_cast(Ref); } void DPCTLBuildOptionList_Append(__dpctl_keep DPCTLBuildOptionListRef Ref, __dpctl_keep const char *Option) { - Ref->options.emplace_back(Option); + reinterpret_cast(Ref)->emplace_back(Option); } +using kernel_name_list_t = std::vector; + __dpctl_give DPCTLKernelNameListRef DPCTLKernelNameList_Create() { - return new DPCTLKernelNameList; + auto KernelNameList = + std::unique_ptr(new kernel_name_list_t()); + auto *RetVal = + reinterpret_cast(KernelNameList.get()); + KernelNameList.release(); + return RetVal; } void DPCTLKernelNameList_Delete(__dpctl_take DPCTLKernelNameListRef Ref) { - delete Ref; + delete reinterpret_cast(Ref); } void DPCTLKernelNameList_Append(__dpctl_keep DPCTLKernelNameListRef Ref, __dpctl_keep const char *Option) { - Ref->names.emplace_back(Option); + reinterpret_cast(Ref)->emplace_back(Option); } +using virtual_header_list_t = std::vector>; + __dpctl_give DPCTLVirtualHeaderListRef DPCTLVirtualHeaderList_Create() { - return new DPCTLVirtualHeaderList; + auto HeaderList = + std::unique_ptr(new virtual_header_list_t()); + auto *RetVal = + reinterpret_cast(HeaderList.get()); + HeaderList.release(); + return RetVal; } void DPCTLVirtualHeaderList_Delete(__dpctl_take DPCTLVirtualHeaderListRef Ref) { - delete Ref; + delete reinterpret_cast(Ref); } void DPCTLVirtualHeaderList_Append(__dpctl_keep DPCTLVirtualHeaderListRef Ref, @@ -809,11 +830,34 @@ void DPCTLVirtualHeaderList_Append(__dpctl_keep DPCTLVirtualHeaderListRef Ref, __dpctl_keep const char *Content) { auto Header = std::make_pair(Name, Content); - Ref->headers.push_back(Header); + reinterpret_cast(Ref)->push_back(Header); } namespace syclex = sycl::ext::oneapi::experimental; +#if defined(SYCL_EXT_ONEAPI_KERNEL_COMPILER) && \ + defined(__SYCL_COMPILER_VERSION) && !defined(SUPPORTS_SYCL_COMPILATION) +// SYCL source code compilation is supported from 2025.1 onwards. +#if __SYCL_COMPILER_VERSION >= 20250317u +#define SUPPORTS_SYCL_COMPILATION 1 +#else +#define SUPPORTS_SYCL_COMPILATION 0 +#endif +#endif + +#if (SUPPORTS_SYCL_COMPILATION > 0) +#ifndef __SYCL_COMPILER_VERSION +#error SYCL compiler version not defined +#else +// The property was renamed to `registered_names` after 2025.1 +#if __SYCL_COMPILER_VERSION > 20250317u +using registered_names_property_t = syclex::registered_names; +#else +using registered_names_property_t = syclex::registered_kernel_names; +#endif +#endif +#endif + __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource( __dpctl_keep const DPCTLSyclContextRef Ctx, __dpctl_keep const DPCTLSyclDeviceRef Dev, @@ -822,7 +866,7 @@ __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource( __dpctl_keep DPCTLKernelNameListRef Names, __dpctl_keep DPCTLBuildOptionListRef BuildOptions) { -#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER +#if (SUPPORTS_SYCL_COMPILATION > 0) context *SyclCtx = unwrap(Ctx); device *SyclDev = unwrap(Dev); if (!SyclDev->ext_oneapi_can_compile(syclex::source_language::sycl)) { @@ -830,7 +874,9 @@ __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource( } try { syclex::include_files IncludeFiles; - for (auto &Include : Headers->headers) { + for (auto &Include : + *reinterpret_cast(Headers)) + { const auto &[Name, Content] = Include; IncludeFiles.add(Name, Content); } @@ -840,12 +886,15 @@ __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource( *SyclCtx, syclex::source_language::sycl, Src, syclex::properties{IncludeFiles}); - syclex::registered_names RegisteredNames; - for (const std::string &Name : Names->names) { + registered_names_property_t RegisteredNames; + for (const std::string &Name : + *reinterpret_cast(Names)) + { RegisteredNames.add(Name); } - syclex::build_options Opts{BuildOptions->options}; + syclex::build_options Opts{ + *reinterpret_cast(BuildOptions)}; std::vector Devices({*SyclDev}); @@ -869,7 +918,7 @@ __dpctl_give DPCTLSyclKernelRef DPCTLKernelBundle_GetSyclKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, __dpctl_keep const char *KernelName) { -#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER +#if (SUPPORTS_SYCL_COMPILATION > 0) try { auto KernelBundle = unwrap>(KBRef); @@ -888,7 +937,7 @@ bool DPCTLKernelBundle_HasSyclKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, __dpctl_keep const char *KernelName) { -#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER +#if (SUPPORTS_SYCL_COMPILATION > 0) try { auto KernelBundle = unwrap>(KBRef); From 92a7fe7b8c87996ac54791ea8190e1c91cc24907 Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Thu, 24 Apr 2025 14:02:08 +0100 Subject: [PATCH 3/5] Work around missing constructor implementation This commit works around a bug in DPC++ version 2025.1. The constructor with no parameter of class `include_files` was only declared, but never defined. Calling it when creating a SYCL source kernel bundle therefore leads to references to undefined symbols with DPC++ version 2025.1. This change works around this issue by calling an alternative constructor, which is defined in the release. Signed-off-by: Lukas Sommer --- .../dpctl_sycl_kernel_bundle_interface.cpp | 43 +++++++++++++------ 1 file changed, 31 insertions(+), 12 deletions(-) diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp index 0a6afacf70..ce9098ccad 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -873,18 +873,37 @@ __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource( return nullptr; } try { - syclex::include_files IncludeFiles; - for (auto &Include : - *reinterpret_cast(Headers)) - { - const auto &[Name, Content] = Include; - IncludeFiles.add(Name, Content); - } - + auto *IncludeFileList = + reinterpret_cast(Headers); + std::unique_ptr> + SrcBundle; std::string Src(Source); - auto SrcBundle = syclex::create_kernel_bundle_from_source( - *SyclCtx, syclex::source_language::sycl, Src, - syclex::properties{IncludeFiles}); + // The following logic is to work around a bug in DPC++ version 2025.1. + // This version declares a constructor with no parameters for the + // `include_files` property, but does not implement it. Therefore, the + // only way to create `include_files` is with the name and content of + // the first virtual header, if any. + if (!IncludeFileList->empty()) { + auto IncludeFileIt = IncludeFileList->begin(); + syclex::include_files IncludeFiles{IncludeFileIt->first, + IncludeFileIt->second}; + for (std::advance(IncludeFileIt, 1); + IncludeFileIt != IncludeFileList->end(); ++IncludeFileIt) + { + IncludeFiles.add(IncludeFileIt->first, IncludeFileIt->second); + } + SrcBundle = std::make_unique< + kernel_bundle>( + syclex::create_kernel_bundle_from_source( + *SyclCtx, syclex::source_language::sycl, Src, + syclex::properties{IncludeFiles})); + } + else { + SrcBundle = std::make_unique< + kernel_bundle>( + syclex::create_kernel_bundle_from_source( + *SyclCtx, syclex::source_language::sycl, Src)); + } registered_names_property_t RegisteredNames; for (const std::string &Name : @@ -899,7 +918,7 @@ __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource( std::vector Devices({*SyclDev}); auto ExeBundle = syclex::build( - SrcBundle, Devices, syclex::properties{RegisteredNames, Opts}); + *SrcBundle, Devices, syclex::properties{RegisteredNames, Opts}); auto ResultBundle = std::make_unique>( ExeBundle); From 588c0db912ba3878d43a1b43d6667bf36730ca17 Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Tue, 6 May 2025 13:34:15 +0100 Subject: [PATCH 4/5] Make the property name detection logic more robust Instead of relying on the SYCL compiler version macro, we use C++ type_traits to detect which of the two names for the property actually refers to a fully defined type. Signed-off-by: Lukas Sommer --- .../dpctl_sycl_kernel_bundle_interface.cpp | 36 +++++++++++++------ 1 file changed, 26 insertions(+), 10 deletions(-) diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp index ce9098ccad..d8121da1d8 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -846,16 +846,32 @@ namespace syclex = sycl::ext::oneapi::experimental; #endif #if (SUPPORTS_SYCL_COMPILATION > 0) -#ifndef __SYCL_COMPILER_VERSION -#error SYCL compiler version not defined -#else -// The property was renamed to `registered_names` after 2025.1 -#if __SYCL_COMPILER_VERSION > 20250317u -using registered_names_property_t = syclex::registered_names; -#else -using registered_names_property_t = syclex::registered_kernel_names; -#endif -#endif +// The property for registering names was renamed between DPC++ versions 2025.1 +// and 2025.2. The original name was `registered_kernel_names`, the new name is +// `registered_names`. To select the correct name without being overly reliant +// on the SYCL compiler version definition, we forward declare both names and +// then select the new name if it is defined (i.e., not only declared). +namespace sycl::ext::oneapi::experimental +{ +struct registered_names; +struct registered_kernel_names; +} // namespace sycl::ext::oneapi::experimental + +template +struct new_type_if_defined +{ + using type = FallbackT; +}; + +template +struct new_type_if_defined> +{ + using type = NewT; +}; + +using registered_names_property_t = + new_type_if_defined::type; #endif __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource( From 153ec69eb12bca62c7a45412ec34e1b53506666d Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Fri, 9 May 2025 13:07:55 +0100 Subject: [PATCH 5/5] Adapt test for DPC++ 2025.1 Signed-off-by: Lukas Sommer --- dpctl/tests/test_sycl_program.py | 18 +++++++++++++++--- .../test_sycl_kernel_bundle_interface.cpp | 13 ++++++++++++- 2 files changed, 27 insertions(+), 4 deletions(-) diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index 8cf8f06215..4df119c87b 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -330,10 +330,22 @@ def test_create_program_from_sycl_source(): assert type(prog.addressof_ref()) is int assert prog.has_sycl_kernel("vector_add") - assert prog.has_sycl_kernel("vector_add_template") - regularKernel = prog.get_sycl_kernel("vector_add") - templateKernel = prog.get_sycl_kernel("vector_add_template") + + # DPC++ version 2025.1 supports compilation of SYCL template kernels, but + # does not yet support referencing them with the unmangled name. + hasTemplateName = prog.has_sycl_kernel("vector_add_template") + hasMangledName = prog.has_sycl_kernel( + "_Z33__sycl_kernel_vector_add_templateIiEvPT_S1_S1_" + ) + assert hasTemplateName or hasMangledName + + if hasTemplateName: + templateKernel = prog.get_sycl_kernel("vector_add_template") + else: + templateKernel = prog.get_sycl_kernel( + "_Z33__sycl_kernel_vector_add_templateIiEvPT_S1_S1_" + ) assert "vector_add" == regularKernel.get_function_name() assert type(regularKernel.addressof_ref()) is int diff --git a/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp index 252b3260ed..902d63a3d2 100644 --- a/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp @@ -371,8 +371,12 @@ TEST_P(TestSYCLKernelBundleFromSource, CheckCreateFromSYCLSource) { ASSERT_TRUE(KBRef != nullptr); ASSERT_TRUE(DPCTLKernelBundle_HasSyclKernel(KBRef, "vector_add")); + // DPC++ version 2025.1 supports compilation of SYCL template kernels, + // but does not yet support referencing them with the unmangled name. ASSERT_TRUE( - DPCTLKernelBundle_HasSyclKernel(KBRef, "vector_add_template")); + DPCTLKernelBundle_HasSyclKernel(KBRef, "vector_add_template") || + DPCTLKernelBundle_HasSyclKernel( + KBRef, "_Z33__sycl_kernel_vector_add_templateIiEvPT_S1_S1_")); } TEST_P(TestSYCLKernelBundleFromSource, CheckGetKernelSYCLSource) @@ -380,6 +384,13 @@ TEST_P(TestSYCLKernelBundleFromSource, CheckGetKernelSYCLSource) auto AddKernel = DPCTLKernelBundle_GetSyclKernel(KBRef, "vector_add"); auto AxpyKernel = DPCTLKernelBundle_GetSyclKernel(KBRef, "vector_add_template"); + if (AxpyKernel == nullptr) { + // DPC++ version 2025.1 supports compilation of SYCL template kernels, + // but does not yet support referencing them with the unmangled name. + AxpyKernel = DPCTLKernelBundle_GetSyclKernel( + KBRef, "_Z33__sycl_kernel_vector_add_templateIiEvPT_S1_S1_"); + } + ASSERT_TRUE(AddKernel != nullptr); ASSERT_TRUE(AxpyKernel != nullptr); DPCTLKernel_Delete(AddKernel);