From 1db2bfa036d0e9e4d5c8f408c30faef6d3c0ffa0 Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Thu, 10 Apr 2025 10:39:23 +0100 Subject: [PATCH 1/2] 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 | 7 +- dpctl/program/_program.pyx | 148 ++++++++++- 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, 993 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 96efc310c3..08a6fa0dd6 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, @@ -2367,6 +2370,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 a96d33f04a..e209b68b40 100644 --- a/dpctl/program/__init__.py +++ b/dpctl/program/__init__.py @@ -26,11 +26,13 @@ SyclProgramCompilationError, create_program_from_source, create_program_from_spirv, + create_program_from_sycl_source, ) __all__ = [ "create_program_from_source", "create_program_from_spirv", + "create_program_from_sycl_source", "SyclKernel", "SyclProgram", "SyclProgramCompilationError", diff --git a/dpctl/program/_program.pxd b/dpctl/program/_program.pxd index dc4208a29b..880843c27f 100644 --- a/dpctl/program/_program.pxd +++ b/dpctl/program/_program.pxd @@ -49,9 +49,11 @@ cdef api class SyclProgram [object PySyclProgramObject, type PySyclProgramType]: binary file. """ cdef DPCTLSyclKernelBundleRef _program_ref + cdef bint _is_sycl_source @staticmethod - cdef SyclProgram _create (DPCTLSyclKernelBundleRef pref) + cdef SyclProgram _create (DPCTLSyclKernelBundleRef pref, + bint _is_sycl_source) cdef DPCTLSyclKernelBundleRef get_program_ref (self) cpdef SyclKernel get_sycl_kernel(self, str kernel_name) @@ -59,3 +61,6 @@ cdef api class SyclProgram [object PySyclProgramObject, type PySyclProgramType]: cpdef create_program_from_source (SyclQueue q, unicode source, unicode copts=*) cpdef create_program_from_spirv (SyclQueue q, const unsigned char[:] IL, unicode copts=*) +cpdef create_program_from_sycl_source(SyclQueue q, unicode source, + list headers=*, list registered_names=*, + list copts=*) diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index 3859314505..64433947ac 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -28,6 +28,10 @@ a OpenCL source string or a SPIR-V binary file. from libc.stdint cimport uint32_t from dpctl._backend cimport ( # noqa: E211, E402; + DPCTLBuildOptionList_Append, + DPCTLBuildOptionList_Create, + DPCTLBuildOptionList_Delete, + DPCTLBuildOptionListRef, DPCTLKernel_Copy, DPCTLKernel_Delete, DPCTLKernel_GetCompileNumSubGroups, @@ -38,16 +42,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__ = [ @@ -196,9 +215,11 @@ cdef class SyclProgram: """ @staticmethod - cdef SyclProgram _create(DPCTLSyclKernelBundleRef KBRef): + cdef SyclProgram _create(DPCTLSyclKernelBundleRef KBRef, + bint is_sycl_source): cdef SyclProgram ret = SyclProgram.__new__(SyclProgram) ret._program_ref = KBRef + ret._is_sycl_source = is_sycl_source return ret def __dealloc__(self): @@ -209,6 +230,10 @@ cdef class SyclProgram: cpdef SyclKernel get_sycl_kernel(self, str kernel_name): name = kernel_name.encode("utf8") + if self._is_sycl_source: + return SyclKernel._create( + DPCTLKernelBundle_GetSyclKernel(self._program_ref, name), + kernel_name) return SyclKernel._create( DPCTLKernelBundle_GetKernel(self._program_ref, name), kernel_name @@ -216,6 +241,8 @@ cdef class SyclProgram: def has_sycl_kernel(self, str kernel_name): name = kernel_name.encode("utf8") + if self._is_sycl_source: + return DPCTLKernelBundle_HasSyclKernel(self._program_ref, name) return DPCTLKernelBundle_HasKernel(self._program_ref, name) def addressof_ref(self): @@ -271,7 +298,7 @@ cpdef create_program_from_source(SyclQueue q, str src, str copts=""): if KBref is NULL: raise SyclProgramCompilationError() - return SyclProgram._create(KBref) + return SyclProgram._create(KBref, False) cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL, @@ -317,7 +344,120 @@ cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL, if KBref is NULL: raise SyclProgramCompilationError() - return SyclProgram._create(KBref) + return SyclProgram._create(KBref, False) + + +cpdef create_program_from_sycl_source(SyclQueue q, unicode source, + list headers=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:`.SyclProgram` is going to be built. + source (unicode) + SYCL source code string. + headers (list) + Optional list of virtual headers, where each entry in the list + needs to be a tuple of header name and header content. See the + documentation of the ``include_files`` property in the DPC++ + ``kernel_compiler`` extension for more information. + Default: [] + registered_names (list, optional) + Optional list of kernel names to register. See the + documentation of the ``registered_names`` property in the DPC++ + ``kernel_compiler`` extension for more information. + Default: [] + copts (list) + Optional list of compilation flags that will be used + when compiling the program. Default: ``""``. + + Returns: + program (:class:`.SyclProgram`) + A :class:`.SyclProgram` object wrapping the + ``sycl::kernel_bundle`` + returned by the C API. + + Raises: + SyclProgramCompilationError + If a SYCL kernel bundle could not be created. 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 SyclProgramCompilationError() + bOpt = opt.encode("utf8") + sOpt = bOpt + DPCTLBuildOptionList_Append(BuildOpts, sOpt) + + cdef DPCTLKernelNameListRef KernelNames = DPCTLKernelNameList_Create() + for name in registered_names: + if not isinstance(name, unicode): + DPCTLBuildOptionList_Delete(BuildOpts) + DPCTLKernelNameList_Delete(KernelNames) + raise SyclProgramCompilationError() + bName = name.encode("utf8") + sName = bName + DPCTLKernelNameList_Append(KernelNames, sName) + + cdef DPCTLVirtualHeaderListRef VirtualHeaders + VirtualHeaders = DPCTLVirtualHeaderList_Create() + + for name, content in headers: + if not isinstance(name, unicode) or not isinstance(content, unicode): + DPCTLBuildOptionList_Delete(BuildOpts) + DPCTLKernelNameList_Delete(KernelNames) + DPCTLVirtualHeaderList_Delete(VirtualHeaders) + raise SyclProgramCompilationError() + bName = name.encode("utf8") + sName = bName + bContent = content.encode("utf8") + sContent = bContent + DPCTLVirtualHeaderList_Append(VirtualHeaders, sName, sContent) + + 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 SyclProgramCompilationError(buildLogStr) + + DPCTLBuildOptionList_Delete(BuildOpts) + DPCTLKernelNameList_Delete(KernelNames) + DPCTLVirtualHeaderList_Delete(VirtualHeaders) + DPCTLKernelBuildLog_Delete(BuildLog) + + return SyclProgram._create(KBref, True) cdef api DPCTLSyclKernelBundleRef SyclProgram_GetKernelBundleRef( @@ -336,4 +476,4 @@ cdef api SyclProgram SyclProgram_Make(DPCTLSyclKernelBundleRef KBRef): reference. """ cdef DPCTLSyclKernelBundleRef copied_KBRef = DPCTLKernelBundle_Copy(KBRef) - return SyclProgram._create(copied_KBRef) + return SyclProgram._create(copied_KBRef, False) diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index a791a59daa..d0d1348da5 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -81,8 +81,7 @@ def _check_cpython_api_SyclProgram_Make(sycl_prog): make_prog_fn = callable_maker(make_prog_fn_ptr) p2 = make_prog_fn(sycl_prog.addressof_ref()) - assert p2.has_sycl_kernel("add") - assert p2.has_sycl_kernel("axpy") + return p2 def _check_cpython_api_SyclKernel_GetKernelRef(krn): @@ -187,7 +186,9 @@ def _check_multi_kernel_program(prog): assert type(cmsgsz) is int _check_cpython_api_SyclProgram_GetKernelBundleRef(prog) - _check_cpython_api_SyclProgram_Make(prog) + p2 = _check_cpython_api_SyclProgram_Make(prog) + assert p2.has_sycl_kernel("add") + assert p2.has_sycl_kernel("axpy") def test_create_program_from_source_ocl(): @@ -263,3 +264,150 @@ def test_create_program_from_invalid_src_ocl(): }" with pytest.raises(dpctl_prog.SyclProgramCompilationError): dpctl_prog.create_program_from_source(q, invalid_oclSrc) + + +def test_create_program_from_sycl_source(): + try: + q = dpctl.SyclQueue("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_program_from_sycl_source( + q, + sycl_source, + headers=[ + ("math_ops.hpp", header_content), + ("math_template_ops.hpp", header2_content), + ], + registered_names=["vector_add_template"], + copts=["-fno-fast-math"], + ) + + assert type(prog) is dpctl_prog.SyclProgram + + assert type(prog.addressof_ref()) is int + assert prog.has_sycl_kernel("vector_add") + 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_SyclProgram_GetKernelBundleRef(prog) + + +def test_create_program_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_program_from_sycl_source( + q, + sycl_source, + headers=[], + registered_names=[], + copts=[], + ) + assert False + except dpctl_prog.SyclProgramCompilationError 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 529bc3cca1..1ddfbe95f8 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 d136c700b6..740385283b 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 de9866128cdacf96c4d21e04f0a0a8beb4548be4 Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Wed, 16 Jul 2025 15:10:19 +0100 Subject: [PATCH 2/2] Skip test if oneAPI Base Toolkit isn't available As the library aren't included in .bc format in version 2025.2 and earlier, we need to skip the test if no oneAPI Base Toolkit is installed. Signed-off-by: Lukas Sommer --- dpctl/tests/test_sycl_program.py | 13 +++++++++++++ 1 file changed, 13 insertions(+) diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index d0d1348da5..88cba25170 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -18,6 +18,7 @@ """ import os +import shutil import pytest @@ -272,6 +273,12 @@ def test_create_program_from_sycl_source(): except dpctl.SyclQueueCreationError: pytest.skip("No OpenCL queue is available") + if not shutil.which("icpx"): + # In version 2025.2 and before, the packages do not contain the + # libraries in the .bc format necessary for RTC. Therefore, + # installation of the base toolkit is required. + pytest.skip("oneAPI Base Toolkit not installed") + if not q.get_sycl_device().can_compile("sycl"): pytest.skip("SYCL source compilation not supported") @@ -382,6 +389,12 @@ def test_create_program_from_invalid_src_sycl(): except dpctl.SyclQueueCreationError: pytest.skip("No OpenCL queue is available") + if not shutil.which("icpx"): + # In version 2025.2 and before, the packages do not contain the + # libraries in the .bc format necessary for RTC. Therefore, + # installation of the base toolkit is required. + pytest.skip("oneAPI Base Toolkit not installed") + if not q.get_sycl_device().can_compile("sycl"): pytest.skip("SYCL source compilation not supported")