diff --git a/dpctl-capi/dbg_build.sh b/dpctl-capi/dbg_build.sh index 60c8674a42..5b42f82273 100755 --- a/dpctl-capi/dbg_build.sh +++ b/dpctl-capi/dbg_build.sh @@ -9,15 +9,16 @@ rm -rf ${INSTALL_PREFIX} export ONEAPI_ROOT=/opt/intel/oneapi DPCPP_ROOT=${ONEAPI_ROOT}/compiler/latest/linux -cmake \ - -DCMAKE_BUILD_TYPE=Debug \ - -DCMAKE_INSTALL_PREFIX=${INSTALL_PREFIX} \ - -DCMAKE_PREFIX_PATH=${INSTALL_PREFIX} \ - -DDPCPP_INSTALL_DIR=${DPCPP_ROOT} \ - -DCMAKE_C_COMPILER:PATH=${DPCPP_ROOT}/bin/clang \ - -DCMAKE_CXX_COMPILER:PATH=${DPCPP_ROOT}/bin/dpcpp \ - -DDPCTL_BUILD_CAPI_TESTS=ON \ - -DDPCTL_GENERATE_COVERAGE=ON \ +cmake \ + -DCMAKE_BUILD_TYPE=Debug \ + -DCMAKE_INSTALL_PREFIX=${INSTALL_PREFIX} \ + -DCMAKE_PREFIX_PATH=${INSTALL_PREFIX} \ + -DDPCPP_INSTALL_DIR=${DPCPP_ROOT} \ + -DCMAKE_C_COMPILER:PATH=${DPCPP_ROOT}/bin/clang \ + -DCMAKE_CXX_COMPILER:PATH=${DPCPP_ROOT}/bin/dpcpp \ + -DDPCTL_ENABLE_LO_PROGRAM_CREATION=ON \ + -DDPCTL_BUILD_CAPI_TESTS=ON \ + -DDPCTL_GENERATE_COVERAGE=ON \ .. make V=1 -n -j 4 && make check && make install diff --git a/dpctl-capi/include/dpctl_sycl_program_interface.h b/dpctl-capi/include/dpctl_sycl_program_interface.h index c931b4b166..8e25d1df40 100644 --- a/dpctl-capi/include/dpctl_sycl_program_interface.h +++ b/dpctl-capi/include/dpctl_sycl_program_interface.h @@ -52,14 +52,17 @@ DPCTL_C_EXTERN_C_BEGIN * @param Ctx An opaque pointer to a sycl::context * @param IL SPIR-V binary * @param Length The size of the IL binary in bytes. + * @param CompileOpts Optional compiler flags used when compiling th + * SPIR-V binary. * @return A new SyclProgramRef pointer if the program creation succeeded, * else returns NULL. */ DPCTL_API __dpctl_give DPCTLSyclProgramRef -DPCTLProgram_CreateFromOCLSpirv (__dpctl_keep const DPCTLSyclContextRef Ctx, - __dpctl_keep const void *IL, - size_t Length); +DPCTLProgram_CreateFromSpirv (__dpctl_keep const DPCTLSyclContextRef Ctx, + __dpctl_keep const void *IL, + size_t Length, + const char *CompileOpts); /*! * @brief Create a Sycl program from an OpenCL kernel source string. diff --git a/dpctl-capi/source/dpctl_sycl_program_interface.cpp b/dpctl-capi/source/dpctl_sycl_program_interface.cpp index 7f8b75cf63..ae735abbe2 100644 --- a/dpctl-capi/source/dpctl_sycl_program_interface.cpp +++ b/dpctl-capi/source/dpctl_sycl_program_interface.cpp @@ -25,10 +25,15 @@ //===----------------------------------------------------------------------===// #include "dpctl_sycl_program_interface.h" +#include "Config/dpctl_config.h" #include "Support/CBindingWrapping.h" -#include /* Sycl headers */ -#include /* OpenCL headers */ +#include /* Sycl headers */ +#include /* OpenCL headers */ +#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION +#include /* Level Zero headers */ +#include +#endif using namespace cl::sycl; @@ -41,7 +46,8 @@ DEFINE_SIMPLE_CONVERSION_FUNCTIONS(kernel, DPCTLSyclKernelRef) __dpctl_give DPCTLSyclProgramRef createOpenCLInterOpProgram (const context &SyclCtx, __dpctl_keep const void *IL, - size_t length) + size_t length, + const char * /* */) { cl_int err; auto CLCtx = SyclCtx.get(); @@ -83,12 +89,63 @@ createOpenCLInterOpProgram (const context &SyclCtx, } } +#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION +__dpctl_give DPCTLSyclProgramRef +createLevelZeroInterOpProgram (const context &SyclCtx, + const void *IL, + size_t length, + const char *CompileOpts) +{ + auto ZeCtx = SyclCtx.get_native(); + auto SyclDevices = SyclCtx.get_devices(); + if(SyclDevices.size() > 1) { + // We only support build to one device with Level Zero now. + // TODO: log error + 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.format = ZE_MODULE_FORMAT_IL_SPIRV; + ZeModuleDesc.inputSize = length; + ZeModuleDesc.pInputModule = (uint8_t*)IL; + ZeModuleDesc.pBuildFlags = CompileOpts; + ZeModuleDesc.pConstants = &ZeSpecConstants; + + auto ZeDevice = SyclDevices[0].get_native(); + ze_module_handle_t ZeModule; + auto ret = zeModuleCreate(ZeCtx, ZeDevice, &ZeModuleDesc, &ZeModule, + nullptr); + if(ret != ZE_RESULT_SUCCESS) { + // TODO: handle error + return nullptr; + } + + // Create the Sycl program from the ZeModule + try { + auto ZeProgram = new program(sycl::level_zero::make_program( + SyclCtx, reinterpret_cast(ZeModule)) + ); + return wrap(ZeProgram); + } catch (invalid_object_error &e) { + // \todo record error + std::cerr << e.what() << '\n'; + return nullptr; + } +} +#endif } /* end of anonymous namespace */ __dpctl_give DPCTLSyclProgramRef -DPCTLProgram_CreateFromOCLSpirv (__dpctl_keep const DPCTLSyclContextRef CtxRef, - __dpctl_keep const void *IL, - size_t length) +DPCTLProgram_CreateFromSpirv (__dpctl_keep const DPCTLSyclContextRef CtxRef, + __dpctl_keep const void *IL, + size_t length, + const char *CompileOpts) { DPCTLSyclProgramRef Pref = nullptr; context *SyclCtx = nullptr; @@ -96,21 +153,23 @@ DPCTLProgram_CreateFromOCLSpirv (__dpctl_keep const DPCTLSyclContextRef CtxRef, // \todo handle error 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); + Pref = createOpenCLInterOpProgram(*SyclCtx, IL, length, CompileOpts); break; case backend::level_zero: +#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION + Pref = createLevelZeroInterOpProgram(*SyclCtx, IL, length, + CompileOpts); +#endif break; default: break; } - return Pref; } diff --git a/dpctl-capi/tests/test_sycl_program_interface.cpp b/dpctl-capi/tests/test_sycl_program_interface.cpp index 0e0061b674..791907c1c7 100644 --- a/dpctl-capi/tests/test_sycl_program_interface.cpp +++ b/dpctl-capi/tests/test_sycl_program_interface.cpp @@ -1,4 +1,4 @@ -//===---------- test_sycl_program_interface.cpp - dpctl-C_API --*-- C++ -*-===// +//===---------- test_sycl_program_interface.cpp - dpctl-C_API ---*- C++ -*-===// // // Data Parallel Control Library (dpCtl) // @@ -29,7 +29,7 @@ #include "dpctl_sycl_program_interface.h" #include "dpctl_sycl_queue_interface.h" #include "dpctl_sycl_queue_manager.h" - +#include "Config/dpctl_config.h" #include #include #include @@ -127,12 +127,16 @@ struct TestDPCTLSyclProgramInterface : public ::testing::Test size_t spirvFileSize = 0; std::vector spirvBuffer; size_t nOpenCLGpuQ = 0; +#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION + size_t nL0GpuQ = 0; +#endif TestDPCTLSyclProgramInterface () : spirvFile{"./multi_kernel.spv", std::ios::binary | std::ios::ate}, spirvFileSize(std::filesystem::file_size("./multi_kernel.spv")), spirvBuffer(spirvFileSize), - nOpenCLGpuQ(DPCTLQueueMgr_GetNumQueues(DPCTL_OPENCL, DPCTL_GPU)) + nOpenCLGpuQ(DPCTLQueueMgr_GetNumQueues(DPCTL_OPENCL, DPCTL_GPU)), + nL0GpuQ(DPCTLQueueMgr_GetNumQueues(DPCTL_LEVEL_ZERO, DPCTL_GPU)) { spirvFile.seekg(0, std::ios::beg); spirvFile.read(spirvBuffer.data(), spirvFileSize); @@ -152,7 +156,7 @@ TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromOCLSource) auto QueueRef = DPCTLQueueMgr_GetQueue(DPCTL_OPENCL, DPCTL_GPU, 0); auto CtxRef = DPCTLQueue_GetContext(QueueRef); auto PRef = DPCTLProgram_CreateFromOCLSource(CtxRef, CLProgramStr, - CompileOpts); + CompileOpts); ASSERT_TRUE(PRef != nullptr); ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "add")); ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "axpy")); @@ -162,15 +166,36 @@ TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromOCLSource) DPCTLProgram_Delete(PRef); } -TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromOCLSpirv) +TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromSpirvOCL) { if(!nOpenCLGpuQ) GTEST_SKIP_("Skipping as no OpenCL GPU device found.\n"); auto QueueRef = DPCTLQueueMgr_GetQueue(DPCTL_OPENCL, DPCTL_GPU, 0); auto CtxRef = DPCTLQueue_GetContext(QueueRef); - auto PRef = DPCTLProgram_CreateFromOCLSpirv(CtxRef, spirvBuffer.data(), - spirvFileSize); + auto PRef = DPCTLProgram_CreateFromSpirv(CtxRef, spirvBuffer.data(), + spirvFileSize, + nullptr); + ASSERT_TRUE(PRef != nullptr); + ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "add")); + ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "axpy")); + + DPCTLQueue_Delete(QueueRef); + DPCTLContext_Delete(CtxRef); + DPCTLProgram_Delete(PRef); +} + +#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION +TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromSpirvL0) +{ + if(!nL0GpuQ) + GTEST_SKIP_("Skipping as no OpenCL GPU device found.\n"); + + auto QueueRef = DPCTLQueueMgr_GetQueue(DPCTL_LEVEL_ZERO, DPCTL_GPU, 0); + auto CtxRef = DPCTLQueue_GetContext(QueueRef); + auto PRef = DPCTLProgram_CreateFromSpirv(CtxRef, spirvBuffer.data(), + spirvFileSize, + nullptr); ASSERT_TRUE(PRef != nullptr); ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "add")); ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "axpy")); @@ -179,6 +204,7 @@ TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromOCLSpirv) DPCTLContext_Delete(CtxRef); DPCTLProgram_Delete(PRef); } +#endif TEST_F (TestDPCTLSyclProgramInterface, CheckGetKernelOCLSource) { @@ -207,15 +233,15 @@ TEST_F (TestDPCTLSyclProgramInterface, CheckGetKernelOCLSource) DPCTLProgram_Delete(PRef); } -TEST_F (TestDPCTLSyclProgramInterface, CheckGetKernelOCLSpirv) +TEST_F (TestDPCTLSyclProgramInterface, CheckGetKernelSpirv) { if(!nOpenCLGpuQ) GTEST_SKIP_("Skipping as no OpenCL GPU device found.\n"); auto QueueRef = DPCTLQueueMgr_GetQueue(DPCTL_OPENCL, DPCTL_GPU, 0); auto CtxRef = DPCTLQueue_GetContext(QueueRef); - auto PRef = DPCTLProgram_CreateFromOCLSpirv(CtxRef, spirvBuffer.data(), - spirvFileSize); + auto PRef = DPCTLProgram_CreateFromSpirv(CtxRef, spirvBuffer.data(), + spirvFileSize, nullptr); auto AddKernel = DPCTLProgram_GetKernel(PRef, "add"); auto AxpyKernel = DPCTLProgram_GetKernel(PRef, "axpy"); diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index ff6d0c536f..9e9dbca948 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -140,14 +140,15 @@ cdef extern from "dpctl_sycl_context_interface.h": cdef extern from "dpctl_sycl_program_interface.h": - cdef DPCTLSyclProgramRef DPCTLProgram_CreateFromOCLSpirv ( + cdef DPCTLSyclProgramRef DPCTLProgram_CreateFromSpirv ( const DPCTLSyclContextRef Ctx, const void *IL, - size_t Length) + size_t Length, + const char *CompileOpts) cdef DPCTLSyclProgramRef DPCTLProgram_CreateFromOCLSource ( const DPCTLSyclContextRef Ctx, - const char* Source, - const char* CompileOpts) + const char *Source, + const char *CompileOpts) cdef DPCTLSyclKernelRef DPCTLProgram_GetKernel (DPCTLSyclProgramRef PRef, const char *KernelName) cdef bool DPCTLProgram_HasKernel (DPCTLSyclProgramRef PRef, diff --git a/dpctl/_sycl_core.pyx b/dpctl/_sycl_core.pyx index 25d01330b2..835df44db1 100644 --- a/dpctl/_sycl_core.pyx +++ b/dpctl/_sycl_core.pyx @@ -678,7 +678,7 @@ cdef class _SyclRTManager: raise UnsupportedDeviceError("Device can only be gpu or cpu") except KeyError: raise UnsupportedBackendError("Backend can only be opencl or " - "level-0") + "level0") def _remove_current_queue(self): DPCTLQueueMgr_PopQueue() @@ -970,7 +970,7 @@ cdef class _SyclRTManager: raise UnsupportedDeviceError("Device can only be gpu or cpu") except KeyError: raise UnsupportedBackendError("Backend can only be opencl or " - "level-0") + "level0") # This private instance of the _SyclQueueManager should not be directly diff --git a/dpctl/program/_program.pxd b/dpctl/program/_program.pxd index 27f0716791..1190d5a12c 100644 --- a/dpctl/program/_program.pxd +++ b/dpctl/program/_program.pxd @@ -60,4 +60,5 @@ cdef class SyclProgram: cpdef create_program_from_source (SyclQueue q, unicode source, unicode copts=*) -cpdef create_program_from_spirv (SyclQueue q, const unsigned char[:] IL) +cpdef create_program_from_spirv (SyclQueue q, const unsigned char[:] IL, + unicode copts=*) diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index efa47f7649..2ea8ef6ea8 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -162,7 +162,8 @@ cpdef create_program_from_source(SyclQueue q, unicode src, unicode copts=""): cimport cython.array -cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL): +cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL, + unicode copts=""): """ Creates a Sycl interoperability program from an SPIR-V binary. @@ -173,6 +174,8 @@ cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL): q (SyclQueue): The :class:`SyclQueue` for which the :class:`SyclProgram` is going to be built. IL (const char[:]) : SPIR-V binary IL file for an OpenCL program. + copts (unicode) : Optional compilation flags that will be used + when compiling the program. Returns: program (SyclProgram): A :class:`SyclProgram` object wrapping the sycl::program returned by the C API. @@ -185,7 +188,9 @@ cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL): cdef const unsigned char *dIL = &IL[0] cdef DPCTLSyclContextRef CRef = q.get_sycl_context().get_context_ref() cdef size_t length = IL.shape[0] - Pref = DPCTLProgram_CreateFromOCLSpirv(CRef, dIL, length) + cdef bytes bCOpts = copts.encode('utf8') + cdef const char *COpts = bCOpts + Pref = DPCTLProgram_CreateFromSpirv(CRef, dIL, length, COpts) if Pref is NULL: raise SyclProgramCompilationError() diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index 5eb1c2b66b..ce7f6051f2 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -87,9 +87,16 @@ def test_create_program_from_spirv(self): "No Level0 GPU queues available", ) class TestProgramForLevel0GPU(unittest.TestCase): - @unittest.expectedFailure - def test_create_program_from_spirv(self): + import sys + + # Level zero program creation from a SPIR-V binary is not supported + # on Windows. + @unittest.skipIf( + sys.platform in ["win32", "cygwin"], + "Level Zero module creation unsupported on Windows.", + ) + def test_create_program_from_spirv(self): CURR_DIR = os.path.dirname(os.path.abspath(__file__)) spirv_file = os.path.join(CURR_DIR, "input_files/multi_kernel.spv") with open(spirv_file, "rb") as fin: diff --git a/scripts/build_backend.py b/scripts/build_backend.py index 2708094dd6..c68dd37b40 100644 --- a/scripts/build_backend.py +++ b/scripts/build_backend.py @@ -43,6 +43,7 @@ "-DDPCPP_INSTALL_DIR=" + DPCPP_ROOT, "-DCMAKE_C_COMPILER:PATH=" + os.path.join(DPCPP_ROOT, "bin", "clang"), "-DCMAKE_CXX_COMPILER:PATH=" + os.path.join(DPCPP_ROOT, "bin", "clang++"), + "-DDPCTL_ENABLE_LO_PROGRAM_CREATION=ON", backends, ] subprocess.check_call(cmake_args, stderr=subprocess.STDOUT, shell=False)