Skip to content

Commit 7a3d1b2

Browse files
diptorupddiptorupd
authored
Level zero codegen (#241)
* Add support to build a Level Zero interoperability SYCL program from SPIR-V binary on Linux.
1 parent 91ef1f5 commit 7a3d1b2

File tree

10 files changed

+146
-42
lines changed

10 files changed

+146
-42
lines changed

‎dpctl-capi/dbg_build.sh‎

Lines changed: 10 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -9,15 +9,16 @@ rm -rf ${INSTALL_PREFIX}
99
export ONEAPI_ROOT=/opt/intel/oneapi
1010
DPCPP_ROOT=${ONEAPI_ROOT}/compiler/latest/linux
1111

12-
cmake \
13-
-DCMAKE_BUILD_TYPE=Debug \
14-
-DCMAKE_INSTALL_PREFIX=${INSTALL_PREFIX} \
15-
-DCMAKE_PREFIX_PATH=${INSTALL_PREFIX} \
16-
-DDPCPP_INSTALL_DIR=${DPCPP_ROOT} \
17-
-DCMAKE_C_COMPILER:PATH=${DPCPP_ROOT}/bin/clang \
18-
-DCMAKE_CXX_COMPILER:PATH=${DPCPP_ROOT}/bin/dpcpp \
19-
-DDPCTL_BUILD_CAPI_TESTS=ON \
20-
-DDPCTL_GENERATE_COVERAGE=ON \
12+
cmake \
13+
-DCMAKE_BUILD_TYPE=Debug \
14+
-DCMAKE_INSTALL_PREFIX=${INSTALL_PREFIX} \
15+
-DCMAKE_PREFIX_PATH=${INSTALL_PREFIX} \
16+
-DDPCPP_INSTALL_DIR=${DPCPP_ROOT} \
17+
-DCMAKE_C_COMPILER:PATH=${DPCPP_ROOT}/bin/clang \
18+
-DCMAKE_CXX_COMPILER:PATH=${DPCPP_ROOT}/bin/dpcpp \
19+
-DDPCTL_ENABLE_LO_PROGRAM_CREATION=ON \
20+
-DDPCTL_BUILD_CAPI_TESTS=ON \
21+
-DDPCTL_GENERATE_COVERAGE=ON \
2122
..
2223

2324
make V=1 -n -j 4 && make check && make install

‎dpctl-capi/include/dpctl_sycl_program_interface.h‎

Lines changed: 6 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -52,14 +52,17 @@ DPCTL_C_EXTERN_C_BEGIN
5252
* @param Ctx An opaque pointer to a sycl::context
5353
* @param IL SPIR-V binary
5454
* @param Length The size of the IL binary in bytes.
55+
* @param CompileOpts Optional compiler flags used when compiling th
56+
* SPIR-V binary.
5557
* @return A new SyclProgramRef pointer if the program creation succeeded,
5658
* else returns NULL.
5759
*/
5860
DPCTL_API
5961
__dpctl_give DPCTLSyclProgramRef
60-
DPCTLProgram_CreateFromOCLSpirv (__dpctl_keep const DPCTLSyclContextRef Ctx,
61-
__dpctl_keep const void *IL,
62-
size_t Length);
62+
DPCTLProgram_CreateFromSpirv (__dpctl_keep const DPCTLSyclContextRef Ctx,
63+
__dpctl_keep const void *IL,
64+
size_t Length,
65+
const char *CompileOpts);
6366

6467
/*!
6568
* @brief Create a Sycl program from an OpenCL kernel source string.

‎dpctl-capi/source/dpctl_sycl_program_interface.cpp‎

Lines changed: 68 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -25,10 +25,15 @@
2525
//===----------------------------------------------------------------------===//
2626

2727
#include "dpctl_sycl_program_interface.h"
28+
#include "Config/dpctl_config.h"
2829
#include "Support/CBindingWrapping.h"
2930

30-
#include <CL/sycl.hpp> /* Sycl headers */
31-
#include <CL/cl.h> /* OpenCL headers */
31+
#include <CL/sycl.hpp> /* Sycl headers */
32+
#include <CL/cl.h> /* OpenCL headers */
33+
#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION
34+
#include <level_zero/zet_api.h> /* Level Zero headers */
35+
#include <CL/sycl/backend/level_zero.hpp>
36+
#endif
3237

3338
using namespace cl::sycl;
3439

@@ -41,7 +46,8 @@ DEFINE_SIMPLE_CONVERSION_FUNCTIONS(kernel, DPCTLSyclKernelRef)
4146
__dpctl_give DPCTLSyclProgramRef
4247
createOpenCLInterOpProgram (const context &SyclCtx,
4348
__dpctl_keep const void *IL,
44-
size_t length)
49+
size_t length,
50+
const char * /* */)
4551
{
4652
cl_int err;
4753
auto CLCtx = SyclCtx.get();
@@ -83,34 +89,87 @@ createOpenCLInterOpProgram (const context &SyclCtx,
8389
}
8490
}
8591

