215 lines · cpp
1//===- SyclRuntimeWrappers.cpp - MLIR SYCL wrapper library ------------===//2//3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.4// See https://llvm.org/LICENSE.txt for license information.5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception6//7//===----------------------------------------------------------------------===//8//9// Implements wrappers around the sycl runtime library with C linkage10//11//===----------------------------------------------------------------------===//12 13#include <level_zero/ze_api.h>14#include <sycl/ext/oneapi/backend/level_zero.hpp>15#include <sycl/sycl.hpp>16 17#ifdef _WIN3218#define SYCL_RUNTIME_EXPORT __declspec(dllexport)19#else20#define SYCL_RUNTIME_EXPORT21#endif // _WIN3222 23namespace {24 25template <typename F>26auto catchAll(F &&func) {27 try {28 return func();29 } catch (const std::exception &e) {30 fprintf(stdout, "An exception was thrown: %s\n", e.what());31 fflush(stdout);32 abort();33 } catch (...) {34 fprintf(stdout, "An unknown exception was thrown\n");35 fflush(stdout);36 abort();37 }38}39 40#define L0_SAFE_CALL(call) \41 { \42 ze_result_t status = (call); \43 if (status != ZE_RESULT_SUCCESS) { \44 fprintf(stdout, "L0 error %d\n", status); \45 fflush(stdout); \46 abort(); \47 } \48 }49 50} // namespace51 52static sycl::device getDefaultDevice() {53 static sycl::device syclDevice;54 static bool isDeviceInitialised = false;55 if (!isDeviceInitialised) {56 auto platformList = sycl::platform::get_platforms();57 for (const auto &platform : platformList) {58 auto platformName = platform.get_info<sycl::info::platform::name>();59 bool isLevelZero = platformName.find("Level-Zero") != std::string::npos;60 if (!isLevelZero)61 continue;62 63 syclDevice = platform.get_devices()[0];64 isDeviceInitialised = true;65 return syclDevice;66 }67 throw std::runtime_error("getDefaultDevice failed");68 } else {69 return syclDevice;70 }71}72 73static sycl::context getDefaultContext() {74 static sycl::context syclContext{getDefaultDevice()};75 return syclContext;76}77 78static void *allocDeviceMemory(sycl::queue *queue, size_t size, bool isShared) {79 void *memPtr = nullptr;80 if (isShared) {81 memPtr = sycl::aligned_alloc_shared(64, size, getDefaultDevice(),82 getDefaultContext());83 } else {84 memPtr = sycl::aligned_alloc_device(64, size, *queue);85 }86 if (memPtr == nullptr) {87 throw std::runtime_error("mem allocation failed!");88 }89 return memPtr;90}91 92static void deallocDeviceMemory(sycl::queue *queue, void *ptr) {93 sycl::free(ptr, *queue);94}95 96static ze_module_handle_t loadModule(const void *data, size_t dataSize) {97 assert(data);98 ze_module_handle_t zeModule;99 ze_module_desc_t desc = {ZE_STRUCTURE_TYPE_MODULE_DESC,100 nullptr,101 ZE_MODULE_FORMAT_IL_SPIRV,102 dataSize,103 (const uint8_t *)data,104 nullptr,105 nullptr};106 auto zeDevice = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(107 getDefaultDevice());108 auto zeContext = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(109 getDefaultContext());110 L0_SAFE_CALL(zeModuleCreate(zeContext, zeDevice, &desc, &zeModule, nullptr));111 return zeModule;112}113 114static sycl::kernel *getKernel(ze_module_handle_t zeModule, const char *name) {115 assert(zeModule);116 assert(name);117 ze_kernel_handle_t zeKernel;118 ze_kernel_desc_t desc = {};119 desc.pKernelName = name;120 121 L0_SAFE_CALL(zeKernelCreate(zeModule, &desc, &zeKernel));122 sycl::kernel_bundle<sycl::bundle_state::executable> kernelBundle =123 sycl::make_kernel_bundle<sycl::backend::ext_oneapi_level_zero,124 sycl::bundle_state::executable>(125 {zeModule}, getDefaultContext());126 127 auto kernel = sycl::make_kernel<sycl::backend::ext_oneapi_level_zero>(128 {kernelBundle, zeKernel}, getDefaultContext());129 return new sycl::kernel(kernel);130}131 132static void launchKernel(sycl::queue *queue, sycl::kernel *kernel, size_t gridX,133 size_t gridY, size_t gridZ, size_t blockX,134 size_t blockY, size_t blockZ, size_t sharedMemBytes,135 void **params, size_t paramsCount) {136 auto syclGlobalRange =137 sycl::range<3>(blockZ * gridZ, blockY * gridY, blockX * gridX);138 auto syclLocalRange = sycl::range<3>(blockZ, blockY, blockX);139 sycl::nd_range<3> syclNdRange(syclGlobalRange, syclLocalRange);140 141 queue->submit([&](sycl::handler &cgh) {142 for (size_t i = 0; i < paramsCount; i++) {143 cgh.set_arg(static_cast<uint32_t>(i), *(static_cast<void **>(params[i])));144 }145 cgh.parallel_for(syclNdRange, *kernel);146 });147}148 149// Wrappers150 151extern "C" SYCL_RUNTIME_EXPORT sycl::queue *mgpuStreamCreate() {152 153 return catchAll([&]() {154 sycl::queue *queue =155 new sycl::queue(getDefaultContext(), getDefaultDevice());156 return queue;157 });158}159 160extern "C" SYCL_RUNTIME_EXPORT void mgpuStreamDestroy(sycl::queue *queue) {161 catchAll([&]() { delete queue; });162}163 164extern "C" SYCL_RUNTIME_EXPORT void *165mgpuMemAlloc(uint64_t size, sycl::queue *queue, bool isShared) {166 return catchAll([&]() {167 return allocDeviceMemory(queue, static_cast<size_t>(size), true);168 });169}170 171extern "C" SYCL_RUNTIME_EXPORT void mgpuMemFree(void *ptr, sycl::queue *queue) {172 catchAll([&]() {173 if (ptr) {174 deallocDeviceMemory(queue, ptr);175 }176 });177}178 179extern "C" SYCL_RUNTIME_EXPORT ze_module_handle_t180mgpuModuleLoad(const void *data, size_t gpuBlobSize) {181 return catchAll([&]() { return loadModule(data, gpuBlobSize); });182}183 184extern "C" SYCL_RUNTIME_EXPORT sycl::kernel *185mgpuModuleGetFunction(ze_module_handle_t module, const char *name) {186 return catchAll([&]() { return getKernel(module, name); });187}188 189extern "C" SYCL_RUNTIME_EXPORT void190mgpuLaunchKernel(sycl::kernel *kernel, size_t gridX, size_t gridY, size_t gridZ,191 size_t blockX, size_t blockY, size_t blockZ,192 size_t sharedMemBytes, sycl::queue *queue, void **params,193 void ** /*extra*/, size_t paramsCount) {194 return catchAll([&]() {195 launchKernel(queue, kernel, gridX, gridY, gridZ, blockX, blockY, blockZ,196 sharedMemBytes, params, paramsCount);197 });198}199 200extern "C" SYCL_RUNTIME_EXPORT void mgpuStreamSynchronize(sycl::queue *queue) {201 202 catchAll([&]() { queue->wait(); });203}204 205extern "C" SYCL_RUNTIME_EXPORT void206mgpuModuleUnload(ze_module_handle_t module) {207 208 catchAll([&]() { L0_SAFE_CALL(zeModuleDestroy(module)); });209}210 211extern "C" SYCL_RUNTIME_EXPORT void212mgpuMemcpy(void *dst, void *src, size_t sizeBytes, sycl::queue *queue) {213 catchAll([&]() { queue->memcpy(dst, src, sizeBytes).wait(); });214}215