From a8e316ed0bc9d15c787ee2940a68c7c6290cb177 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Fri, 20 May 2022 12:42:02 -0500 Subject: [PATCH 01/14] Changes to delete DPCTLKernel_GetFunctionName --- dpctl/_backend.pxd | 1 - dpctl/program/_program.pxd | 4 ++-- dpctl/program/_program.pyx | 10 ++++------ 3 files changed, 6 insertions(+), 9 deletions(-) diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index eb9124ffa5..e3fbcb7c10 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -250,7 +250,6 @@ cdef extern from "syclinterface/dpctl_sycl_event_interface.h": cdef extern from "syclinterface/dpctl_sycl_kernel_interface.h": - cdef const char* DPCTLKernel_GetFunctionName(const DPCTLSyclKernelRef KRef) cdef size_t DPCTLKernel_GetNumArgs(const DPCTLSyclKernelRef KRef) cdef void DPCTLKernel_Delete(DPCTLSyclKernelRef KRef) diff --git a/dpctl/program/_program.pxd b/dpctl/program/_program.pxd index e0b035793c..1e3b86ffd8 100644 --- a/dpctl/program/_program.pxd +++ b/dpctl/program/_program.pxd @@ -33,11 +33,11 @@ cdef class SyclKernel: kernel. ''' cdef DPCTLSyclKernelRef _kernel_ref - cdef const char *_function_name + cdef str _function_name cdef DPCTLSyclKernelRef get_kernel_ref (self) @staticmethod - cdef SyclKernel _create (DPCTLSyclKernelRef kref) + cdef SyclKernel _create (DPCTLSyclKernelRef kref, str name) cdef class SyclProgram: diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index 923deef367..ce42ea786a 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -30,7 +30,6 @@ cimport cython.array from dpctl._backend cimport ( # noqa: E211, E402 DPCTLCString_Delete, DPCTLKernel_Delete, - DPCTLKernel_GetFunctionName, DPCTLKernel_GetNumArgs, DPCTLProgram_CreateFromOCLSource, DPCTLProgram_CreateFromSpirv, @@ -61,20 +60,19 @@ cdef class SyclKernel: """ """ @staticmethod - cdef SyclKernel _create(DPCTLSyclKernelRef kref): + cdef SyclKernel _create(DPCTLSyclKernelRef kref, str name): cdef SyclKernel ret = SyclKernel.__new__(SyclKernel) ret._kernel_ref = kref - ret._function_name = DPCTLKernel_GetFunctionName(kref) + ret._function_name = name return ret def __dealloc__(self): DPCTLKernel_Delete(self._kernel_ref) - DPCTLCString_Delete(self._function_name) def get_function_name(self): """ Returns the name of the ``sycl::kernel`` function. """ - return self._function_name.decode() + return self._function_name def get_num_args(self): """ Returns the number of arguments for this kernel function. @@ -121,7 +119,7 @@ cdef class SyclProgram: cpdef SyclKernel get_sycl_kernel(self, str kernel_name): name = kernel_name.encode('utf8') return SyclKernel._create(DPCTLProgram_GetKernel(self._program_ref, - name)) + name), kernel_name) def has_sycl_kernel(self, str kernel_name): name = kernel_name.encode('utf8') From b5b06eccbeb36db0f55f1e4a97b7648f1670a4c6 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Fri, 20 May 2022 12:45:02 -0500 Subject: [PATCH 02/14] Removed DPCTLKernel_GetFunctionName --- .../include/dpctl_sycl_kernel_interface.h | 13 ---------- .../source/dpctl_sycl_kernel_interface.cpp | 17 ------------ .../tests/test_sycl_kernel_interface.cpp | 26 ------------------- 3 files changed, 56 deletions(-) diff --git a/libsyclinterface/include/dpctl_sycl_kernel_interface.h b/libsyclinterface/include/dpctl_sycl_kernel_interface.h index 582b19b6ef..3bc0470bd7 100644 --- a/libsyclinterface/include/dpctl_sycl_kernel_interface.h +++ b/libsyclinterface/include/dpctl_sycl_kernel_interface.h @@ -38,19 +38,6 @@ DPCTL_C_EXTERN_C_BEGIN * @defgroup KernelInterface Kernel class C wrapper */ -/*! - * @brief Returns a C string for the kernel name. - * - * @param KRef DPCTLSyclKernelRef pointer to an OpenCL - * interoperability kernel. - * @return If a kernel name exists then returns it as a C string, else - * returns a nullptr. - * @ingroup KernelInterface - */ -DPCTL_API -__dpctl_give const char * -DPCTLKernel_GetFunctionName(__dpctl_keep const DPCTLSyclKernelRef KRef); - /*! * @brief Returns the number of arguments for the OpenCL kernel. * diff --git a/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp index 48d8ff919e..e2fc31bba1 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp @@ -39,23 +39,6 @@ DEFINE_SIMPLE_CONVERSION_FUNCTIONS(kernel, DPCTLSyclKernelRef) } /* end of anonymous namespace */ -__dpctl_give const char * -DPCTLKernel_GetFunctionName(__dpctl_keep const DPCTLSyclKernelRef Kernel) -{ - if (!Kernel) { - error_handler("Cannot get the number of arguments " - "as input is a nullptr.", - __FILE__, __func__, __LINE__); - return nullptr; - } - - auto SyclKernel = unwrap(Kernel); - auto kernel_name = SyclKernel->get_info(); - if (kernel_name.empty()) - return nullptr; - return dpctl::helper::cstring_from_string(kernel_name); -} - size_t DPCTLKernel_GetNumArgs(__dpctl_keep const DPCTLSyclKernelRef Kernel) { if (!Kernel) { diff --git a/libsyclinterface/tests/test_sycl_kernel_interface.cpp b/libsyclinterface/tests/test_sycl_kernel_interface.cpp index 3a5f5b5f0f..ca53c130fb 100644 --- a/libsyclinterface/tests/test_sycl_kernel_interface.cpp +++ b/libsyclinterface/tests/test_sycl_kernel_interface.cpp @@ -81,31 +81,6 @@ struct TestDPCTLSyclKernelInterface }; } // namespace -TEST_P(TestDPCTLSyclKernelInterface, CheckGetFunctionName) -{ - auto QueueRef = DPCTLQueue_CreateForDevice(DRef, nullptr, 0); - auto CtxRef = DPCTLQueue_GetContext(QueueRef); - auto PRef = - DPCTLProgram_CreateFromOCLSource(CtxRef, CLProgramStr, CompileOpts); - auto AddKernel = DPCTLProgram_GetKernel(PRef, "add"); - auto AxpyKernel = DPCTLProgram_GetKernel(PRef, "axpy"); - - auto fnName1 = DPCTLKernel_GetFunctionName(AddKernel); - auto fnName2 = DPCTLKernel_GetFunctionName(AxpyKernel); - - ASSERT_STREQ("add", fnName1); - ASSERT_STREQ("axpy", fnName2); - - DPCTLCString_Delete(fnName1); - DPCTLCString_Delete(fnName2); - - DPCTLQueue_Delete(QueueRef); - DPCTLContext_Delete(CtxRef); - DPCTLProgram_Delete(PRef); - DPCTLKernel_Delete(AddKernel); - DPCTLKernel_Delete(AxpyKernel); -} - TEST_P(TestDPCTLSyclKernelInterface, CheckGetNumArgs) { auto QueueRef = DPCTLQueue_CreateForDevice(DRef, nullptr, 0); @@ -130,7 +105,6 @@ TEST_P(TestDPCTLSyclKernelInterface, CheckNullPtrArg) DPCTLSyclKernelRef AddKernel = nullptr; ASSERT_EQ(DPCTLKernel_GetNumArgs(AddKernel), -1); - ASSERT_EQ(DPCTLKernel_GetFunctionName(AddKernel), nullptr); } INSTANTIATE_TEST_SUITE_P(TestKernelInterfaceFunctions, From c8dc538eb63d5beba0fc1383cef01e2fe17967a4 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Mon, 23 May 2022 18:04:56 -0500 Subject: [PATCH 03/14] Refactoring of program support in dpctl and SyclInterface Removed uses of `sycl::program` throughout the code. In SyclInterface replace DPCTLSyclProgramRef with DPCTLSyclKernelBundleRef which is reference to `sycl::kernel_bundle`. Functions `DPCTLProgram_*` were replaced with `DCPTLKernelBundle_*`. Functions to create program now take both context and device. Tests were modified. dpctl.SyclProgram stays with this name, but it now encapsulates DPCTLSyclKernelBundleRef instead of removed DPCTLSyclProgramRef. OpenCL functions are no longer directly used (instead of used via loader, like in the case of level-zero backend), hence removed linkage to OpenCL library. --- dpctl/_backend.pxd | 18 +- dpctl/program/_program.pxd | 12 +- dpctl/program/_program.pyx | 83 +- libsyclinterface/CMakeLists.txt | 18 +- .../helper/include/dpctl_dynamic_lib_helper.h | 4 +- .../include/Config/dpctl_config.h.in | 4 +- .../include/dpctl_sycl_program_interface.h | 70 +- libsyclinterface/include/dpctl_sycl_types.h | 8 +- .../source/dpctl_sycl_program_interface.cpp | 721 ++++++++++++++---- .../tests/test_sycl_kernel_interface.cpp | 10 +- .../tests/test_sycl_program_interface.cpp | 65 +- .../tests/test_sycl_queue_submit.cpp | 24 +- 12 files changed, 747 insertions(+), 290 deletions(-) diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index e3fbcb7c10..7781d9bf73 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -120,7 +120,7 @@ cdef extern from "syclinterface/dpctl_sycl_types.h": cdef struct DPCTLOpaqueSyclEvent cdef struct DPCTLOpaqueSyclKernel cdef struct DPCTLOpaqueSyclPlatform - cdef struct DPCTLOpaqueSyclProgram + cdef struct DPCTLOpaqueSyclKernelBundle cdef struct DPCTLOpaqueSyclQueue cdef struct DPCTLOpaqueSyclUSM @@ -130,7 +130,7 @@ cdef extern from "syclinterface/dpctl_sycl_types.h": ctypedef DPCTLOpaqueSyclEvent *DPCTLSyclEventRef ctypedef DPCTLOpaqueSyclKernel *DPCTLSyclKernelRef ctypedef DPCTLOpaqueSyclPlatform *DPCTLSyclPlatformRef - ctypedef DPCTLOpaqueSyclProgram *DPCTLSyclProgramRef + ctypedef DPCTLOpaqueSyclKernelBundle *DPCTLSyclKernelBundleRef ctypedef DPCTLOpaqueSyclQueue *DPCTLSyclQueueRef ctypedef DPCTLOpaqueSyclUSM *DPCTLSyclUSMRef @@ -305,21 +305,23 @@ cdef extern from "syclinterface/dpctl_sycl_context_interface.h": cdef extern from "syclinterface/dpctl_sycl_program_interface.h": - cdef DPCTLSyclProgramRef DPCTLProgram_CreateFromSpirv( + cdef DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSpirv( const DPCTLSyclContextRef Ctx, + const DPCTLSyclDeviceRef Dev, const void *IL, size_t Length, const char *CompileOpts) - cdef DPCTLSyclProgramRef DPCTLProgram_CreateFromOCLSource( + cdef DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromOCLSource( const DPCTLSyclContextRef Ctx, + const DPCTLSyclDeviceRef Dev, const char *Source, const char *CompileOpts) - cdef DPCTLSyclKernelRef DPCTLProgram_GetKernel( - DPCTLSyclProgramRef PRef, + cdef DPCTLSyclKernelRef DPCTLKernelBundle_GetKernel( + DPCTLSyclKernelBundleRef KBRef, const char *KernelName) - cdef bool DPCTLProgram_HasKernel(DPCTLSyclProgramRef PRef, + cdef bool DPCTLKernelBundle_HasKernel(DPCTLSyclKernelBundleRef KBRef, const char *KernelName) - cdef void DPCTLProgram_Delete(DPCTLSyclProgramRef PRef) + cdef void DPCTLKernelBundle_Delete(DPCTLSyclKernelBundleRef KBRef) cdef extern from "syclinterface/dpctl_sycl_queue_interface.h": diff --git a/dpctl/program/_program.pxd b/dpctl/program/_program.pxd index 1e3b86ffd8..dda9e2662c 100644 --- a/dpctl/program/_program.pxd +++ b/dpctl/program/_program.pxd @@ -22,7 +22,7 @@ """ -from .._backend cimport DPCTLSyclKernelRef, DPCTLSyclProgramRef +from .._backend cimport DPCTLSyclKernelBundleRef, DPCTLSyclKernelRef from .._sycl_context cimport SyclContext from .._sycl_device cimport SyclDevice from .._sycl_queue cimport SyclQueue @@ -41,18 +41,18 @@ cdef class SyclKernel: cdef class SyclProgram: - ''' Wraps a sycl::program object created from an OpenCL interoperability - program. + ''' Wraps a sycl::kernel_bundle object created from + using SYCL interoperability layer for OpenCL and Level-Zero backends. SyclProgram exposes the C API from dpctl_sycl_program_interface.h. A SyclProgram can be created from either a source string or a SPIR-V binary file. ''' - cdef DPCTLSyclProgramRef _program_ref + cdef DPCTLSyclKernelBundleRef _program_ref @staticmethod - cdef SyclProgram _create (DPCTLSyclProgramRef pref) - cdef DPCTLSyclProgramRef get_program_ref (self) + cdef SyclProgram _create (DPCTLSyclKernelBundleRef pref) + cdef DPCTLSyclKernelBundleRef get_program_ref (self) cpdef SyclKernel get_sycl_kernel(self, str kernel_name) diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index ce42ea786a..b953a1cf3e 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -31,14 +31,15 @@ from dpctl._backend cimport ( # noqa: E211, E402 DPCTLCString_Delete, DPCTLKernel_Delete, DPCTLKernel_GetNumArgs, - DPCTLProgram_CreateFromOCLSource, - DPCTLProgram_CreateFromSpirv, - DPCTLProgram_Delete, - DPCTLProgram_GetKernel, - DPCTLProgram_HasKernel, + DPCTLKernelBundle_CreateFromOCLSource, + DPCTLKernelBundle_CreateFromSpirv, + DPCTLKernelBundle_Delete, + DPCTLKernelBundle_GetKernel, + DPCTLKernelBundle_HasKernel, DPCTLSyclContextRef, + DPCTLSyclDeviceRef, + DPCTLSyclKernelBundleRef, DPCTLSyclKernelRef, - DPCTLSyclProgramRef, ) __all__ = [ @@ -50,8 +51,8 @@ __all__ = [ ] cdef class SyclProgramCompilationError(Exception): - """This exception is raised when a ``sycl::program`` could not be built from - either a SPIR-V binary file or a string source. + """This exception is raised when a ``sycl::kernel_bundle`` could not be + built from either a SPIR-V binary file or a string source. """ pass @@ -105,33 +106,35 @@ cdef class SyclProgram: """ @staticmethod - cdef SyclProgram _create(DPCTLSyclProgramRef pref): + cdef SyclProgram _create(DPCTLSyclKernelBundleRef KBRef): cdef SyclProgram ret = SyclProgram.__new__(SyclProgram) - ret._program_ref = pref + ret._program_ref = KBRef return ret def __dealloc__(self): - DPCTLProgram_Delete(self._program_ref) + DPCTLKernelBundle_Delete(self._program_ref) - cdef DPCTLSyclProgramRef get_program_ref(self): + cdef DPCTLSyclKernelBundleRef get_program_ref(self): return self._program_ref cpdef SyclKernel get_sycl_kernel(self, str kernel_name): name = kernel_name.encode('utf8') - return SyclKernel._create(DPCTLProgram_GetKernel(self._program_ref, - name), kernel_name) + return SyclKernel._create( + DPCTLKernelBundle_GetKernel(self._program_ref, name), + kernel_name + ) def has_sycl_kernel(self, str kernel_name): name = kernel_name.encode('utf8') - return DPCTLProgram_HasKernel(self._program_ref, name) + return DPCTLKernelBundle_HasKernel(self._program_ref, name) def addressof_ref(self): - """Returns the address of the C API DPCTLSyclProgramRef pointer + """Returns the address of the C API DPCTLSyclKernelBundleRef pointer as a long. Returns: - The address of the ``DPCTLSyclProgramRef`` pointer used to create - this :class:`dpctl.SyclProgram` object cast to a ``size_t``. + The address of the ``DPCTLSyclKernelBundleRef`` pointer used to + create this :class:`dpctl.SyclProgram` object cast to a ``size_t``. """ return int(self._program_ref) @@ -140,9 +143,10 @@ cpdef create_program_from_source(SyclQueue q, unicode src, unicode copts=""): """ Creates a Sycl interoperability program from an OpenCL source string. - We use the ``DPCTLProgram_CreateFromOCLSource()`` C API function to - create a ``sycl::program`` from an OpenCL source program that can - contain multiple kernels. Note currently only supported for OpenCL. + We use the ``DPCTLKernelBundle_CreateFromOCLSource()`` C API function + to create a ``sycl::kernel_bundle`` + from an OpenCL source program that can contain multiple kernels. + Note: This function is currently only supported for the OpenCL backend. Parameters: q (SyclQueue) : The :class:`SyclQueue` for which the @@ -153,24 +157,27 @@ cpdef create_program_from_source(SyclQueue q, unicode src, unicode copts=""): Returns: program (SyclProgram): A :class:`SyclProgram` object wrapping the - ``sycl::program`` returned by the C API. + ``sycl::kernel_bundle`` returned + by the C API. Raises: - SyclProgramCompilationError: If a SYCL program could not be created. + SyclProgramCompilationError: If a SYCL kernel bundle could not be + created. """ - cdef DPCTLSyclProgramRef Pref + cdef DPCTLSyclKernelBundleRef KBref cdef bytes bSrc = src.encode('utf8') cdef bytes bCOpts = copts.encode('utf8') cdef const char *Src = bSrc cdef const char *COpts = bCOpts cdef DPCTLSyclContextRef CRef = q.get_sycl_context().get_context_ref() - Pref = DPCTLProgram_CreateFromOCLSource(CRef, Src, COpts) + cdef DPCTLSyclDeviceRef DRef = q.get_sycl_device().get_device_ref() + KBref = DPCTLKernelBundle_CreateFromOCLSource(CRef, DRef, Src, COpts) - if Pref is NULL: + if KBref is NULL: raise SyclProgramCompilationError() - return SyclProgram._create(Pref) + return SyclProgram._create(KBref) cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL, @@ -178,8 +185,9 @@ cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL, """ Creates a Sycl interoperability program from an SPIR-V binary. - We use the ``DPCTLProgram_CreateFromOCLSpirv()`` C API function to - create a ``sycl::program`` object from an compiled SPIR-V binary file. + We use the ``DPCTLKernelBundle_CreateFromOCLSpirv()`` C API function to + create a ``sycl::kernel_bundle`` object + from an compiled SPIR-V binary file. Parameters: q (SyclQueue): The :class:`SyclQueue` for which the @@ -190,20 +198,25 @@ cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL, Returns: program (SyclProgram): A :class:`SyclProgram` object wrapping the - ``sycl::program`` returned by the C API. + ``sycl::kernel_bundle`` returned by + the C API. Raises: - SyclProgramCompilationError: If a SYCL program could not be created. + SyclProgramCompilationError: If a SYCL kernel bundle could not be + created. """ - cdef DPCTLSyclProgramRef Pref + cdef DPCTLSyclKernelBundleRef KBref cdef const unsigned char *dIL = &IL[0] cdef DPCTLSyclContextRef CRef = q.get_sycl_context().get_context_ref() + cdef DPCTLSyclDeviceRef DRef = q.get_sycl_device().get_device_ref() cdef size_t length = IL.shape[0] cdef bytes bCOpts = copts.encode('utf8') cdef const char *COpts = bCOpts - Pref = DPCTLProgram_CreateFromSpirv(CRef, dIL, length, COpts) - if Pref is NULL: + KBref = DPCTLKernelBundle_CreateFromSpirv( + CRef, DRef, dIL, length, COpts + ) + if KBref is NULL: raise SyclProgramCompilationError() - return SyclProgram._create(Pref) + return SyclProgram._create(KBref) diff --git a/libsyclinterface/CMakeLists.txt b/libsyclinterface/CMakeLists.txt index 695975f7ae..05af34940b 100644 --- a/libsyclinterface/CMakeLists.txt +++ b/libsyclinterface/CMakeLists.txt @@ -71,6 +71,23 @@ if(DPCTL_ENABLE_L0_PROGRAM_CREATION) endif() endif() +if (UNIX) + find_library(PI_OPENCL_LIB + NAMES pi_opencl + HINTS ${IntelSycl_LIBRARY_DIR} + ) + find_program(READELF_PROG readelf) + find_program(GREP_PROG grep) + execute_process( + COMMAND ${READELF_PROG} -d ${PI_OPENCL_LIB} + COMMAND ${GREP_PROG} OpenCL + COMMAND ${GREP_PROG} -Po "libOpenCL[^\]]*" + OUTPUT_VARIABLE LIBCL_LOADER_FILENAME + OUTPUT_STRIP_TRAILING_WHITESPACE + ERROR_STRIP_TRAILING_WHITESPACE + ) +endif() + configure_file( ${CMAKE_CURRENT_SOURCE_DIR}/include/Config/dpctl_config.h.in ${CMAKE_CURRENT_SOURCE_DIR}/include/Config/dpctl_config.h @@ -193,7 +210,6 @@ target_include_directories(DPCTLSyclInterface ) target_link_libraries(DPCTLSyclInterface PRIVATE ${IntelSycl_SYCL_LIBRARY} - PRIVATE ${IntelSycl_OPENCL_LIBRARY} ) if(DPCTL_ENABLE_GLOG) diff --git a/libsyclinterface/helper/include/dpctl_dynamic_lib_helper.h b/libsyclinterface/helper/include/dpctl_dynamic_lib_helper.h index 864ef326a8..c1a725f8d3 100644 --- a/libsyclinterface/helper/include/dpctl_dynamic_lib_helper.h +++ b/libsyclinterface/helper/include/dpctl_dynamic_lib_helper.h @@ -75,13 +75,13 @@ class DynamicLibHelper final void *sym = dlsym(_handle, symName); char *error = dlerror(); - if (NULL != error) { + if (nullptr != error) { return nullptr; } #elif defined(_WIN32) || defined(_WIN64) void *sym = (void *)GetProcAddress((HMODULE)_handle, symName); - if (NULL == sym) { + if (nullptr == sym) { return nullptr; } #endif diff --git a/libsyclinterface/include/Config/dpctl_config.h.in b/libsyclinterface/include/Config/dpctl_config.h.in index c50b22db87..d630060b9e 100644 --- a/libsyclinterface/include/Config/dpctl_config.h.in +++ b/libsyclinterface/include/Config/dpctl_config.h.in @@ -26,9 +26,11 @@ #pragma once /* Defined when dpctl was built with level zero program creation enabled. */ -#cmakedefine DPCTL_ENABLE_L0_PROGRAM_CREATION @DPCTL_ENABLE_L0_PROGRAM_CREATION@ +#cmakedefine DPCTL_ENABLE_L0_PROGRAM_CREATION \ + @DPCTL_ENABLE_L0_PROGRAM_CREATION @ /* The DPCPP version used to build dpctl */ #define DPCTL_DPCPP_VERSION "@IntelSycl_VERSION@" #define DPCTL_LIBZE_LOADER_FILENAME "@LIBZE_LOADER_FILENAME@" +#define DPCTL_LIBCL_LOADER_FILENAME "@LIBCL_LOADER_FILENAME@" diff --git a/libsyclinterface/include/dpctl_sycl_program_interface.h b/libsyclinterface/include/dpctl_sycl_program_interface.h index 4667631a98..335d274140 100644 --- a/libsyclinterface/include/dpctl_sycl_program_interface.h +++ b/libsyclinterface/include/dpctl_sycl_program_interface.h @@ -35,52 +35,50 @@ DPCTL_C_EXTERN_C_BEGIN /** - * @defgroup ProgramInterface Program class C wrapper + * @defgroup KernelBundleInterface Kernel_bundle class C wrapper */ /*! - * @brief Create a Sycl program from an OpenCL SPIR-V binary file. + * @brief Create a Sycl kernel_bundle from an OpenCL SPIR-V binary file. * - * Sycl 1.2 does not expose any method to create a sycl::program from a SPIR-V - * IL file. To get around this limitation, we first creare a SYCL - * interoperability program and then create a SYCL program from the - * interoperability program. Currently, interoperability programs can be created - * for OpenCL and Level-0 backends. - * - * The feature to create a Sycl kernel from a SPIR-V IL binary will be available - * in Sycl 2.0 at which point this function may become deprecated. + * Uses SYCL2020 interoperability layer to create sycl::kernel_bundle object + * in executable state for OpenCL and Level-Zero backends from SPIR-V binary. * * @param Ctx An opaque pointer to a sycl::context + * @param Dev An opaque pointer to a sycl::device * @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. - * @return A new SyclProgramRef pointer if the program creation succeeded, - * else returns NULL. - * @ingroup ProgramInterface + * @return A new SyclKernelBundleRef pointer if the kernel_bundle creation + * succeeded, else returns NULL. + * @ingroup KernelBundleInterface */ DPCTL_API -__dpctl_give DPCTLSyclProgramRef -DPCTLProgram_CreateFromSpirv(__dpctl_keep const DPCTLSyclContextRef Ctx, - __dpctl_keep const void *IL, - size_t Length, - const char *CompileOpts); +__dpctl_give DPCTLSyclKernelBundleRef +DPCTLKernelBundle_CreateFromSpirv(__dpctl_keep const DPCTLSyclContextRef Ctx, + __dpctl_keep const DPCTLSyclDeviceRef Dev, + __dpctl_keep const void *IL, + size_t Length, + const char *CompileOpts); /*! - * @brief Create a Sycl program from an OpenCL kernel source string. + * @brief Create a Sycl kernel bundle from an OpenCL kernel source string. * * @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.) - * @return A new SyclProgramRef pointer if the program creation succeeded, - * else returns NULL. - * @ingroup ProgramInterface + * @return A new SyclKernelBundleRef pointer if the program creation + * succeeded, else returns NULL. + * @ingroup KernelBundleInterface */ DPCTL_API -__dpctl_give DPCTLSyclProgramRef -DPCTLProgram_CreateFromOCLSource(__dpctl_keep const DPCTLSyclContextRef Ctx, - __dpctl_keep const char *Source, - __dpctl_keep const char *CompileOpts); +__dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromOCLSource( + __dpctl_keep const DPCTLSyclContextRef Ctx, + __dpctl_keep const DPCTLSyclDeviceRef Dev, + __dpctl_keep const char *Source, + __dpctl_keep const char *CompileOpts); /*! * @brief Returns the SyclKernel with given name from the program, if not found @@ -89,12 +87,12 @@ DPCTLProgram_CreateFromOCLSource(__dpctl_keep const DPCTLSyclContextRef Ctx, * @param PRef Opaque pointer to a sycl::program * @param KernelName Name of kernel * @return A SyclKernel reference if the kernel exists, else NULL - * @ingroup ProgramInterface + * @ingroup KernelBundleInterface */ DPCTL_API __dpctl_give DPCTLSyclKernelRef -DPCTLProgram_GetKernel(__dpctl_keep DPCTLSyclProgramRef PRef, - __dpctl_keep const char *KernelName); +DPCTLKernelBundle_GetKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, + __dpctl_keep const char *KernelName); /*! * @brief Return True if a SyclKernel with given name exists in the program, if @@ -103,19 +101,19 @@ DPCTLProgram_GetKernel(__dpctl_keep DPCTLSyclProgramRef PRef, * @param PRef Opaque pointer to a sycl::program * @param KernelName Name of kernel * @return True if the kernel exists, else False - * @ingroup ProgramInterface + * @ingroup KernelBundleInterface */ DPCTL_API -bool DPCTLProgram_HasKernel(__dpctl_keep DPCTLSyclProgramRef PRef, - __dpctl_keep const char *KernelName); +bool DPCTLKernelBundle_HasKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, + __dpctl_keep const char *KernelName); /*! - * @brief Frees the DPCTLSyclProgramRef pointer. + * @brief Frees the DPCTLSyclKernelBundleRef pointer. * - * @param PRef Opaque pointer to a sycl::program - * @ingroup ProgramInterface + * @param PRef Opaque pointer to a sycl::kernel_bundle + * @ingroup KernelBundleInterface */ DPCTL_API -void DPCTLProgram_Delete(__dpctl_take DPCTLSyclProgramRef PRef); +void DPCTLKernelBundle_Delete(__dpctl_take DPCTLSyclKernelBundleRef KBRef); DPCTL_C_EXTERN_C_END diff --git a/libsyclinterface/include/dpctl_sycl_types.h b/libsyclinterface/include/dpctl_sycl_types.h index 12fadb901c..b00cab4702 100644 --- a/libsyclinterface/include/dpctl_sycl_types.h +++ b/libsyclinterface/include/dpctl_sycl_types.h @@ -60,16 +60,16 @@ typedef struct DPCTLOpaqueSyclEvent *DPCTLSyclEventRef; typedef struct DPCTLOpaqueSyclKernel *DPCTLSyclKernelRef; /*! - * @brief Opaque pointer to a ``sycl::platform`` + * @brief Opaque pointer to a ``sycl::kernel_bundle`` * */ -typedef struct DPCTLOpaqueSyclPlatform *DPCTLSyclPlatformRef; +typedef struct DPCTLOpaqueSyclKernelBundle *DPCTLSyclKernelBundleRef; /*! - * @brief Opaque pointer to a ``sycl::program`` + * @brief Opaque pointer to a ``sycl::platform`` * */ -typedef struct DPCTLOpaqueSyclProgram *DPCTLSyclProgramRef; +typedef struct DPCTLOpaqueSyclPlatform *DPCTLSyclPlatformRef; /*! * @brief Opaque pointer to a ``sycl::queue`` diff --git a/libsyclinterface/source/dpctl_sycl_program_interface.cpp b/libsyclinterface/source/dpctl_sycl_program_interface.cpp index 6a62e948c0..c2774515dd 100644 --- a/libsyclinterface/source/dpctl_sycl_program_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_program_interface.cpp @@ -24,14 +24,10 @@ /// //===----------------------------------------------------------------------===// -#ifndef __SYCL_INTERNAL_API -// make sure that sycl::program is defined and implemented -#define __SYCL_INTERNAL_API -#endif - #include "dpctl_sycl_program_interface.h" #include "Config/dpctl_config.h" #include "Support/CBindingWrapping.h" +#include "dpctl_dynamic_lib_helper.h" #include "dpctl_error_handlers.h" #include /* OpenCL headers */ #include /* Sycl headers */ @@ -39,7 +35,6 @@ #include #ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION -#include "dpctl_dynamic_lib_helper.h" // Note: include ze_api.h before level_zero.hpp. Make sure clang-format does // not reorder the includes. // clang-format off @@ -52,112 +47,442 @@ using namespace cl::sycl; namespace { -#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION +DEFINE_SIMPLE_CONVERSION_FUNCTIONS(context, DPCTLSyclContextRef) +DEFINE_SIMPLE_CONVERSION_FUNCTIONS(device, DPCTLSyclDeviceRef) +DEFINE_SIMPLE_CONVERSION_FUNCTIONS(kernel_bundle, + DPCTLSyclKernelBundleRef) +DEFINE_SIMPLE_CONVERSION_FUNCTIONS(kernel, DPCTLSyclKernelRef) #ifdef __linux__ -static const char *zeLoaderName = DPCTL_LIBZE_LOADER_FILENAME; -static const int libLoadFlags = RTLD_NOLOAD | RTLD_NOW | RTLD_LOCAL; +static const char *clLoaderName = DPCTL_LIBCL_LOADER_FILENAME; +static const int clLibLoadFlags = RTLD_NOLOAD | RTLD_NOW | RTLD_LOCAL; #elif defined(_WIN64) -static const char *zeLoaderName = "ze_loader.dll"; -static const int libLoadFlags = 0; +static const char *clLoaderName = "OpenCL.dll"; +static const int clLibLoadFlags = 0; #else -#error "Level Zero program compilation is unavailable for this platform" +#error "OpenCL program compilation is unavailable for this platform" #endif -typedef ze_result_t (*zeModuleCreateFT)(ze_context_handle_t, - ze_device_handle_t, - const ze_module_desc_t *, - ze_module_handle_t *, - ze_module_build_log_handle_t *); +typedef cl_program (*clCreateProgramWithSourceFT)(cl_context, + cl_uint, + const char **, + const size_t *, + cl_int *); +const char *clCreateProgramWithSource_Name = "clCreateProgramWithSource"; +clCreateProgramWithSourceFT get_clCreateProgramWithSource() +{ + static dpctl::DynamicLibHelper clLib(clLoaderName, clLibLoadFlags); + if (!clLib.opened()) { + error_handler("The OpenCL loader dynamic library could not " + "be opened.", + __FILE__, __func__, __LINE__); + return nullptr; + } -const char *zeModuleCreateFuncName = "zeModuleCreate"; + static auto st_clCreateProgramWithSourceF = + clLib.getSymbol( + clCreateProgramWithSource_Name); -#endif // #ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION + return st_clCreateProgramWithSourceF; +} -DEFINE_SIMPLE_CONVERSION_FUNCTIONS(context, DPCTLSyclContextRef) -DEFINE_SIMPLE_CONVERSION_FUNCTIONS(program, DPCTLSyclProgramRef) -DEFINE_SIMPLE_CONVERSION_FUNCTIONS(kernel, DPCTLSyclKernelRef) +typedef cl_program (*clCreateProgramWithILFT)(cl_context, + const void *, + size_t, + cl_int *); +const char *clCreateProgramWithIL_Name = "clCreateProgramWithIL"; +clCreateProgramWithILFT get_clCreateProgramWithIL() +{ + static dpctl::DynamicLibHelper clLib(clLoaderName, clLibLoadFlags); + if (!clLib.opened()) { + error_handler("The OpenCL loader dynamic library could not " + "be opened.", + __FILE__, __func__, __LINE__); + return nullptr; + } + static auto st_clCreateProgramWithILF = + clLib.getSymbol(clCreateProgramWithIL_Name); + + return st_clCreateProgramWithILF; +} +typedef cl_int (*clBuildProgramFT)(cl_program, + cl_uint, + const cl_device_id *, + const char *, + void (*)(cl_program, void *), + void *); +const char *clBuildProgram_Name = "clBuildProgram"; +clBuildProgramFT get_clBuldProgram() +{ + static dpctl::DynamicLibHelper clLib(clLoaderName, clLibLoadFlags); + if (!clLib.opened()) { + error_handler("The OpenCL loader dynamic library could not " + "be opened.", + __FILE__, __func__, __LINE__); + return nullptr; + } + static auto st_clBuildProgramF = + clLib.getSymbol(clBuildProgram_Name); + + return st_clBuildProgramF; +} -__dpctl_give DPCTLSyclProgramRef -createOpenCLInterOpProgram(const context &SyclCtx, - __dpctl_keep const void *IL, - size_t length, - const char *CompileOpts) +typedef cl_kernel (*clCreateKernelFT)(cl_program, const char *, cl_int *); +const char *clCreateKernel_Name = "clCreateKernel"; +clCreateKernelFT get_clCreateKernel() { - cl_int err; - auto CLCtx = get_native(SyclCtx); - auto CLProgram = clCreateProgramWithIL(CLCtx, IL, length, &err); - if (err) { - std::stringstream ss; - ss << "OpenCL program could not be created from the SPIR-V " - "binary. OpenCL Error " - << err << "."; - error_handler(ss.str(), __FILE__, __func__, __LINE__); + static dpctl::DynamicLibHelper clLib(clLoaderName, clLibLoadFlags); + if (!clLib.opened()) { + error_handler("The OpenCL loader dynamic library could not " + "be opened.", + __FILE__, __func__, __LINE__); return nullptr; } - auto SyclDevices = SyclCtx.get_devices(); + static auto st_clCreateKernelF = + clLib.getSymbol(clCreateKernel_Name); - // Get a list of CL Devices from the Sycl devices - auto CLDevices = new cl_device_id[SyclDevices.size()]; - for (auto i = 0ul; i < SyclDevices.size(); ++i) - CLDevices[i] = get_native(SyclDevices[i]); + return st_clCreateKernelF; +} - // Build the OpenCL interoperability program - err = clBuildProgram(CLProgram, (cl_uint)(SyclDevices.size()), CLDevices, - CompileOpts, nullptr, nullptr); - // free the CLDevices array - delete[] CLDevices; +std::string _GetErrorCode_ocl_impl(cl_int code) +{ + if (code == CL_BUILD_PROGRAM_FAILURE) { + return "CL_BUILD_PROGRAM_FAILURE (code=" + + std::to_string(static_cast(code)) + ")"; + } + else if (code == CL_INVALID_CONTEXT) { + return "CL_INVALID_CONTEXT (code=" + + std::to_string(static_cast(code)) + ")"; + } + else if (code == CL_INVALID_DEVICE) { + return "CL_INVALID_DEVICE (code=" + + std::to_string(static_cast(code)) + ")"; + } + else if (code == CL_INVALID_VALUE) { + return "CL_INVALID_VALUE (code=" + + std::to_string(static_cast(code)) + ")"; + } + else if (code == CL_OUT_OF_RESOURCES) { + return "CL_OUT_OF_RESOURCES (code=" + + std::to_string(static_cast(code)) + ")"; + } + else if (code == CL_OUT_OF_HOST_MEMORY) { + return "CL_OUT_OF_HOST_MEMORY (code=" + + std::to_string(static_cast(code)) + ")"; + } + else if (code == CL_INVALID_OPERATION) { + return "CL_INVALID_OPERATION (code=" + + std::to_string(static_cast(code)) + ")"; + } + else if (code == CL_INVALID_BINARY) { + return "CL_INVALID_BINARY (code=" + + std::to_string(static_cast(code)) + ")"; + } + + return "<< ERROR CODE UNRECOGNIZED >> (code=" + + std::to_string(static_cast(code)) + ")"; +} - if (err) { - std::stringstream ss; - ss << "OpenCL program could not be built. OpenCL Error " << err << "."; - error_handler(ss.str(), __FILE__, __func__, __LINE__); +constexpr backend cl_be = backend::opencl; + +DPCTLSyclKernelBundleRef +_CreateKernelBundle_common_ocl_impl(cl_program clProgram, + const context &ctx, + const device &dev, + const char *CompileOpts) +{ + backend_traits::return_type clDevice; + clDevice = get_native(dev); + + // Last to pointers are notification function pointer and user-data pointer + // that can be passed to the notification function. + auto clBuildProgramF = get_clBuldProgram(); + if (clBuildProgramF == nullptr) { return nullptr; } + cl_int build_status = + clBuildProgramF(clProgram, 1, &clDevice, CompileOpts, nullptr, nullptr); - // Create the Sycl program from OpenCL program - try { - auto SyclProgram = new program(SyclCtx, CLProgram); - return wrap(SyclProgram); - } catch (std::exception const &e) { - error_handler(e, __FILE__, __func__, __LINE__); + if (build_status != CL_SUCCESS) { + error_handler("clBuildProgram failed: " + + _GetErrorCode_ocl_impl(build_status), + __FILE__, __func__, __LINE__); + return nullptr; + } + + kernel_bundle kb = + make_kernel_bundle(clProgram, ctx); + return wrap(new kernel_bundle(kb)); +} + +DPCTLSyclKernelBundleRef +_CreateKernelBundleWithOCLSource_ocl_impl(const context &ctx, + const device &dev, + const char *oclSrc, + const char *CompileOpts) +{ + auto clCreateProgramWithSourceF = get_clCreateProgramWithSource(); + if (clCreateProgramWithSourceF == nullptr) { + return nullptr; + } + + backend_traits::return_type clContext; + clContext = get_native(ctx); + + cl_int build_with_source_err_code = CL_SUCCESS; + cl_program clProgram = clCreateProgramWithSourceF( + clContext, 1, &oclSrc, nullptr, &build_with_source_err_code); + + if (build_with_source_err_code != CL_SUCCESS) { + error_handler("clPCreateProgramWithSource failed with " + + _GetErrorCode_ocl_impl(build_with_source_err_code), + __FILE__, __func__, __LINE__); + return nullptr; + } + + return _CreateKernelBundle_common_ocl_impl(clProgram, ctx, dev, + CompileOpts); +} + +DPCTLSyclKernelBundleRef +_CreateKernelBundleWithIL_ocl_impl(const context &ctx, + const device &dev, + const void *IL, + size_t il_length, + const char *CompileOpts) +{ + auto clCreateProgramWithILF = get_clCreateProgramWithIL(); + if (clCreateProgramWithILF == nullptr) { + return nullptr; + } + + backend_traits::return_type clContext; + clContext = get_native(ctx); + + cl_int create_err_code = CL_SUCCESS; + cl_program clProgram = + clCreateProgramWithILF(clContext, IL, il_length, &create_err_code); + + if (create_err_code != CL_SUCCESS) { + error_handler("OpenCL program could not be created from the SPIR-V " + "binary. OpenCL Error " + + _GetErrorCode_ocl_impl(create_err_code), + __FILE__, __func__, __LINE__); + return nullptr; + } + + return _CreateKernelBundle_common_ocl_impl(clProgram, ctx, dev, + CompileOpts); +} + +bool _HasKernel_ocl_impl(const kernel_bundle &kb, + const char *kernel_name) +{ + auto clCreateKernelF = get_clCreateKernel(); + if (clCreateKernelF == nullptr) { + return false; + } + + std::vector oclKB = get_native(kb); + + bool found = false; + for (auto &cl_pr : oclKB) { + cl_int create_kernel_err_code = CL_SUCCESS; + [[maybe_unused]] cl_kernel try_kern = + clCreateKernelF(cl_pr, kernel_name, &create_kernel_err_code); + if (create_kernel_err_code == CL_SUCCESS) { + found = true; + break; + } + } + return found; +} + +__dpctl_give DPCTLSyclKernelRef +_GetKernel_ocl_impl(const kernel_bundle &kb, + const char *kernel_name) +{ + auto clCreateKernelF = get_clCreateKernel(); + if (clCreateKernelF == nullptr) { + return nullptr; + } + + std::vector oclKB = get_native(kb); + + bool found = false; + cl_kernel ocl_kernel_from_kb; + for (auto &cl_pr : oclKB) { + cl_int create_kernel_err_code = CL_SUCCESS; + cl_kernel try_kern = + clCreateKernelF(cl_pr, kernel_name, &create_kernel_err_code); + if (create_kernel_err_code == CL_SUCCESS) { + found = true; + ocl_kernel_from_kb = try_kern; + break; + } + } + if (found) { + try { + context ctx = kb.get_context(); + + kernel interop_kernel = make_kernel(ocl_kernel_from_kb, ctx); + + return wrap(new kernel(interop_kernel)); + } catch (std::exception const &e) { + error_handler(e, __FILE__, __func__, __LINE__); + return nullptr; + } + } + else { + error_handler("Kernel " + std::string(kernel_name) + " not found.", + __FILE__, __func__, __LINE__); return nullptr; } } #ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION -zeModuleCreateFT getZeModuleCreateFn() +#ifdef __linux__ +static const char *zeLoaderName = DPCTL_LIBZE_LOADER_FILENAME; +static const int zeLibLoadFlags = RTLD_NOLOAD | RTLD_NOW | RTLD_LOCAL; +#elif defined(_WIN64) +static const char *zeLoaderName = "ze_loader.dll"; +static const int zeLibLoadFlags = 0; +#else +#error "Level Zero program compilation is unavailable for this platform" +#endif + +constexpr sycl::backend ze_be = sycl::backend::ext_oneapi_level_zero; + +typedef ze_result_t (*zeModuleCreateFT)(ze_context_handle_t, + ze_device_handle_t, + const ze_module_desc_t *, + ze_module_handle_t *, + ze_module_build_log_handle_t *); +const char *zeModuleCreate_Name = "zeModuleCreate"; +zeModuleCreateFT get_zeModuleCreate() { - static dpctl::DynamicLibHelper zeLib(zeLoaderName, libLoadFlags); + static dpctl::DynamicLibHelper zeLib(zeLoaderName, zeLibLoadFlags); if (!zeLib.opened()) { error_handler("The level zero loader dynamic library could not " "be opened.", __FILE__, __func__, __LINE__); return nullptr; } - static auto stZeModuleCreateF = - zeLib.getSymbol(zeModuleCreateFuncName); + static auto st_zeModuleCreateF = + zeLib.getSymbol(zeModuleCreate_Name); - return stZeModuleCreateF; + return st_zeModuleCreateF; } -__dpctl_give DPCTLSyclProgramRef -createLevelZeroInterOpProgram(const context &SyclCtx, - const void *IL, - size_t length, - const char *CompileOpts) +typedef ze_result_t (*zeModuleDestroyFT)(ze_module_handle_t); +const char *zeModuleDestroy_Name = "zeModuleDestroy"; +zeModuleDestroyFT get_zeModuleDestroy() { - auto ZeCtx = get_native(SyclCtx); - auto SyclDevices = SyclCtx.get_devices(); - if (SyclDevices.size() > 1) { - error_handler("Level zero program can be created for only one device.", + static dpctl::DynamicLibHelper zeLib(zeLoaderName, zeLibLoadFlags); + if (!zeLib.opened()) { + error_handler("The level zero loader dynamic library could not " + "be opened.", __FILE__, __func__, __LINE__); return nullptr; } + static auto st_zeModuleDestroyF = + zeLib.getSymbol(zeModuleDestroy_Name); - // Specialization constants are not yet supported. - // Refer https://bit.ly/33UEDYN for details on specialization constants. + return st_zeModuleDestroyF; +} + +typedef ze_result_t (*zeKernelCreateFT)(ze_module_handle_t, + const ze_kernel_desc_t *, + ze_kernel_handle_t *); +const char *zeKernelCreate_Name = "zeKernelCreate"; +zeKernelCreateFT get_zeKernelCreate() +{ + static dpctl::DynamicLibHelper zeLib(zeLoaderName, zeLibLoadFlags); + if (!zeLib.opened()) { + error_handler("The level zero loader dynamic library could not " + "be opened.", + __FILE__, __func__, __LINE__); + return nullptr; + } + static auto st_zeKernelCreateF = + zeLib.getSymbol(zeKernelCreate_Name); + + return st_zeKernelCreateF; +} + +std::string _GetErrorCode_ze_impl(ze_result_t code) +{ + if (code == ZE_RESULT_ERROR_UNINITIALIZED) { + return "ZE_RESULT_ERROR_UNINITIALIZED (code=" + + std::to_string(static_cast(code)) + ")"; + } + else if (code == ZE_RESULT_ERROR_DEVICE_LOST) { + return "ZE_RESULT_ERROR_DEVICE_LOST (code=" + + std::to_string(static_cast(code)) + ")"; + } + else if (code == ZE_RESULT_ERROR_INVALID_NULL_HANDLE) { + return "ZE_RESULT_ERROR_INVALID_NULL_HANDLE (code=" + + std::to_string(static_cast(code)) + ")"; + } + else if (code == ZE_RESULT_ERROR_INVALID_NULL_POINTER) { + return "ZE_RESULT_ERROR_INVALID_NULL_POINTER (code=" + + std::to_string(static_cast(code)) + ")"; + } + else if (code == ZE_RESULT_ERROR_INVALID_ENUMERATION) { + return "ZE_RESULT_ERROR_INVALID_ENUMERATION (code=" + + std::to_string(static_cast(code)) + ")"; + } + else if (code == ZE_RESULT_ERROR_INVALID_NATIVE_BINARY) { + return "ZE_RESULT_ERROR_INVALID_NATIVE_BINARY (code=" + + std::to_string(static_cast(code)) + ")"; + } + else if (code == ZE_RESULT_ERROR_INVALID_SIZE) { + return "ZE_RESULT_ERROR_INVALID_SIZE (code=" + + std::to_string(static_cast(code)) + ")"; + } + else if (code == ZE_RESULT_ERROR_OUT_OF_HOST_MEMORY) { + return "ZE_RESULT_ERROR_OUT_OF_HOST_MEMORY (code=" + + std::to_string(static_cast(code)) + ")"; + } + else if (code == ZE_RESULT_ERROR_OUT_OF_DEVICE_MEMORY) { + return "ZE_RESULT_ERROR_OUT_OF_DEVICE_MEMORY (code=" + + std::to_string(static_cast(code)) + ")"; + } + else if (code == ZE_RESULT_ERROR_MODULE_BUILD_FAILURE) { + return "ZE_RESULT_ERROR_MODULE_BUILD_FAILURE (code=" + + std::to_string(static_cast(code)) + ")"; + } + else if (code == ZE_RESULT_ERROR_INVALID_MODULE_UNLINKED) { + return "ZE_RESULT_ERROR_INVALID_MODULE_UNLINKED (code=" + + std::to_string(static_cast(code)) + ")"; + } + + return "<< UNRECOGNIZE ZE_RESULT_T CODE >> (code=" + + std::to_string(static_cast(code)) + ")"; +} + +__dpctl_give DPCTLSyclKernelBundleRef +_CreateKernelBundleWithIL_ze_impl(const context &SyclCtx, + const device &SyclDev, + const void *IL, + size_t il_length, + const char *CompileOpts) +{ + auto zeModuleCreateFn = get_zeModuleCreate(); + if (zeModuleCreateFn == nullptr) { + error_handler("ZeModuleCreateFn is invalid.", __FILE__, __func__, + __LINE__); + return nullptr; + } + + backend_traits::return_type ZeContext; + ZeContext = get_native(SyclCtx); + + backend_traits::return_type ZeDevice; + ZeDevice = get_native(SyclDev); + + // Specialization constants are not supported by DPCTL at the moment ze_module_constants_t ZeSpecConstants = {}; ZeSpecConstants.numConstants = 0; @@ -165,118 +490,210 @@ createLevelZeroInterOpProgram(const context &SyclCtx, ze_module_desc_t ZeModuleDesc = {}; ZeModuleDesc.stype = ZE_STRUCTURE_TYPE_MODULE_DESC; ZeModuleDesc.format = ZE_MODULE_FORMAT_IL_SPIRV; - ZeModuleDesc.inputSize = length; + ZeModuleDesc.inputSize = il_length; ZeModuleDesc.pInputModule = (uint8_t *)IL; ZeModuleDesc.pBuildFlags = CompileOpts; ZeModuleDesc.pConstants = &ZeSpecConstants; - auto ZeDevice = get_native(SyclDevices[0]); ze_module_handle_t ZeModule; - auto stZeModuleCreateF = getZeModuleCreateFn(); + auto ret_code = zeModuleCreateFn(ZeContext, ZeDevice, &ZeModuleDesc, + &ZeModule, nullptr); + if (ret_code != ZE_RESULT_SUCCESS) { + error_handler("Module creation failed " + + _GetErrorCode_ze_impl(ret_code), + __FILE__, __func__, __LINE__); + return nullptr; + } + + try { + auto kb = make_kernel_bundle( + {ZeModule, ext::oneapi::level_zero::ownership::keep}, SyclCtx); - if (!stZeModuleCreateF) { - error_handler("ZeModuleCreateFn is invalid.", __FILE__, __func__, - __LINE__); + return wrap(new kernel_bundle(kb)); + } catch (std::exception const &e) { + error_handler(e, __FILE__, __func__, __LINE__); + auto zeModuleDestroyFn = get_zeModuleDestroy(); + if (zeModuleDestroyFn) { + zeModuleDestroyFn(ZeModule); + } return nullptr; } +} - auto ret = - stZeModuleCreateF(ZeCtx, ZeDevice, &ZeModuleDesc, &ZeModule, nullptr); - if (ret != ZE_RESULT_SUCCESS) { - error_handler("ZeModule creation failed.", __FILE__, __func__, - __LINE__); +__dpctl_give DPCTLSyclKernelRef +_GetKernel_ze_impl(const kernel_bundle &kb, + const char *kernel_name) +{ + auto zeKernelCreateFn = get_zeKernelCreate(); + if (zeKernelCreateFn == nullptr) { + error_handler("Could not load zeKernelCreate function.", __FILE__, + __func__, __LINE__); return nullptr; } - // Create the Sycl program from the ZeModule - try { - auto ZeProgram = - new program(sycl::ext::oneapi::level_zero::make_program( - SyclCtx, reinterpret_cast(ZeModule))); - return wrap(ZeProgram); - } catch (std::exception const &e) { - error_handler(e, __FILE__, __func__, __LINE__); + auto ZeKernelBundle = sycl::get_native(kb); + bool found = false; + + // Populate the Level Zero kernel descriptions + ze_kernel_desc_t ZeKernelDescr = {ZE_STRUCTURE_TYPE_KERNEL_DESC, nullptr, + 0, // flags + kernel_name}; + + std::unique_ptr syclInteropKern_ptr; + ze_kernel_handle_t ZeKern; + for (auto &ZeM : ZeKernelBundle) { + ze_result_t ze_status = zeKernelCreateFn(ZeM, &ZeKernelDescr, &ZeKern); + + if (ze_status == ZE_RESULT_SUCCESS) { + found = true; + auto ctx = kb.get_context(); + auto k = make_kernel( + {kb, ZeKern, ext::oneapi::level_zero::ownership::keep}, ctx); + syclInteropKern_ptr = std::unique_ptr(new kernel(k)); + break; + } + else { + if (ze_status != ZE_RESULT_ERROR_INVALID_KERNEL_NAME) { + error_handler("zeKernelCreate failed: " + + _GetErrorCode_ze_impl(ze_status), + __FILE__, __func__, __LINE__); + return nullptr; + } + } + } + + if (found) { + return wrap(new kernel(*syclInteropKern_ptr)); + } + else { + error_handler("Kernel named " + std::string(kernel_name) + + " could not be found.", + __FILE__, __func__, __LINE__); return nullptr; } } + +bool _HasKernel_ze_impl(const kernel_bundle &kb, + const char *kernel_name) +{ + auto zeKernelCreateFn = get_zeKernelCreate(); + if (zeKernelCreateFn == nullptr) { + error_handler("Could not load zeKernelCreate function.", __FILE__, + __func__, __LINE__); + return false; + } + + auto ZeKernelBundle = sycl::get_native(kb); + + // Populate the Level Zero kernel descriptions + ze_kernel_desc_t ZeKernelDescr = {ZE_STRUCTURE_TYPE_KERNEL_DESC, nullptr, + 0, // flags + kernel_name}; + + std::unique_ptr syclInteropKern_ptr; + ze_kernel_handle_t ZeKern; + for (auto &ZeM : ZeKernelBundle) { + ze_result_t ze_status = zeKernelCreateFn(ZeM, &ZeKernelDescr, &ZeKern); + + if (ze_status == ZE_RESULT_SUCCESS) { + return true; + } + else { + if (ze_status != ZE_RESULT_ERROR_INVALID_KERNEL_NAME) { + error_handler("zeKernelCreate failed: " + + _GetErrorCode_ze_impl(ze_status), + __FILE__, __func__, __LINE__); + return false; + } + } + } + + return false; +} + #endif /* #ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION */ } /* end of anonymous namespace */ -__dpctl_give DPCTLSyclProgramRef -DPCTLProgram_CreateFromSpirv(__dpctl_keep const DPCTLSyclContextRef CtxRef, - __dpctl_keep const void *IL, - size_t length, - const char *CompileOpts) +__dpctl_give DPCTLSyclKernelBundleRef +DPCTLKernelBundle_CreateFromSpirv(__dpctl_keep const DPCTLSyclContextRef CtxRef, + __dpctl_keep const DPCTLSyclDeviceRef DevRef, + __dpctl_keep const void *IL, + size_t length, + const char *CompileOpts) { - DPCTLSyclProgramRef Pref = nullptr; - context *SyclCtx = nullptr; + DPCTLSyclKernelBundleRef KBRef = nullptr; if (!CtxRef) { error_handler("Cannot create program from SPIR-V as the supplied SYCL " "context is NULL.", __FILE__, __func__, __LINE__); - return Pref; + return KBRef; + } + if (!DevRef) { + error_handler("Cannot create program from SPIR-V as the supplied SYCL " + "device is NULL.", + __FILE__, __func__, __LINE__); + return KBRef; } - SyclCtx = unwrap(CtxRef); + + context *SyclCtx = unwrap(CtxRef); + device *SyclDev = unwrap(DevRef); // get the backend type auto BE = SyclCtx->get_platform().get_backend(); switch (BE) { case backend::opencl: - Pref = createOpenCLInterOpProgram(*SyclCtx, IL, length, CompileOpts); + KBRef = _CreateKernelBundleWithIL_ocl_impl(*SyclCtx, *SyclDev, IL, + length, CompileOpts); break; case backend::ext_oneapi_level_zero: #ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION - Pref = createLevelZeroInterOpProgram(*SyclCtx, IL, length, CompileOpts); + KBRef = _CreateKernelBundleWithIL_ze_impl(*SyclCtx, *SyclDev, IL, + length, CompileOpts); #endif break; default: + error_handler("Backend " + std::to_string(static_cast(BE)) + + " is not supported", + __FILE__, __func__, __LINE__); break; } - return Pref; + return KBRef; } -__dpctl_give DPCTLSyclProgramRef -DPCTLProgram_CreateFromOCLSource(__dpctl_keep const DPCTLSyclContextRef Ctx, - __dpctl_keep const char *Source, - __dpctl_keep const char *CompileOpts) +__dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromOCLSource( + __dpctl_keep const DPCTLSyclContextRef Ctx, + __dpctl_keep const DPCTLSyclDeviceRef Dev, + __dpctl_keep const char *Source, + __dpctl_keep const char *CompileOpts) { - std::string compileOpts; context *SyclCtx = nullptr; - program *SyclProgram = nullptr; + device *SyclDev = nullptr; if (!Ctx) { error_handler("Input Ctx is nullptr.", __FILE__, __func__, __LINE__); return nullptr; } - + if (!Dev) { + error_handler("Input Dev is nullptr.", __FILE__, __func__, __LINE__); + return nullptr; + } if (!Source) { error_handler("Input Source is nullptr.", __FILE__, __func__, __LINE__); return nullptr; } SyclCtx = unwrap(Ctx); - try { - SyclProgram = new program(*SyclCtx); - } catch (std::exception const &e) { - error_handler(e, __FILE__, __func__, __LINE__); - return nullptr; - } - std::string source = Source; - - if (CompileOpts) { - compileOpts = CompileOpts; - } + SyclDev = unwrap(Dev); // get the backend type auto BE = SyclCtx->get_platform().get_backend(); switch (BE) { case backend::opencl: try { - SyclProgram->build_with_source(source, compileOpts); - return wrap(SyclProgram); + return _CreateKernelBundleWithOCLSource_ocl_impl( + *SyclCtx, *SyclDev, Source, CompileOpts); } catch (std::exception const &e) { - delete SyclProgram; error_handler(e, __FILE__, __func__, __LINE__); return nullptr; } @@ -284,45 +701,47 @@ DPCTLProgram_CreateFromOCLSource(__dpctl_keep const DPCTLSyclContextRef Ctx, case backend::ext_oneapi_level_zero: error_handler("CreateFromSource is not supported in Level Zero.", __FILE__, __func__, __LINE__); - delete SyclProgram; return nullptr; default: error_handler("CreateFromSource is not supported in unknown backend.", __FILE__, __func__, __LINE__); - delete SyclProgram; return nullptr; } } __dpctl_give DPCTLSyclKernelRef -DPCTLProgram_GetKernel(__dpctl_keep DPCTLSyclProgramRef PRef, - __dpctl_keep const char *KernelName) +DPCTLKernelBundle_GetKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, + __dpctl_keep const char *KernelName) { - if (!PRef) { - error_handler("Input PRef is nullptr", __FILE__, __func__, __LINE__); + if (!KBRef) { + error_handler("Input KBRef is nullptr", __FILE__, __func__, __LINE__); return nullptr; } - auto SyclProgram = unwrap(PRef); if (!KernelName) { error_handler("Input KernelName is nullptr", __FILE__, __func__, __LINE__); return nullptr; } - std::string name = KernelName; - try { - auto SyclKernel = new kernel(SyclProgram->get_kernel(name)); - return wrap(SyclKernel); - } catch (std::exception const &e) { - error_handler(e, __FILE__, __func__, __LINE__); + auto SyclKB = unwrap(KBRef); + sycl::backend be = SyclKB->get_backend(); + switch (be) { + case sycl::backend::opencl: + return _GetKernel_ocl_impl(*SyclKB, KernelName); + case sycl::backend::ext_oneapi_level_zero: + return _GetKernel_ze_impl(*SyclKB, KernelName); + default: + error_handler("Backend " + std::to_string(static_cast(be)) + + " is not supported.", + __FILE__, __func__, __LINE__); return nullptr; } } -bool DPCTLProgram_HasKernel(__dpctl_keep DPCTLSyclProgramRef PRef, - __dpctl_keep const char *KernelName) +bool DPCTLKernelBundle_HasKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, + __dpctl_keep const char *KernelName) { - if (!PRef) { - error_handler("Input PRef is nullptr", __FILE__, __func__, __LINE__); + if (!KBRef) { + error_handler("Input KBRef is nullptr", __FILE__, __func__, __LINE__); return false; } if (!KernelName) { @@ -331,16 +750,22 @@ bool DPCTLProgram_HasKernel(__dpctl_keep DPCTLSyclProgramRef PRef, return false; } - auto SyclProgram = unwrap(PRef); - try { - return SyclProgram->has_kernel(KernelName); - } catch (std::exception const &e) { - error_handler(e, __FILE__, __func__, __LINE__); + auto SyclKB = unwrap(KBRef); + sycl::backend be = SyclKB->get_backend(); + switch (be) { + case sycl::backend::opencl: + return _HasKernel_ocl_impl(*SyclKB, KernelName); + case sycl::backend::ext_oneapi_level_zero: + return _HasKernel_ze_impl(*SyclKB, KernelName); + default: + error_handler("Backend " + std::to_string(static_cast(be)) + + " is not supported.", + __FILE__, __func__, __LINE__); return false; } } -void DPCTLProgram_Delete(__dpctl_take DPCTLSyclProgramRef PRef) +void DPCTLKernelBundle_Delete(__dpctl_take DPCTLSyclKernelBundleRef KBRef) { - delete unwrap(PRef); + delete unwrap(KBRef); } diff --git a/libsyclinterface/tests/test_sycl_kernel_interface.cpp b/libsyclinterface/tests/test_sycl_kernel_interface.cpp index ca53c130fb..0db8c227bf 100644 --- a/libsyclinterface/tests/test_sycl_kernel_interface.cpp +++ b/libsyclinterface/tests/test_sycl_kernel_interface.cpp @@ -85,17 +85,17 @@ TEST_P(TestDPCTLSyclKernelInterface, CheckGetNumArgs) { auto QueueRef = DPCTLQueue_CreateForDevice(DRef, nullptr, 0); auto CtxRef = DPCTLQueue_GetContext(QueueRef); - auto PRef = - DPCTLProgram_CreateFromOCLSource(CtxRef, CLProgramStr, CompileOpts); - auto AddKernel = DPCTLProgram_GetKernel(PRef, "add"); - auto AxpyKernel = DPCTLProgram_GetKernel(PRef, "axpy"); + auto KBRef = DPCTLKernelBundle_CreateFromOCLSource( + CtxRef, DRef, CLProgramStr, CompileOpts); + auto AddKernel = DPCTLKernelBundle_GetKernel(KBRef, "add"); + auto AxpyKernel = DPCTLKernelBundle_GetKernel(KBRef, "axpy"); ASSERT_EQ(DPCTLKernel_GetNumArgs(AddKernel), 3ul); ASSERT_EQ(DPCTLKernel_GetNumArgs(AxpyKernel), 4ul); DPCTLQueue_Delete(QueueRef); DPCTLContext_Delete(CtxRef); - DPCTLProgram_Delete(PRef); + DPCTLKernelBundle_Delete(KBRef); DPCTLKernel_Delete(AddKernel); DPCTLKernel_Delete(AxpyKernel); } diff --git a/libsyclinterface/tests/test_sycl_program_interface.cpp b/libsyclinterface/tests/test_sycl_program_interface.cpp index 9824e09088..a6ffe0783b 100644 --- a/libsyclinterface/tests/test_sycl_program_interface.cpp +++ b/libsyclinterface/tests/test_sycl_program_interface.cpp @@ -47,7 +47,7 @@ struct TestDPCTLSyclProgramInterface DPCTLSyclDeviceRef DRef = nullptr; DPCTLSyclContextRef CRef = nullptr; DPCTLSyclQueueRef QRef = nullptr; - DPCTLSyclProgramRef PRef = nullptr; + DPCTLSyclKernelBundleRef KBRef = nullptr; std::ifstream spirvFile; size_t spirvFileSize; std::vector spirvBuffer; @@ -67,8 +67,8 @@ struct TestDPCTLSyclProgramInterface spirvBuffer.reserve(spirvFileSize); spirvFile.seekg(0, std::ios::beg); spirvFile.read(spirvBuffer.data(), spirvFileSize); - PRef = DPCTLProgram_CreateFromSpirv(CRef, spirvBuffer.data(), - spirvFileSize, nullptr); + KBRef = DPCTLKernelBundle_CreateFromSpirv( + CRef, DRef, spirvBuffer.data(), spirvFileSize, nullptr); } } @@ -88,41 +88,42 @@ struct TestDPCTLSyclProgramInterface DPCTLDevice_Delete(DRef); DPCTLQueue_Delete(QRef); DPCTLContext_Delete(CRef); - DPCTLProgram_Delete(PRef); + DPCTLKernelBundle_Delete(KBRef); } }; TEST_P(TestDPCTLSyclProgramInterface, ChkCreateFromSpirv) { - ASSERT_TRUE(PRef != nullptr); - ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "add")); - ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "axpy")); - ASSERT_FALSE(DPCTLProgram_HasKernel(PRef, nullptr)); + ASSERT_TRUE(KBRef != nullptr); + ASSERT_TRUE(DPCTLKernelBundle_HasKernel(KBRef, "add")); + ASSERT_TRUE(DPCTLKernelBundle_HasKernel(KBRef, "axpy")); + ASSERT_FALSE(DPCTLKernelBundle_HasKernel(KBRef, nullptr)); } TEST_P(TestDPCTLSyclProgramInterface, ChkCreateFromSpirvNull) { DPCTLSyclContextRef Null_CRef = nullptr; + DPCTLSyclDeviceRef Null_DRef = nullptr; const void *null_spirv = nullptr; - DPCTLSyclProgramRef PRef = nullptr; - EXPECT_NO_FATAL_FAILURE( - PRef = DPCTLProgram_CreateFromSpirv(Null_CRef, null_spirv, 0, nullptr)); - ASSERT_TRUE(PRef == nullptr); + DPCTLSyclKernelBundleRef KBRef = nullptr; + EXPECT_NO_FATAL_FAILURE(KBRef = DPCTLKernelBundle_CreateFromSpirv( + Null_CRef, Null_DRef, null_spirv, 0, nullptr)); + ASSERT_TRUE(KBRef == nullptr); } TEST_P(TestDPCTLSyclProgramInterface, ChkHasKernelNullProgram) { - DPCTLSyclProgramRef NullRef = nullptr; - ASSERT_FALSE(DPCTLProgram_HasKernel(NullRef, "add")); + DPCTLSyclKernelBundleRef NullRef = nullptr; + ASSERT_FALSE(DPCTLKernelBundle_HasKernel(NullRef, "add")); } TEST_P(TestDPCTLSyclProgramInterface, ChkGetKernel) { - auto AddKernel = DPCTLProgram_GetKernel(PRef, "add"); - auto AxpyKernel = DPCTLProgram_GetKernel(PRef, "axpy"); - auto NullKernel = DPCTLProgram_GetKernel(PRef, nullptr); + auto AddKernel = DPCTLKernelBundle_GetKernel(KBRef, "add"); + auto AxpyKernel = DPCTLKernelBundle_GetKernel(KBRef, "axpy"); + auto NullKernel = DPCTLKernelBundle_GetKernel(KBRef, nullptr); ASSERT_TRUE(AddKernel != nullptr); ASSERT_TRUE(AxpyKernel != nullptr); @@ -134,10 +135,10 @@ TEST_P(TestDPCTLSyclProgramInterface, ChkGetKernel) TEST_P(TestDPCTLSyclProgramInterface, ChkGetKernelNullProgram) { - DPCTLSyclProgramRef NullRef = nullptr; + DPCTLSyclKernelBundleRef NullRef = nullptr; DPCTLSyclKernelRef KRef = nullptr; - EXPECT_NO_FATAL_FAILURE(KRef = DPCTLProgram_GetKernel(NullRef, "add")); + EXPECT_NO_FATAL_FAILURE(KRef = DPCTLKernelBundle_GetKernel(NullRef, "add")); EXPECT_TRUE(KRef == nullptr); } @@ -159,7 +160,7 @@ struct TestOCLProgramFromSource : public ::testing::Test DPCTLSyclDeviceRef DRef = nullptr; DPCTLSyclContextRef CRef = nullptr; DPCTLSyclQueueRef QRef = nullptr; - DPCTLSyclProgramRef PRef = nullptr; + DPCTLSyclKernelBundleRef KBRef = nullptr; TestOCLProgramFromSource() { @@ -170,8 +171,8 @@ struct TestOCLProgramFromSource : public ::testing::Test QRef = DPCTLQueue_Create(CRef, DRef, nullptr, DPCTL_DEFAULT_PROPERTY); if (DRef) - PRef = DPCTLProgram_CreateFromOCLSource(CRef, CLProgramStr, - CompileOpts); + KBRef = DPCTLKernelBundle_CreateFromOCLSource( + CRef, DRef, CLProgramStr, CompileOpts); } ~TestOCLProgramFromSource() @@ -179,7 +180,7 @@ struct TestOCLProgramFromSource : public ::testing::Test DPCTLDevice_Delete(DRef); DPCTLQueue_Delete(QRef); DPCTLContext_Delete(CRef); - DPCTLProgram_Delete(PRef); + DPCTLKernelBundle_Delete(KBRef); } }; @@ -188,9 +189,9 @@ TEST_F(TestOCLProgramFromSource, CheckCreateFromOCLSource) if (!DRef) GTEST_SKIP_("Skipping as no OpenCL GPU device found.\n"); - ASSERT_TRUE(PRef != nullptr); - ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "add")); - ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "axpy")); + ASSERT_TRUE(KBRef != nullptr); + ASSERT_TRUE(DPCTLKernelBundle_HasKernel(KBRef, "add")); + ASSERT_TRUE(DPCTLKernelBundle_HasKernel(KBRef, "axpy")); } TEST_F(TestOCLProgramFromSource, CheckCreateFromOCLSourceNull) @@ -201,14 +202,14 @@ TEST_F(TestOCLProgramFromSource, CheckCreateFromOCLSourceNull) b[index] = a[index]; } )CLC"; - DPCTLSyclProgramRef PRef = nullptr; + DPCTLSyclKernelBundleRef KBRef = nullptr; if (!DRef) GTEST_SKIP_("Skipping as no OpenCL GPU device found.\n"); - EXPECT_NO_FATAL_FAILURE(PRef = DPCTLProgram_CreateFromOCLSource( - CRef, InvalidCLProgramStr, CompileOpts);); - ASSERT_TRUE(PRef == nullptr); + EXPECT_NO_FATAL_FAILURE(KBRef = DPCTLKernelBundle_CreateFromOCLSource( + CRef, DRef, InvalidCLProgramStr, CompileOpts);); + ASSERT_TRUE(KBRef == nullptr); } TEST_F(TestOCLProgramFromSource, CheckGetKernelOCLSource) @@ -216,8 +217,8 @@ TEST_F(TestOCLProgramFromSource, CheckGetKernelOCLSource) if (!DRef) GTEST_SKIP_("Skipping as no OpenCL GPU device found.\n"); - auto AddKernel = DPCTLProgram_GetKernel(PRef, "add"); - auto AxpyKernel = DPCTLProgram_GetKernel(PRef, "axpy"); + auto AddKernel = DPCTLKernelBundle_GetKernel(KBRef, "add"); + auto AxpyKernel = DPCTLKernelBundle_GetKernel(KBRef, "axpy"); ASSERT_TRUE(AddKernel != nullptr); ASSERT_TRUE(AxpyKernel != nullptr); DPCTLKernel_Delete(AddKernel); diff --git a/libsyclinterface/tests/test_sycl_queue_submit.cpp b/libsyclinterface/tests/test_sycl_queue_submit.cpp index b0892ccb9c..efec07e7d0 100644 --- a/libsyclinterface/tests/test_sycl_queue_submit.cpp +++ b/libsyclinterface/tests/test_sycl_queue_submit.cpp @@ -79,11 +79,11 @@ TEST_F(TestQueueSubmit, CheckSubmitRange_saxpy) ASSERT_TRUE(QRef); auto CRef = DPCTLQueue_GetContext(QRef); ASSERT_TRUE(CRef); - auto PRef = DPCTLProgram_CreateFromSpirv(CRef, spirvBuffer.data(), - spirvFileSize, nullptr); - ASSERT_TRUE(PRef != nullptr); - ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "axpy")); - auto AxpyKernel = DPCTLProgram_GetKernel(PRef, "axpy"); + auto KBRef = DPCTLKernelBundle_CreateFromSpirv( + CRef, DRef, spirvBuffer.data(), spirvFileSize, nullptr); + ASSERT_TRUE(KBRef != nullptr); + ASSERT_TRUE(DPCTLKernelBundle_HasKernel(KBRef, "axpy")); + auto AxpyKernel = DPCTLKernelBundle_GetKernel(KBRef, "axpy"); // Create the input args auto a = DPCTLmalloc_shared(SIZE * sizeof(float), QRef); @@ -120,7 +120,7 @@ TEST_F(TestQueueSubmit, CheckSubmitRange_saxpy) DPCTLfree_with_queue((DPCTLSyclUSMRef)c, QRef); DPCTLQueue_Delete(QRef); DPCTLContext_Delete(CRef); - DPCTLProgram_Delete(PRef); + DPCTLKernelBundle_Delete(KBRef); DPCTLDevice_Delete(DRef); DPCTLDeviceSelector_Delete(DSRef); } @@ -139,11 +139,11 @@ TEST_F(TestQueueSubmit, CheckSubmitNDRange_saxpy) ASSERT_TRUE(QRef); auto CRef = DPCTLQueue_GetContext(QRef); ASSERT_TRUE(CRef); - auto PRef = DPCTLProgram_CreateFromSpirv(CRef, spirvBuffer.data(), - spirvFileSize, nullptr); - ASSERT_TRUE(PRef != nullptr); - ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "axpy")); - auto AxpyKernel = DPCTLProgram_GetKernel(PRef, "axpy"); + auto KBRef = DPCTLKernelBundle_CreateFromSpirv( + CRef, DRef, spirvBuffer.data(), spirvFileSize, nullptr); + ASSERT_TRUE(KBRef != nullptr); + ASSERT_TRUE(DPCTLKernelBundle_HasKernel(KBRef, "axpy")); + auto AxpyKernel = DPCTLKernelBundle_GetKernel(KBRef, "axpy"); // Create the input args auto a = DPCTLmalloc_shared(SIZE * sizeof(float), QRef); @@ -185,7 +185,7 @@ TEST_F(TestQueueSubmit, CheckSubmitNDRange_saxpy) DPCTLfree_with_queue((DPCTLSyclUSMRef)c, QRef); DPCTLQueue_Delete(QRef); DPCTLContext_Delete(CRef); - DPCTLProgram_Delete(PRef); + DPCTLKernelBundle_Delete(KBRef); DPCTLDevice_Delete(DRef); DPCTLDeviceSelector_Delete(DSRef); } From cdaf00335f3d7614768fceb167a059f7b2f7ed13 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Tue, 24 May 2022 06:09:40 -0500 Subject: [PATCH 04/14] Modularized symbol getters --- .../source/dpctl_sycl_program_interface.cpp | 117 ++++++++++-------- 1 file changed, 65 insertions(+), 52 deletions(-) diff --git a/libsyclinterface/source/dpctl_sycl_program_interface.cpp b/libsyclinterface/source/dpctl_sycl_program_interface.cpp index c2774515dd..c91aadbfb1 100644 --- a/libsyclinterface/source/dpctl_sycl_program_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_program_interface.cpp @@ -1,4 +1,5 @@ -//===- dpctl_sycl_program_interface.cpp - Implements C API for sycl::program =// +//===- dpctl_sycl_program_interface.cpp - Implements C API for +// sycl::kernel_bundle =// // // Data Parallel Control (dpctl) // @@ -63,6 +64,35 @@ static const int clLibLoadFlags = 0; #error "OpenCL program compilation is unavailable for this platform" #endif +constexpr backend cl_be = backend::opencl; + +struct cl_loader +{ +public: + static cl_loader &get() + { + static cl_loader _loader; + return _loader; + } + + template retTy getSymbol(const char *name) + { + if (!opened) { + error_handler("The OpenCL loader dynamic library could not " + "be opened.", + __FILE__, __func__, __LINE__); + + return nullptr; + } + return clLib.getSymbol(name); + } + +private: + dpctl::DynamicLibHelper clLib; + bool opened; + cl_loader() : clLib(clLoaderName, clLibLoadFlags), opened(clLib.opened()) {} +}; + typedef cl_program (*clCreateProgramWithSourceFT)(cl_context, cl_uint, const char **, @@ -71,16 +101,8 @@ typedef cl_program (*clCreateProgramWithSourceFT)(cl_context, const char *clCreateProgramWithSource_Name = "clCreateProgramWithSource"; clCreateProgramWithSourceFT get_clCreateProgramWithSource() { - static dpctl::DynamicLibHelper clLib(clLoaderName, clLibLoadFlags); - if (!clLib.opened()) { - error_handler("The OpenCL loader dynamic library could not " - "be opened.", - __FILE__, __func__, __LINE__); - return nullptr; - } - static auto st_clCreateProgramWithSourceF = - clLib.getSymbol( + cl_loader::get().getSymbol( clCreateProgramWithSource_Name); return st_clCreateProgramWithSourceF; @@ -93,15 +115,9 @@ typedef cl_program (*clCreateProgramWithILFT)(cl_context, const char *clCreateProgramWithIL_Name = "clCreateProgramWithIL"; clCreateProgramWithILFT get_clCreateProgramWithIL() { - static dpctl::DynamicLibHelper clLib(clLoaderName, clLibLoadFlags); - if (!clLib.opened()) { - error_handler("The OpenCL loader dynamic library could not " - "be opened.", - __FILE__, __func__, __LINE__); - return nullptr; - } static auto st_clCreateProgramWithILF = - clLib.getSymbol(clCreateProgramWithIL_Name); + cl_loader::get().getSymbol( + clCreateProgramWithIL_Name); return st_clCreateProgramWithILF; } @@ -114,15 +130,8 @@ typedef cl_int (*clBuildProgramFT)(cl_program, const char *clBuildProgram_Name = "clBuildProgram"; clBuildProgramFT get_clBuldProgram() { - static dpctl::DynamicLibHelper clLib(clLoaderName, clLibLoadFlags); - if (!clLib.opened()) { - error_handler("The OpenCL loader dynamic library could not " - "be opened.", - __FILE__, __func__, __LINE__); - return nullptr; - } static auto st_clBuildProgramF = - clLib.getSymbol(clBuildProgram_Name); + cl_loader::get().getSymbol(clBuildProgram_Name); return st_clBuildProgramF; } @@ -183,8 +192,6 @@ std::string _GetErrorCode_ocl_impl(cl_int code) std::to_string(static_cast(code)) + ")"; } -constexpr backend cl_be = backend::opencl; - DPCTLSyclKernelBundleRef _CreateKernelBundle_common_ocl_impl(cl_program clProgram, const context &ctx, @@ -354,6 +361,33 @@ static const int zeLibLoadFlags = 0; constexpr sycl::backend ze_be = sycl::backend::ext_oneapi_level_zero; +struct ze_loader +{ +public: + static ze_loader &get() + { + static ze_loader _loader; + return _loader; + } + + template retTy getSymbol(const char *name) + { + if (!opened) { + error_handler("The Level-Zero loader dynamic library could not " + "be opened.", + __FILE__, __func__, __LINE__); + + return nullptr; + } + return zeLib.getSymbol(name); + } + +private: + dpctl::DynamicLibHelper zeLib; + bool opened; + ze_loader() : zeLib(zeLoaderName, zeLibLoadFlags), opened(zeLib.opened()) {} +}; + typedef ze_result_t (*zeModuleCreateFT)(ze_context_handle_t, ze_device_handle_t, const ze_module_desc_t *, @@ -362,15 +396,8 @@ typedef ze_result_t (*zeModuleCreateFT)(ze_context_handle_t, const char *zeModuleCreate_Name = "zeModuleCreate"; zeModuleCreateFT get_zeModuleCreate() { - static dpctl::DynamicLibHelper zeLib(zeLoaderName, zeLibLoadFlags); - if (!zeLib.opened()) { - error_handler("The level zero loader dynamic library could not " - "be opened.", - __FILE__, __func__, __LINE__); - return nullptr; - } static auto st_zeModuleCreateF = - zeLib.getSymbol(zeModuleCreate_Name); + ze_loader::get().getSymbol(zeModuleCreate_Name); return st_zeModuleCreateF; } @@ -379,15 +406,8 @@ typedef ze_result_t (*zeModuleDestroyFT)(ze_module_handle_t); const char *zeModuleDestroy_Name = "zeModuleDestroy"; zeModuleDestroyFT get_zeModuleDestroy() { - static dpctl::DynamicLibHelper zeLib(zeLoaderName, zeLibLoadFlags); - if (!zeLib.opened()) { - error_handler("The level zero loader dynamic library could not " - "be opened.", - __FILE__, __func__, __LINE__); - return nullptr; - } static auto st_zeModuleDestroyF = - zeLib.getSymbol(zeModuleDestroy_Name); + ze_loader::get().getSymbol(zeModuleDestroy_Name); return st_zeModuleDestroyF; } @@ -398,15 +418,8 @@ typedef ze_result_t (*zeKernelCreateFT)(ze_module_handle_t, const char *zeKernelCreate_Name = "zeKernelCreate"; zeKernelCreateFT get_zeKernelCreate() { - static dpctl::DynamicLibHelper zeLib(zeLoaderName, zeLibLoadFlags); - if (!zeLib.opened()) { - error_handler("The level zero loader dynamic library could not " - "be opened.", - __FILE__, __func__, __LINE__); - return nullptr; - } static auto st_zeKernelCreateF = - zeLib.getSymbol(zeKernelCreate_Name); + ze_loader::get().getSymbol(zeKernelCreate_Name); return st_zeKernelCreateF; } From 89ac958fcebd31f71fb4907d9dec3d9a361d7d34 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Tue, 24 May 2022 09:36:26 -0500 Subject: [PATCH 05/14] Renamed dpctl_sycl_program_interface to dpctl_sycl_kernel_bundle_interface --- dpctl/_backend.pxd | 2 +- dpctl/program/_program.pxd | 2 +- dpctl/program/_program.pyx | 11 ++++++----- ...terface.h => dpctl_sycl_kernel_bundle_interface.h} | 7 ++++--- ...ace.cpp => dpctl_sycl_kernel_bundle_interface.cpp} | 8 ++++---- libsyclinterface/tests/test_sycl_kernel_interface.cpp | 4 ++-- .../tests/test_sycl_program_interface.cpp | 5 +++-- libsyclinterface/tests/test_sycl_queue_submit.cpp | 2 +- 8 files changed, 22 insertions(+), 19 deletions(-) rename libsyclinterface/include/{dpctl_sycl_program_interface.h => dpctl_sycl_kernel_bundle_interface.h} (93%) rename libsyclinterface/source/{dpctl_sycl_program_interface.cpp => dpctl_sycl_kernel_bundle_interface.cpp} (99%) diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index 7781d9bf73..7d46ad0b77 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -304,7 +304,7 @@ cdef extern from "syclinterface/dpctl_sycl_context_interface.h": cdef void DPCTLContext_Delete(DPCTLSyclContextRef CtxRef) -cdef extern from "syclinterface/dpctl_sycl_program_interface.h": +cdef extern from "syclinterface/dpctl_sycl_kernel_bundle_interface.h": cdef DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSpirv( const DPCTLSyclContextRef Ctx, const DPCTLSyclDeviceRef Dev, diff --git a/dpctl/program/_program.pxd b/dpctl/program/_program.pxd index dda9e2662c..8880b853ce 100644 --- a/dpctl/program/_program.pxd +++ b/dpctl/program/_program.pxd @@ -44,7 +44,7 @@ cdef class SyclProgram: ''' Wraps a sycl::kernel_bundle object created from using SYCL interoperability layer for OpenCL and Level-Zero backends. - SyclProgram exposes the C API from dpctl_sycl_program_interface.h. A + SyclProgram exposes the C API from dpctl_sycl_kernel_bundle_interface.h. A SyclProgram can be created from either a source string or a SPIR-V binary file. ''' diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index b953a1cf3e..f638087793 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -97,12 +97,13 @@ cdef class SyclKernel: cdef class SyclProgram: - """ Wraps a ``sycl::program`` object created from an OpenCL interoperability - program. + """ Wraps a ``sycl::kernel_bundle`` object + created using SYCL interoperability layer with underlying backends. Only the + OpenCL and Level-Zero backends are currently supported. - SyclProgram exposes the C API from ``dpctl_sycl_program_interface.h``. A - SyclProgram can be created from either a source string or a SPIR-V - binary file. + SyclProgram exposes the C API from ``dpctl_sycl_kernel_bundle_interface.h``. + A SyclProgram can be created from either a source string or a SPIR-V + binary file. """ @staticmethod diff --git a/libsyclinterface/include/dpctl_sycl_program_interface.h b/libsyclinterface/include/dpctl_sycl_kernel_bundle_interface.h similarity index 93% rename from libsyclinterface/include/dpctl_sycl_program_interface.h rename to libsyclinterface/include/dpctl_sycl_kernel_bundle_interface.h index 335d274140..8dacfbf581 100644 --- a/libsyclinterface/include/dpctl_sycl_program_interface.h +++ b/libsyclinterface/include/dpctl_sycl_kernel_bundle_interface.h @@ -1,4 +1,5 @@ -//===- dpctl_sycl_program_interface.h - C API for sycl::program -*-C++-*- ===// +//===- dpctl_sycl_kernel_bundle_interface.h - C API for +// sycl::kernel_bundle -*-C++-*- ===// // // Data Parallel Control (dpctl) // @@ -84,7 +85,7 @@ __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromOCLSource( * @brief Returns the SyclKernel with given name from the program, if not found * then return NULL. * - * @param PRef Opaque pointer to a sycl::program + * @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 @@ -98,7 +99,7 @@ DPCTLKernelBundle_GetKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, * @brief Return True if a SyclKernel with given name exists in the program, if * not found then returns False. * - * @param PRef Opaque pointer to a sycl::program + * @param KBRef Opaque pointer to a sycl::kernel_bundle * @param KernelName Name of kernel * @return True if the kernel exists, else False * @ingroup KernelBundleInterface diff --git a/libsyclinterface/source/dpctl_sycl_program_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp similarity index 99% rename from libsyclinterface/source/dpctl_sycl_program_interface.cpp rename to libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp index c91aadbfb1..112894f078 100644 --- a/libsyclinterface/source/dpctl_sycl_program_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -1,5 +1,5 @@ -//===- dpctl_sycl_program_interface.cpp - Implements C API for -// sycl::kernel_bundle =// +//===- dpctl_sycl_kernel_bundle_interface.cpp - Implements C API for +// sycl::kernel_bundle ---------------===// // // Data Parallel Control (dpctl) // @@ -21,11 +21,11 @@ /// /// \file /// This file implements the functions declared in -/// dpctl_sycl_program_interface.h. +/// dpctl_sycl_kernel_bundle_interface.h. /// //===----------------------------------------------------------------------===// -#include "dpctl_sycl_program_interface.h" +#include "dpctl_sycl_kernel_bundle_interface.h" #include "Config/dpctl_config.h" #include "Support/CBindingWrapping.h" #include "dpctl_dynamic_lib_helper.h" diff --git a/libsyclinterface/tests/test_sycl_kernel_interface.cpp b/libsyclinterface/tests/test_sycl_kernel_interface.cpp index 0db8c227bf..89cc586aab 100644 --- a/libsyclinterface/tests/test_sycl_kernel_interface.cpp +++ b/libsyclinterface/tests/test_sycl_kernel_interface.cpp @@ -1,4 +1,4 @@ -//===-- test_sycl_program_interface.cpp - Test cases for kernel interface ===// +//===-- test_sycl_kernel_interface.cpp - Test cases for kernel interface ===// // // Data Parallel Control (dpctl) // @@ -27,8 +27,8 @@ #include "dpctl_sycl_context_interface.h" #include "dpctl_sycl_device_interface.h" #include "dpctl_sycl_device_selector_interface.h" +#include "dpctl_sycl_kernel_bundle_interface.h" #include "dpctl_sycl_kernel_interface.h" -#include "dpctl_sycl_program_interface.h" #include "dpctl_sycl_queue_interface.h" #include "dpctl_sycl_queue_manager.h" #include "dpctl_utils.h" diff --git a/libsyclinterface/tests/test_sycl_program_interface.cpp b/libsyclinterface/tests/test_sycl_program_interface.cpp index a6ffe0783b..54f655e752 100644 --- a/libsyclinterface/tests/test_sycl_program_interface.cpp +++ b/libsyclinterface/tests/test_sycl_program_interface.cpp @@ -1,4 +1,5 @@ -//===-- test_sycl_program_interface.cpp - Test cases for module interface -===// +//===- test_sycl_kernel_bundle_interface.cpp - +// Test cases for module interface -===// // // Data Parallel Control (dpctl) // @@ -29,8 +30,8 @@ #include "dpctl_sycl_device_interface.h" #include "dpctl_sycl_device_manager.h" #include "dpctl_sycl_device_selector_interface.h" +#include "dpctl_sycl_kernel_bundle_interface.h" #include "dpctl_sycl_kernel_interface.h" -#include "dpctl_sycl_program_interface.h" #include "dpctl_sycl_queue_interface.h" #include "dpctl_sycl_queue_manager.h" #include diff --git a/libsyclinterface/tests/test_sycl_queue_submit.cpp b/libsyclinterface/tests/test_sycl_queue_submit.cpp index efec07e7d0..87a3a91b55 100644 --- a/libsyclinterface/tests/test_sycl_queue_submit.cpp +++ b/libsyclinterface/tests/test_sycl_queue_submit.cpp @@ -28,8 +28,8 @@ #include "dpctl_sycl_device_interface.h" #include "dpctl_sycl_device_selector_interface.h" #include "dpctl_sycl_event_interface.h" +#include "dpctl_sycl_kernel_bundle_interface.h" #include "dpctl_sycl_kernel_interface.h" -#include "dpctl_sycl_program_interface.h" #include "dpctl_sycl_queue_interface.h" #include "dpctl_sycl_usm_interface.h" #include From 97d6912c08df09720b3702af653be8faea694718 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Tue, 24 May 2022 12:13:24 -0500 Subject: [PATCH 06/14] Expanded program test suite Make sure to exercise addressof_ref() methods of SyclProgram and SyclKernel. Add an example of invalid source code to trip throwing of an exception. --- dpctl/tests/test_sycl_program.py | 61 +++++++++++++++++++------------- 1 file changed, 36 insertions(+), 25 deletions(-) diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index f69fb8c410..bd55e2b4cf 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -31,6 +31,24 @@ def get_spirv_abspath(fn): return spirv_file +def _check_multi_kernel_program(prog): + assert type(prog) is dpctl_prog.SyclProgram + + assert type(prog.addressof_ref()) is int + assert prog.has_sycl_kernel("add") + assert prog.has_sycl_kernel("axpy") + + addKernel = prog.get_sycl_kernel("add") + axpyKernel = prog.get_sycl_kernel("axpy") + + assert "add" == addKernel.get_function_name() + assert "axpy" == axpyKernel.get_function_name() + assert 3 == addKernel.get_num_args() + assert 4 == axpyKernel.get_num_args() + assert type(addKernel.addressof_ref()) is int + assert type(axpyKernel.addressof_ref()) is int + + def test_create_program_from_source_ocl(): oclSrc = " \ kernel void add(global int* a, global int* b, global int* c) { \ @@ -46,18 +64,7 @@ def test_create_program_from_source_ocl(): except dpctl.SyclQueueCreationError: pytest.skip("No OpenCL queue is available") prog = dpctl_prog.create_program_from_source(q, oclSrc) - assert prog is not None - - assert prog.has_sycl_kernel("add") - assert prog.has_sycl_kernel("axpy") - - addKernel = prog.get_sycl_kernel("add") - axpyKernel = prog.get_sycl_kernel("axpy") - - assert "add" == addKernel.get_function_name() - assert "axpy" == axpyKernel.get_function_name() - assert 3 == addKernel.get_num_args() - assert 4 == axpyKernel.get_num_args() + _check_multi_kernel_program(prog) def test_create_program_from_spirv_ocl(): @@ -69,17 +76,7 @@ def test_create_program_from_spirv_ocl(): with open(spirv_file, "rb") as fin: spirv = fin.read() prog = dpctl_prog.create_program_from_spirv(q, spirv) - assert prog is not None - assert prog.has_sycl_kernel("add") - assert prog.has_sycl_kernel("axpy") - - addKernel = prog.get_sycl_kernel("add") - axpyKernel = prog.get_sycl_kernel("axpy") - - assert "add" == addKernel.get_function_name() - assert "axpy" == axpyKernel.get_function_name() - assert 3 == addKernel.get_num_args() - assert 4 == axpyKernel.get_num_args() + _check_multi_kernel_program(prog) def test_create_program_from_spirv_l0(): @@ -90,7 +87,8 @@ def test_create_program_from_spirv_l0(): spirv_file = get_spirv_abspath("multi_kernel.spv") with open(spirv_file, "rb") as fin: spirv = fin.read() - dpctl_prog.create_program_from_spirv(q, spirv) + prog = dpctl_prog.create_program_from_spirv(q, spirv) + _check_multi_kernel_program(prog) @pytest.mark.xfail( @@ -110,4 +108,17 @@ def test_create_program_from_source_l0(): size_t index = get_global_id(0); \ c[index] = a[index] + d*b[index]; \ }" - dpctl_prog.create_program_from_source(q, oclSrc) + prog = dpctl_prog.create_program_from_source(q, oclSrc) + _check_multi_kernel_program(prog) + + +def test_create_program_from_invalid_src_ocl(): + try: + q = dpctl.SyclQueue("opencl") + except dpctl.SyclQueueCreationError: + pytest.skip("No OpenCL queue is available") + invalid_oclSrc = " \ + kernel void add( \ + }" + with pytest.raises(dpctl_prog.SyclProgramCompilationError): + dpctl_prog.create_program_from_source(q, invalid_oclSrc) From ca9517c7c9e1296bb0edc3d27915b3183c5ceab4 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Tue, 24 May 2022 12:59:41 -0500 Subject: [PATCH 07/14] backend::cuda -> backend::ext_oneapi_cuda With this change, removed use of -Wno-deprecated-declarations, since code no longer uses any. --- libsyclinterface/CMakeLists.txt | 2 -- libsyclinterface/helper/source/dpctl_utils_helper.cpp | 4 ++-- libsyclinterface/source/dpctl_sycl_context_interface.cpp | 2 +- libsyclinterface/tests/test_helper.cpp | 6 +++--- libsyclinterface/tests/test_sycl_queue_interface.cpp | 2 +- 5 files changed, 7 insertions(+), 9 deletions(-) diff --git a/libsyclinterface/CMakeLists.txt b/libsyclinterface/CMakeLists.txt index 05af34940b..a3ceca466d 100644 --- a/libsyclinterface/CMakeLists.txt +++ b/libsyclinterface/CMakeLists.txt @@ -104,7 +104,6 @@ if(WIN32) "-Wunused-function " "-Wuninitialized " "-Wmissing-declarations " - "-Wno-deprecated-declarations " ) set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${WARNING_FLAGS}") set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${WARNING_FLAGS}") @@ -123,7 +122,6 @@ elseif(UNIX) "-Wuninitialized " "-Wmissing-declarations " "-fdiagnostics-color=auto " - "-Wno-deprecated-declarations " ) string(CONCAT SDL_FLAGS "-fstack-protector " diff --git a/libsyclinterface/helper/source/dpctl_utils_helper.cpp b/libsyclinterface/helper/source/dpctl_utils_helper.cpp index 9aeec30d63..7a10f5f894 100644 --- a/libsyclinterface/helper/source/dpctl_utils_helper.cpp +++ b/libsyclinterface/helper/source/dpctl_utils_helper.cpp @@ -89,7 +89,7 @@ backend DPCTL_DPCTLBackendTypeToSyclBackend(DPCTLSyclBackendType BeTy) { switch (BeTy) { case DPCTLSyclBackendType::DPCTL_CUDA: - return backend::cuda; + return backend::ext_oneapi_cuda; case DPCTLSyclBackendType::DPCTL_HOST: return backend::host; case DPCTLSyclBackendType::DPCTL_LEVEL_ZERO: @@ -106,7 +106,7 @@ backend DPCTL_DPCTLBackendTypeToSyclBackend(DPCTLSyclBackendType BeTy) DPCTLSyclBackendType DPCTL_SyclBackendToDPCTLBackendType(backend B) { switch (B) { - case backend::cuda: + case backend::ext_oneapi_cuda: return DPCTLSyclBackendType::DPCTL_CUDA; case backend::host: return DPCTLSyclBackendType::DPCTL_HOST; diff --git a/libsyclinterface/source/dpctl_sycl_context_interface.cpp b/libsyclinterface/source/dpctl_sycl_context_interface.cpp index 157b122236..9709e7574e 100644 --- a/libsyclinterface/source/dpctl_sycl_context_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_context_interface.cpp @@ -196,7 +196,7 @@ DPCTLContext_GetBackend(__dpctl_keep const DPCTLSyclContextRef CtxRef) return DPCTL_OPENCL; case backend::ext_oneapi_level_zero: return DPCTL_LEVEL_ZERO; - case backend::cuda: + case backend::ext_oneapi_cuda: return DPCTL_CUDA; default: return DPCTL_UNKNOWN_BACKEND; diff --git a/libsyclinterface/tests/test_helper.cpp b/libsyclinterface/tests/test_helper.cpp index ea7d638aab..dedaeae10b 100644 --- a/libsyclinterface/tests/test_helper.cpp +++ b/libsyclinterface/tests/test_helper.cpp @@ -90,7 +90,7 @@ TEST_F(TestHelperFns, ChkDPCTLBackendTypeToSyclBackend) EXPECT_NO_FATAL_FAILURE(res = DPCTL_DPCTLBackendTypeToSyclBackend( DPCTLSyclBackendType::DPCTL_CUDA)); - ASSERT_TRUE(res == sycl::backend::cuda); + ASSERT_TRUE(res == sycl::backend::ext_oneapi_cuda); EXPECT_NO_FATAL_FAILURE(res = DPCTL_DPCTLBackendTypeToSyclBackend( DPCTLSyclBackendType::DPCTL_HOST)); @@ -125,8 +125,8 @@ TEST_F(TestHelperFns, ChkSyclBackendToDPCTLBackendType) DTy = DPCTL_SyclBackendToDPCTLBackendType(sycl::backend::host)); ASSERT_TRUE(DTy == DPCTLSyclBackendType::DPCTL_HOST); - EXPECT_NO_FATAL_FAILURE( - DTy = DPCTL_SyclBackendToDPCTLBackendType(sycl::backend::cuda)); + EXPECT_NO_FATAL_FAILURE(DTy = DPCTL_SyclBackendToDPCTLBackendType( + sycl::backend::ext_oneapi_cuda)); ASSERT_TRUE(DTy == DPCTLSyclBackendType::DPCTL_CUDA); EXPECT_NO_FATAL_FAILURE( diff --git a/libsyclinterface/tests/test_sycl_queue_interface.cpp b/libsyclinterface/tests/test_sycl_queue_interface.cpp index 9653735f21..25cb4e5481 100644 --- a/libsyclinterface/tests/test_sycl_queue_interface.cpp +++ b/libsyclinterface/tests/test_sycl_queue_interface.cpp @@ -353,7 +353,7 @@ TEST_P(TestDPCTLQueueMemberFunctions, CheckGetBackend) auto Bty = DPCTLQueue_GetBackend(QRef); switch (Bty) { case DPCTL_CUDA: - EXPECT_TRUE(Backend == backend::cuda); + EXPECT_TRUE(Backend == backend::ext_oneapi_cuda); break; case DPCTL_HOST: EXPECT_TRUE(Backend == backend::host); From b4d2d8b38817eb5c7d26ad17aa3f979c2fde0ce2 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Tue, 24 May 2022 13:17:26 -0500 Subject: [PATCH 08/14] Remove setting -Wno-deprecated-declarations in examples --- examples/pybind11/external_usm_allocation/CMakeLists.txt | 1 - examples/pybind11/onemkl_gemv/CMakeLists.txt | 2 +- examples/pybind11/use_dpctl_syclqueue/CMakeLists.txt | 1 - 3 files changed, 1 insertion(+), 3 deletions(-) diff --git a/examples/pybind11/external_usm_allocation/CMakeLists.txt b/examples/pybind11/external_usm_allocation/CMakeLists.txt index 6f42a3e6d3..d3ba8f4dd9 100644 --- a/examples/pybind11/external_usm_allocation/CMakeLists.txt +++ b/examples/pybind11/external_usm_allocation/CMakeLists.txt @@ -28,7 +28,6 @@ pybind11_add_module(${py_module_name} external_usm_allocation/_usm_alloc_example.cpp ) target_include_directories(${py_module_name} PUBLIC ${Dpctl_INCLUDE_DIRS}) -target_compile_options(${py_module_name} PRIVATE -Wno-deprecated-declarations) install(TARGETS ${py_module_name} DESTINATION external_usm_allocation ) diff --git a/examples/pybind11/onemkl_gemv/CMakeLists.txt b/examples/pybind11/onemkl_gemv/CMakeLists.txt index 848e8c7727..b9c4b087ed 100644 --- a/examples/pybind11/onemkl_gemv/CMakeLists.txt +++ b/examples/pybind11/onemkl_gemv/CMakeLists.txt @@ -50,7 +50,7 @@ target_include_directories(${py_module_name} PUBLIC ${Dpctl_INCLUDE_DIRS}) get_target_property(_sycl_gemm_sources ${py_module_name} SOURCES) set_source_files_properties(${_sycl_gemm_sources} PROPERTIES - COMPILE_OPTIONS "-O3;-Wno-deprecated-declarations" + COMPILE_OPTIONS "-O3" ) add_executable(standalone_cpp diff --git a/examples/pybind11/use_dpctl_syclqueue/CMakeLists.txt b/examples/pybind11/use_dpctl_syclqueue/CMakeLists.txt index a788c56bae..0d4e262d1e 100644 --- a/examples/pybind11/use_dpctl_syclqueue/CMakeLists.txt +++ b/examples/pybind11/use_dpctl_syclqueue/CMakeLists.txt @@ -28,7 +28,6 @@ pybind11_add_module(${py_module_name} use_queue_device/_example.cpp ) target_include_directories(${py_module_name} PUBLIC ${Dpctl_INCLUDE_DIRS}) -target_compile_options(${py_module_name} PRIVATE -Wno-deprecated-declarations) install(TARGETS ${py_module_name} DESTINATION use_queue_device ) From 1f6e0a8364cb3b6e44fc47dfad5e92a9d2dabcc3 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Wed, 25 May 2022 11:27:13 -0500 Subject: [PATCH 09/14] Cleaned up _GetErrorCode_*_impl --- .../dpctl_sycl_kernel_bundle_interface.cpp | 112 +++++------------- 1 file changed, 32 insertions(+), 80 deletions(-) diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp index 112894f078..afe62d1e51 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -64,6 +64,13 @@ static const int clLibLoadFlags = 0; #error "OpenCL program compilation is unavailable for this platform" #endif +#define CodeStringSuffix(code) \ + std::string(" (code=") + std::to_string(static_cast(code)) + ")" + +#define EnumCaseString(code) \ + case code: \ + return std::string(#code) + CodeStringSuffix(code) + constexpr backend cl_be = backend::opencl; struct cl_loader @@ -155,41 +162,18 @@ clCreateKernelFT get_clCreateKernel() std::string _GetErrorCode_ocl_impl(cl_int code) { - if (code == CL_BUILD_PROGRAM_FAILURE) { - return "CL_BUILD_PROGRAM_FAILURE (code=" + - std::to_string(static_cast(code)) + ")"; - } - else if (code == CL_INVALID_CONTEXT) { - return "CL_INVALID_CONTEXT (code=" + - std::to_string(static_cast(code)) + ")"; - } - else if (code == CL_INVALID_DEVICE) { - return "CL_INVALID_DEVICE (code=" + - std::to_string(static_cast(code)) + ")"; - } - else if (code == CL_INVALID_VALUE) { - return "CL_INVALID_VALUE (code=" + - std::to_string(static_cast(code)) + ")"; - } - else if (code == CL_OUT_OF_RESOURCES) { - return "CL_OUT_OF_RESOURCES (code=" + - std::to_string(static_cast(code)) + ")"; - } - else if (code == CL_OUT_OF_HOST_MEMORY) { - return "CL_OUT_OF_HOST_MEMORY (code=" + - std::to_string(static_cast(code)) + ")"; - } - else if (code == CL_INVALID_OPERATION) { - return "CL_INVALID_OPERATION (code=" + - std::to_string(static_cast(code)) + ")"; - } - else if (code == CL_INVALID_BINARY) { - return "CL_INVALID_BINARY (code=" + - std::to_string(static_cast(code)) + ")"; + switch (code) { + EnumCaseString(CL_BUILD_PROGRAM_FAILURE); + EnumCaseString(CL_INVALID_CONTEXT); + EnumCaseString(CL_INVALID_DEVICE); + EnumCaseString(CL_INVALID_VALUE); + EnumCaseString(CL_OUT_OF_RESOURCES); + EnumCaseString(CL_OUT_OF_HOST_MEMORY); + EnumCaseString(CL_INVALID_OPERATION); + EnumCaseString(CL_INVALID_BINARY); + default: + return "<< ERROR CODE UNRECOGNIZED >>" + CodeStringSuffix(code); } - - return "<< ERROR CODE UNRECOGNIZED >> (code=" + - std::to_string(static_cast(code)) + ")"; } DPCTLSyclKernelBundleRef @@ -426,53 +410,21 @@ zeKernelCreateFT get_zeKernelCreate() std::string _GetErrorCode_ze_impl(ze_result_t code) { - if (code == ZE_RESULT_ERROR_UNINITIALIZED) { - return "ZE_RESULT_ERROR_UNINITIALIZED (code=" + - std::to_string(static_cast(code)) + ")"; - } - else if (code == ZE_RESULT_ERROR_DEVICE_LOST) { - return "ZE_RESULT_ERROR_DEVICE_LOST (code=" + - std::to_string(static_cast(code)) + ")"; - } - else if (code == ZE_RESULT_ERROR_INVALID_NULL_HANDLE) { - return "ZE_RESULT_ERROR_INVALID_NULL_HANDLE (code=" + - std::to_string(static_cast(code)) + ")"; - } - else if (code == ZE_RESULT_ERROR_INVALID_NULL_POINTER) { - return "ZE_RESULT_ERROR_INVALID_NULL_POINTER (code=" + - std::to_string(static_cast(code)) + ")"; - } - else if (code == ZE_RESULT_ERROR_INVALID_ENUMERATION) { - return "ZE_RESULT_ERROR_INVALID_ENUMERATION (code=" + - std::to_string(static_cast(code)) + ")"; - } - else if (code == ZE_RESULT_ERROR_INVALID_NATIVE_BINARY) { - return "ZE_RESULT_ERROR_INVALID_NATIVE_BINARY (code=" + - std::to_string(static_cast(code)) + ")"; - } - else if (code == ZE_RESULT_ERROR_INVALID_SIZE) { - return "ZE_RESULT_ERROR_INVALID_SIZE (code=" + - std::to_string(static_cast(code)) + ")"; - } - else if (code == ZE_RESULT_ERROR_OUT_OF_HOST_MEMORY) { - return "ZE_RESULT_ERROR_OUT_OF_HOST_MEMORY (code=" + - std::to_string(static_cast(code)) + ")"; - } - else if (code == ZE_RESULT_ERROR_OUT_OF_DEVICE_MEMORY) { - return "ZE_RESULT_ERROR_OUT_OF_DEVICE_MEMORY (code=" + - std::to_string(static_cast(code)) + ")"; - } - else if (code == ZE_RESULT_ERROR_MODULE_BUILD_FAILURE) { - return "ZE_RESULT_ERROR_MODULE_BUILD_FAILURE (code=" + - std::to_string(static_cast(code)) + ")"; - } - else if (code == ZE_RESULT_ERROR_INVALID_MODULE_UNLINKED) { - return "ZE_RESULT_ERROR_INVALID_MODULE_UNLINKED (code=" + - std::to_string(static_cast(code)) + ")"; + switch (code) { + EnumCaseString(ZE_RESULT_ERROR_UNINITIALIZED); + EnumCaseString(ZE_RESULT_ERROR_DEVICE_LOST); + EnumCaseString(ZE_RESULT_ERROR_INVALID_NULL_HANDLE); + EnumCaseString(ZE_RESULT_ERROR_INVALID_NULL_POINTER); + EnumCaseString(ZE_RESULT_ERROR_INVALID_ENUMERATION); + EnumCaseString(ZE_RESULT_ERROR_INVALID_NATIVE_BINARY); + EnumCaseString(ZE_RESULT_ERROR_INVALID_SIZE); + EnumCaseString(ZE_RESULT_ERROR_OUT_OF_HOST_MEMORY); + EnumCaseString(ZE_RESULT_ERROR_OUT_OF_DEVICE_MEMORY); + EnumCaseString(ZE_RESULT_ERROR_MODULE_BUILD_FAILURE); + EnumCaseString(ZE_RESULT_ERROR_INVALID_MODULE_UNLINKED); + default: + return "<< UNRECOGNIZED ZE_RESULT_T CODE >> " + CodeStringSuffix(code); } - - return "<< UNRECOGNIZE ZE_RESULT_T CODE >> (code=" + - std::to_string(static_cast(code)) + ")"; } __dpctl_give DPCTLSyclKernelBundleRef From fc400fdd125520cabc3a7fed336462bd342ed4c1 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Wed, 25 May 2022 13:05:24 -0500 Subject: [PATCH 10/14] Fixed get_clCreateKernel implementation to use cl_loader --- .../source/dpctl_sycl_kernel_bundle_interface.cpp | 9 +-------- 1 file changed, 1 insertion(+), 8 deletions(-) diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp index afe62d1e51..ca1fbed8bf 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -147,15 +147,8 @@ typedef cl_kernel (*clCreateKernelFT)(cl_program, const char *, cl_int *); const char *clCreateKernel_Name = "clCreateKernel"; clCreateKernelFT get_clCreateKernel() { - static dpctl::DynamicLibHelper clLib(clLoaderName, clLibLoadFlags); - if (!clLib.opened()) { - error_handler("The OpenCL loader dynamic library could not " - "be opened.", - __FILE__, __func__, __LINE__); - return nullptr; - } static auto st_clCreateKernelF = - clLib.getSymbol(clCreateKernel_Name); + cl_loader::get().getSymbol(clCreateKernel_Name); return st_clCreateKernelF; } From 65a5d4f556ca9da5a195be65e1ea226571d2a265 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Wed, 25 May 2022 15:05:23 -0500 Subject: [PATCH 11/14] Added tests for null-dref, null-spirv pointer Added a test for non-supported backend (host). This only runs if SYCL_ENABLE_HOST_DEVICE environment variable is set. --- .../dpctl_sycl_kernel_bundle_interface.cpp | 5 + ... => test_sycl_kernel_bundle_interface.cpp} | 148 ++++++++++++++---- 2 files changed, 124 insertions(+), 29 deletions(-) rename libsyclinterface/tests/{test_sycl_program_interface.cpp => test_sycl_kernel_bundle_interface.cpp} (62%) diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp index ca1fbed8bf..625a303a0b 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -594,6 +594,11 @@ DPCTLKernelBundle_CreateFromSpirv(__dpctl_keep const DPCTLSyclContextRef CtxRef, __FILE__, __func__, __LINE__); return KBRef; } + if ((!IL) || (length == 0)) { + error_handler("Cannot create program from null SPIR-V buffer.", + __FILE__, __func__, __LINE__); + return KBRef; + } context *SyclCtx = unwrap(CtxRef); device *SyclDev = unwrap(DevRef); diff --git a/libsyclinterface/tests/test_sycl_program_interface.cpp b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp similarity index 62% rename from libsyclinterface/tests/test_sycl_program_interface.cpp rename to libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp index 54f655e752..7a8b3fdf38 100644 --- a/libsyclinterface/tests/test_sycl_program_interface.cpp +++ b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp @@ -42,24 +42,22 @@ using namespace cl::sycl; -struct TestDPCTLSyclProgramInterface +struct TestDPCTLSyclKernelBundleInterface : public ::testing::TestWithParam { DPCTLSyclDeviceRef DRef = nullptr; DPCTLSyclContextRef CRef = nullptr; - DPCTLSyclQueueRef QRef = nullptr; DPCTLSyclKernelBundleRef KBRef = nullptr; std::ifstream spirvFile; size_t spirvFileSize; std::vector spirvBuffer; - TestDPCTLSyclProgramInterface() + TestDPCTLSyclKernelBundleInterface() { auto DS = DPCTLFilterSelector_Create(GetParam()); DRef = DPCTLDevice_CreateFromSelector(DS); DPCTLDeviceSelector_Delete(DS); CRef = DPCTLDeviceMgr_GetCachedContext(DRef); - QRef = DPCTLQueue_Create(CRef, DRef, nullptr, DPCTL_DEFAULT_PROPERTY); if (DRef) { spirvFile.open("./multi_kernel.spv", @@ -82,18 +80,20 @@ struct TestDPCTLSyclProgramInterface } } - ~TestDPCTLSyclProgramInterface() + ~TestDPCTLSyclKernelBundleInterface() { - if (DRef) + if (DRef) { spirvFile.close(); - DPCTLDevice_Delete(DRef); - DPCTLQueue_Delete(QRef); - DPCTLContext_Delete(CRef); - DPCTLKernelBundle_Delete(KBRef); + DPCTLDevice_Delete(DRef); + } + if (CRef) + DPCTLContext_Delete(CRef); + if (KBRef) + DPCTLKernelBundle_Delete(KBRef); } }; -TEST_P(TestDPCTLSyclProgramInterface, ChkCreateFromSpirv) +TEST_P(TestDPCTLSyclKernelBundleInterface, ChkCreateFromSpirv) { ASSERT_TRUE(KBRef != nullptr); @@ -102,25 +102,36 @@ TEST_P(TestDPCTLSyclProgramInterface, ChkCreateFromSpirv) ASSERT_FALSE(DPCTLKernelBundle_HasKernel(KBRef, nullptr)); } -TEST_P(TestDPCTLSyclProgramInterface, ChkCreateFromSpirvNull) +TEST_P(TestDPCTLSyclKernelBundleInterface, ChkCreateFromSpirvNull) { DPCTLSyclContextRef Null_CRef = nullptr; DPCTLSyclDeviceRef Null_DRef = nullptr; const void *null_spirv = nullptr; DPCTLSyclKernelBundleRef KBRef = nullptr; + // Null context EXPECT_NO_FATAL_FAILURE(KBRef = DPCTLKernelBundle_CreateFromSpirv( Null_CRef, Null_DRef, null_spirv, 0, nullptr)); ASSERT_TRUE(KBRef == nullptr); + + // Null device + EXPECT_NO_FATAL_FAILURE(KBRef = DPCTLKernelBundle_CreateFromSpirv( + CRef, Null_DRef, null_spirv, 0, nullptr)); + ASSERT_TRUE(KBRef == nullptr); + + // Null IL + EXPECT_NO_FATAL_FAILURE(KBRef = DPCTLKernelBundle_CreateFromSpirv( + CRef, DRef, null_spirv, 0, nullptr)); + ASSERT_TRUE(KBRef == nullptr); } -TEST_P(TestDPCTLSyclProgramInterface, ChkHasKernelNullProgram) +TEST_P(TestDPCTLSyclKernelBundleInterface, ChkHasKernelNullProgram) { DPCTLSyclKernelBundleRef NullRef = nullptr; ASSERT_FALSE(DPCTLKernelBundle_HasKernel(NullRef, "add")); } -TEST_P(TestDPCTLSyclProgramInterface, ChkGetKernel) +TEST_P(TestDPCTLSyclKernelBundleInterface, ChkGetKernel) { auto AddKernel = DPCTLKernelBundle_GetKernel(KBRef, "add"); auto AxpyKernel = DPCTLKernelBundle_GetKernel(KBRef, "axpy"); @@ -134,7 +145,7 @@ TEST_P(TestDPCTLSyclProgramInterface, ChkGetKernel) EXPECT_NO_FATAL_FAILURE(DPCTLKernel_Delete(NullKernel)); } -TEST_P(TestDPCTLSyclProgramInterface, ChkGetKernelNullProgram) +TEST_P(TestDPCTLSyclKernelBundleInterface, ChkGetKernelNullProgram) { DPCTLSyclKernelBundleRef NullRef = nullptr; DPCTLSyclKernelRef KRef = nullptr; @@ -143,7 +154,7 @@ TEST_P(TestDPCTLSyclProgramInterface, ChkGetKernelNullProgram) EXPECT_TRUE(KRef == nullptr); } -struct TestOCLProgramFromSource : public ::testing::Test +struct TestOCLKernelBundleFromSource : public ::testing::Test { const char *CLProgramStr = R"CLC( kernel void add(global int* a, global int* b, global int* c) { @@ -160,32 +171,32 @@ struct TestOCLProgramFromSource : public ::testing::Test const char *CompileOpts = "-cl-fast-relaxed-math"; DPCTLSyclDeviceRef DRef = nullptr; DPCTLSyclContextRef CRef = nullptr; - DPCTLSyclQueueRef QRef = nullptr; DPCTLSyclKernelBundleRef KBRef = nullptr; - TestOCLProgramFromSource() + TestOCLKernelBundleFromSource() { auto DS = DPCTLFilterSelector_Create("opencl:gpu"); DRef = DPCTLDevice_CreateFromSelector(DS); DPCTLDeviceSelector_Delete(DS); CRef = DPCTLDeviceMgr_GetCachedContext(DRef); - QRef = DPCTLQueue_Create(CRef, DRef, nullptr, DPCTL_DEFAULT_PROPERTY); if (DRef) KBRef = DPCTLKernelBundle_CreateFromOCLSource( CRef, DRef, CLProgramStr, CompileOpts); } - ~TestOCLProgramFromSource() + ~TestOCLKernelBundleFromSource() { - DPCTLDevice_Delete(DRef); - DPCTLQueue_Delete(QRef); - DPCTLContext_Delete(CRef); - DPCTLKernelBundle_Delete(KBRef); + if (DRef) + DPCTLDevice_Delete(DRef); + if (CRef) + DPCTLContext_Delete(CRef); + if (KBRef) + DPCTLKernelBundle_Delete(KBRef); } }; -TEST_F(TestOCLProgramFromSource, CheckCreateFromOCLSource) +TEST_F(TestOCLKernelBundleFromSource, CheckCreateFromOCLSource) { if (!DRef) GTEST_SKIP_("Skipping as no OpenCL GPU device found.\n"); @@ -195,7 +206,7 @@ TEST_F(TestOCLProgramFromSource, CheckCreateFromOCLSource) ASSERT_TRUE(DPCTLKernelBundle_HasKernel(KBRef, "axpy")); } -TEST_F(TestOCLProgramFromSource, CheckCreateFromOCLSourceNull) +TEST_F(TestOCLKernelBundleFromSource, CheckCreateFromOCLSourceNull) { const char *InvalidCLProgramStr = R"CLC( kernel void invalid(global foo* a, global bar* b) { @@ -211,9 +222,21 @@ TEST_F(TestOCLProgramFromSource, CheckCreateFromOCLSourceNull) EXPECT_NO_FATAL_FAILURE(KBRef = DPCTLKernelBundle_CreateFromOCLSource( CRef, DRef, InvalidCLProgramStr, CompileOpts);); ASSERT_TRUE(KBRef == nullptr); + + DPCTLSyclContextRef Null_CRef = nullptr; + EXPECT_NO_FATAL_FAILURE( + KBRef = DPCTLKernelBundle_CreateFromOCLSource( + Null_CRef, DRef, InvalidCLProgramStr, CompileOpts);); + ASSERT_TRUE(KBRef == nullptr); + + DPCTLSyclDeviceRef Null_DRef = nullptr; + EXPECT_NO_FATAL_FAILURE( + KBRef = DPCTLKernelBundle_CreateFromOCLSource( + CRef, Null_DRef, InvalidCLProgramStr, CompileOpts);); + ASSERT_TRUE(KBRef == nullptr); } -TEST_F(TestOCLProgramFromSource, CheckGetKernelOCLSource) +TEST_F(TestOCLKernelBundleFromSource, CheckGetKernelOCLSource) { if (!DRef) GTEST_SKIP_("Skipping as no OpenCL GPU device found.\n"); @@ -226,8 +249,8 @@ TEST_F(TestOCLProgramFromSource, CheckGetKernelOCLSource) DPCTLKernel_Delete(AxpyKernel); } -INSTANTIATE_TEST_SUITE_P(ProgramCreationFromSpriv, - TestDPCTLSyclProgramInterface, +INSTANTIATE_TEST_SUITE_P(KernelBundleCreationFromSpirv, + TestDPCTLSyclKernelBundleInterface, ::testing::Values("opencl", "opencl:gpu", "opencl:cpu", @@ -237,3 +260,70 @@ INSTANTIATE_TEST_SUITE_P(ProgramCreationFromSpriv, "level_zero:gpu", #endif "opencl:cpu:0")); + +struct TestKernelBundleUnsupportedBackend : public ::testing::Test +{ + DPCTLSyclDeviceRef DRef = nullptr; + DPCTLSyclContextRef CRef = nullptr; + + TestKernelBundleUnsupportedBackend() + { + auto DS = DPCTLFilterSelector_Create("host:host"); + DRef = DPCTLDevice_CreateFromSelector(DS); + DPCTLDeviceSelector_Delete(DS); + if (DRef) + CRef = DPCTLDeviceMgr_GetCachedContext(DRef); + } + + void SetUp() + { + if (!DRef) { + std::string message = "Skipping as host device is not enabled."; + GTEST_SKIP_(message.c_str()); + } + } + + ~TestKernelBundleUnsupportedBackend() + { + if (DRef) + DPCTLDevice_Delete(DRef); + if (CRef) + DPCTLContext_Delete(CRef); + } +}; + +TEST_F(TestKernelBundleUnsupportedBackend, CheckCreateFromSource) +{ + const char *src = R"CLC( + kernel void set(global int* a, int v) { + size_t index = get_global_id(0); + a[index] = v; + } + )CLC"; + const char *opts = ""; + + DPCTLSyclKernelBundleRef KBRef = nullptr; + EXPECT_NO_FATAL_FAILURE( + KBRef = DPCTLKernelBundle_CreateFromOCLSource(CRef, DRef, src, opts)); + ASSERT_TRUE(KBRef == nullptr); +} + +TEST_F(TestKernelBundleUnsupportedBackend, CheckCreateFromSpirv) +{ + std::ifstream spirvFile; + size_t spirvFileSize; + std::vector spirvBuffer; + + spirvFile.open("./multi_kernel.spv", std::ios::binary | std::ios::ate); + spirvFileSize = std::filesystem::file_size("./multi_kernel.spv"); + spirvBuffer.reserve(spirvFileSize); + spirvFile.seekg(0, std::ios::beg); + spirvFile.read(spirvBuffer.data(), spirvFileSize); + spirvFile.close(); + + DPCTLSyclKernelBundleRef KBRef = nullptr; + EXPECT_NO_FATAL_FAILURE( + KBRef = DPCTLKernelBundle_CreateFromSpirv( + CRef, DRef, spirvBuffer.data(), spirvFileSize, nullptr)); + ASSERT_TRUE(KBRef == nullptr); +} From 8a7bae8642c34851d3be167c9ed31ed0ec892fb2 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Wed, 25 May 2022 17:47:42 -0500 Subject: [PATCH 12/14] Run coverage workflow with host device enabled Also parametrized test for compiling from source to run on OCL gpu and OCL cpu if found. --- .github/workflows/generate-coverage.yaml | 2 +- .../test_sycl_kernel_bundle_interface.cpp | 33 +++++++++++-------- 2 files changed, 20 insertions(+), 15 deletions(-) diff --git a/.github/workflows/generate-coverage.yaml b/.github/workflows/generate-coverage.yaml index d1ef2618cf..7a1183e29b 100644 --- a/.github/workflows/generate-coverage.yaml +++ b/.github/workflows/generate-coverage.yaml @@ -85,7 +85,7 @@ jobs: shell: bash -l {0} run: | source /opt/intel/oneapi/setvars.sh - python scripts/gen_coverage.py + SYCL_ENABLE_HOST_DEVICE=1 python scripts/gen_coverage.py - name: Install coverall dependencies shell: bash -l {0} diff --git a/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp index 7a8b3fdf38..9bbd55f009 100644 --- a/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp @@ -154,7 +154,8 @@ TEST_P(TestDPCTLSyclKernelBundleInterface, ChkGetKernelNullProgram) EXPECT_TRUE(KRef == nullptr); } -struct TestOCLKernelBundleFromSource : public ::testing::Test +struct TestOCLKernelBundleFromSource + : public ::testing::TestWithParam { const char *CLProgramStr = R"CLC( kernel void add(global int* a, global int* b, global int* c) { @@ -175,7 +176,7 @@ struct TestOCLKernelBundleFromSource : public ::testing::Test TestOCLKernelBundleFromSource() { - auto DS = DPCTLFilterSelector_Create("opencl:gpu"); + auto DS = DPCTLFilterSelector_Create(GetParam()); DRef = DPCTLDevice_CreateFromSelector(DS); DPCTLDeviceSelector_Delete(DS); CRef = DPCTLDeviceMgr_GetCachedContext(DRef); @@ -185,6 +186,15 @@ struct TestOCLKernelBundleFromSource : public ::testing::Test CRef, DRef, CLProgramStr, CompileOpts); } + void SetUp() + { + if (!DRef) { + auto message = "Skipping as no device of type " + + std::string(GetParam()) + "."; + GTEST_SKIP_(message.c_str()); + } + } + ~TestOCLKernelBundleFromSource() { if (DRef) @@ -196,17 +206,14 @@ struct TestOCLKernelBundleFromSource : public ::testing::Test } }; -TEST_F(TestOCLKernelBundleFromSource, CheckCreateFromOCLSource) +TEST_P(TestOCLKernelBundleFromSource, CheckCreateFromOCLSource) { - if (!DRef) - GTEST_SKIP_("Skipping as no OpenCL GPU device found.\n"); - ASSERT_TRUE(KBRef != nullptr); ASSERT_TRUE(DPCTLKernelBundle_HasKernel(KBRef, "add")); ASSERT_TRUE(DPCTLKernelBundle_HasKernel(KBRef, "axpy")); } -TEST_F(TestOCLKernelBundleFromSource, CheckCreateFromOCLSourceNull) +TEST_P(TestOCLKernelBundleFromSource, CheckCreateFromOCLSourceNull) { const char *InvalidCLProgramStr = R"CLC( kernel void invalid(global foo* a, global bar* b) { @@ -216,9 +223,6 @@ TEST_F(TestOCLKernelBundleFromSource, CheckCreateFromOCLSourceNull) )CLC"; DPCTLSyclKernelBundleRef KBRef = nullptr; - if (!DRef) - GTEST_SKIP_("Skipping as no OpenCL GPU device found.\n"); - EXPECT_NO_FATAL_FAILURE(KBRef = DPCTLKernelBundle_CreateFromOCLSource( CRef, DRef, InvalidCLProgramStr, CompileOpts);); ASSERT_TRUE(KBRef == nullptr); @@ -236,11 +240,8 @@ TEST_F(TestOCLKernelBundleFromSource, CheckCreateFromOCLSourceNull) ASSERT_TRUE(KBRef == nullptr); } -TEST_F(TestOCLKernelBundleFromSource, CheckGetKernelOCLSource) +TEST_P(TestOCLKernelBundleFromSource, CheckGetKernelOCLSource) { - if (!DRef) - GTEST_SKIP_("Skipping as no OpenCL GPU device found.\n"); - auto AddKernel = DPCTLKernelBundle_GetKernel(KBRef, "add"); auto AxpyKernel = DPCTLKernelBundle_GetKernel(KBRef, "axpy"); ASSERT_TRUE(AddKernel != nullptr); @@ -261,6 +262,10 @@ INSTANTIATE_TEST_SUITE_P(KernelBundleCreationFromSpirv, #endif "opencl:cpu:0")); +INSTANTIATE_TEST_SUITE_P(KernelBundleCreationFromSource, + TestOCLKernelBundleFromSource, + ::testing::Values("opencl:gpu", "opencl:cpu")); + struct TestKernelBundleUnsupportedBackend : public ::testing::Test { DPCTLSyclDeviceRef DRef = nullptr; From 49000a3bfc79e2414268b5e66bd1064981546e73 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Tue, 31 May 2022 10:58:11 -0500 Subject: [PATCH 13/14] Fixed grammar in SyclProgram class docstring --- dpctl/program/_program.pxd | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dpctl/program/_program.pxd b/dpctl/program/_program.pxd index 8880b853ce..dffaea8c5f 100644 --- a/dpctl/program/_program.pxd +++ b/dpctl/program/_program.pxd @@ -41,7 +41,7 @@ cdef class SyclKernel: cdef class SyclProgram: - ''' Wraps a sycl::kernel_bundle object created from + ''' Wraps a sycl::kernel_bundle object created by using SYCL interoperability layer for OpenCL and Level-Zero backends. SyclProgram exposes the C API from dpctl_sycl_kernel_bundle_interface.h. A From afd12ef39bc33adc1b7ff78bcd21c8e0f084a260 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Thu, 2 Jun 2022 07:01:52 -0500 Subject: [PATCH 14/14] Update libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp Fixed typo in comment --- libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp index 625a303a0b..47cda39092 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -178,7 +178,7 @@ _CreateKernelBundle_common_ocl_impl(cl_program clProgram, backend_traits::return_type clDevice; clDevice = get_native(dev); - // Last to pointers are notification function pointer and user-data pointer + // Last two pointers are notification function pointer and user-data pointer // that can be passed to the notification function. auto clBuildProgramF = get_clBuldProgram(); if (clBuildProgramF == nullptr) {