92+
#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION
93+
__dpctl_give DPCTLSyclProgramRef
94+
createLevelZeroInterOpProgram (const context &SyclCtx,
95+
const void *IL,
96+
size_t length,
97+
const char *CompileOpts)
98+
{
99+
auto ZeCtx = SyclCtx.get_native<backend::level_zero>();
100+
auto SyclDevices = SyclCtx.get_devices();
101+
if(SyclDevices.size() > 1) {
102+
// We only support build to one device with Level Zero now.
103+
// TODO: log error
104+
return nullptr;
105+
}
106+
107+
// Specialization constants are not yet supported.
108+
// Refer https://bit.ly/33UEDYN for details on specialization constants.
109+
ze_module_constants_t ZeSpecConstants = {};
110+
ZeSpecConstants.numConstants = 0;
111+
112+
// Populate the Level Zero module descriptions
113+
ze_module_desc_t ZeModuleDesc = {};
114+
ZeModuleDesc.format = ZE_MODULE_FORMAT_IL_SPIRV;
115+
ZeModuleDesc.inputSize = length;
116+
ZeModuleDesc.pInputModule = (uint8_t*)IL;
117+
ZeModuleDesc.pBuildFlags = CompileOpts;
118+
ZeModuleDesc.pConstants = &ZeSpecConstants;
119+
120+
auto ZeDevice = SyclDevices[0].get_native<backend::level_zero>();
121+
ze_module_handle_t ZeModule;
122+
auto ret = zeModuleCreate(ZeCtx, ZeDevice, &ZeModuleDesc, &ZeModule,
123+
nullptr);
124+
if(ret != ZE_RESULT_SUCCESS) {
125+
// TODO: handle error
126+
return nullptr;
127+
}
128+
129+
// Create the Sycl program from the ZeModule
130+
try {
131+
auto ZeProgram = new program(sycl::level_zero::make_program(
132+
SyclCtx, reinterpret_cast<uintptr_t>(ZeModule))
133+
);
134+
return wrap(ZeProgram);
135+
} catch (invalid_object_error &e) {
136+
// \todo record error
137+
std::cerr << e.what() << '\n';
138+
return nullptr;
139+
}
140+
}
141+
#endif
86142
} /* end of anonymous namespace */
87143

88144
__dpctl_give DPCTLSyclProgramRef
89-
DPCTLProgram_CreateFromOCLSpirv (__dpctl_keep const DPCTLSyclContextRef CtxRef,
90-
__dpctl_keep const void *IL,
91-
size_t length)
145+
DPCTLProgram_CreateFromSpirv (__dpctl_keep const DPCTLSyclContextRef CtxRef,
146+
__dpctl_keep const void *IL,
147+
size_t length,
148+
const char *CompileOpts)
92149
{
93150
DPCTLSyclProgramRef Pref = nullptr;
94151
context *SyclCtx = nullptr;
95152
if(!CtxRef) {
96153
// \todo handle error
97154
return Pref;
98155
}
99-
100156
SyclCtx = unwrap(CtxRef);
101157
// get the backend type
102158
auto BE = SyclCtx->get_platform().get_backend();
103159
switch (BE)
104160
{
105161
case backend::opencl:
106-
Pref = createOpenCLInterOpProgram(*SyclCtx, IL, length);
162+
Pref = createOpenCLInterOpProgram(*SyclCtx, IL, length, CompileOpts);
107163
break;
108164
case backend::level_zero:
165+
#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION
166+
Pref = createLevelZeroInterOpProgram(*SyclCtx, IL, length,
167+
CompileOpts);
168+
#endif
109169
break;
110170
default:
111171
break;
112172
}
113-
114173
return Pref;
115174
}
116175

