Added dynamic linking in ispcrt

ispcrt/detail/Device.h

@@ -24,6 +24,8 @@ struct Device : public RefCounted {
 
     virtual Module *newModule(const char *moduleFile, const ISPCRTModuleOptions &opts) const = 0;
 
+    virtual void linkModules(Module **modules, uint32_t numModules) const = 0;
+
     virtual Kernel *newKernel(const Module &module, const char *name) const = 0;
 
     virtual void *platformNativeHandle() const = 0;

ispcrt/detail/cpu/CPUDevice.cpp

@@ -221,6 +221,8 @@ ispcrt::base::Module *CPUDevice::newModule(const char *moduleFile, const ISPCRTM
     return new cpu::Module(moduleFile);
 }
 
+void CPUDevice::linkModules(base::Module **modules, const uint32_t numModules) const {}
+
 ispcrt::base::Kernel *CPUDevice::newKernel(const ispcrt::base::Module &module, const char *name) const {
     return new cpu::Kernel(module, name);
 }

ispcrt/detail/cpu/CPUDevice.h

@@ -24,6 +24,8 @@ struct CPUDevice : public base::Device {
 
     base::Module *newModule(const char *moduleFile, const ISPCRTModuleOptions &moduleOpts) const override;
 
+    void linkModules(base::Module **modules, const uint32_t numModules) const override;
+
     base::Kernel *newKernel(const base::Module &module, const char *name) const override;
 
     void *platformNativeHandle() const override;

ispcrt/detail/gpu/GPUDevice.cpp

@@ -560,12 +560,21 @@ struct Module : public ispcrt::base::Module {
         // + or = sign. '+' means that the content of the variable should
         // be added to the default igc options, while '=' will replace
         // the options with the content of the env var.
-        std::string igcOptions = "-vc-codegen -no-optimize -Xfinalizer '-presched'";
+        std::string igcOptions;
+        // If scalar module is passed to ISPC Runtime, do not use VC backend
+        // options on it
+        if (opts.moduleType != ISPCRTModuleType::ISPCRT_SCALAR_MODULE) {
+            igcOptions += "-vc-codegen -no-optimize -Xfinalizer '-presched'";
+        }
         // If stackSize has default value 0, do not set -stateless-stack-mem-size,
         // it will be set to 8192 in VC backend by default.
         if (opts.stackSize > 0) {
             igcOptions += " -stateless-stack-mem-size=" + std::to_string(opts.stackSize);
         }
+        // If module is a library for the kernel, add " -library-compilation"
+        if (opts.libraryCompilation) {
+            igcOptions += " -library-compilation";
+        }
         constexpr auto MAX_ISPCRT_IGC_OPTIONS = 2000UL;
 #if defined(_WIN32) || defined(_WIN64)
         char *userIgcOptionsEnv = nullptr;
@@ -1027,6 +1036,30 @@ ISPCRTDeviceInfo deviceInfo(uint32_t deviceIdx) {
     return info;
 }
 
+void linkModules(gpu::Module **modules, const uint32_t numModules) {
+    std::vector<ze_module_handle_t> moduleHandles;
+    for (int i = 0; i< numModules; i++) {
+        moduleHandles.push_back(modules[i]->handle());
+    }
+
+    ze_module_build_log_handle_t linkLog;
+    L0_SAFE_CALL(zeModuleDynamicLink(numModules, moduleHandles.data(), &linkLog));
+
+    size_t buildLogSize;
+    L0_SAFE_CALL(zeModuleBuildLogGetString(linkLog, &buildLogSize, nullptr));
+    char *dynLogBuffer = new char[buildLogSize]();
+    L0_SAFE_CALL(zeModuleBuildLogGetString(linkLog, &buildLogSize, dynLogBuffer));
+
+    // For now always print dynamic linking log.
+    // TODO: introduce verbose mode to ISPCRT
+    std::cout << dynLogBuffer << "\n";
+    delete[] dynLogBuffer;
+    if (linkLog)
+        L0_SAFE_CALL_NOEXCEPT(zeModuleBuildLogDestroy(linkLog));
+
+}
+
+
 } // namespace gpu
 
 // Use the first available device by default for now.
@@ -1082,6 +1115,10 @@ base::Module *GPUDevice::newModule(const char *moduleFile, const ISPCRTModuleOpt
     return new gpu::Module((ze_device_handle_t)m_device, (ze_context_handle_t)m_context, moduleFile, m_is_mock, opts);
 }
 
+void GPUDevice::linkModules(base::Module **modules, const uint32_t numModules) const {
+    gpu::linkModules((gpu::Module **)modules, numModules);
+}
+
 base::Kernel *GPUDevice::newKernel(const base::Module &module, const char *name) const {
     return new gpu::Kernel(module, name);
 }

ispcrt/detail/gpu/GPUDevice.h

@@ -30,6 +30,8 @@ struct GPUDevice : public base::Device {
 
     base::Module *newModule(const char *moduleFile, const ISPCRTModuleOptions &opts) const override;
 
+    void linkModules(base::Module **modules, const uint32_t numModules) const override;
+
     base::Kernel *newKernel(const base::Module &module, const char *name) const override;
 
     void *platformNativeHandle() const override;

ispcrt/ispcrt.cpp

@@ -256,16 +256,24 @@ void *ispcrtSharedPtr(ISPCRTMemoryView h) ISPCRT_CATCH_BEGIN {
 ISPCRT_CATCH_END(nullptr)
 
 ///////////////////////////////////////////////////////////////////////////////
-// Kernels ////////////////////////////////////////////////////////////////////
+// Modules ////////////////////////////////////////////////////////////////////
 ///////////////////////////////////////////////////////////////////////////////
 
 ISPCRTModule ispcrtLoadModule(ISPCRTDevice d, const char *moduleFile,
                               ISPCRTModuleOptions moduleOpts) ISPCRT_CATCH_BEGIN {
     const auto &device = referenceFromHandle<ispcrt::base::Device>(d);
-    return (ISPCRTModule)device.newModule(moduleFile, moduleOpts);
+    auto module = (ISPCRTModule)device.newModule(moduleFile, moduleOpts);
+    return module;
+
 }
 ISPCRT_CATCH_END(nullptr)
 
+void ispcrtLinkModules(ISPCRTDevice d, ISPCRTModule *modules, const uint32_t numModules) ISPCRT_CATCH_BEGIN {
+    const auto &device = referenceFromHandle<ispcrt::base::Device>(d);
+    device.linkModules((ispcrt::base::Module **)modules, numModules);
+}
+ISPCRT_CATCH_END()
+
 ISPCRTKernel ispcrtNewKernel(ISPCRTDevice d, ISPCRTModule m, const char *name) ISPCRT_CATCH_BEGIN {
     const auto &device = referenceFromHandle<ispcrt::base::Device>(d);
     const auto &module = referenceFromHandle<ispcrt::base::Module>(m);

ispcrt/ispcrt.h

@@ -113,12 +113,23 @@ size_t ispcrtSize(ISPCRTMemoryView);
 ISPCRTAllocationType ispcrtGetMemoryViewAllocType(ISPCRTMemoryView);
 ISPCRTAllocationType ispcrtGetMemoryAllocType(ISPCRTDevice d, void* memBuffer);
 
-// Kernels ////////////////////////////////////////////////////////////////////
-typedef struct {
-    uint32_t stackSize;
-} ISPCRTModuleOptions;
+// Modules ////////////////////////////////////////////////////////////////////
+typedef enum {
+    // Module using IGC VC backend
+    ISPCRT_VECTOR_MODULE = 0,
+    // Module using IGC scalar backend
+    ISPCRT_SCALAR_MODULE,
+} ISPCRTModuleType;
+
+struct ISPCRTModuleOptions_ {
+    uint32_t stackSize{0};
+    bool libraryCompilation{false};
+    ISPCRTModuleType moduleType{ISPCRTModuleType::ISPCRT_VECTOR_MODULE};
+};
+typedef struct ISPCRTModuleOptions_ ISPCRTModuleOptions;
 
 ISPCRTModule ispcrtLoadModule(ISPCRTDevice, const char *moduleFile, ISPCRTModuleOptions);
+void ispcrtLinkModules(ISPCRTDevice, ISPCRTModule *modules, uint32_t numModules);
 ISPCRTKernel ispcrtNewKernel(ISPCRTDevice, ISPCRTModule, const char *name);
 
 // Task queues ////////////////////////////////////////////////////////////////

ispcrt/ispcrt.hpp

@@ -102,7 +102,6 @@ inline uint64_t Future::time() const { return ispcrtFutureGetTimeNs(handle()); }
 /////////////////////////////////////////////////////////////////////////////
 // Device wrapper ///////////////////////////////////////////////////////////
 /////////////////////////////////////////////////////////////////////////////
-
 class Device : public GenericObject<ISPCRTDevice> {
   public:
     Device() = default;
@@ -119,6 +118,8 @@ class Device : public GenericObject<ISPCRTDevice> {
     static uint32_t deviceCount(ISPCRTDeviceType type);
     static ISPCRTDeviceInfo deviceInformation(ISPCRTDeviceType type, uint32_t deviceIdx);
     static std::vector<ISPCRTDeviceInfo> allDevicesInformation(ISPCRTDeviceType type);
+    // link modules
+    void linkModules(ISPCRTModule* modules, const uint32_t num);
 };
 
 // Inlined definitions //
@@ -149,6 +150,10 @@ inline std::vector<ISPCRTDeviceInfo> Device::allDevicesInformation(ISPCRTDeviceT
     return devInfo;
 }
 
+inline void Device::linkModules(ISPCRTModule* modules, const uint32_t num) {
+    ispcrtLinkModules(handle(), (ISPCRTModule*)modules, num);
+}
+
 /////////////////////////////////////////////////////////////////////////////
 // Arrays (MemoryView wrapper w/ element type) //////////////////////////////
 /////////////////////////////////////////////////////////////////////////////

ispcrt/tests/level_zero_mock/ze_mock_driver.cpp

@@ -1,4 +1,4 @@
-// Copyright 2020 Intel Corporation
+// Copyright 2020-2022 Intel Corporation
 // SPDX-License-Identifier: BSD-3-Clause
 
 #include "ze_mock.h"
@@ -265,6 +265,30 @@ ze_result_t zeModuleDestroy(ze_module_handle_t hModule) {
     MOCK_RET;
 }
 
+ze_result_t zeModuleDynamicLink(uint32_t numModules,
+                                ze_module_handle_t* phModules,
+                                ze_module_build_log_handle_t* phLinkLog) {
+    MOCK_CNT_CALL;
+    if (phModules == NULL)
+        return ZE_RESULT_ERROR_INVALID_NULL_POINTER;
+    MOCK_RET;
+}
+
+ze_result_t zeModuleBuildLogGetString(ze_module_build_log_handle_t hModuleBuildLog,
+                                      size_t* pSize,
+                                      char* pBuildLog) {
+    MOCK_CNT_CALL;
+    *pSize = 0;
+    MOCK_RET;
+}
+
+ze_result_t zeModuleBuildLogDestroy(ze_module_build_log_handle_t hModuleBuildLog) {
+    MOCK_CNT_CALL;
+    if (hModuleBuildLog == NULL)
+        return ZE_RESULT_ERROR_UNINITIALIZED;
+    MOCK_RET;
+}
+
 ze_result_t zeKernelCreate(ze_module_handle_t hModule, const ze_kernel_desc_t *desc, ze_kernel_handle_t *phKernel) {
     MOCK_CNT_CALL;
     if (hModule != ModuleHandle.get())
@@ -414,6 +438,13 @@ ze_result_t zeGetMemProcAddrTable(ze_api_version_t version, ze_mem_dditable_t *p
 ze_result_t zeGetModuleProcAddrTable(ze_api_version_t version, ze_module_dditable_t *pDdiTable) {
     pDdiTable->pfnCreate = ispcrt::testing::mock::driver::zeModuleCreate;
     pDdiTable->pfnDestroy = ispcrt::testing::mock::driver::zeModuleDestroy;
+    pDdiTable->pfnDynamicLink = ispcrt::testing::mock::driver::zeModuleDynamicLink;
+    return ZE_RESULT_SUCCESS;
+}
+
+ze_result_t zeGetModuleBuildLogProcAddrTable(ze_api_version_t version, ze_module_build_log_dditable_t *pDdiTable) {
+    pDdiTable->pfnGetString = ispcrt::testing::mock::driver::zeModuleBuildLogGetString;
+    pDdiTable->pfnDestroy = ispcrt::testing::mock::driver::zeModuleBuildLogDestroy;
     return ZE_RESULT_SUCCESS;
 }
 
@@ -422,7 +453,6 @@ ze_result_t zeGetModuleProcAddrTable(ze_api_version_t version, ze_module_dditabl
 
 MOCK_DDI_FUN(zeGetFenceProcAddrTable, ze_fence_dditable_t)
 MOCK_DDI_FUN(zeGetImageProcAddrTable, ze_image_dditable_t)
-MOCK_DDI_FUN(zeGetModuleBuildLogProcAddrTable, ze_module_build_log_dditable_t)
 MOCK_DDI_FUN(zeGetPhysicalMemProcAddrTable, ze_physical_mem_dditable_t)
 MOCK_DDI_FUN(zeGetSamplerProcAddrTable, ze_sampler_dditable_t)
 MOCK_DDI_FUN(zeGetVirtualMemProcAddrTable, ze_virtual_mem_dditable_t)

ispcrt/tests/mock_tests/ispcrt_mock_main.cpp

@@ -1,4 +1,4 @@
-// Copyright 2020-2021 Intel Corporation
+// Copyright 2020-2022 Intel Corporation
 // SPDX-License-Identifier: BSD-3-Clause
 
 #include "ispcrt.hpp"
@@ -177,6 +177,33 @@ TEST_F(MockTestWithDevice, Module_Constructor_zeModuleCreate) {
     ASSERT_EQ(sm_rt_error, ISPCRT_DEVICE_LOST);
 }
 
+/////////////////////////////////////////////////////////////////////
+// Dynamic binary linking tests
+
+TEST_F(MockTestWithDevice, Module_DynamicLink) {
+    // Create 2 modules and link them
+    ASSERT_NE(m_device, 0);
+    ispcrt::Module m1(m_device, "");
+    ispcrt::Module m2(m_device, "");
+    Config::setRetValue("zeModuleBuildLogDestroy", ZE_RESULT_SUCCESS);
+    std::array<ISPCRTModule, 2> modules = {
+        (ISPCRTModule)m1.handle(), (ISPCRTModule)m2.handle()};
+    m_device.linkModules(modules.data(), modules.size());
+    ASSERT_EQ(sm_rt_error, ISPCRT_NO_ERROR);
+}
+
+TEST_F(MockTestWithDevice, Module_DynamicLink_zeModuleDynamicLink) {
+    // Check if error is reported when zeModuleDynamicLink is not successful
+    ASSERT_NE(m_device, 0);
+    ispcrt::Module m1(m_device, "");
+    ispcrt::Module m2(m_device, "");
+    Config::setRetValue("zeModuleDynamicLink", ZE_RESULT_ERROR_DEVICE_LOST);
+    std::array<ISPCRTModule, 2> modules = {
+        (ISPCRTModule)m1.handle(), (ISPCRTModule)m2.handle()};
+    m_device.linkModules(modules.data(), modules.size());
+    ASSERT_EQ(sm_rt_error, ISPCRT_DEVICE_LOST);
+}
+
 /////////////////////////////////////////////////////////////////////
 // Kernel tests