From 686861e05d03f04e63fd649a25aa6023ddc0a5bb Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Thu, 10 Apr 2025 10:39:23 +0100 Subject: [PATCH 1/7] Support compilation from SYCL source code Enable SYCL source compilation, but only for DPC++ versions that actually support the compilation, based on the __SYCL_COMPILER_VERSION reported. Uses the correct naming for the property based on DPC++ version, detected through C++ type traits to check which property actually refers to a fully defined type. This commit also 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/_backend.pxd | 50 +++- dpctl/_sycl_device.pxd | 1 + dpctl/_sycl_device.pyx | 31 +++ dpctl/program/__init__.py | 2 + dpctl/program/_program.pxd | 8 +- dpctl/program/_program.pyx | 154 ++++++++++- dpctl/tests/test_sycl_program.py | 154 ++++++++++- .../dpctl_sycl_device_interface.h | 36 +++ .../dpctl_sycl_kernel_bundle_interface.h | 169 ++++++++++++ .../source/dpctl_sycl_device_interface.cpp | 25 ++ .../dpctl_sycl_kernel_bundle_interface.cpp | 247 ++++++++++++++++++ .../test_sycl_kernel_bundle_interface.cpp | 132 ++++++++++ 12 files changed, 1000 insertions(+), 9 deletions(-) diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index 93d9b5ef97..069f9a33fc 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -287,9 +287,12 @@ cdef extern from "syclinterface/dpctl_sycl_device_interface.h": _peer_access PT) cdef void DPCTLDevice_EnablePeerAccess(const DPCTLSyclDeviceRef DRef, const DPCTLSyclDeviceRef PDRef) - cdef void DPCTLDevice_DisablePeerAccess(const DPCTLSyclDeviceRef DRef, const DPCTLSyclDeviceRef PDRef) + 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": cdef DPCTLDeviceVectorRef DPCTLDeviceVector_CreateFromArray( @@ -452,6 +455,51 @@ 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 + cdef struct DPCTLKernelBuildLog + ctypedef DPCTLBuildOptionList* DPCTLBuildOptionListRef + ctypedef DPCTLKernelNameList* DPCTLKernelNameListRef + ctypedef DPCTLVirtualHeaderList* DPCTLVirtualHeaderListRef + ctypedef DPCTLKernelBuildLog* DPCTLKernelBuildLogRef + + 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 DPCTLKernelBuildLogRef DPCTLKernelBuildLog_Create() + cdef void DPCTLKernelBuildLog_Delete(DPCTLKernelBuildLogRef Ref) + cdef const char *DPCTLKernelBuildLog_Get(DPCTLKernelBuildLogRef) + + cdef DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource( + const DPCTLSyclContextRef Ctx, + const DPCTLSyclDeviceRef Dev, + const char *Source, + DPCTLVirtualHeaderListRef Headers, + DPCTLKernelNameListRef Names, + DPCTLBuildOptionListRef BuildOptions, + DPCTLKernelBuildLogRef BuildLog) + + 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 6c6638e2e3..d4bb39690f 100644 --- a/dpctl/_sycl_device.pyx +++ b/dpctl/_sycl_device.pyx @@ -26,6 +26,9 @@ from ._backend cimport ( # noqa: E211 DPCTLDefaultSelector_Create, DPCTLDevice_AreEq, DPCTLDevice_CanAccessPeer, + DPCTLDevice_CanCompileOpenCL, + DPCTLDevice_CanCompileSPIRV, + DPCTLDevice_CanCompileSYCL, DPCTLDevice_Copy, DPCTLDevice_CreateFromSelector, DPCTLDevice_CreateSubDevicesByAffinity, @@ -2363,6 +2366,34 @@ 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_CanCompileSPIRV(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 71302e4186..d47f7ac2d3 100644 --- a/dpctl/program/__init__.py +++ b/dpctl/program/__init__.py @@ -27,6 +27,7 @@ SyclKernelBundleCompilationError, create_kernel_bundle_from_source, create_kernel_bundle_from_spirv, + create_kernel_bundle_from_sycl_source, create_program_from_source, create_program_from_spirv, ) @@ -36,6 +37,7 @@ "create_kernel_bundle_from_spirv", "create_program_from_source", "create_program_from_spirv", + "create_kernel_bundle_from_sycl_source", "SyclKernel", "SyclKernelBundle", "SyclKernelBundleCompilationError", diff --git a/dpctl/program/_program.pxd b/dpctl/program/_program.pxd index 435ef68521..5e09d05f34 100644 --- a/dpctl/program/_program.pxd +++ b/dpctl/program/_program.pxd @@ -52,9 +52,11 @@ cdef api class SyclKernelBundle [ binary file. """ cdef DPCTLSyclKernelBundleRef _kernel_bundle_ref + cdef bint _is_sycl_source @staticmethod - cdef SyclKernelBundle _create (DPCTLSyclKernelBundleRef kbref) + cdef SyclKernelBundle _create (DPCTLSyclKernelBundleRef kbref, + bint _is_sycl_source) cdef DPCTLSyclKernelBundleRef get_kernel_bundle_ref (self) cpdef SyclKernel get_sycl_kernel(self, str kernel_name) @@ -69,3 +71,7 @@ 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_kernel_bundle_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 8737be4762..84dd10fc32 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -31,6 +31,10 @@ from libc.stdint cimport uint32_t import warnings from dpctl._backend cimport ( # noqa: E211, E402; + DPCTLBuildOptionList_Append, + DPCTLBuildOptionList_Create, + DPCTLBuildOptionList_Delete, + DPCTLBuildOptionListRef, DPCTLKernel_Copy, DPCTLKernel_Delete, DPCTLKernel_GetCompileNumSubGroups, @@ -41,16 +45,31 @@ from dpctl._backend cimport ( # noqa: E211, E402; DPCTLKernel_GetPreferredWorkGroupSizeMultiple, DPCTLKernel_GetPrivateMemSize, DPCTLKernel_GetWorkGroupSize, + DPCTLKernelBuildLog_Create, + DPCTLKernelBuildLog_Delete, + DPCTLKernelBuildLog_Get, + DPCTLKernelBuildLogRef, 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__ = [ @@ -199,9 +218,11 @@ cdef class SyclKernelBundle: """ @staticmethod - cdef SyclKernelBundle _create(DPCTLSyclKernelBundleRef KBRef): + cdef SyclKernelBundle _create(DPCTLSyclKernelBundleRef KBRef, + bint is_sycl_source): cdef SyclKernelBundle ret = SyclKernelBundle.__new__(SyclKernelBundle) ret._kernel_bundle_ref = KBRef + ret._is_sycl_source = is_sycl_source return ret def __dealloc__(self): @@ -212,6 +233,13 @@ cdef class SyclKernelBundle: 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._kernel_bundle_ref, name + ), + kernel_name + ) return SyclKernel._create( DPCTLKernelBundle_GetKernel(self._kernel_bundle_ref, name), kernel_name @@ -219,6 +247,10 @@ cdef class SyclKernelBundle: def has_sycl_kernel(self, str kernel_name): name = kernel_name.encode("utf8") + if self._is_sycl_source: + return DPCTLKernelBundle_HasSyclKernel( + self._kernel_bundle_ref, name + ) return DPCTLKernelBundle_HasKernel(self._kernel_bundle_ref, name) def addressof_ref(self): @@ -249,7 +281,7 @@ cdef api SyclKernelBundle SyclKernelBundle_Make(DPCTLSyclKernelBundleRef KBRef): reference. """ cdef DPCTLSyclKernelBundleRef copied_KBRef = DPCTLKernelBundle_Copy(KBRef) - return SyclKernelBundle._create(copied_KBRef) + return SyclKernelBundle._create(copied_KBRef, False) cpdef create_kernel_bundle_from_source(SyclQueue q, str src, str copts=""): @@ -295,7 +327,7 @@ cpdef create_kernel_bundle_from_source(SyclQueue q, str src, str copts=""): if KBref is NULL: raise SyclKernelBundleCompilationError() - return SyclKernelBundle._create(KBref) + return SyclKernelBundle._create(KBref, False) cpdef create_kernel_bundle_from_spirv( @@ -342,7 +374,121 @@ cpdef create_kernel_bundle_from_spirv( if KBref is NULL: raise SyclKernelBundleCompilationError() - return SyclKernelBundle._create(KBref) + return SyclKernelBundle._create(KBref, False) + + +cpdef create_kernel_bundle_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. + + 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:`.SyclKernelBundle` 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: + kernel_bundle (:class:`.SyclKernelBundle`) + A :class:`.SyclKernelBundle` object wrapping the + ``sycl::kernel_bundle`` + returned by the C API. + + Raises: + SyclKernelBundleCompilationError + If a SYCL kernel bundle could not be created. The exception + message contains the build log for more details. + """ + 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 + cdef const char* buildLogContent + for opt in copts: + if not isinstance(opt, unicode): + DPCTLBuildOptionList_Delete(BuildOpts) + raise SyclKernelBundleCompilationError() + 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 SyclKernelBundleCompilationError() + bName = name.encode("utf8") + sName = bName + DPCTLKernelNameList_Append(KernelNames, sName) + + cdef DPCTLVirtualHeaderListRef VirtualHeaders + 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 SyclKernelBundleCompilationError() + bName = name.encode("utf8") + sName = bName + bContent = content.encode("utf8") + sContent = bContent + DPCTLVirtualHeaderList_Append(VirtualHeaders, sName, sContent) + + cdef DPCTLKernelBuildLogRef BuildLog + BuildLog = DPCTLKernelBuildLog_Create() + + KBref = DPCTLKernelBundle_CreateFromSYCLSource(CRef, DRef, Src, + VirtualHeaders, KernelNames, + BuildOpts, BuildLog) + + if KBref is NULL: + buildLogContent = DPCTLKernelBuildLog_Get(BuildLog) + buildLogStr = str(buildLogContent, "utf-8") + DPCTLBuildOptionList_Delete(BuildOpts) + DPCTLKernelNameList_Delete(KernelNames) + DPCTLVirtualHeaderList_Delete(VirtualHeaders) + DPCTLKernelBuildLog_Delete(BuildLog) + raise SyclKernelBundleCompilationError(buildLogStr) + + DPCTLBuildOptionList_Delete(BuildOpts) + DPCTLKernelNameList_Delete(KernelNames) + DPCTLVirtualHeaderList_Delete(VirtualHeaders) + DPCTLKernelBuildLog_Delete(BuildLog) + + return SyclKernelBundle._create(KBref, True) cpdef create_program_from_source(SyclQueue q, str src, str copts=""): diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index 1c09adc28e..4dddc0794b 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -80,8 +80,7 @@ def _check_cpython_api_SyclKernelBundle_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): @@ -186,7 +185,9 @@ def _check_multi_kernel_program(kb): assert type(cmsgsz) is int _check_cpython_api_SyclKernelBundle_GetKernelBundleRef(kb) - _check_cpython_api_SyclKernelBundle_Make(kb) + p2 = _check_cpython_api_SyclKernelBundle_Make(kb) + assert p2.has_sycl_kernel("add") + assert p2.has_sycl_kernel("axpy") def test_create_kernel_bundle_from_source_ocl(): @@ -262,3 +263,150 @@ def test_create_kernel_bundle_from_invalid_src_ocl(): }" with pytest.raises(dpctl_prog.SyclKernelBundleCompilationError): dpctl_prog.create_kernel_bundle_from_source(q, invalid_oclSrc) + + +def test_create_kernel_bundle_from_sycl_source(): + try: + q = dpctl.SyclQueue("opencl") + except dpctl.SyclQueueCreationError: + pytest.skip("No OpenCL 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_kernel_bundle_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.SyclKernelBundle + + assert type(prog.addressof_ref()) is int + assert prog.has_sycl_kernel("vector_add") + regularKernel = prog.get_sycl_kernel("vector_add") + + # 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 + 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_SyclKernelBundle_GetKernelBundleRef(prog) + + +def test_create_kernel_bundle_from_invalid_src_sycl(): + try: + q = dpctl.SyclQueue("opencl") + except dpctl.SyclQueueCreationError: + pytest.skip("No OpenCL queue is available") + + if not q.get_sycl_device().can_compile("sycl"): + pytest.skip("SYCL source compilation not supported") + + sycl_source = """ + #include + + 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] = in1[globalID] + in2[globalID]; + } + """ + try: + _ = dpctl.program.create_kernel_bundle_from_sycl_source( + q, + sycl_source, + headers=[], + registered_names=[], + copts=[], + ) + assert False + except dpctl_prog.SyclKernelBundleCompilationError as prog_error: + print(str(prog_error)) + assert "error: expected ';' at end of declaration" in str(prog_error) diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_device_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_device_interface.h index 72b0261e1f..95e30ac03f 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_device_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_device_interface.h @@ -828,4 +828,40 @@ DPCTL_API void DPCTLDevice_DisablePeerAccess(__dpctl_keep const DPCTLSyclDeviceRef DRef, __dpctl_keep const DPCTLSyclDeviceRef PDRef); +/*! + * @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 07a76c3fd8..d0f270db26 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h @@ -129,4 +129,173 @@ DPCTL_API __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_Copy(__dpctl_keep const DPCTLSyclKernelBundleRef KBRef); +typedef struct DPCTLBuildOptionList *DPCTLBuildOptionListRef; +typedef struct DPCTLKernelNameList *DPCTLKernelNameListRef; +typedef struct DPCTLVirtualHeaderList *DPCTLVirtualHeaderListRef; +typedef struct DPCTLKernelBuildLog *DPCTLKernelBuildLogRef; + +/*! + * @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 Ref 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 Ref 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 Ref 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 an empty kernel build log. + * + * @return Opaque pointer to the kernel build log. + * @ingroup KernelBundleInterface + */ +DPCTL_API __dpctl_give DPCTLKernelBuildLogRef DPCTLKernelBuildLog_Create(); + +/*! + * @brief Frees the DPCTLKernelBuildLogRef pointer. + * + * @param Ref Opaque pointer to a kernel build log. + * @ingroup KernelBundleInterface + */ +DPCTL_API +void DPCTLKernelBuildLog_Delete(__dpctl_take DPCTLKernelBuildLogRef Ref); + +/*! + * @brief Get the content of the build log. + * + * @param Ref Opaque pointer to the kernel build log. + * @return Content of the build log + * @ingroup KernelBundleInterface + */ +DPCTL_API const char * +DPCTLKernelBuildLog_Get(__dpctl_keep DPCTLKernelBuildLogRef); + +/*! + * @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, + __dpctl_keep DPCTLKernelBuildLogRef BuildLog); + +/*! + * @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 1378f6f818..4bf5c6293b 100644 --- a/libsyclinterface/source/dpctl_sycl_device_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_device_interface.cpp @@ -982,3 +982,28 @@ void DPCTLDevice_DisablePeerAccess(__dpctl_keep const DPCTLSyclDeviceRef DRef, } return; } + +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 78c714ecbb..94a0e3099e 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -761,3 +761,250 @@ DPCTLKernelBundle_Copy(__dpctl_keep const DPCTLSyclKernelBundleRef KBRef) return nullptr; } } + +using build_option_list_t = std::vector; + +__dpctl_give DPCTLBuildOptionListRef DPCTLBuildOptionList_Create() +{ + 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 reinterpret_cast(Ref); +} + +void DPCTLBuildOptionList_Append(__dpctl_keep DPCTLBuildOptionListRef Ref, + __dpctl_keep const char *Option) +{ + reinterpret_cast(Ref)->emplace_back(Option); +} + +using kernel_name_list_t = std::vector; + +__dpctl_give DPCTLKernelNameListRef DPCTLKernelNameList_Create() +{ + 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 reinterpret_cast(Ref); +} + +void DPCTLKernelNameList_Append(__dpctl_keep DPCTLKernelNameListRef Ref, + __dpctl_keep const char *Option) +{ + reinterpret_cast(Ref)->emplace_back(Option); +} + +using virtual_header_list_t = std::vector>; + +__dpctl_give DPCTLVirtualHeaderListRef DPCTLVirtualHeaderList_Create() +{ + 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 reinterpret_cast(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); + reinterpret_cast(Ref)->push_back(Header); +} + +using kernel_build_log_t = std::string; + +__dpctl_give DPCTLKernelBuildLogRef DPCTLKernelBuildLog_Create() +{ + auto BuildLog = + std::unique_ptr(new kernel_build_log_t("")); + auto *RetVal = reinterpret_cast(BuildLog.get()); + BuildLog.release(); + return RetVal; +} + +void DPCTLKernelBuildLog_Delete(__dpctl_take DPCTLKernelBuildLogRef Ref) +{ + delete reinterpret_cast(Ref); +} + +const char *DPCTLKernelBuildLog_Get(__dpctl_keep DPCTLKernelBuildLogRef Ref) +{ + return reinterpret_cast(Ref)->data(); +} + +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) +// 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( + __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, + __dpctl_keep DPCTLKernelBuildLogRef BuildLog) +{ +#if (SUPPORTS_SYCL_COMPILATION > 0) + context *SyclCtx = unwrap(Ctx); + device *SyclDev = unwrap(Dev); + if (!SyclDev->ext_oneapi_can_compile(syclex::source_language::sycl)) { + return nullptr; + } + try { + auto *IncludeFileList = + reinterpret_cast(Headers); + std::unique_ptr> + SrcBundle; + std::string Src(Source); + // 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 : + *reinterpret_cast(Names)) + { + RegisteredNames.add(Name); + } + + syclex::build_options Opts{ + *reinterpret_cast(BuildOptions)}; + + 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) { + auto *RawBuildLog = reinterpret_cast(BuildLog); + *RawBuildLog = e.what(); + return nullptr; + } +#else + return nullptr; +#endif +} + +__dpctl_give DPCTLSyclKernelRef +DPCTLKernelBundle_GetSyclKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, + __dpctl_keep const char *KernelName) +{ +#if (SUPPORTS_SYCL_COMPILATION > 0) + 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) +{ +#if (SUPPORTS_SYCL_COMPILATION > 0) + 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 a835c277b9..3c41662f39 100644 --- a/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp @@ -273,6 +273,132 @@ 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); + DPCTLKernelBuildLogRef KBLRef = DPCTLKernelBuildLog_Create(); + KBRef = DPCTLKernelBundle_CreateFromSYCLSource( + CRef, DRef, sycl_source, VHRef, KNRef, BORef, KBLRef); + DPCTLVirtualHeaderList_Delete(VHRef); + DPCTLKernelNameList_Delete(KNRef); + DPCTLBuildOptionList_Delete(BORef); + DPCTLKernelBuildLog_Delete(KBLRef); + } + } + + 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")); + // 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, "_Z33__sycl_kernel_vector_add_templateIiEvPT_S1_S1_")); +} + +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); + DPCTLKernel_Delete(AxpyKernel); +} + INSTANTIATE_TEST_SUITE_P(KernelBundleCreationFromSpirv, TestDPCTLSyclKernelBundleInterface, ::testing::Values("opencl", @@ -289,6 +415,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 3b9321e3fcd0bd612d86c336e6cdd977a95998ff Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Wed, 11 Feb 2026 20:56:53 -0800 Subject: [PATCH 2/7] adds utility to dpctl.program --- dpctl/_backend.pxd | 2 ++ dpctl/program/__init__.py | 2 ++ dpctl/program/_program.pyx | 13 +++++++++++++ .../source/dpctl_sycl_kernel_bundle_interface.cpp | 9 +++++++++ 4 files changed, 26 insertions(+) diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index 069f9a33fc..951743110b 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -484,6 +484,8 @@ cdef extern from "syclinterface/dpctl_sycl_kernel_bundle_interface.h": cdef void DPCTLKernelBuildLog_Delete(DPCTLKernelBuildLogRef Ref) cdef const char *DPCTLKernelBuildLog_Get(DPCTLKernelBuildLogRef) + cdef bool DPCTLKernelBundle_CreateFromSYCLSource_Available() + cdef DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource( const DPCTLSyclContextRef Ctx, const DPCTLSyclDeviceRef Dev, diff --git a/dpctl/program/__init__.py b/dpctl/program/__init__.py index d47f7ac2d3..2d447c2fef 100644 --- a/dpctl/program/__init__.py +++ b/dpctl/program/__init__.py @@ -30,6 +30,7 @@ create_kernel_bundle_from_sycl_source, create_program_from_source, create_program_from_spirv, + is_sycl_source_compilation_available, ) __all__ = [ @@ -38,6 +39,7 @@ "create_program_from_source", "create_program_from_spirv", "create_kernel_bundle_from_sycl_source", + "is_sycl_source_compilation_available", "SyclKernel", "SyclKernelBundle", "SyclKernelBundleCompilationError", diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index 84dd10fc32..f7ffc8ffe0 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -53,6 +53,7 @@ from dpctl._backend cimport ( # noqa: E211, E402; DPCTLKernelBundle_CreateFromOCLSource, DPCTLKernelBundle_CreateFromSpirv, DPCTLKernelBundle_CreateFromSYCLSource, + DPCTLKernelBundle_CreateFromSYCLSource_Available, DPCTLKernelBundle_Delete, DPCTLKernelBundle_GetKernel, DPCTLKernelBundle_GetSyclKernel, @@ -75,6 +76,7 @@ from dpctl._backend cimport ( # noqa: E211, E402; __all__ = [ "create_kernel_bundle_from_source", "create_kernel_bundle_from_spirv", + "is_sycl_source_compilation_available", "SyclKernel", "SyclKernelBundle", "SyclKernelBundleCompilationError", @@ -87,6 +89,17 @@ cdef class SyclKernelBundleCompilationError(Exception): pass +cpdef bint is_sycl_source_compilation_available(): + """Returns True if dpctl was built with compiler that supports the DPC++ + `kernel_compiler` extension API used by + :func:`create_kernel_bundle_from_sycl_source`. + + Device support is separate; callers should also check + ``q.sycl_device.can_compile('sycl')`` (or similar) for specific devices. + """ + return DPCTLKernelBundle_CreateFromSYCLSource_Available() + + cdef class SyclKernel: """ """ diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp index 94a0e3099e..685c6cfa4e 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -866,6 +866,15 @@ namespace syclex = sycl::ext::oneapi::experimental; #endif #endif +bool DPCTLKernelBundle_CreateFromSYCLSource_Available() +{ +#if (SUPPORTS_SYCL_COMPILATION > 0) + return true; +#else + return false; +#endif +} + #if (SUPPORTS_SYCL_COMPILATION > 0) // 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 From c521084030b56b86594d2e440a3cc67750fc2f1b Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Wed, 8 Apr 2026 02:38:53 -0700 Subject: [PATCH 3/7] adds tests for SYCL source compilation --- dpctl/tests/test_sycl_program.py | 106 +++++++++++++++++++++++++++---- 1 file changed, 92 insertions(+), 14 deletions(-) diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index 4dddc0794b..b05c07af0b 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -18,12 +18,28 @@ import os +import numpy as np import pytest import dpctl +import dpctl.memory as dpm import dpctl.program as dpctl_prog +def _get_opencl_queue_or_skip(): + try: + return dpctl.SyclQueue("opencl") + except dpctl.SyclQueueCreationError: + pytest.skip("No OpenCL queue is available") + + +def _skip_if_no_sycl_source_compilation(q): + if not dpctl.program.is_sycl_source_compilation_available(): + pytest.skip("SYCL source compilation extension not available") + if not q.get_sycl_device().can_compile("sycl"): + pytest.skip("SYCL source compilation not supported") + + def get_spirv_abspath(fn): curr_dir = os.path.dirname(os.path.abspath(__file__)) spirv_file = os.path.join(curr_dir, "input_files", fn) @@ -266,13 +282,8 @@ def test_create_kernel_bundle_from_invalid_src_ocl(): def test_create_kernel_bundle_from_sycl_source(): - try: - q = dpctl.SyclQueue("opencl") - except dpctl.SyclQueueCreationError: - pytest.skip("No OpenCL queue is available") - - if not q.get_sycl_device().can_compile("sycl"): - pytest.skip("SYCL source compilation not supported") + q = _get_opencl_queue_or_skip() + _skip_if_no_sycl_source_compilation(q) sycl_source = """ #include @@ -376,13 +387,8 @@ def test_create_kernel_bundle_from_sycl_source(): def test_create_kernel_bundle_from_invalid_src_sycl(): - try: - q = dpctl.SyclQueue("opencl") - except dpctl.SyclQueueCreationError: - pytest.skip("No OpenCL queue is available") - - if not q.get_sycl_device().can_compile("sycl"): - pytest.skip("SYCL source compilation not supported") + q = _get_opencl_queue_or_skip() + _skip_if_no_sycl_source_compilation(q) sycl_source = """ #include @@ -410,3 +416,75 @@ def test_create_kernel_bundle_from_invalid_src_sycl(): except dpctl_prog.SyclKernelBundleCompilationError as prog_error: print(str(prog_error)) assert "error: expected ';' at end of declaration" in str(prog_error) + + +def test_sycl_source_compilation_is_available_returns_bool(): + v = dpctl.program.is_sycl_source_compilation_available() + assert type(v) is bool + + +def test_sycl_source_vector_add_correctness(): + q = _get_opencl_queue_or_skip() + _skip_if_no_sycl_source_compilation(q) + + sycl_source = """ + #include + #include "math_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]); + } + """ + + header_content = """ + int math_op(int a, int b){ + return a + b; + } + """ + + prog = dpctl.program.create_kernel_bundle_from_sycl_source( + q, + sycl_source, + headers=[("math_ops.hpp", header_content)], + registered_names=[], + copts=["-fno-fast-math"], + ) + + kernel = prog.get_sycl_kernel("vector_add") + + local_size = 16 + global_size = local_size * 8 + + in1 = np.arange(global_size, dtype=np.int32) + in2 = (np.arange(global_size, dtype=np.int32) * 3 - 7).astype(np.int32) + out = np.empty(global_size, dtype=np.int32) + expected = (in1 + in2).astype(np.int32) + + in1_usm = dpm.MemoryUSMDevice(in1.nbytes, queue=q) + in2_usm = dpm.MemoryUSMDevice(in2.nbytes, queue=q) + out_usm = dpm.MemoryUSMDevice(out.nbytes, queue=q) + + ev1 = q.memcpy_async(dest=in1_usm, src=in1, count=in1.nbytes) + ev2 = q.memcpy_async(dest=in2_usm, src=in2, count=in2.nbytes) + + try: + ev3 = q.submit( + kernel, + [in1_usm, in2_usm, out_usm], + [global_size], + [local_size], + dEvents=[ev1, ev2], + ) + except dpctl._sycl_queue.SyclKernelSubmitError: + pytest.skip(f"Kernel submission to {q.sycl_device} failed") + + ev4 = q.memcpy_async(dest=out, src=out_usm, count=out.nbytes, dEvents=[ev3]) + ev4.wait() + assert np.array_equal(out, expected) From 0f40fa025cca42cef57a9515bfdb982c16f7dcf3 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Wed, 8 Apr 2026 02:39:13 -0700 Subject: [PATCH 4/7] add bool DPCTLKernelBundle_CreateFromSYCLSource_Available to header --- .../syclinterface/dpctl_sycl_kernel_bundle_interface.h | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h index d0f270db26..a901e5f8e2 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h @@ -246,6 +246,16 @@ void DPCTLKernelBuildLog_Delete(__dpctl_take DPCTLKernelBuildLogRef Ref); DPCTL_API const char * DPCTLKernelBuildLog_Get(__dpctl_keep DPCTLKernelBuildLogRef); +/*! + * @brief Return True if the DPCTLKernelBundle_CreateFromSYCLSource function is + * available, else False. + * + * @ingroup KernelBundleInterface + */ + +DPCTL_API +bool DPCTLKernelBundle_CreateFromSYCLSource_Available(); + /*! * @brief Create a SYCL kernel bundle from an SYCL kernel source string. * From 9e0b72a41448fb219ae030e5c1f94e81e921ea3a Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Tue, 26 May 2026 07:54:26 -0700 Subject: [PATCH 5/7] add SYCL source creation function to __all__ --- dpctl/program/_program.pyx | 1 + 1 file changed, 1 insertion(+) diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index f7ffc8ffe0..ddde4e124b 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -76,6 +76,7 @@ from dpctl._backend cimport ( # noqa: E211, E402; __all__ = [ "create_kernel_bundle_from_source", "create_kernel_bundle_from_spirv", + "create_kernel_bundle_from_sycl_source", "is_sycl_source_compilation_available", "SyclKernel", "SyclKernelBundle", From ece9fa0d7fbe803bc114eb39d9385ed24e3b82d3 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Tue, 26 May 2026 08:26:01 -0700 Subject: [PATCH 6/7] add create_kernel_bundle_from_sycl_source to docs --- docs/doc_sources/api_reference/dpctl/program.rst | 1 + 1 file changed, 1 insertion(+) diff --git a/docs/doc_sources/api_reference/dpctl/program.rst b/docs/doc_sources/api_reference/dpctl/program.rst index aee4b574a8..8e05ea42f5 100644 --- a/docs/doc_sources/api_reference/dpctl/program.rst +++ b/docs/doc_sources/api_reference/dpctl/program.rst @@ -22,6 +22,7 @@ execution via :py:meth:`dpctl.SyclQueue.submit`. create_kernel_bundle_from_source create_kernel_bundle_from_spirv + create_kernel_bundle_from_sycl_source create_program_from_source create_program_from_spirv From 8d662b5ee04d12c48073160df421791b2a4ac31b Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Tue, 26 May 2026 08:37:11 -0700 Subject: [PATCH 7/7] run SYCL source compilation tests on Level Zero if possible --- dpctl/tests/test_sycl_program.py | 55 ++++++++++++++++---------------- 1 file changed, 28 insertions(+), 27 deletions(-) diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index b05c07af0b..aa8ec1310b 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -33,6 +33,13 @@ def _get_opencl_queue_or_skip(): pytest.skip("No OpenCL queue is available") +def _get_level_zero_queue_or_skip(): + try: + return dpctl.SyclQueue("level_zero") + except dpctl.SyclQueueCreationError: + pytest.skip("No Level Zero queue is available") + + def _skip_if_no_sycl_source_compilation(q): if not dpctl.program.is_sycl_source_compilation_available(): pytest.skip("SYCL source compilation extension not available") @@ -216,19 +223,13 @@ def test_create_kernel_bundle_from_source_ocl(): size_t index = get_global_id(0); \ c[index] = a[index] + d*b[index]; \ }" - try: - q = dpctl.SyclQueue("opencl") - except dpctl.SyclQueueCreationError: - pytest.skip("No OpenCL queue is available") + q = _get_opencl_queue_or_skip() kb = dpctl_prog.create_kernel_bundle_from_source(q, oclSrc) _check_multi_kernel_program(kb) def test_create_kernel_bundle_from_spirv_ocl(): - try: - q = dpctl.SyclQueue("opencl") - except dpctl.SyclQueueCreationError: - pytest.skip("No OpenCL queue is available") + q = _get_opencl_queue_or_skip() spirv_file = get_spirv_abspath("multi_kernel.spv") with open(spirv_file, "rb") as fin: spirv = fin.read() @@ -237,10 +238,7 @@ def test_create_kernel_bundle_from_spirv_ocl(): def test_create_kernel_bundle_from_spirv_l0(): - try: - q = dpctl.SyclQueue("level_zero") - except dpctl.SyclQueueCreationError: - pytest.skip("No Level-zero queue is available") + q = _get_level_zero_queue_or_skip() spirv_file = get_spirv_abspath("multi_kernel.spv") with open(spirv_file, "rb") as fin: spirv = fin.read() @@ -249,13 +247,10 @@ def test_create_kernel_bundle_from_spirv_l0(): @pytest.mark.xfail( - reason="Level-zero backend does not support compilation from source" + reason="Level Zero backend does not support compilation from source" ) def test_create_kernel_bundle_from_source_l0(): - try: - q = dpctl.SyclQueue("level_zero") - except dpctl.SyclQueueCreationError: - pytest.skip("No Level-zero queue is available") + q = _get_level_zero_queue_or_skip() oclSrc = " \ kernel void add(global int* a, global int* b, global int* c) { \ size_t index = get_global_id(0); \ @@ -270,10 +265,7 @@ def test_create_kernel_bundle_from_source_l0(): def test_create_kernel_bundle_from_invalid_src_ocl(): - try: - q = dpctl.SyclQueue("opencl") - except dpctl.SyclQueueCreationError: - pytest.skip("No OpenCL queue is available") + q = _get_opencl_queue_or_skip() invalid_oclSrc = " \ kernel void add( \ }" @@ -281,8 +273,11 @@ def test_create_kernel_bundle_from_invalid_src_ocl(): dpctl_prog.create_kernel_bundle_from_source(q, invalid_oclSrc) -def test_create_kernel_bundle_from_sycl_source(): - q = _get_opencl_queue_or_skip() +@pytest.mark.parametrize( + "queue_selector", [_get_opencl_queue_or_skip, _get_level_zero_queue_or_skip] +) +def test_create_kernel_bundle_from_sycl_source(queue_selector): + q = queue_selector() _skip_if_no_sycl_source_compilation(q) sycl_source = """ @@ -386,8 +381,11 @@ def test_create_kernel_bundle_from_sycl_source(): _check_cpython_api_SyclKernelBundle_GetKernelBundleRef(prog) -def test_create_kernel_bundle_from_invalid_src_sycl(): - q = _get_opencl_queue_or_skip() +@pytest.mark.parametrize( + "queue_selector", [_get_opencl_queue_or_skip, _get_level_zero_queue_or_skip] +) +def test_create_kernel_bundle_from_invalid_src_sycl(queue_selector): + q = queue_selector() _skip_if_no_sycl_source_compilation(q) sycl_source = """ @@ -423,8 +421,11 @@ def test_sycl_source_compilation_is_available_returns_bool(): assert type(v) is bool -def test_sycl_source_vector_add_correctness(): - q = _get_opencl_queue_or_skip() +@pytest.mark.parametrize( + "queue_selector", [_get_opencl_queue_or_skip, _get_level_zero_queue_or_skip] +) +def test_sycl_source_vector_add_correctness(queue_selector): + q = queue_selector() _skip_if_no_sycl_source_compilation(q) sycl_source = """