‎dpctl-capi/tests/test_sycl_program_interface.cpp‎

Lines changed: 36 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
//===---------- test_sycl_program_interface.cpp - dpctl-C_API --*-- C++ -*-===//
1+
//===---------- test_sycl_program_interface.cpp - dpctl-C_API ---*- C++ -*-===//
22
//
33
// Data Parallel Control Library (dpCtl)
44
//
@@ -29,7 +29,7 @@
2929
#include "dpctl_sycl_program_interface.h"
3030
#include "dpctl_sycl_queue_interface.h"
3131
#include "dpctl_sycl_queue_manager.h"
32-
32+
#include "Config/dpctl_config.h"
3333
#include <array>
3434
#include <fstream>
3535
#include <filesystem>
@@ -127,12 +127,16 @@ struct TestDPCTLSyclProgramInterface : public ::testing::Test
127127
size_t spirvFileSize = 0;
128128
std::vector<char> spirvBuffer;
129129
size_t nOpenCLGpuQ = 0;
130+
#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION
131+
size_t nL0GpuQ = 0;
132+
#endif
130133

131134
TestDPCTLSyclProgramInterface () :
132135
spirvFile{"./multi_kernel.spv", std::ios::binary | std::ios::ate},
133136
spirvFileSize(std::filesystem::file_size("./multi_kernel.spv")),
134137
spirvBuffer(spirvFileSize),
135-
nOpenCLGpuQ(DPCTLQueueMgr_GetNumQueues(DPCTL_OPENCL, DPCTL_GPU))
138+
nOpenCLGpuQ(DPCTLQueueMgr_GetNumQueues(DPCTL_OPENCL, DPCTL_GPU)),
139+
nL0GpuQ(DPCTLQueueMgr_GetNumQueues(DPCTL_LEVEL_ZERO, DPCTL_GPU))
136140
{
137141
spirvFile.seekg(0, std::ios::beg);
138142
spirvFile.read(spirvBuffer.data(), spirvFileSize);
@@ -152,7 +156,7 @@ TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromOCLSource)
152156
auto QueueRef = DPCTLQueueMgr_GetQueue(DPCTL_OPENCL, DPCTL_GPU, 0);
153157
auto CtxRef = DPCTLQueue_GetContext(QueueRef);
154158
auto PRef = DPCTLProgram_CreateFromOCLSource(CtxRef, CLProgramStr,
155-
CompileOpts);
159+
CompileOpts);
156160
ASSERT_TRUE(PRef != nullptr);
157161
ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "add"));
158162
ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "axpy"));
@@ -162,15 +166,36 @@ TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromOCLSource)
162166
DPCTLProgram_Delete(PRef);
163167
}
164168

