diff --git a/CHANGELOG.md b/CHANGELOG.md index 686d06eed7..7dafc1fde2 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -9,6 +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`, `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/docs/doc_sources/api_reference/dpctl/program.rst b/docs/doc_sources/api_reference/dpctl/program.rst index 6b2bf22bad..f8c9433eab 100644 --- a/docs/doc_sources/api_reference/dpctl/program.rst +++ b/docs/doc_sources/api_reference/dpctl/program.rst @@ -22,8 +22,10 @@ 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 + is_sycl_source_compilation_available .. autosummary:: :toctree: generated diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index a85565b729..6ef017fbe4 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -331,7 +331,6 @@ 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 uint32_t DPCTLDevice_GetVendorId(const DPCTLSyclDeviceRef DRef) @@ -374,6 +373,10 @@ cdef extern from "syclinterface/dpctl_sycl_device_interface.h": const DPCTLSyclDeviceRef DRef, size_t *res_len) cdef int *DPCTLDevice_GetPartitionAffinityDomains( const DPCTLSyclDeviceRef DRef, size_t *res_len) + 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( @@ -542,6 +545,53 @@ 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 bool DPCTLKernelBundle_CreateFromSYCLSource_Available() + + 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 1bc00fdd2a..5d78424193 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, @@ -2869,6 +2872,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..0c09226f99 100644 --- a/dpctl/program/__init__.py +++ b/dpctl/program/__init__.py @@ -29,8 +29,10 @@ 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, + is_sycl_source_compilation_available, ) __all__ = [ @@ -38,6 +40,8 @@ "create_kernel_bundle_from_spirv", "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.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..de7c375704 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,32 @@ 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_CreateFromSYCLSource_Available, 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, ) @@ -73,6 +93,8 @@ 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", "SyclKernelBundleCompilationError", @@ -86,6 +108,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: """ """ @@ -217,9 +250,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 +265,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 +279,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 +313,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 +564,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 +650,143 @@ 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) + 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: ``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: ``None`` + copts (list, optional) + Optional list of compilation flags that will be used + when compiling the program. + Default: ``None`` + + 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 = 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 + # 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 + cdef const char* sName + 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 = [] + + try: + BuildOpts = DPCTLBuildOptionList_Create() + KernelNames = DPCTLKernelNameList_Create() + VirtualHeaders = DPCTLVirtualHeaderList_Create() + 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) + + 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) + + 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) + + 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..8f74b747a1 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -22,10 +22,32 @@ 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 _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") + 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) @@ -82,8 +104,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 +209,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(): @@ -201,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() @@ -222,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() @@ -234,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); \ @@ -255,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( \ }" @@ -412,3 +420,290 @@ def test_spec_const_equality(): assert sp1 == sp2 assert sp1 != sp3 assert sp1 != sp4 + + +@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 = """ + #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) + + +@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 = """ + #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) + + +def test_sycl_source_compilation_is_available_returns_bool(): + v = dpctl.program.is_sycl_source_compilation_available() + 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( + "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] +) +def test_sycl_source_vector_add_correctness(queue_selector): + q = queue_selector() + _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) diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_device_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_device_interface.h index 2d8bfad0c6..83174cec91 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_device_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_device_interface.h @@ -1094,4 +1094,40 @@ __dpctl_give int *DPCTLDevice_GetPartitionAffinityDomains( __dpctl_keep const DPCTLSyclDeviceRef DRef, size_t *res_len); +/*! + * @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..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 @@ -140,4 +144,187 @@ 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 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. + * + * @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, 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 +__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 d7e1b092ef..5029e60650 100644 --- a/libsyclinterface/source/dpctl_sycl_device_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_device_interface.cpp @@ -1282,3 +1282,52 @@ __dpctl_give int *DPCTLDevice_GetPartitionAffinityDomains( return get_info_enum_array( DRef, res_len, DPCTL_SyclPartitionAffinityDomainToDPCTLType); } + +bool DPCTLDevice_CanCompileSPIRV(__dpctl_keep const DPCTLSyclDeviceRef DRef) +{ + bool canCompile = false; + auto Dev = unwrap(DRef); + 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); + 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); + 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 +} diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp index ec94e5a077..c984123628 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -872,3 +872,308 @@ 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 + +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 +// `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) + // 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 { + 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 (!IncludeFiles.empty()) { + auto IncludeFileIt = IncludeFiles.begin(); + syclex::include_files IncludeFilesProp{IncludeFileIt->first, + IncludeFileIt->second}; + for (std::advance(IncludeFileIt, 1); + IncludeFileIt != IncludeFiles.end(); ++IncludeFileIt) + { + 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{IncludeFilesProp})); + } + 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 : KernelNames) { + RegisteredNames.add(Name); + } + + syclex::build_options Opts{Options}; + + std::vector Devices({*SyclDev}); + + auto ExeBundle = syclex::build( + *SrcBundle, Devices, syclex::properties{RegisteredNames, Opts}); + auto ResultBundle = + std::make_unique>( + ExeBundle); + return wrap>( + ResultBundle.release()); + } catch (const std::exception &e) { + set_build_log(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) + 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); + 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) + 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); + 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..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"); @@ -277,6 +320,206 @@ 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]); + } + )==="; + + // 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; + } + )==="; + + 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); +} + +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", @@ -293,6 +536,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;