diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index ca9e9ccb9f..351c9bf8eb 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -278,6 +278,9 @@ cdef extern from "syclinterface/dpctl_sycl_device_interface.h": cdef DPCTLDeviceVectorRef DPCTLDevice_GetComponentDevices( const DPCTLSyclDeviceRef DRef ) + cdef bool DPCTLDevice_CanCompileSPIRV(const DPCTLSyclDeviceRef DRef) + cdef bool DPCTLDevice_CanCompileOpenCL(const DPCTLSyclDeviceRef DRef) + cdef bool DPCTLDevice_CanCompileSYCL(const DPCTLSyclDeviceRef DRef) cdef extern from "syclinterface/dpctl_sycl_device_manager.h": @@ -441,6 +444,44 @@ cdef extern from "syclinterface/dpctl_sycl_kernel_bundle_interface.h": cdef DPCTLSyclKernelBundleRef DPCTLKernelBundle_Copy( const DPCTLSyclKernelBundleRef KBRef) + cdef struct DPCTLBuildOptionList + cdef struct DPCTLKernelNameList + cdef struct DPCTLVirtualHeaderList + ctypedef DPCTLBuildOptionList* DPCTLBuildOptionListRef + ctypedef DPCTLKernelNameList* DPCTLKernelNameListRef + ctypedef DPCTLVirtualHeaderList* DPCTLVirtualHeaderListRef + + cdef DPCTLBuildOptionListRef DPCTLBuildOptionList_Create() + cdef void DPCTLBuildOptionList_Delete(DPCTLBuildOptionListRef Ref) + cdef void DPCTLBuildOptionList_Append(DPCTLBuildOptionListRef Ref, + const char *Option) + + cdef DPCTLKernelNameListRef DPCTLKernelNameList_Create() + cdef void DPCTLKernelNameList_Delete(DPCTLKernelNameListRef Ref) + cdef void DPCTLKernelNameList_Append(DPCTLKernelNameListRef Ref, + const char *Option) + + cdef DPCTLVirtualHeaderListRef DPCTLVirtualHeaderList_Create() + cdef void DPCTLVirtualHeaderList_Delete(DPCTLVirtualHeaderListRef Ref) + cdef void DPCTLVirtualHeaderList_Append(DPCTLVirtualHeaderListRef Ref, + const char *Name, + const char *Content) + + cdef DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource( + const DPCTLSyclContextRef Ctx, + const DPCTLSyclDeviceRef Dev, + const char *Source, + DPCTLVirtualHeaderListRef Headers, + DPCTLKernelNameListRef Names, + DPCTLBuildOptionListRef BuildOptions) + + cdef DPCTLSyclKernelRef DPCTLKernelBundle_GetSyclKernel( + DPCTLSyclKernelBundleRef KBRef, + const char *KernelName) + + cdef bool DPCTLKernelBundle_HasSyclKernel(DPCTLSyclKernelBundleRef KBRef, + const char *KernelName) + cdef extern from "syclinterface/dpctl_sycl_queue_interface.h": ctypedef struct _md_local_accessor "MDLocalAccessor": diff --git a/dpctl/_sycl_device.pxd b/dpctl/_sycl_device.pxd index 190d981cd0..d9378f0897 100644 --- a/dpctl/_sycl_device.pxd +++ b/dpctl/_sycl_device.pxd @@ -61,3 +61,4 @@ cdef public api class SyclDevice(_SyclDevice) [ cdef int get_overall_ordinal(self) cdef int get_backend_ordinal(self) cdef int get_backend_and_device_type_ordinal(self) + cpdef bint can_compile(self, str language) diff --git a/dpctl/_sycl_device.pyx b/dpctl/_sycl_device.pyx index 5b43ffed1a..b5be0af1ee 100644 --- a/dpctl/_sycl_device.pyx +++ b/dpctl/_sycl_device.pyx @@ -25,6 +25,8 @@ from ._backend cimport ( # noqa: E211 DPCTLCString_Delete, DPCTLDefaultSelector_Create, DPCTLDevice_AreEq, + DPCTLDevice_CanCompileOpenCL, + DPCTLDevice_CanCompileSYCL, DPCTLDevice_Copy, DPCTLDevice_CreateFromSelector, DPCTLDevice_CreateSubDevicesByAffinity, @@ -2160,6 +2162,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_CanCompileSYCL(self._device_ref) + if language == "opencl" or language == "ocl": + return DPCTLDevice_CanCompileOpenCL(self._device_ref) + if language == "sycl": + return DPCTLDevice_CanCompileSYCL(self._device_ref) + + raise ValueError(f"Unknown source language {language}") + cdef api DPCTLSyclDeviceRef SyclDevice_GetDeviceRef(SyclDevice dev): """ diff --git a/dpctl/program/__init__.py b/dpctl/program/__init__.py index a96d33f04a..e209b68b40 100644 --- a/dpctl/program/__init__.py +++ b/dpctl/program/__init__.py @@ -26,11 +26,13 @@ SyclProgramCompilationError, create_program_from_source, create_program_from_spirv, + create_program_from_sycl_source, ) __all__ = [ "create_program_from_source", "create_program_from_spirv", + "create_program_from_sycl_source", "SyclKernel", "SyclProgram", "SyclProgramCompilationError", diff --git a/dpctl/program/_program.pxd b/dpctl/program/_program.pxd index dc4208a29b..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..f371149bfd 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -28,6 +28,10 @@ a OpenCL source string or a SPIR-V binary file. from libc.stdint cimport uint32_t from dpctl._backend cimport ( # noqa: E211, E402; + DPCTLBuildOptionList_Append, + DPCTLBuildOptionList_Create, + DPCTLBuildOptionList_Delete, + DPCTLBuildOptionListRef, DPCTLKernel_Copy, DPCTLKernel_Delete, DPCTLKernel_GetCompileNumSubGroups, @@ -41,13 +45,24 @@ from dpctl._backend cimport ( # noqa: E211, E402; DPCTLKernelBundle_Copy, DPCTLKernelBundle_CreateFromOCLSource, DPCTLKernelBundle_CreateFromSpirv, + DPCTLKernelBundle_CreateFromSYCLSource, DPCTLKernelBundle_Delete, DPCTLKernelBundle_GetKernel, + DPCTLKernelBundle_GetSyclKernel, DPCTLKernelBundle_HasKernel, + DPCTLKernelBundle_HasSyclKernel, + DPCTLKernelNameList_Append, + DPCTLKernelNameList_Create, + DPCTLKernelNameList_Delete, + DPCTLKernelNameListRef, DPCTLSyclContextRef, DPCTLSyclDeviceRef, DPCTLSyclKernelBundleRef, DPCTLSyclKernelRef, + DPCTLVirtualHeaderList_Append, + DPCTLVirtualHeaderList_Create, + DPCTLVirtualHeaderList_Delete, + DPCTLVirtualHeaderListRef, ) __all__ = [ @@ -196,9 +211,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 +226,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 +237,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 +294,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 +340,111 @@ 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. + """ + cdef DPCTLSyclKernelBundleRef KBref + cdef DPCTLSyclContextRef CRef = q.get_sycl_context().get_context_ref() + cdef DPCTLSyclDeviceRef DRef = q.get_sycl_device().get_device_ref() + cdef bytes bSrc = source.encode("utf8") + cdef const char *Src = bSrc + cdef DPCTLBuildOptionListRef BuildOpts = DPCTLBuildOptionList_Create() + cdef bytes bOpt + cdef const char* sOpt + cdef bytes bName + cdef const char* sName + cdef bytes bContent + cdef const char* sContent + for opt in copts: + if not isinstance(opt, unicode): + DPCTLBuildOptionList_Delete(BuildOpts) + raise SyclProgramCompilationError() + bOpt = opt.encode("utf8") + sOpt = bOpt + DPCTLBuildOptionList_Append(BuildOpts, sOpt) + + cdef DPCTLKernelNameListRef KernelNames = DPCTLKernelNameList_Create() + for name in registered_names: + if not isinstance(name, unicode): + DPCTLBuildOptionList_Delete(BuildOpts) + DPCTLKernelNameList_Delete(KernelNames) + raise SyclProgramCompilationError() + bName = name.encode("utf8") + sName = bName + DPCTLKernelNameList_Append(KernelNames, sName) + + cdef DPCTLVirtualHeaderListRef VirtualHeaders + VirtualHeaders = DPCTLVirtualHeaderList_Create() + + for name, content in headers: + if not isinstance(name, unicode) or not isinstance(content, unicode): + DPCTLBuildOptionList_Delete(BuildOpts) + DPCTLKernelNameList_Delete(KernelNames) + DPCTLVirtualHeaderList_Delete(VirtualHeaders) + raise SyclProgramCompilationError() + bName = name.encode("utf8") + sName = bName + bContent = content.encode("utf8") + sContent = bContent + DPCTLVirtualHeaderList_Append(VirtualHeaders, sName, sContent) + + KBref = DPCTLKernelBundle_CreateFromSYCLSource(CRef, DRef, Src, + VirtualHeaders, KernelNames, + BuildOpts) + + if KBref is NULL: + DPCTLBuildOptionList_Delete(BuildOpts) + DPCTLKernelNameList_Delete(KernelNames) + DPCTLVirtualHeaderList_Delete(VirtualHeaders) + raise SyclProgramCompilationError() + + DPCTLBuildOptionList_Delete(BuildOpts) + DPCTLKernelNameList_Delete(KernelNames) + DPCTLVirtualHeaderList_Delete(VirtualHeaders) + + return SyclProgram._create(KBref, True) cdef api DPCTLSyclKernelBundleRef SyclProgram_GetKernelBundleRef( @@ -336,4 +463,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..4df119c87b 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,113 @@ def test_create_program_from_invalid_src_ocl(): }" with pytest.raises(dpctl_prog.SyclProgramCompilationError): dpctl_prog.create_program_from_source(q, invalid_oclSrc) + + +def test_create_program_from_sycl_source(): + try: + q = dpctl.SyclQueue("level_zero") + except dpctl.SyclQueueCreationError: + pytest.skip("No Level-zero queue is available") + + if not q.get_sycl_device().can_compile("sycl"): + pytest.skip("SYCL source compilation not supported") + + sycl_source = """ + #include + #include "math_ops.hpp" + #include "math_template_ops.hpp" + + namespace syclext = sycl::ext::oneapi::experimental; + + extern "C" SYCL_EXTERNAL + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclext::nd_range_kernel<1>)) + void vector_add(int* in1, int* in2, int* out){ + sycl::nd_item<1> item = + sycl::ext::oneapi::this_work_item::get_nd_item<1>(); + size_t globalID = item.get_global_linear_id(); + out[globalID] = math_op(in1[globalID],in2[globalID]); + } + + template + SYCL_EXTERNAL + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclext::nd_range_kernel<1>)) + void vector_add_template(T* in1, T* in2, T* out){ + sycl::nd_item<1> item = + sycl::ext::oneapi::this_work_item::get_nd_item<1>(); + size_t globalID = item.get_global_linear_id(); + out[globalID] = math_op_template(in1[globalID], in2[globalID]); + } + """ + + header_content = """ + int math_op(int a, int b){ + return a + b; + } + """ + + header2_content = """ + template + T math_op_template(T a, T b){ + return a + b; + } + """ + + prog = dpctl.program.create_program_from_sycl_source( + q, + sycl_source, + headers=[ + ("math_ops.hpp", header_content), + ("math_template_ops.hpp", header2_content), + ], + registered_names=["vector_add_template"], + copts=["-fno-fast-math"], + ) + + assert type(prog) is dpctl_prog.SyclProgram + + assert type(prog.addressof_ref()) is int + assert prog.has_sycl_kernel("vector_add") + 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) diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_device_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_device_interface.h index 6fddb2967f..b712e7cb41 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_device_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_device_interface.h @@ -792,4 +792,40 @@ DPCTL_API __dpctl_give DPCTLDeviceVectorRef DPCTLDevice_GetComponentDevices(__dpctl_keep const DPCTLSyclDeviceRef DRef); +/*! + * @brief Checks whether it is possible to create executables kernel bundles + * from SPIR-V binaries on this device. + * + * @param DRef Opaque pointer to a ``sycl::device``. + * @return True if creation is supported. + * #DPCTLSyclDeviceRef objects + * @ingroup DeviceInterface + */ +DPCTL_API +bool DPCTLDevice_CanCompileSPIRV(__dpctl_keep const DPCTLSyclDeviceRef DRef); + +/*! + * @brief Checks whether it is possible to create executables kernel bundles + * from OpenCL source code on this device. + * + * @param DRef Opaque pointer to a ``sycl::device``. + * @return True if creation is supported. + * #DPCTLSyclDeviceRef objects + * @ingroup DeviceInterface + */ +DPCTL_API +bool DPCTLDevice_CanCompileOpenCL(__dpctl_keep const DPCTLSyclDeviceRef DRef); + +/*! + * @brief Checks whether it is possible to create executables kernel bundles + * from SYCL source code on this device. + * + * @param DRef Opaque pointer to a ``sycl::device``. + * @return True if creation is supported. + * #DPCTLSyclDeviceRef objects + * @ingroup DeviceInterface + */ +DPCTL_API +bool DPCTLDevice_CanCompileSYCL(__dpctl_keep const DPCTLSyclDeviceRef DRef); + DPCTL_C_EXTERN_C_END diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h index 529bc3cca1..32cd289f20 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h @@ -129,4 +129,144 @@ DPCTL_API __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_Copy(__dpctl_keep const DPCTLSyclKernelBundleRef KBRef); +typedef struct DPCTLBuildOptionList *DPCTLBuildOptionListRef; +typedef struct DPCTLKernelNameList *DPCTLKernelNameListRef; +typedef struct DPCTLVirtualHeaderList *DPCTLVirtualHeaderListRef; + +/*! + * @brief Create an empty list of build options. + * + * @return Opaque pointer to the build option file list. + * @ingroup KernelBundleInterface + */ +DPCTL_API +__dpctl_give DPCTLBuildOptionListRef DPCTLBuildOptionList_Create(); + +/*! + * @brief Frees the DPCTLBuildOptionListRef pointer. + * + * @param KBRef Opaque pointer to a list of build options + * @ingroup KernelBundleInterface + */ +DPCTL_API void +DPCTLBuildOptionList_Delete(__dpctl_take DPCTLBuildOptionListRef Ref); + +/*! + * @brief Append a build option to the list of build options + * + * @param Ref Opaque pointer to the list of build options + * @param Option Option to append + */ +DPCTL_API +void DPCTLBuildOptionList_Append(__dpctl_keep DPCTLBuildOptionListRef Ref, + __dpctl_keep const char *Option); + +/*! + * @brief Create an empty list of kernel names to register. + * + * @return Opaque pointer to the list of kernel names to register. + * @ingroup KernelBundleInterface + */ +DPCTL_API +__dpctl_give DPCTLKernelNameListRef DPCTLKernelNameList_Create(); + +/*! + * @brief Frees the DPCTLKernelNameListRef pointer. + * + * @param KBRef Opaque pointer to a list of kernels to register + * @ingroup KernelBundleInterface + */ +DPCTL_API void +DPCTLKernelNameList_Delete(__dpctl_take DPCTLKernelNameListRef Ref); + +/*! + * @brief Append a kernel name to register to the list of build options + * + * @param Ref Opaque pointer to the list of kernel names + * @param Option Kernel name to append + */ +DPCTL_API +void DPCTLKernelNameList_Append(__dpctl_keep DPCTLKernelNameListRef Ref, + __dpctl_keep const char *Option); +/*! + * @brief Create an empty list of virtual header files. + * + * @return Opaque pointer to the virtual header file list. + * @ingroup KernelBundleInterface + */ +DPCTL_API +__dpctl_give DPCTLVirtualHeaderListRef DPCTLVirtualHeaderList_Create(); + +/*! + * @brief Frees the DPCTLVirtualHeaderListRef pointer. + * + * @param KBRef Opaque pointer to a list of virtual headers + * @ingroup KernelBundleInterface + */ +DPCTL_API void +DPCTLVirtualHeaderList_Delete(__dpctl_take DPCTLVirtualHeaderListRef Ref); + +/*! + * @brief Append a kernel name to register to the list of virtual header files + * + * @param Ref Opaque pointer to the list of header files + * @param Name Name of the virtual header file + * @param Content Content of the virtual header + */ +DPCTL_API +void DPCTLVirtualHeaderList_Append(__dpctl_keep DPCTLVirtualHeaderListRef Ref, + __dpctl_keep const char *Name, + __dpctl_keep const char *Content); + +/*! + * @brief Create a SYCL kernel bundle from an SYCL kernel source string. + * + * @param Ctx An opaque pointer to a sycl::context + * @param Dev An opaque pointer to a sycl::device + * @param Source SYCL source string + * @param Headers List of virtual headers + * @param Names List of kernel names to register + * @param CompileOpts List of extra compiler flags (refer Sycl spec.) + * @return A new SyclKernelBundleRef pointer if the program creation + * succeeded, else returns NULL. + * @ingroup KernelBundleInterface + */ +DPCTL_API +__dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource( + __dpctl_keep const DPCTLSyclContextRef Ctx, + __dpctl_keep const DPCTLSyclDeviceRef Dev, + __dpctl_keep const char *Source, + __dpctl_keep DPCTLVirtualHeaderListRef Headers, + __dpctl_keep DPCTLKernelNameListRef Names, + __dpctl_keep DPCTLBuildOptionListRef BuildOptions); + +/*! + * @brief Returns the SyclKernel with given name from the program compiled from + * SYCL source code, if not found then return NULL. + * + * @param KBRef Opaque pointer to a sycl::kernel_bundle + * @param KernelName Name of kernel + * @return A SyclKernel reference if the kernel exists, else NULL + * @ingroup KernelBundleInterface + */ +DPCTL_API +__dpctl_give DPCTLSyclKernelRef +DPCTLKernelBundle_GetSyclKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, + __dpctl_keep const char *KernelName); + +/*! + * @brief Return True if a SyclKernel with given name exists in the program + * compiled from SYCL source code, if not found then returns False. + * + * @param KBRef Opaque pointer to a sycl::kernel_bundle + * @param KernelName Name of kernel + * @return True if the kernel exists, else False + * @ingroup KernelBundleInterface + */ + +DPCTL_API +bool DPCTLKernelBundle_HasSyclKernel(__dpctl_keep DPCTLSyclKernelBundleRef + KBRef, + __dpctl_keep const char *KernelName); + DPCTL_C_EXTERN_C_END diff --git a/libsyclinterface/source/dpctl_sycl_device_interface.cpp b/libsyclinterface/source/dpctl_sycl_device_interface.cpp index 7b1e900b58..7eff0b242d 100644 --- a/libsyclinterface/source/dpctl_sycl_device_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_device_interface.cpp @@ -903,3 +903,28 @@ DPCTLDevice_GetCompositeDevice(__dpctl_keep const DPCTLSyclDeviceRef DRef) else return nullptr; } + +bool DPCTLDevice_CanCompileSPIRV(__dpctl_keep const DPCTLSyclDeviceRef DRef) +{ + auto Dev = unwrap(DRef); + auto Backend = Dev->get_platform().get_backend(); + return Backend == backend::opencl || + Backend == backend::ext_oneapi_level_zero; +} + +bool DPCTLDevice_CanCompileOpenCL(__dpctl_keep const DPCTLSyclDeviceRef DRef) +{ + auto Dev = unwrap(DRef); + return Dev->get_platform().get_backend() == backend::opencl; +} + +bool DPCTLDevice_CanCompileSYCL(__dpctl_keep const DPCTLSyclDeviceRef DRef) +{ +#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER + auto Dev = unwrap(DRef); + return Dev->ext_oneapi_can_compile( + ext::oneapi::experimental::source_language::sycl); +#else + return false; +#endif +} diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp index c702018687..d8121da1d8 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -761,3 +761,227 @@ 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); +} + +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) +{ +#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) { + error_handler(e, __FILE__, __func__, __LINE__); + return nullptr; + } +#else + return nullptr; +#endif +} + +__dpctl_give DPCTLSyclKernelRef +DPCTLKernelBundle_GetSyclKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, + __dpctl_keep const char *KernelName) +{ +#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..902d63a3d2 100644 --- a/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp @@ -273,6 +273,130 @@ TEST_P(TestOCLKernelBundleFromSource, CheckGetKernelOCLSource) DPCTLKernel_Delete(AxpyKernel); } +struct TestSYCLKernelBundleFromSource + : public ::testing::TestWithParam +{ + const char *sycl_source = R"===( + #include + #include "math_ops.hpp" + #include "math_template_ops.hpp" + + namespace syclext = sycl::ext::oneapi::experimental; + + extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclext::nd_range_kernel<1>)) + void vector_add(int* in1, int* in2, int* out){ + sycl::nd_item<1> item = sycl::ext::oneapi::this_work_item::get_nd_item<1>(); + size_t globalID = item.get_global_linear_id(); + out[globalID] = math_op(in1[globalID],in2[globalID]); + } + + template + SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclext::nd_range_kernel<1>)) + void vector_add_template(T* in1, T* in2, T* out){ + sycl::nd_item<1> item = sycl::ext::oneapi::this_work_item::get_nd_item<1>(); + size_t globalID = item.get_global_linear_id(); + out[globalID] = math_op_template(in1[globalID], in2[globalID]); + } + )==="; + + const char *header1_content = R"===( + int math_op(int a, int b){ + return a + b; + } + )==="; + + const char *header2_content = R"===( + template + T math_op_template(T a, T b){ + return a + b; + } + )==="; + + const char *CompileOpt = "-fno-fast-math"; + const char *KernelName = "vector_add_template"; + const char *Header1Name = "math_ops.hpp"; + const char *Header2Name = "math_template_ops.hpp"; + DPCTLSyclDeviceRef DRef = nullptr; + DPCTLSyclContextRef CRef = nullptr; + DPCTLSyclKernelBundleRef KBRef = nullptr; + + TestSYCLKernelBundleFromSource() + { + auto DS = DPCTLFilterSelector_Create(GetParam()); + DRef = DPCTLDevice_CreateFromSelector(DS); + DPCTLDeviceSelector_Delete(DS); + CRef = DPCTLDeviceMgr_GetCachedContext(DRef); + + if (DRef) { + DPCTLBuildOptionListRef BORef = DPCTLBuildOptionList_Create(); + DPCTLBuildOptionList_Append(BORef, CompileOpt); + DPCTLKernelNameListRef KNRef = DPCTLKernelNameList_Create(); + DPCTLKernelNameList_Append(KNRef, KernelName); + DPCTLVirtualHeaderListRef VHRef = DPCTLVirtualHeaderList_Create(); + DPCTLVirtualHeaderList_Append(VHRef, Header1Name, header1_content); + DPCTLVirtualHeaderList_Append(VHRef, Header2Name, header2_content); + KBRef = DPCTLKernelBundle_CreateFromSYCLSource( + CRef, DRef, sycl_source, VHRef, KNRef, BORef); + DPCTLVirtualHeaderList_Delete(VHRef); + DPCTLKernelNameList_Delete(KNRef); + DPCTLBuildOptionList_Delete(BORef); + } + } + + void SetUp() + { + if (!DRef) { + auto message = "Skipping as no device of type " + + std::string(GetParam()) + "."; + GTEST_SKIP_(message.c_str()); + } + if (!DPCTLDevice_CanCompileSYCL(DRef)) { + const char *message = "Skipping as SYCL compilation not supported"; + GTEST_SKIP_(message); + } + } + + ~TestSYCLKernelBundleFromSource() + { + if (DRef) + DPCTLDevice_Delete(DRef); + if (CRef) + DPCTLContext_Delete(CRef); + if (KBRef) + DPCTLKernelBundle_Delete(KBRef); + } +}; + +TEST_P(TestSYCLKernelBundleFromSource, CheckCreateFromSYCLSource) +{ + ASSERT_TRUE(KBRef != nullptr); + ASSERT_TRUE(DPCTLKernelBundle_HasSyclKernel(KBRef, "vector_add")); + // 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 +413,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;