165-
TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromOCLSpirv)
169+
TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromSpirvOCL)
166170
{
167171
if(!nOpenCLGpuQ)
168172
GTEST_SKIP_("Skipping as no OpenCL GPU device found.\n");
169173

170174
auto QueueRef = DPCTLQueueMgr_GetQueue(DPCTL_OPENCL, DPCTL_GPU, 0);
171175
auto CtxRef = DPCTLQueue_GetContext(QueueRef);
172-
auto PRef = DPCTLProgram_CreateFromOCLSpirv(CtxRef, spirvBuffer.data(),
173-
spirvFileSize);
176+
auto PRef = DPCTLProgram_CreateFromSpirv(CtxRef, spirvBuffer.data(),
177+
spirvFileSize,
178+
nullptr);
179+
ASSERT_TRUE(PRef != nullptr);
180+
ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "add"));
181+
ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "axpy"));
182+
183+
DPCTLQueue_Delete(QueueRef);
184+
DPCTLContext_Delete(CtxRef);
185+
DPCTLProgram_Delete(PRef);
186+
}
187+
188+
#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION
189+
TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromSpirvL0)
190+
{
191+
if(!nL0GpuQ)
192+
GTEST_SKIP_("Skipping as no OpenCL GPU device found.\n");
193+
194+
auto QueueRef = DPCTLQueueMgr_GetQueue(DPCTL_LEVEL_ZERO, DPCTL_GPU, 0);
195+
auto CtxRef = DPCTLQueue_GetContext(QueueRef);
196+
auto PRef = DPCTLProgram_CreateFromSpirv(CtxRef, spirvBuffer.data(),
197+
spirvFileSize,
198+
nullptr);
174199
ASSERT_TRUE(PRef != nullptr);
175200
ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "add"));
176201
ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "axpy"));
@@ -179,6 +204,7 @@ TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromOCLSpirv)
179204
DPCTLContext_Delete(CtxRef);
180205
DPCTLProgram_Delete(PRef);
181206
}
207+
#endif
182208

183209
TEST_F (TestDPCTLSyclProgramInterface, CheckGetKernelOCLSource)
184210
{
@@ -207,15 +233,15 @@ TEST_F (TestDPCTLSyclProgramInterface, CheckGetKernelOCLSource)
207233
DPCTLProgram_Delete(PRef);
208234
}
209235

210-
TEST_F (TestDPCTLSyclProgramInterface, CheckGetKernelOCLSpirv)
236+
TEST_F (TestDPCTLSyclProgramInterface, CheckGetKernelSpirv)
211237
{
212238
if(!nOpenCLGpuQ)
213239
GTEST_SKIP_("Skipping as no OpenCL GPU device found.\n");
214240

215241
auto QueueRef = DPCTLQueueMgr_GetQueue(DPCTL_OPENCL, DPCTL_GPU, 0);
216242
auto CtxRef = DPCTLQueue_GetContext(QueueRef);
217-
auto PRef = DPCTLProgram_CreateFromOCLSpirv(CtxRef, spirvBuffer.data(),
218-
spirvFileSize);
243+
auto PRef = DPCTLProgram_CreateFromSpirv(CtxRef, spirvBuffer.data(),
244+
spirvFileSize, nullptr);
219245
auto AddKernel = DPCTLProgram_GetKernel(PRef, "add");
220246
auto AxpyKernel = DPCTLProgram_GetKernel(PRef, "axpy");
221247

‎dpctl/_backend.pxd‎

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -140,14 +140,15 @@ cdef extern from "dpctl_sycl_context_interface.h":
140140

141141

142142
cdef extern from "dpctl_sycl_program_interface.h":
143-
cdef DPCTLSyclProgramRef DPCTLProgram_CreateFromOCLSpirv (
143+
cdef DPCTLSyclProgramRef DPCTLProgram_CreateFromSpirv (
144144
const DPCTLSyclContextRef Ctx,
145145
const void *IL,
146-
size_t Length)
146+
size_t Length,
147+
const char *CompileOpts)
147148
cdef DPCTLSyclProgramRef DPCTLProgram_CreateFromOCLSource (
148149
const DPCTLSyclContextRef Ctx,
149-
const char* Source,
150-
const char* CompileOpts)
150+
const char *Source,
151+
const char *CompileOpts)
151152
cdef DPCTLSyclKernelRef DPCTLProgram_GetKernel (DPCTLSyclProgramRef PRef,
152153
const char *KernelName)
153154
cdef bool DPCTLProgram_HasKernel (DPCTLSyclProgramRef PRef,

‎dpctl/_sycl_core.pyx‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -678,7 +678,7 @@ cdef class _SyclRTManager:
678678
raise UnsupportedDeviceError("Device can only be gpu or cpu")
679679
except KeyError:
680680
raise UnsupportedBackendError("Backend can only be opencl or "
681-
"level-0")
681+
"level0")
682682

