From d1ef4becd25088ff037da78054095b4f0a7baa8c Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Thu, 10 Apr 2025 10:39:23 +0100 Subject: [PATCH 01/13] 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 452ff8f255..03621ed612 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -288,9 +288,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( @@ -459,6 +462,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 69ca6d0887..55d14f93f1 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 506664ab18..b4dc360a92 100644 --- a/dpctl/_sycl_device.pyx +++ b/dpctl/_sycl_device.pyx @@ -27,6 +27,9 @@ from ._backend cimport ( # noqa: E211 DPCTLDefaultSelector_Create, DPCTLDevice_AreEq, DPCTLDevice_CanAccessPeer, + DPCTLDevice_CanCompileOpenCL, + DPCTLDevice_CanCompileSPIRV, + DPCTLDevice_CanCompileSYCL, DPCTLDevice_Copy, DPCTLDevice_CreateFromSelector, DPCTLDevice_CreateSubDevicesByAffinity, @@ -2376,6 +2379,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 4eedbee7a8..f3f76261df 100644 --- a/dpctl/program/__init__.py +++ b/dpctl/program/__init__.py @@ -29,6 +29,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, ) @@ -38,6 +39,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 de75753495..f2feb6beca 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) @@ -72,3 +74,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 1c9227ec02..269dfe3700 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -43,6 +43,10 @@ from libc.string cimport memcmp, memcpy import warnings from dpctl._backend cimport ( # noqa: E211, E402; + DPCTLBuildOptionList_Append, + DPCTLBuildOptionList_Create, + DPCTLBuildOptionList_Delete, + DPCTLBuildOptionListRef, DPCTLKernel_Copy, DPCTLKernel_Delete, DPCTLKernel_GetCompileNumSubGroups, @@ -53,16 +57,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, _spec_const, ) @@ -217,9 +236,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): @@ -230,6 +251,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 @@ -237,6 +265,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): @@ -267,7 +299,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) cdef class SpecializationConstant: @@ -518,7 +550,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( @@ -604,7 +636,121 @@ cpdef create_kernel_bundle_from_spirv( if spconsts != NULL: free(spconsts) - 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 bbf16df98c..ecc4d5dbc9 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -82,8 +82,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): @@ -188,7 +187,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(): @@ -412,3 +413,150 @@ def test_spec_const_equality(): assert sp1 == sp2 assert sp1 != sp3 assert sp1 != sp4 + + +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 c07a0f8f46..ceb4e33aa3 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 7b8a11c06d..7c131b7a2b 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h @@ -140,4 +140,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 65be01a0bb..fdda1d5740 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 ec94e5a077..3f10d88319 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -872,3 +872,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 7917157c28..23c81aad3b 100644 --- a/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp @@ -277,6 +277,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", @@ -293,6 +419,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 f7e8163fa2f298636cfd597f0347f0fca4155ad9 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Wed, 11 Feb 2026 20:56:53 -0800 Subject: [PATCH 02/13] 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 03621ed612..5dfd8db4b4 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -491,6 +491,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 f3f76261df..0c09226f99 100644 --- a/dpctl/program/__init__.py +++ b/dpctl/program/__init__.py @@ -32,6 +32,7 @@ create_kernel_bundle_from_sycl_source, create_program_from_source, create_program_from_spirv, + is_sycl_source_compilation_available, ) __all__ = [ @@ -40,6 +41,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 269dfe3700..19ee9771da 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -65,6 +65,7 @@ from dpctl._backend cimport ( # noqa: E211, E402; DPCTLKernelBundle_CreateFromOCLSource, DPCTLKernelBundle_CreateFromSpirv, DPCTLKernelBundle_CreateFromSYCLSource, + DPCTLKernelBundle_CreateFromSYCLSource_Available, DPCTLKernelBundle_Delete, DPCTLKernelBundle_GetKernel, DPCTLKernelBundle_GetSyclKernel, @@ -92,6 +93,7 @@ import numpy as np __all__ = [ "create_kernel_bundle_from_source", "create_kernel_bundle_from_spirv", + "is_sycl_source_compilation_available", "SyclKernel", "SyclKernelBundle", "SyclKernelBundleCompilationError", @@ -105,6 +107,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 3f10d88319..6008fca48d 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -977,6 +977,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 44ee08b68f460acad2b5e10e440b085cb6f1b8f8 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Wed, 8 Apr 2026 02:38:53 -0700 Subject: [PATCH 03/13] adds tests for SYCL source compilation --- dpctl/tests/test_sycl_program.py | 105 ++++++++++++++++++++++++++----- 1 file changed, 91 insertions(+), 14 deletions(-) diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index ecc4d5dbc9..cb0589f0f4 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -22,10 +22,25 @@ import pytest import dpctl +import dpctl.memory as dpm import dpctl.program as dpctl_prog from dpctl.program.utils import parse_spirv_specializations +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) @@ -416,13 +431,8 @@ def test_spec_const_equality(): 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 @@ -526,13 +536,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 @@ -560,3 +565,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 705a859d5f51f9e115c31322efbe8ca09bfc44f9 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Wed, 8 Apr 2026 02:39:13 -0700 Subject: [PATCH 04/13] 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 7c131b7a2b..aee14b8deb 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h @@ -257,6 +257,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 0082c9f89be83ab5140bb6ecd46e4b5d10719b97 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Tue, 26 May 2026 07:54:26 -0700 Subject: [PATCH 05/13] 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 19ee9771da..6c5b3f421d 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -93,6 +93,7 @@ import numpy as np __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 a892c8f1cd8c923790a26a6fadae8a7ece35f281 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Tue, 26 May 2026 08:26:01 -0700 Subject: [PATCH 06/13] 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 6b2bf22bad..a37aec5531 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 76bb60d1957800c5a47a9fd3dbb2e9618f9e8bb1 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Tue, 26 May 2026 08:37:11 -0700 Subject: [PATCH 07/13] 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 cb0589f0f4..48fdfb7a59 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -34,6 +34,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") @@ -217,19 +224,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() @@ -238,10 +239,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() @@ -250,13 +248,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); \ @@ -271,10 +266,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( \ }" @@ -430,8 +422,11 @@ def test_spec_const_equality(): assert sp1 != sp4 -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 = """ @@ -535,8 +530,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 = """ @@ -572,8 +570,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 = """ From 0a7524b9fad7346c85e208ad7046120c16d02175 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Wed, 5 Aug 2026 11:37:42 -0700 Subject: [PATCH 08/13] fix default values --- dpctl/program/_program.pyx | 19 ++++++++++++++----- dpctl/tests/test_sycl_program.py | 28 ++++++++++++++++++++++++++++ 2 files changed, 42 insertions(+), 5 deletions(-) diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index 6c5b3f421d..453d7e26bf 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -671,20 +671,21 @@ cpdef create_kernel_bundle_from_sycl_source(SyclQueue q, :class:`.SyclKernelBundle` is going to be built. source (unicode) SYCL source code string. - headers (list) + headers (list, optional) 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: [] + Default: ``None`` 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) + Default: ``None`` + copts (list, optional) Optional list of compilation flags that will be used - when compiling the program. Default: ``""``. + when compiling the program. + Default: ``None`` Returns: kernel_bundle (:class:`.SyclKernelBundle`) @@ -710,6 +711,14 @@ cpdef create_kernel_bundle_from_sycl_source(SyclQueue q, cdef bytes bContent cdef const char* sContent cdef const char* buildLogContent + + if headers is None: + headers = [] + if registered_names is None: + registered_names = [] + if copts is None: + copts = [] + for opt in copts: if not isinstance(opt, unicode): DPCTLBuildOptionList_Delete(BuildOpts) diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index 48fdfb7a59..c045b2aafe 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -570,6 +570,34 @@ def test_sycl_source_compilation_is_available_returns_bool(): assert type(v) is bool +@pytest.mark.parametrize( + "queue_selector", [_get_opencl_queue_or_skip, _get_level_zero_queue_or_skip] +) +def test_create_kernel_bundle_from_sycl_source_defaults(queue_selector): + q = queue_selector() + _skip_if_no_sycl_source_compilation(q) + + 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]; + } + """ + + prog = dpctl.program.create_kernel_bundle_from_sycl_source(q, sycl_source) + + assert type(prog) is dpctl_prog.SyclKernelBundle + assert prog.has_sycl_kernel("vector_add") + + @pytest.mark.parametrize( "queue_selector", [_get_opencl_queue_or_skip, _get_level_zero_queue_or_skip] ) From 8a8358caac3f048649cec5158559df7f7196b3b0 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Wed, 5 Aug 2026 14:03:47 -0700 Subject: [PATCH 09/13] handle null arguments in source compilation APIs --- .../dpctl_sycl_kernel_bundle_interface.h | 20 ++- .../dpctl_sycl_kernel_bundle_interface.cpp | 81 +++++++++--- .../test_sycl_kernel_bundle_interface.cpp | 117 ++++++++++++++++++ 3 files changed, 196 insertions(+), 22 deletions(-) diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h index aee14b8deb..b65303d094 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h @@ -57,11 +57,13 @@ typedef struct DPCTLSpecConstTy * @param IL SPIR-V binary * @param Length The size of the IL binary in bytes. * @param CompileOpts Optional compiler flags used when compiling the - * SPIR-V binary. + * SPIR-V binary, or NULL for none. * @param NumSpecConsts The number of specialization constants. - * @param SpecConsts An array of specialization constants. + * @param SpecConsts An array of specialization constants, or NULL if + * NumSpecConsts is 0. * @return A new SyclKernelBundleRef pointer if the kernel_bundle creation * succeeded, else returns NULL. + * * @ingroup KernelBundleInterface */ DPCTL_API @@ -80,9 +82,11 @@ __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSpirv( * @param Ctx An opaque pointer to a sycl::context * @param Dev An opaque pointer to a sycl::device * @param Source OpenCL source string - * @param CompileOpts Extra compiler flags (refer Sycl spec.) + * @param CompileOpts Extra compiler flags (refer Sycl spec.), or NULL + * for none * @return A new SyclKernelBundleRef pointer if the kernel bundle creation * succeeded, else returns NULL. + * * @ingroup KernelBundleInterface */ DPCTL_API @@ -273,11 +277,15 @@ bool DPCTLKernelBundle_CreateFromSYCLSource_Available(); * @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.) + * @param Headers List of virtual headers, or NULL for none + * @param Names List of kernel names to register, or NULL for none + * @param BuildOptions List of extra compiler flags (refer Sycl spec.), + * or NULL for none + * @param BuildLog Build log to write diagnostics to, or NULL to + * discard the log * @return A new SyclKernelBundleRef pointer if the program creation * succeeded, else returns NULL. + * * @ingroup KernelBundleInterface */ DPCTL_API diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp index 6008fca48d..c984123628 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -1025,14 +1025,48 @@ __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource( __dpctl_keep DPCTLKernelBuildLogRef BuildLog) { #if (SUPPORTS_SYCL_COMPILATION > 0) + // optional arguments and BuildLog may be NULL, which is equivalent to + // passing an empty list or discarding the log + // Writing to the log is routed through a helper + auto *RawBuildLog = reinterpret_cast(BuildLog); + auto set_build_log = [RawBuildLog](const std::string &Msg) { + if (RawBuildLog) + *RawBuildLog = Msg; + }; + auto fail = [&set_build_log](const char *Msg, int Line) { + set_build_log(Msg); + error_handler(Msg, __FILE__, "DPCTLKernelBundle_CreateFromSYCLSource", + Line); + return nullptr; + }; + + if (!Ctx) + return fail("Input Ctx is nullptr.", __LINE__); + if (!Dev) + return fail("Input Dev is nullptr.", __LINE__); + if (!Source) + return fail("Input Source is nullptr.", __LINE__); + + static const virtual_header_list_t EmptyHeaders{}; + static const kernel_name_list_t EmptyNames{}; + static const build_option_list_t EmptyOptions{}; + + const auto &IncludeFiles = + Headers ? *reinterpret_cast(Headers) + : EmptyHeaders; + const auto &KernelNames = + Names ? *reinterpret_cast(Names) : EmptyNames; + const auto &Options = + BuildOptions ? *reinterpret_cast(BuildOptions) + : EmptyOptions; + context *SyclCtx = unwrap(Ctx); device *SyclDev = unwrap(Dev); if (!SyclDev->ext_oneapi_can_compile(syclex::source_language::sycl)) { + set_build_log("Device does not support compilation of SYCL source."); return nullptr; } try { - auto *IncludeFileList = - reinterpret_cast(Headers); std::unique_ptr> SrcBundle; std::string Src(Source); @@ -1041,20 +1075,21 @@ __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource( // `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}; + if (!IncludeFiles.empty()) { + auto IncludeFileIt = IncludeFiles.begin(); + syclex::include_files IncludeFilesProp{IncludeFileIt->first, + IncludeFileIt->second}; for (std::advance(IncludeFileIt, 1); - IncludeFileIt != IncludeFileList->end(); ++IncludeFileIt) + IncludeFileIt != IncludeFiles.end(); ++IncludeFileIt) { - IncludeFiles.add(IncludeFileIt->first, IncludeFileIt->second); + IncludeFilesProp.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})); + syclex::properties{IncludeFilesProp})); } else { SrcBundle = std::make_unique< @@ -1064,14 +1099,11 @@ __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource( } registered_names_property_t RegisteredNames; - for (const std::string &Name : - *reinterpret_cast(Names)) - { + for (const std::string &Name : KernelNames) { RegisteredNames.add(Name); } - syclex::build_options Opts{ - *reinterpret_cast(BuildOptions)}; + syclex::build_options Opts{Options}; std::vector Devices({*SyclDev}); @@ -1083,8 +1115,7 @@ __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource( return wrap>( ResultBundle.release()); } catch (const std::exception &e) { - auto *RawBuildLog = reinterpret_cast(BuildLog); - *RawBuildLog = e.what(); + set_build_log(e.what()); return nullptr; } #else @@ -1097,6 +1128,15 @@ DPCTLKernelBundle_GetSyclKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, __dpctl_keep const char *KernelName) { #if (SUPPORTS_SYCL_COMPILATION > 0) + if (!KBRef) { + error_handler("Input KBRef is nullptr", __FILE__, __func__, __LINE__); + return nullptr; + } + if (!KernelName) { + error_handler("Input KernelName is nullptr", __FILE__, __func__, + __LINE__); + return nullptr; + } try { auto KernelBundle = unwrap>(KBRef); @@ -1116,6 +1156,15 @@ bool DPCTLKernelBundle_HasSyclKernel(__dpctl_keep DPCTLSyclKernelBundleRef __dpctl_keep const char *KernelName) { #if (SUPPORTS_SYCL_COMPILATION > 0) + if (!KBRef) { + error_handler("Input KBRef is nullptr", __FILE__, __func__, __LINE__); + return false; + } + if (!KernelName) { + error_handler("Input KernelName is nullptr", __FILE__, __func__, + __LINE__); + return false; + } try { auto KernelBundle = unwrap>(KBRef); diff --git a/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp index 23c81aad3b..329c33fe9b 100644 --- a/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp @@ -151,6 +151,36 @@ TEST_P(TestDPCTLSyclKernelBundleInterface, ChkCreateFromSpirvNull) ASSERT_TRUE(KBRef == nullptr); } +TEST_P(TestDPCTLSyclKernelBundleInterface, ChkCreateFromSpirvNullCompileOpts) +{ + // CompileOpts is optional and NULL means "no build options". It is + // forwarded to clBuildProgram (whose `options` argument is nullable) for + // the OpenCL backend, and to ze_module_desc_t::pBuildFlags (documented + // as [in][optional]) for the Level Zero backend. + DPCTLSyclKernelBundleRef KB = nullptr; + + EXPECT_NO_FATAL_FAILURE(KB = DPCTLKernelBundle_CreateFromSpirv( + CRef, DRef, spirvBuffer.data(), spirvFileSize, + nullptr, 0, nullptr)); + ASSERT_TRUE(KB != nullptr); + EXPECT_TRUE(DPCTLKernelBundle_HasKernel(KB, "add")); + EXPECT_TRUE(DPCTLKernelBundle_HasKernel(KB, "axpy")); + DPCTLKernelBundle_Delete(KB); +} + +TEST_P(TestDPCTLSyclKernelBundleInterface, ChkCreateFromSpirvWithCompileOpts) +{ + DPCTLSyclKernelBundleRef KB = nullptr; + + EXPECT_NO_FATAL_FAILURE(KB = DPCTLKernelBundle_CreateFromSpirv( + CRef, DRef, spirvBuffer.data(), spirvFileSize, + "-cl-fast-relaxed-math", 0, nullptr)); + ASSERT_TRUE(KB != nullptr); + EXPECT_TRUE(DPCTLKernelBundle_HasKernel(KB, "add")); + EXPECT_TRUE(DPCTLKernelBundle_HasKernel(KB, "axpy")); + DPCTLKernelBundle_Delete(KB); +} + TEST_P(TestDPCTLSyclKernelBundleInterface, ChkHasKernelNullKernelBundle) { @@ -267,6 +297,19 @@ TEST_P(TestOCLKernelBundleFromSource, CheckCreateFromOCLSourceNull) ASSERT_TRUE(KBRef == nullptr); } +TEST_P(TestOCLKernelBundleFromSource, CheckCreateFromOCLSourceNullCompileOpts) +{ + // CompileOpts is optional: it is forwarded to clBuildProgram, whose + // `options` argument accepts NULL to mean "no build options". + DPCTLSyclKernelBundleRef KB = nullptr; + EXPECT_NO_FATAL_FAILURE(KB = DPCTLKernelBundle_CreateFromOCLSource( + CRef, DRef, CLProgramStr, nullptr)); + ASSERT_TRUE(KB != nullptr); + EXPECT_TRUE(DPCTLKernelBundle_HasKernel(KB, "add")); + EXPECT_TRUE(DPCTLKernelBundle_HasKernel(KB, "axpy")); + DPCTLKernelBundle_Delete(KB); +} + TEST_P(TestOCLKernelBundleFromSource, CheckGetKernelOCLSource) { auto AddKernel = DPCTLKernelBundle_GetKernel(KBRef, "add"); @@ -303,6 +346,21 @@ struct TestSYCLKernelBundleFromSource } )==="; + // A source string that needs no virtual headers, so that it can be built + // while passing NULL for the optional arguments. + const char *standalone_sycl_source = R"===( + #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]; + } + )==="; + const char *header1_content = R"===( int math_op(int a, int b){ return a + b; @@ -403,6 +461,65 @@ TEST_P(TestSYCLKernelBundleFromSource, CheckGetKernelSYCLSource) DPCTLKernel_Delete(AxpyKernel); } +TEST_P(TestSYCLKernelBundleFromSource, CheckCreateFromSYCLSourceNullArgs) +{ + DPCTLBuildOptionListRef BORef = DPCTLBuildOptionList_Create(); + DPCTLKernelNameListRef KNRef = DPCTLKernelNameList_Create(); + DPCTLVirtualHeaderListRef VHRef = DPCTLVirtualHeaderList_Create(); + DPCTLKernelBuildLogRef KBLRef = DPCTLKernelBuildLog_Create(); + + // Ctx, Dev and Source are required; passing NULL must return NULL rather + // than dereference it. + auto create = [&](DPCTLSyclContextRef C, DPCTLSyclDeviceRef D, + const char *S, DPCTLVirtualHeaderListRef H, + DPCTLKernelNameListRef N, DPCTLBuildOptionListRef B, + DPCTLKernelBuildLogRef L) { + return DPCTLKernelBundle_CreateFromSYCLSource(C, D, S, H, N, B, L); + }; + + EXPECT_EQ(create(nullptr, DRef, sycl_source, VHRef, KNRef, BORef, KBLRef), + nullptr); + EXPECT_EQ(create(CRef, nullptr, sycl_source, VHRef, KNRef, BORef, KBLRef), + nullptr); + EXPECT_EQ(create(CRef, DRef, nullptr, VHRef, KNRef, BORef, KBLRef), + nullptr); + + DPCTLVirtualHeaderList_Delete(VHRef); + DPCTLKernelNameList_Delete(KNRef); + DPCTLBuildOptionList_Delete(BORef); + DPCTLKernelBuildLog_Delete(KBLRef); +} + +TEST_P(TestSYCLKernelBundleFromSource, CheckCreateFromSYCLSourceNullOptionals) +{ + // Headers, Names, BuildOptions and BuildLog are optional: passing NULL is + // equivalent to passing an empty list / discarding the log. + DPCTLSyclKernelBundleRef KB = DPCTLKernelBundle_CreateFromSYCLSource( + CRef, DRef, standalone_sycl_source, nullptr, nullptr, nullptr, nullptr); + + ASSERT_TRUE(KB != nullptr); + EXPECT_TRUE(DPCTLKernelBundle_HasSyclKernel(KB, "vector_add")); + DPCTLKernelBundle_Delete(KB); +} + +TEST_P(TestSYCLKernelBundleFromSource, CheckNullBuildLogOnFailure) +{ + // A compilation failure with a NULL BuildLog must return NULL rather than + // write the diagnostics through a null pointer. + const char *bad_source = "this is not valid SYCL"; + EXPECT_EQ(DPCTLKernelBundle_CreateFromSYCLSource( + CRef, DRef, bad_source, nullptr, nullptr, nullptr, nullptr), + nullptr); +} + +TEST_P(TestSYCLKernelBundleFromSource, CheckSyclKernelQueriesNullArgs) +{ + EXPECT_EQ(DPCTLKernelBundle_GetSyclKernel(nullptr, "vector_add"), nullptr); + EXPECT_EQ(DPCTLKernelBundle_GetSyclKernel(KBRef, nullptr), nullptr); + EXPECT_FALSE(DPCTLKernelBundle_HasSyclKernel(nullptr, "vector_add")); + EXPECT_FALSE(DPCTLKernelBundle_HasSyclKernel(KBRef, nullptr)); +} + INSTANTIATE_TEST_SUITE_P(KernelBundleCreationFromSpirv, TestDPCTLSyclKernelBundleInterface, ::testing::Values("opencl", From 4453c9dbfce89fd443b0522af72015cb132b6b3d Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Wed, 5 Aug 2026 14:04:20 -0700 Subject: [PATCH 10/13] prevent leaks when creating kernel bundle from SYCL source --- dpctl/program/_program.pyx | 109 +++++++++++++++++-------------- dpctl/tests/test_sycl_program.py | 41 ++++++++++++ 2 files changed, 102 insertions(+), 48 deletions(-) diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index 453d7e26bf..de7c375704 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -698,12 +698,17 @@ cpdef create_kernel_bundle_from_sycl_source(SyclQueue q, If a SYCL kernel bundle could not be created. The exception message contains the build log for more details. """ - cdef DPCTLSyclKernelBundleRef KBref + cdef DPCTLSyclKernelBundleRef KBref = NULL 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() + # initialized to NULL so that the `finally` clause below can release + # whichever of these were created before an exception was raised + cdef DPCTLBuildOptionListRef BuildOpts = NULL + cdef DPCTLKernelNameListRef KernelNames = NULL + cdef DPCTLVirtualHeaderListRef VirtualHeaders = NULL + cdef DPCTLKernelBuildLogRef BuildLog = NULL cdef bytes bOpt cdef const char* sOpt cdef bytes bName @@ -719,59 +724,67 @@ cpdef create_kernel_bundle_from_sycl_source(SyclQueue q, if copts is None: copts = [] - 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) + try: + BuildOpts = DPCTLBuildOptionList_Create() + KernelNames = DPCTLKernelNameList_Create() + VirtualHeaders = DPCTLVirtualHeaderList_Create() + BuildLog = DPCTLKernelBuildLog_Create() - cdef DPCTLKernelBuildLogRef BuildLog - BuildLog = DPCTLKernelBuildLog_Create() + for opt in copts: + if not isinstance(opt, unicode): + raise TypeError( + "Every element of `copts` must be a string, got " + f"{type(opt)}" + ) + bOpt = opt.encode("utf8") + sOpt = bOpt + DPCTLBuildOptionList_Append(BuildOpts, sOpt) - KBref = DPCTLKernelBundle_CreateFromSYCLSource(CRef, DRef, Src, - VirtualHeaders, KernelNames, - BuildOpts, BuildLog) + for name in registered_names: + if not isinstance(name, unicode): + raise TypeError( + "Every element of `registered_names` must be a string, " + f"got {type(name)}" + ) + bName = name.encode("utf8") + sName = bName + DPCTLKernelNameList_Append(KernelNames, sName) - if KBref is NULL: - buildLogContent = DPCTLKernelBuildLog_Get(BuildLog) - buildLogStr = str(buildLogContent, "utf-8") + for header in headers: + if not isinstance(header, tuple) or len(header) != 2: + raise TypeError( + "Every element of `headers` must be a 2-tuple of header " + "name and header content" + ) + name, content = header + if ( + not isinstance(name, unicode) + or not isinstance(content, unicode) + ): + raise TypeError( + "Header names and header content must be strings, got " + f"({type(name)}, {type(content)})" + ) + 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, BuildLog + ) + + if KBref is NULL: + buildLogContent = DPCTLKernelBuildLog_Get(BuildLog) + raise SyclKernelBundleCompilationError( + str(buildLogContent, "utf-8") + ) + finally: 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) diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index c045b2aafe..8f74b747a1 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -598,6 +598,47 @@ def test_create_kernel_bundle_from_sycl_source_defaults(queue_selector): assert prog.has_sycl_kernel("vector_add") +@pytest.mark.parametrize( + "kwargs", + [ + {"copts": [1]}, + {"copts": ["-fno-fast-math", None]}, + {"registered_names": [1]}, + {"registered_names": ["vector_add", None]}, + {"headers": [1]}, + {"headers": [("math_ops.hpp",)]}, + {"headers": [("math_ops.hpp", "int f(){return 0;}", "extra")]}, + {"headers": [(1, "int f(){return 0;}")]}, + {"headers": [("math_ops.hpp", 1)]}, + ], +) +def test_create_kernel_bundle_from_sycl_source_bad_args(kwargs): + q = _get_opencl_queue_or_skip() + _skip_if_no_sycl_source_compilation(q) + + 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]; + } + """ + + # Malformed arguments must raise rather than crash, and the C API handles + # allocated for the argument lists must be released on the way out. + with pytest.raises(TypeError): + dpctl.program.create_kernel_bundle_from_sycl_source( + q, sycl_source, **kwargs + ) + + @pytest.mark.parametrize( "queue_selector", [_get_opencl_queue_or_skip, _get_level_zero_queue_or_skip] ) From 7601d3129cd489f0bf6427a1b469cb0959ea3bd3 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Wed, 5 Aug 2026 14:04:55 -0700 Subject: [PATCH 11/13] document is_sycl_source_compilation_available --- 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 a37aec5531..f8c9433eab 100644 --- a/docs/doc_sources/api_reference/dpctl/program.rst +++ b/docs/doc_sources/api_reference/dpctl/program.rst @@ -25,6 +25,7 @@ execution via :py:meth:`dpctl.SyclQueue.submit`. create_kernel_bundle_from_sycl_source create_program_from_source create_program_from_spirv + is_sycl_source_compilation_available .. autosummary:: :toctree: generated From 57cd0477d362ebb855c62e77616a7519713bff61 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Wed, 5 Aug 2026 14:10:08 -0700 Subject: [PATCH 12/13] add gh-2206 to changelog --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index a183f475d1..c058ca81a1 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -8,6 +8,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 ### Added * `dpctl.SyclQueue.copy` and `dpctl.SyclQueue.copy_async` methods [gh-2273](https://github.com/IntelPython/dpctl/pull/2273) +* Added `create_kernel_bundle_from_sycl_source` for creating `dpctl.SyclKernelBundle` from SYCL source strings via DPC++ extension, as well as corresponding C-API functions to support it [gh-2206](https://github.com/IntelPython/dpctl/pull/2206) ### Changed * Bump minimum NumPy version to 1.26 [gh-2192](https://github.com/IntelPython/dpctl/pull/2192) From eaac81eab0cd5c29548b305b8c10d6648c418840 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Thu, 6 Aug 2026 06:28:02 -0700 Subject: [PATCH 13/13] address review feedback --- CHANGELOG.md | 2 +- .../source/dpctl_sycl_device_interface.cpp | 36 +++++++++++++++---- 2 files changed, 31 insertions(+), 7 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index e330a540ed..7dafc1fde2 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -9,7 +9,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 ### Added * `dpctl.SyclQueue.copy` and `dpctl.SyclQueue.copy_async` methods [gh-2273](https://github.com/IntelPython/dpctl/pull/2273) * Added a number of `sycl::device` info queries to `dpctl.SyclDevice` [gh-2324](https://github.com/IntelPython/dpctl/pull/2324) -* Added `create_kernel_bundle_from_sycl_source` for creating `dpctl.SyclKernelBundle` from SYCL source strings via DPC++ extension, as well as corresponding C-API functions to support it [gh-2206](https://github.com/IntelPython/dpctl/pull/2206) +* Added `create_kernel_bundle_from_sycl_source`, `is_sycl_source_compilation_available`, and `dpctl.SyclDevice.can_compile` for supporting the creation of `dpctl.SyclKernelBundle`s from SYCL source strings via DPC++ extension, as well as corresponding C-API functions to support it [gh-2206](https://github.com/IntelPython/dpctl/pull/2206) ### Changed * Bump minimum NumPy version to 1.26 [gh-2192](https://github.com/IntelPython/dpctl/pull/2192) diff --git a/libsyclinterface/source/dpctl_sycl_device_interface.cpp b/libsyclinterface/source/dpctl_sycl_device_interface.cpp index a28d61f3ee..5029e60650 100644 --- a/libsyclinterface/source/dpctl_sycl_device_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_device_interface.cpp @@ -1285,24 +1285,48 @@ __dpctl_give int *DPCTLDevice_GetPartitionAffinityDomains( bool DPCTLDevice_CanCompileSPIRV(__dpctl_keep const DPCTLSyclDeviceRef DRef) { + bool canCompile = false; auto Dev = unwrap(DRef); - auto Backend = Dev->get_platform().get_backend(); - return Backend == backend::opencl || - Backend == backend::ext_oneapi_level_zero; + if (Dev) { + try { + auto Backend = Dev->get_platform().get_backend(); + canCompile = Backend == backend::opencl || + Backend == backend::ext_oneapi_level_zero; + } catch (std::exception const &e) { + error_handler(e, __FILE__, __func__, __LINE__); + } + } + return canCompile; } bool DPCTLDevice_CanCompileOpenCL(__dpctl_keep const DPCTLSyclDeviceRef DRef) { + bool canCompile = false; auto Dev = unwrap(DRef); - return Dev->get_platform().get_backend() == backend::opencl; + if (Dev) { + try { + canCompile = Dev->get_platform().get_backend() == backend::opencl; + } catch (std::exception const &e) { + error_handler(e, __FILE__, __func__, __LINE__); + } + } + return canCompile; } bool DPCTLDevice_CanCompileSYCL(__dpctl_keep const DPCTLSyclDeviceRef DRef) { #ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER + bool canCompile = false; auto Dev = unwrap(DRef); - return Dev->ext_oneapi_can_compile( - ext::oneapi::experimental::source_language::sycl); + if (Dev) { + try { + canCompile = Dev->ext_oneapi_can_compile( + ext::oneapi::experimental::source_language::sycl); + } catch (std::exception const &e) { + error_handler(e, __FILE__, __func__, __LINE__); + } + } + return canCompile; #else return false; #endif