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/dpctl/_backend.pxd b/dpctl/_backend.pxd index eb9124ffa5..7d46ad0b77 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 @@ -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) @@ -305,22 +304,24 @@ 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 DPCTLSyclProgramRef DPCTLProgram_CreateFromSpirv( +cdef extern from "syclinterface/dpctl_sycl_kernel_bundle_interface.h": + 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 e0b035793c..dffaea8c5f 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 @@ -33,26 +33,26 @@ 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: - ''' Wraps a sycl::program object created from an OpenCL interoperability - program. + ''' 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_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. ''' - 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 923deef367..f638087793 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -30,16 +30,16 @@ cimport cython.array from dpctl._backend cimport ( # noqa: E211, E402 DPCTLCString_Delete, DPCTLKernel_Delete, - DPCTLKernel_GetFunctionName, 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__ = [ @@ -51,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 @@ -61,20 +61,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. @@ -98,42 +97,45 @@ 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 - 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)) + 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) @@ -142,9 +144,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 @@ -155,24 +158,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, @@ -180,8 +186,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 @@ -192,20 +199,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/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) 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 ) diff --git a/libsyclinterface/CMakeLists.txt b/libsyclinterface/CMakeLists.txt index 695975f7ae..a3ceca466d 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 @@ -87,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}") @@ -106,7 +122,6 @@ elseif(UNIX) "-Wuninitialized " "-Wmissing-declarations " "-fdiagnostics-color=auto " - "-Wno-deprecated-declarations " ) string(CONCAT SDL_FLAGS "-fstack-protector " @@ -193,7 +208,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/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/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_kernel_bundle_interface.h similarity index 51% rename from libsyclinterface/include/dpctl_sycl_program_interface.h rename to libsyclinterface/include/dpctl_sycl_kernel_bundle_interface.h index 4667631a98..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) // @@ -35,87 +36,85 @@ 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 * 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 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 * 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 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_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/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_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/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp new file mode 100644 index 0000000000..47cda39092 --- /dev/null +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -0,0 +1,734 @@ +//===- dpctl_sycl_kernel_bundle_interface.cpp - Implements C API for +// sycl::kernel_bundle ---------------===// +// +// Data Parallel Control (dpctl) +// +// Copyright 2020-2021 Intel Corporation +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// +//===----------------------------------------------------------------------===// +/// +/// \file +/// This file implements the functions declared in +/// dpctl_sycl_kernel_bundle_interface.h. +/// +//===----------------------------------------------------------------------===// + +#include "dpctl_sycl_kernel_bundle_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 */ +#include +#include + +#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION +// Note: include ze_api.h before level_zero.hpp. Make sure clang-format does +// not reorder the includes. +// clang-format off +#include "ze_api.h" /* Level Zero headers */ +#include "sycl/ext/oneapi/backend/level_zero.hpp" +// clang-format on +#endif + +using namespace cl::sycl; + +namespace +{ +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 *clLoaderName = DPCTL_LIBCL_LOADER_FILENAME; +static const int clLibLoadFlags = RTLD_NOLOAD | RTLD_NOW | RTLD_LOCAL; +#elif defined(_WIN64) +static const char *clLoaderName = "OpenCL.dll"; +static const int clLibLoadFlags = 0; +#else +#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 +{ +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 **, + const size_t *, + cl_int *); +const char *clCreateProgramWithSource_Name = "clCreateProgramWithSource"; +clCreateProgramWithSourceFT get_clCreateProgramWithSource() +{ + static auto st_clCreateProgramWithSourceF = + cl_loader::get().getSymbol( + clCreateProgramWithSource_Name); + + return st_clCreateProgramWithSourceF; +} + +typedef cl_program (*clCreateProgramWithILFT)(cl_context, + const void *, + size_t, + cl_int *); +const char *clCreateProgramWithIL_Name = "clCreateProgramWithIL"; +clCreateProgramWithILFT get_clCreateProgramWithIL() +{ + static auto st_clCreateProgramWithILF = + cl_loader::get().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 auto st_clBuildProgramF = + cl_loader::get().getSymbol(clBuildProgram_Name); + + return st_clBuildProgramF; +} + +typedef cl_kernel (*clCreateKernelFT)(cl_program, const char *, cl_int *); +const char *clCreateKernel_Name = "clCreateKernel"; +clCreateKernelFT get_clCreateKernel() +{ + static auto st_clCreateKernelF = + cl_loader::get().getSymbol(clCreateKernel_Name); + + return st_clCreateKernelF; +} + +std::string _GetErrorCode_ocl_impl(cl_int 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); + } +} + +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 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) { + return nullptr; + } + cl_int build_status = + clBuildProgramF(clProgram, 1, &clDevice, CompileOpts, nullptr, nullptr); + + 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 + +#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; + +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 *, + ze_module_handle_t *, + ze_module_build_log_handle_t *); +const char *zeModuleCreate_Name = "zeModuleCreate"; +zeModuleCreateFT get_zeModuleCreate() +{ + static auto st_zeModuleCreateF = + ze_loader::get().getSymbol(zeModuleCreate_Name); + + return st_zeModuleCreateF; +} + +typedef ze_result_t (*zeModuleDestroyFT)(ze_module_handle_t); +const char *zeModuleDestroy_Name = "zeModuleDestroy"; +zeModuleDestroyFT get_zeModuleDestroy() +{ + static auto st_zeModuleDestroyF = + ze_loader::get().getSymbol(zeModuleDestroy_Name); + + 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 auto st_zeKernelCreateF = + ze_loader::get().getSymbol(zeKernelCreate_Name); + + return st_zeKernelCreateF; +} + +std::string _GetErrorCode_ze_impl(ze_result_t 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); + } +} + +__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; + + // Populate the Level Zero module descriptions + ze_module_desc_t ZeModuleDesc = {}; + ZeModuleDesc.stype = ZE_STRUCTURE_TYPE_MODULE_DESC; + ZeModuleDesc.format = ZE_MODULE_FORMAT_IL_SPIRV; + ZeModuleDesc.inputSize = il_length; + ZeModuleDesc.pInputModule = (uint8_t *)IL; + ZeModuleDesc.pBuildFlags = CompileOpts; + ZeModuleDesc.pConstants = &ZeSpecConstants; + + ze_module_handle_t ZeModule; + + 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); + + 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; + } +} + +__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; + } + + 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 DPCTLSyclKernelBundleRef +DPCTLKernelBundle_CreateFromSpirv(__dpctl_keep const DPCTLSyclContextRef CtxRef, + __dpctl_keep const DPCTLSyclDeviceRef DevRef, + __dpctl_keep const void *IL, + size_t length, + const char *CompileOpts) +{ + DPCTLSyclKernelBundleRef KBRef = nullptr; + if (!CtxRef) { + error_handler("Cannot create program from SPIR-V as the supplied SYCL " + "context is NULL.", + __FILE__, __func__, __LINE__); + return KBRef; + } + if (!DevRef) { + error_handler("Cannot create program from SPIR-V as the supplied SYCL " + "device is NULL.", + __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); + // get the backend type + auto BE = SyclCtx->get_platform().get_backend(); + switch (BE) { + case backend::opencl: + KBRef = _CreateKernelBundleWithIL_ocl_impl(*SyclCtx, *SyclDev, IL, + length, CompileOpts); + break; + case backend::ext_oneapi_level_zero: +#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION + 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 KBRef; +} + +__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) +{ + context *SyclCtx = 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); + SyclDev = unwrap(Dev); + + // get the backend type + auto BE = SyclCtx->get_platform().get_backend(); + switch (BE) { + case backend::opencl: + try { + return _CreateKernelBundleWithOCLSource_ocl_impl( + *SyclCtx, *SyclDev, Source, CompileOpts); + } catch (std::exception const &e) { + error_handler(e, __FILE__, __func__, __LINE__); + return nullptr; + } + break; + case backend::ext_oneapi_level_zero: + error_handler("CreateFromSource is not supported in Level Zero.", + __FILE__, __func__, __LINE__); + return nullptr; + default: + error_handler("CreateFromSource is not supported in unknown backend.", + __FILE__, __func__, __LINE__); + return nullptr; + } +} + +__dpctl_give DPCTLSyclKernelRef +DPCTLKernelBundle_GetKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, + __dpctl_keep const char *KernelName) +{ + if (!KBRef) { + error_handler("Input KBRef is nullptr", __FILE__, __func__, __LINE__); + return nullptr; + } + if (!KernelName) { + error_handler("Input KernelName is nullptr", __FILE__, __func__, + __LINE__); + return nullptr; + } + 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 DPCTLKernelBundle_HasKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, + __dpctl_keep const char *KernelName) +{ + if (!KBRef) { + error_handler("Input KBRef is nullptr", __FILE__, __func__, __LINE__); + return false; + } + if (!KernelName) { + error_handler("Input KernelName is nullptr", __FILE__, __func__, + __LINE__); + return false; + } + + 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 DPCTLKernelBundle_Delete(__dpctl_take DPCTLSyclKernelBundleRef KBRef) +{ + delete unwrap(KBRef); +} 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/source/dpctl_sycl_program_interface.cpp b/libsyclinterface/source/dpctl_sycl_program_interface.cpp deleted file mode 100644 index 6a62e948c0..0000000000 --- a/libsyclinterface/source/dpctl_sycl_program_interface.cpp +++ /dev/null @@ -1,346 +0,0 @@ -//===- dpctl_sycl_program_interface.cpp - Implements C API for sycl::program =// -// -// Data Parallel Control (dpctl) -// -// Copyright 2020-2021 Intel Corporation -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -//===----------------------------------------------------------------------===// -/// -/// \file -/// This file implements the functions declared in -/// dpctl_sycl_program_interface.h. -/// -//===----------------------------------------------------------------------===// - -#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_error_handlers.h" -#include /* OpenCL headers */ -#include /* Sycl headers */ -#include -#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 -#include "ze_api.h" /* Level Zero headers */ -#include "sycl/ext/oneapi/backend/level_zero.hpp" -// clang-format on -#endif - -using namespace cl::sycl; - -namespace -{ -#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION - -#ifdef __linux__ -static const char *zeLoaderName = DPCTL_LIBZE_LOADER_FILENAME; -static const int libLoadFlags = RTLD_NOLOAD | RTLD_NOW | RTLD_LOCAL; -#elif defined(_WIN64) -static const char *zeLoaderName = "ze_loader.dll"; -static const int libLoadFlags = 0; -#else -#error "Level Zero 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 *); - -const char *zeModuleCreateFuncName = "zeModuleCreate"; - -#endif // #ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION - -DEFINE_SIMPLE_CONVERSION_FUNCTIONS(context, DPCTLSyclContextRef) -DEFINE_SIMPLE_CONVERSION_FUNCTIONS(program, DPCTLSyclProgramRef) -DEFINE_SIMPLE_CONVERSION_FUNCTIONS(kernel, DPCTLSyclKernelRef) - -__dpctl_give DPCTLSyclProgramRef -createOpenCLInterOpProgram(const context &SyclCtx, - __dpctl_keep const void *IL, - size_t length, - const char *CompileOpts) -{ - 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__); - return nullptr; - } - auto SyclDevices = SyclCtx.get_devices(); - - // 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]); - - // Build the OpenCL interoperability program - err = clBuildProgram(CLProgram, (cl_uint)(SyclDevices.size()), CLDevices, - CompileOpts, nullptr, nullptr); - // free the CLDevices array - delete[] CLDevices; - - if (err) { - std::stringstream ss; - ss << "OpenCL program could not be built. OpenCL Error " << err << "."; - error_handler(ss.str(), __FILE__, __func__, __LINE__); - return 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__); - return nullptr; - } -} - -#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION - -zeModuleCreateFT getZeModuleCreateFn() -{ - static dpctl::DynamicLibHelper zeLib(zeLoaderName, libLoadFlags); - 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); - - return stZeModuleCreateF; -} - -__dpctl_give DPCTLSyclProgramRef -createLevelZeroInterOpProgram(const context &SyclCtx, - const void *IL, - size_t length, - const char *CompileOpts) -{ - 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.", - __FILE__, __func__, __LINE__); - return nullptr; - } - - // Specialization constants are not yet supported. - // Refer https://bit.ly/33UEDYN for details on specialization constants. - ze_module_constants_t ZeSpecConstants = {}; - ZeSpecConstants.numConstants = 0; - - // Populate the Level Zero module descriptions - ze_module_desc_t ZeModuleDesc = {}; - ZeModuleDesc.stype = ZE_STRUCTURE_TYPE_MODULE_DESC; - ZeModuleDesc.format = ZE_MODULE_FORMAT_IL_SPIRV; - ZeModuleDesc.inputSize = 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(); - - if (!stZeModuleCreateF) { - error_handler("ZeModuleCreateFn is invalid.", __FILE__, __func__, - __LINE__); - return nullptr; - } - - auto ret = - stZeModuleCreateF(ZeCtx, ZeDevice, &ZeModuleDesc, &ZeModule, nullptr); - if (ret != ZE_RESULT_SUCCESS) { - error_handler("ZeModule creation failed.", __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__); - return nullptr; - } -} -#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) -{ - DPCTLSyclProgramRef Pref = nullptr; - context *SyclCtx = nullptr; - if (!CtxRef) { - error_handler("Cannot create program from SPIR-V as the supplied SYCL " - "context is NULL.", - __FILE__, __func__, __LINE__); - return Pref; - } - SyclCtx = unwrap(CtxRef); - // get the backend type - auto BE = SyclCtx->get_platform().get_backend(); - switch (BE) { - case backend::opencl: - Pref = createOpenCLInterOpProgram(*SyclCtx, IL, length, CompileOpts); - break; - case backend::ext_oneapi_level_zero: -#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION - Pref = createLevelZeroInterOpProgram(*SyclCtx, IL, length, CompileOpts); -#endif - break; - default: - break; - } - return Pref; -} - -__dpctl_give DPCTLSyclProgramRef -DPCTLProgram_CreateFromOCLSource(__dpctl_keep const DPCTLSyclContextRef Ctx, - __dpctl_keep const char *Source, - __dpctl_keep const char *CompileOpts) -{ - std::string compileOpts; - context *SyclCtx = nullptr; - program *SyclProgram = nullptr; - - if (!Ctx) { - error_handler("Input Ctx 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; - } - - // 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); - } catch (std::exception const &e) { - delete SyclProgram; - error_handler(e, __FILE__, __func__, __LINE__); - return nullptr; - } - break; - 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) -{ - if (!PRef) { - error_handler("Input PRef 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__); - return nullptr; - } -} - -bool DPCTLProgram_HasKernel(__dpctl_keep DPCTLSyclProgramRef PRef, - __dpctl_keep const char *KernelName) -{ - if (!PRef) { - error_handler("Input PRef is nullptr", __FILE__, __func__, __LINE__); - return false; - } - if (!KernelName) { - error_handler("Input KernelName is nullptr", __FILE__, __func__, - __LINE__); - return false; - } - - auto SyclProgram = unwrap(PRef); - try { - return SyclProgram->has_kernel(KernelName); - } catch (std::exception const &e) { - error_handler(e, __FILE__, __func__, __LINE__); - return false; - } -} - -void DPCTLProgram_Delete(__dpctl_take DPCTLSyclProgramRef PRef) -{ - delete unwrap(PRef); -} 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_kernel_bundle_interface.cpp b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp new file mode 100644 index 0000000000..9bbd55f009 --- /dev/null +++ b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp @@ -0,0 +1,334 @@ +//===- test_sycl_kernel_bundle_interface.cpp - +// Test cases for module interface -===// +// +// Data Parallel Control (dpctl) +// +// Copyright 2020-2021 Intel Corporation +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// +//===----------------------------------------------------------------------===// +/// +/// \file +/// This file has unit test cases for functions defined in +/// dpctl_sycl_module_interface.h. +/// +//===----------------------------------------------------------------------===// + +#include "Config/dpctl_config.h" +#include "dpctl_sycl_context_interface.h" +#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_queue_interface.h" +#include "dpctl_sycl_queue_manager.h" +#include +#include +#include +#include +#include + +using namespace cl::sycl; + +struct TestDPCTLSyclKernelBundleInterface + : public ::testing::TestWithParam +{ + DPCTLSyclDeviceRef DRef = nullptr; + DPCTLSyclContextRef CRef = nullptr; + DPCTLSyclKernelBundleRef KBRef = nullptr; + std::ifstream spirvFile; + size_t spirvFileSize; + std::vector spirvBuffer; + + TestDPCTLSyclKernelBundleInterface() + { + auto DS = DPCTLFilterSelector_Create(GetParam()); + DRef = DPCTLDevice_CreateFromSelector(DS); + DPCTLDeviceSelector_Delete(DS); + CRef = DPCTLDeviceMgr_GetCachedContext(DRef); + + if (DRef) { + 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); + KBRef = DPCTLKernelBundle_CreateFromSpirv( + CRef, DRef, spirvBuffer.data(), spirvFileSize, nullptr); + } + } + + void SetUp() + { + if (!DRef) { + auto message = "Skipping as no device of type " + + std::string(GetParam()) + "."; + GTEST_SKIP_(message.c_str()); + } + } + + ~TestDPCTLSyclKernelBundleInterface() + { + if (DRef) { + spirvFile.close(); + DPCTLDevice_Delete(DRef); + } + if (CRef) + DPCTLContext_Delete(CRef); + if (KBRef) + DPCTLKernelBundle_Delete(KBRef); + } +}; + +TEST_P(TestDPCTLSyclKernelBundleInterface, ChkCreateFromSpirv) +{ + + ASSERT_TRUE(KBRef != nullptr); + ASSERT_TRUE(DPCTLKernelBundle_HasKernel(KBRef, "add")); + ASSERT_TRUE(DPCTLKernelBundle_HasKernel(KBRef, "axpy")); + ASSERT_FALSE(DPCTLKernelBundle_HasKernel(KBRef, nullptr)); +} + +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(TestDPCTLSyclKernelBundleInterface, ChkHasKernelNullProgram) +{ + + DPCTLSyclKernelBundleRef NullRef = nullptr; + ASSERT_FALSE(DPCTLKernelBundle_HasKernel(NullRef, "add")); +} + +TEST_P(TestDPCTLSyclKernelBundleInterface, ChkGetKernel) +{ + 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); + ASSERT_TRUE(NullKernel == nullptr); + DPCTLKernel_Delete(AddKernel); + DPCTLKernel_Delete(AxpyKernel); + EXPECT_NO_FATAL_FAILURE(DPCTLKernel_Delete(NullKernel)); +} + +TEST_P(TestDPCTLSyclKernelBundleInterface, ChkGetKernelNullProgram) +{ + DPCTLSyclKernelBundleRef NullRef = nullptr; + DPCTLSyclKernelRef KRef = nullptr; + + EXPECT_NO_FATAL_FAILURE(KRef = DPCTLKernelBundle_GetKernel(NullRef, "add")); + EXPECT_TRUE(KRef == nullptr); +} + +struct TestOCLKernelBundleFromSource + : public ::testing::TestWithParam +{ + const char *CLProgramStr = R"CLC( + kernel void add(global int* a, global int* b, global int* c) { + size_t index = get_global_id(0); + c[index] = a[index] + b[index]; + } + + kernel void axpy(global int* a, global int* b, global int* c, int d) + { + size_t index = get_global_id(0); + c[index] = a[index] + d*b[index]; + } + )CLC"; + const char *CompileOpts = "-cl-fast-relaxed-math"; + DPCTLSyclDeviceRef DRef = nullptr; + DPCTLSyclContextRef CRef = nullptr; + DPCTLSyclKernelBundleRef KBRef = nullptr; + + TestOCLKernelBundleFromSource() + { + auto DS = DPCTLFilterSelector_Create(GetParam()); + DRef = DPCTLDevice_CreateFromSelector(DS); + DPCTLDeviceSelector_Delete(DS); + CRef = DPCTLDeviceMgr_GetCachedContext(DRef); + + if (DRef) + KBRef = DPCTLKernelBundle_CreateFromOCLSource( + 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) + DPCTLDevice_Delete(DRef); + if (CRef) + DPCTLContext_Delete(CRef); + if (KBRef) + DPCTLKernelBundle_Delete(KBRef); + } +}; + +TEST_P(TestOCLKernelBundleFromSource, CheckCreateFromOCLSource) +{ + ASSERT_TRUE(KBRef != nullptr); + ASSERT_TRUE(DPCTLKernelBundle_HasKernel(KBRef, "add")); + ASSERT_TRUE(DPCTLKernelBundle_HasKernel(KBRef, "axpy")); +} + +TEST_P(TestOCLKernelBundleFromSource, CheckCreateFromOCLSourceNull) +{ + const char *InvalidCLProgramStr = R"CLC( + kernel void invalid(global foo* a, global bar* b) { + size_t index = get_global_id(0); + b[index] = a[index]; + } + )CLC"; + DPCTLSyclKernelBundleRef KBRef = nullptr; + + 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_P(TestOCLKernelBundleFromSource, CheckGetKernelOCLSource) +{ + auto AddKernel = DPCTLKernelBundle_GetKernel(KBRef, "add"); + auto AxpyKernel = DPCTLKernelBundle_GetKernel(KBRef, "axpy"); + ASSERT_TRUE(AddKernel != nullptr); + ASSERT_TRUE(AxpyKernel != nullptr); + DPCTLKernel_Delete(AddKernel); + DPCTLKernel_Delete(AxpyKernel); +} + +INSTANTIATE_TEST_SUITE_P(KernelBundleCreationFromSpirv, + TestDPCTLSyclKernelBundleInterface, + ::testing::Values("opencl", + "opencl:gpu", + "opencl:cpu", + "opencl:gpu:0", +#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION + "level_zero", + "level_zero:gpu", +#endif + "opencl:cpu:0")); + +INSTANTIATE_TEST_SUITE_P(KernelBundleCreationFromSource, + TestOCLKernelBundleFromSource, + ::testing::Values("opencl:gpu", "opencl:cpu")); + +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); +} diff --git a/libsyclinterface/tests/test_sycl_kernel_interface.cpp b/libsyclinterface/tests/test_sycl_kernel_interface.cpp index 3a5f5b5f0f..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" @@ -81,46 +81,21 @@ 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); 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); } @@ -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, diff --git a/libsyclinterface/tests/test_sycl_program_interface.cpp b/libsyclinterface/tests/test_sycl_program_interface.cpp deleted file mode 100644 index 9824e09088..0000000000 --- a/libsyclinterface/tests/test_sycl_program_interface.cpp +++ /dev/null @@ -1,237 +0,0 @@ -//===-- test_sycl_program_interface.cpp - Test cases for module interface -===// -// -// Data Parallel Control (dpctl) -// -// Copyright 2020-2021 Intel Corporation -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -//===----------------------------------------------------------------------===// -/// -/// \file -/// This file has unit test cases for functions defined in -/// dpctl_sycl_module_interface.h. -/// -//===----------------------------------------------------------------------===// - -#include "Config/dpctl_config.h" -#include "dpctl_sycl_context_interface.h" -#include "dpctl_sycl_device_interface.h" -#include "dpctl_sycl_device_manager.h" -#include "dpctl_sycl_device_selector_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 -#include -#include -#include -#include - -using namespace cl::sycl; - -struct TestDPCTLSyclProgramInterface - : public ::testing::TestWithParam -{ - DPCTLSyclDeviceRef DRef = nullptr; - DPCTLSyclContextRef CRef = nullptr; - DPCTLSyclQueueRef QRef = nullptr; - DPCTLSyclProgramRef PRef = nullptr; - std::ifstream spirvFile; - size_t spirvFileSize; - std::vector spirvBuffer; - - TestDPCTLSyclProgramInterface() - { - 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", - 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); - PRef = DPCTLProgram_CreateFromSpirv(CRef, spirvBuffer.data(), - spirvFileSize, nullptr); - } - } - - void SetUp() - { - if (!DRef) { - auto message = "Skipping as no device of type " + - std::string(GetParam()) + "."; - GTEST_SKIP_(message.c_str()); - } - } - - ~TestDPCTLSyclProgramInterface() - { - if (DRef) - spirvFile.close(); - DPCTLDevice_Delete(DRef); - DPCTLQueue_Delete(QRef); - DPCTLContext_Delete(CRef); - DPCTLProgram_Delete(PRef); - } -}; - -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)); -} - -TEST_P(TestDPCTLSyclProgramInterface, ChkCreateFromSpirvNull) -{ - DPCTLSyclContextRef Null_CRef = 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); -} - -TEST_P(TestDPCTLSyclProgramInterface, ChkHasKernelNullProgram) -{ - - DPCTLSyclProgramRef NullRef = nullptr; - ASSERT_FALSE(DPCTLProgram_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); - - ASSERT_TRUE(AddKernel != nullptr); - ASSERT_TRUE(AxpyKernel != nullptr); - ASSERT_TRUE(NullKernel == nullptr); - DPCTLKernel_Delete(AddKernel); - DPCTLKernel_Delete(AxpyKernel); - EXPECT_NO_FATAL_FAILURE(DPCTLKernel_Delete(NullKernel)); -} - -TEST_P(TestDPCTLSyclProgramInterface, ChkGetKernelNullProgram) -{ - DPCTLSyclProgramRef NullRef = nullptr; - DPCTLSyclKernelRef KRef = nullptr; - - EXPECT_NO_FATAL_FAILURE(KRef = DPCTLProgram_GetKernel(NullRef, "add")); - EXPECT_TRUE(KRef == nullptr); -} - -struct TestOCLProgramFromSource : public ::testing::Test -{ - const char *CLProgramStr = R"CLC( - kernel void add(global int* a, global int* b, global int* c) { - size_t index = get_global_id(0); - c[index] = a[index] + b[index]; - } - - kernel void axpy(global int* a, global int* b, global int* c, int d) - { - size_t index = get_global_id(0); - c[index] = a[index] + d*b[index]; - } - )CLC"; - const char *CompileOpts = "-cl-fast-relaxed-math"; - DPCTLSyclDeviceRef DRef = nullptr; - DPCTLSyclContextRef CRef = nullptr; - DPCTLSyclQueueRef QRef = nullptr; - DPCTLSyclProgramRef PRef = nullptr; - - TestOCLProgramFromSource() - { - 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) - PRef = DPCTLProgram_CreateFromOCLSource(CRef, CLProgramStr, - CompileOpts); - } - - ~TestOCLProgramFromSource() - { - DPCTLDevice_Delete(DRef); - DPCTLQueue_Delete(QRef); - DPCTLContext_Delete(CRef); - DPCTLProgram_Delete(PRef); - } -}; - -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")); -} - -TEST_F(TestOCLProgramFromSource, CheckCreateFromOCLSourceNull) -{ - const char *InvalidCLProgramStr = R"CLC( - kernel void invalid(global foo* a, global bar* b) { - size_t index = get_global_id(0); - b[index] = a[index]; - } - )CLC"; - DPCTLSyclProgramRef PRef = 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); -} - -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"); - ASSERT_TRUE(AddKernel != nullptr); - ASSERT_TRUE(AxpyKernel != nullptr); - DPCTLKernel_Delete(AddKernel); - DPCTLKernel_Delete(AxpyKernel); -} - -INSTANTIATE_TEST_SUITE_P(ProgramCreationFromSpriv, - TestDPCTLSyclProgramInterface, - ::testing::Values("opencl", - "opencl:gpu", - "opencl:cpu", - "opencl:gpu:0", -#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION - "level_zero", - "level_zero:gpu", -#endif - "opencl:cpu:0")); 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); diff --git a/libsyclinterface/tests/test_sycl_queue_submit.cpp b/libsyclinterface/tests/test_sycl_queue_submit.cpp index b0892ccb9c..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 @@ -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); }