683683
def _remove_current_queue(self):
684684
DPCTLQueueMgr_PopQueue()
@@ -970,7 +970,7 @@ cdef class _SyclRTManager:
970970
raise UnsupportedDeviceError("Device can only be gpu or cpu")
971971
except KeyError:
972972
raise UnsupportedBackendError("Backend can only be opencl or "
973-
"level-0")
973+
"level0")
974974

975975

976976
# This private instance of the _SyclQueueManager should not be directly

‎dpctl/program/_program.pxd‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -60,4 +60,5 @@ cdef class SyclProgram:
6060

6161

6262
cpdef create_program_from_source (SyclQueue q, unicode source, unicode copts=*)
63-
cpdef create_program_from_spirv (SyclQueue q, const unsigned char[:] IL)
63+
cpdef create_program_from_spirv (SyclQueue q, const unsigned char[:] IL,
64+
unicode copts=*)

‎dpctl/program/_program.pyx‎

Lines changed: 7 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -162,7 +162,8 @@ cpdef create_program_from_source(SyclQueue q, unicode src, unicode copts=""):
162162

163163
cimport cython.array
164164

165-
cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL):
165+
cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL,
166+
unicode copts=""):
166167
"""
167168
Creates a Sycl interoperability program from an SPIR-V binary.
168169
@@ -173,6 +174,8 @@ cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL):
173174
q (SyclQueue): The :class:`SyclQueue` for which the
174175
:class:`SyclProgram` is going to be built.
175176
IL (const char[:]) : SPIR-V binary IL file for an OpenCL program.
177+
copts (unicode) : Optional compilation flags that will be used
178+
when compiling the program.
176179
177180
Returns:
178181
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):
185188
cdef const unsigned char *dIL = &IL[0]
186189
cdef DPCTLSyclContextRef CRef = q.get_sycl_context().get_context_ref()
187190
cdef size_t length = IL.shape[0]
188-
Pref = DPCTLProgram_CreateFromOCLSpirv(CRef, <const void*>dIL, length)
191+
cdef bytes bCOpts = copts.encode('utf8')
192+
cdef const char *COpts = <const char*>bCOpts
193+
Pref = DPCTLProgram_CreateFromSpirv(CRef, <const void*>dIL, length, COpts)
189194
if Pref is NULL:
190195
raise SyclProgramCompilationError()
191196

‎dpctl/tests/test_sycl_program.py‎

Lines changed: 9 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -87,9 +87,16 @@ def test_create_program_from_spirv(self):
8787
"No Level0 GPU queues available",
8888
)
8989
class TestProgramForLevel0GPU(unittest.TestCase):
90-
@unittest.expectedFailure
91-
def test_create_program_from_spirv(self):
9290

91+
import sys
92+
93+
# Level zero program creation from a SPIR-V binary is not supported
94+
# on Windows.
95+
@unittest.skipIf(
96+
sys.platform in ["win32", "cygwin"],
97+
"Level Zero module creation unsupported on Windows.",
98+
)
99+
def test_create_program_from_spirv(self):
93100
CURR_DIR = os.path.dirname(os.path.abspath(__file__))
94101
spirv_file = os.path.join(CURR_DIR, "input_files/multi_kernel.spv")
95102
with open(spirv_file, "rb") as fin:

‎scripts/build_backend.py‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -43,6 +43,7 @@
4343
"-DDPCPP_INSTALL_DIR=" + DPCPP_ROOT,
4444
"-DCMAKE_C_COMPILER:PATH=" + os.path.join(DPCPP_ROOT, "bin", "clang"),
4545
"-DCMAKE_CXX_COMPILER:PATH=" + os.path.join(DPCPP_ROOT, "bin", "clang++"),
46+
"-DDPCTL_ENABLE_LO_PROGRAM_CREATION=ON",
4647
backends,
4748
]
4849
subprocess.check_call(cmake_args, stderr=subprocess.STDOUT, shell=False)

0 commit comments

Comments
 (0)