Level zero codegen

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

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.

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 <CL/sycl.hpp> /* Sycl headers */
-#include <CL/cl.h>     /* OpenCL headers */
+#include <CL/sycl.hpp>          /* Sycl headers       */
+#include <CL/cl.h>              /* OpenCL headers     */
+#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION
+#include <level_zero/zet_api.h> /* Level Zero headers */
+#include <CL/sycl/backend/level_zero.hpp>
+#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,34 +89,87 @@ 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<backend::level_zero>();
+    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<backend::level_zero>();
+    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<uintptr_t>(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;
     if(!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;
 }
 

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 <array>
 #include <fstream>
 #include <filesystem>
@@ -127,12 +127,16 @@ struct TestDPCTLSyclProgramInterface : public ::testing::Test
     size_t spirvFileSize = 0;
     std::vector<char> 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");
 

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,

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

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=*)

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, <const void*>dIL, length)
+    cdef bytes bCOpts = copts.encode('utf8')
+    cdef const char *COpts = <const char*>bCOpts
+    Pref = DPCTLProgram_CreateFromSpirv(CRef, <const void*>dIL, length, COpts)
     if Pref is NULL:
         raise SyclProgramCompilationError()
 

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:

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)