• Home
  • Features
  • Pricing
  • Docs
  • Announcements
  • Sign In

IntelPython / dpctl / 27977445675

22 Jun 2026 07:11PM UTC coverage: 75.406% (-0.3%) from 75.677%
27977445675

Pull #2304

github

web-flow
Merge f7daf456f into 61d1fc0cd
Pull Request #2304: Add support for specialization constants

869 of 1218 branches covered (71.35%)

Branch coverage included in aggregate %.

113 of 157 new or added lines in 4 files covered. (71.97%)

110 existing lines in 1 file now uncovered.

3356 of 4385 relevant lines covered (76.53%)

266.27 hits per line

Source File
Press 'n' to go to next uncovered line, 'b' for previous

46.5
/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp
1
//===- dpctl_sycl_kernel_bundle_interface.cpp - Implements C API for
2
//    sycl::kernel_bundle<sycl::bundle_state::executable>  ---------------===//
3
//
4
//                      Data Parallel Control (dpctl)
5
//
6
// Copyright 2020 Intel Corporation
7
//
8
// Licensed under the Apache License, Version 2.0 (the "License");
9
// you may not use this file except in compliance with the License.
10
// You may obtain a copy of the License at
11
//
12
//    http://www.apache.org/licenses/LICENSE-2.0
13
//
14
// Unless required by applicable law or agreed to in writing, software
15
// distributed under the License is distributed on an "AS IS" BASIS,
16
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
17
// See the License for the specific language governing permissions and
18
// limitations under the License.
19
//
20
//===----------------------------------------------------------------------===//
21
///
22
/// \file
23
/// This file implements the functions declared in
24
/// dpctl_sycl_kernel_bundle_interface.h.
25
///
26
//===----------------------------------------------------------------------===//
27

28
#include "dpctl_sycl_kernel_bundle_interface.h"
29
#include "Config/dpctl_config.h"
30
#include "dpctl_dynamic_lib_helper.h"
31
#include "dpctl_error_handlers.h"
32
#include "dpctl_sycl_type_casters.hpp"
33
#include <CL/cl.h> /* OpenCL headers     */
34
#include <cstdint>
35
#include <sstream>
36
#include <stddef.h>
37
#include <sycl/backend/opencl.hpp>
38
#include <sycl/sycl.hpp> /* Sycl headers       */
39
#include <utility>
40

41
#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION
42
// Note: include ze_api.h before level_zero.hpp. Make sure clang-format does
43
// not reorder the includes.
44
// clang-format off
45
#include "ze_api.h" /* Level Zero headers */
46
#include <sycl/ext/oneapi/backend/level_zero.hpp>
47
// clang-format on
48
#endif
49

50
using namespace sycl;
51

52
namespace
53
{
54
static_assert(__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_VERSION_REQUIRED,
55
              "The compiler does not meet minimum version requirement");
56

57
using namespace dpctl::syclinterface;
58

59
#ifdef __linux__
60
static const char *clLoaderName = DPCTL_LIBCL_LOADER_FILENAME;
61
static const int clLibLoadFlags = RTLD_NOLOAD | RTLD_NOW | RTLD_LOCAL;
62
#elif defined(_WIN64)
63
static const char *clLoaderName = "OpenCL.dll";
64
static const int clLibLoadFlags = 0;
65
#else
66
#error "OpenCL program compilation is unavailable for this platform"
67
#endif
68

69
#define CodeStringSuffix(code)                                                 \
70
    std::string(" (code=") + std::to_string(static_cast<int>(code)) + ")"
6✔
71

72
#define EnumCaseString(code)                                                   \
73
    case code:                                                                 \
2✔
74
        return std::string(#code) + CodeStringSuffix(code)
2✔
75

76
static constexpr backend cl_be = backend::opencl;
77

78
struct cl_loader
79
{
80
public:
81
    static cl_loader &get()
82
    {
11✔
83
        static cl_loader _loader;
11✔
84
        return _loader;
11✔
85
    }
11✔
86

87
    template <typename retTy> retTy getSymbol(const char *name)
88
    {
11✔
89
        if (!opened) {
11!
90
            error_handler("The OpenCL loader dynamic library could not "
×
91
                          "be opened.",
×
92
                          __FILE__, __func__, __LINE__);
×
93

94
            return nullptr;
×
95
        }
×
96
        return clLib.getSymbol<retTy>(name);
11✔
97
    }
11✔
98

99
private:
100
    dpctl::DynamicLibHelper clLib;
101
    bool opened;
102
    cl_loader() : clLib(clLoaderName, clLibLoadFlags), opened(clLib.opened()) {}
2✔
103
};
104

105
typedef cl_program (*clCreateProgramWithSourceFT)(cl_context,
106
                                                  cl_uint,
107
                                                  const char **,
108
                                                  const size_t *,
109
                                                  cl_int *);
110
const char *clCreateProgramWithSource_Name = "clCreateProgramWithSource";
111
clCreateProgramWithSourceFT get_clCreateProgramWithSource()
112
{
28✔
113
    static auto st_clCreateProgramWithSourceF =
28✔
114
        cl_loader::get().getSymbol<clCreateProgramWithSourceFT>(
28✔
115
            clCreateProgramWithSource_Name);
28✔
116

117
    return st_clCreateProgramWithSourceF;
28✔
118
}
28✔
119

120
typedef cl_program (*clCreateProgramWithILFT)(cl_context,
121
                                              const void *,
122
                                              size_t,
123
                                              cl_int *);
124
const char *clCreateProgramWithIL_Name = "clCreateProgramWithIL";
125
clCreateProgramWithILFT get_clCreateProgramWithIL()
126
{
66✔
127
    static auto st_clCreateProgramWithILF =
66✔
128
        cl_loader::get().getSymbol<clCreateProgramWithILFT>(
66✔
129
            clCreateProgramWithIL_Name);
66✔
130

131
    return st_clCreateProgramWithILF;
66✔
132
}
66✔
133
typedef cl_int (*clBuildProgramFT)(cl_program,
134
                                   cl_uint,
135
                                   const cl_device_id *,
136
                                   const char *,
137
                                   void (*)(cl_program, void *),
138
                                   void *);
139
const char *clBuildProgram_Name = "clBuildProgram";
140
clBuildProgramFT get_clBuildProgram()
141
{
94✔
142
    static auto st_clBuildProgramF =
94✔
143
        cl_loader::get().getSymbol<clBuildProgramFT>(clBuildProgram_Name);
94✔
144

145
    return st_clBuildProgramF;
94✔
146
}
94✔
147

148
typedef cl_kernel (*clCreateKernelFT)(cl_program, const char *, cl_int *);
149
const char *clCreateKernel_Name = "clCreateKernel";
150
clCreateKernelFT get_clCreateKernel()
151
{
153✔
152
    static auto st_clCreateKernelF =
153✔
153
        cl_loader::get().getSymbol<clCreateKernelFT>(clCreateKernel_Name);
153✔
154

155
    return st_clCreateKernelF;
153✔
156
}
153✔
157

158
std::string _GetErrorCode_ocl_impl(cl_int code)
159
{
2✔
160
    switch (code) {
2✔
161
        EnumCaseString(CL_BUILD_PROGRAM_FAILURE);
2!
162
        EnumCaseString(CL_INVALID_CONTEXT);
×
163
        EnumCaseString(CL_INVALID_DEVICE);
×
164
        EnumCaseString(CL_INVALID_VALUE);
×
165
        EnumCaseString(CL_OUT_OF_RESOURCES);
×
166
        EnumCaseString(CL_OUT_OF_HOST_MEMORY);
×
167
        EnumCaseString(CL_INVALID_OPERATION);
×
168
        EnumCaseString(CL_INVALID_BINARY);
×
169
    default:
×
170
        return "<< ERROR CODE UNRECOGNIZED >>" + CodeStringSuffix(code);
×
171
    }
2✔
172
}
2✔
173

174
typedef cl_int (*clSetProgramSpecializationConstantFT)(cl_program,
175
                                                       cl_uint,
176
                                                       size_t,
177
                                                       const void *);
178
const char *clSetProgramSpecializationConstant_Name =
179
    "clSetProgramSpecializationConstant";
180
clSetProgramSpecializationConstantFT get_clSetProgramSpecializationConstant()
181
{
2✔
182
    static auto st_clSetProgramSpecializationConstantF =
2✔
183
        cl_loader::get().getSymbol<clSetProgramSpecializationConstantFT>(
2✔
184
            clSetProgramSpecializationConstant_Name);
2✔
185

186
    return st_clSetProgramSpecializationConstantF;
2✔
187
}
2✔
188

189
typedef cl_int (*clReleaseProgramFT)(cl_program);
190
const char *clReleaseProgram_Name = "clReleaseProgram";
191
clReleaseProgramFT get_clReleaseProgram()
192
{
2✔
193
    static auto st_clReleaseProgramF =
2✔
194
        cl_loader::get().getSymbol<clReleaseProgramFT>(clReleaseProgram_Name);
2✔
195

196
    return st_clReleaseProgramF;
2✔
197
}
2✔
198

199
DPCTLSyclKernelBundleRef
200
_CreateKernelBundle_common_ocl_impl(cl_program clProgram,
201
                                    const context &ctx,
202
                                    const device &dev,
203
                                    const char *CompileOpts)
204
{
94✔
205
    backend_traits<cl_be>::return_type<device> clDevice;
94✔
206
    clDevice = get_native<cl_be>(dev);
94✔
207

208
    // Last two pointers are notification function pointer and user-data pointer
209
    // that can be passed to the notification function.
210
    auto clBuildProgramF = get_clBuildProgram();
94✔
211
    if (clBuildProgramF == nullptr) {
94!
NEW
212
        auto clReleaseProgramF = get_clReleaseProgram();
×
NEW
213
        if (clReleaseProgramF) {
×
NEW
214
            clReleaseProgramF(clProgram);
×
NEW
215
        }
×
216
        return nullptr;
×
217
    }
×
218
    cl_int build_status =
94✔
219
        clBuildProgramF(clProgram, 1, &clDevice, CompileOpts, nullptr, nullptr);
94✔
220

221
    if (build_status != CL_SUCCESS) {
94✔
222
        error_handler("clBuildProgram failed: " +
2✔
223
                          _GetErrorCode_ocl_impl(build_status),
2✔
224
                      __FILE__, __func__, __LINE__);
2✔
225
        auto clReleaseProgramF = get_clReleaseProgram();
2✔
226
        if (clReleaseProgramF) {
2!
227
            clReleaseProgramF(clProgram);
2✔
228
        }
2✔
229
        return nullptr;
2✔
230
    }
2✔
231

232
    using ekbTy = kernel_bundle<bundle_state::executable>;
92✔
233
    const ekbTy &kb =
92✔
234
        make_kernel_bundle<cl_be, bundle_state::executable>(clProgram, ctx);
92✔
235
    return wrap<ekbTy>(new ekbTy(kb));
92✔
236
}
94✔
237

238
DPCTLSyclKernelBundleRef
239
_CreateKernelBundleWithOCLSource_ocl_impl(const context &ctx,
240
                                          const device &dev,
241
                                          const char *oclSrc,
242
                                          const char *CompileOpts)
243
{
28✔
244
    auto clCreateProgramWithSourceF = get_clCreateProgramWithSource();
28✔
245
    if (clCreateProgramWithSourceF == nullptr) {
28!
246
        return nullptr;
×
247
    }
×
248

249
    backend_traits<cl_be>::return_type<context> clContext;
28✔
250
    clContext = get_native<cl_be>(ctx);
28✔
251

252
    cl_int build_with_source_err_code = CL_SUCCESS;
28✔
253
    cl_program clProgram = clCreateProgramWithSourceF(
28✔
254
        clContext, 1, &oclSrc, nullptr, &build_with_source_err_code);
28✔
255

256
    if (build_with_source_err_code != CL_SUCCESS) {
28!
257
        error_handler("clPCreateProgramWithSource failed with " +
×
258
                          _GetErrorCode_ocl_impl(build_with_source_err_code),
×
259
                      __FILE__, __func__, __LINE__);
×
260
        return nullptr;
×
261
    }
×
262

263
    return _CreateKernelBundle_common_ocl_impl(clProgram, ctx, dev,
28✔
264
                                               CompileOpts);
28✔
265
}
28✔
266

267
DPCTLSyclKernelBundleRef
268
_CreateKernelBundleWithIL_ocl_impl(const context &ctx,
269
                                   const device &dev,
270
                                   const void *IL,
271
                                   size_t il_length,
272
                                   const char *CompileOpts,
273
                                   size_t NumSpecConsts,
274
                                   const DPCTLSpecConst *SpecConsts)
275
{
66✔
276
    auto clCreateProgramWithILF = get_clCreateProgramWithIL();
66✔
277
    if (clCreateProgramWithILF == nullptr) {
66!
278
        return nullptr;
×
279
    }
×
280

281
    backend_traits<cl_be>::return_type<context> clContext;
66✔
282
    clContext = get_native<cl_be>(ctx);
66✔
283

284
    cl_int create_err_code = CL_SUCCESS;
66✔
285
    cl_program clProgram =
66✔
286
        clCreateProgramWithILF(clContext, IL, il_length, &create_err_code);
66✔
287

288
    if (create_err_code != CL_SUCCESS) {
66!
289
        error_handler("OpenCL program could not be created from the SPIR-V "
×
290
                      "binary. OpenCL Error " +
×
291
                          _GetErrorCode_ocl_impl(create_err_code),
×
292
                      __FILE__, __func__, __LINE__);
×
293
        return nullptr;
×
294
    }
×
295

296
    if (SpecConsts != nullptr && NumSpecConsts > 0) {
66!
297
        auto clSetProgramSpecConstF = get_clSetProgramSpecializationConstant();
2✔
298
        if (clSetProgramSpecConstF) {
2!
299
            for (size_t i = 0; i < NumSpecConsts; ++i) {
6✔
300
                clSetProgramSpecConstF(clProgram, SpecConsts[i].id,
4✔
301
                                       SpecConsts[i].size, SpecConsts[i].value);
4✔
302
            }
4✔
303
        }
2✔
NEW
304
        else {
×
NEW
305
            error_handler("clSetProgramSpecializationConstant is not available "
×
NEW
306
                          "in the OpenCL implementation.",
×
NEW
307
                          __FILE__, __func__, __LINE__);
×
308

NEW
309
            auto clReleaseProgramF = get_clReleaseProgram();
×
NEW
310
            if (clReleaseProgramF) {
×
NEW
311
                clReleaseProgramF(clProgram);
×
NEW
312
            }
×
313

NEW
314
            return nullptr;
×
NEW
315
        }
×
316
    }
2✔
317

318
    return _CreateKernelBundle_common_ocl_impl(clProgram, ctx, dev,
66✔
319
                                               CompileOpts);
66✔
320
}
66✔
321

322
bool _HasKernel_ocl_impl(const kernel_bundle<bundle_state::executable> &kb,
323
                         const char *kernel_name)
324
{
62✔
325
    auto clCreateKernelF = get_clCreateKernel();
62✔
326
    if (clCreateKernelF == nullptr) {
62!
327
        return false;
×
328
    }
×
329

330
    std::vector<cl_program> oclKB = get_native<cl_be>(kb);
62✔
331

332
    bool found = false;
62✔
333
    for (auto &cl_pr : oclKB) {
62!
334
        cl_int create_kernel_err_code = CL_SUCCESS;
62✔
335
        [[maybe_unused]] cl_kernel try_kern =
62✔
336
            clCreateKernelF(cl_pr, kernel_name, &create_kernel_err_code);
62✔
337
        if (create_kernel_err_code == CL_SUCCESS) {
62!
338
            found = true;
62✔
339
            break;
62✔
340
        }
62✔
341
    }
62✔
342
    return found;
62✔
343
}
62✔
344

345
__dpctl_give DPCTLSyclKernelRef
346
_GetKernel_ocl_impl(const kernel_bundle<bundle_state::executable> &kb,
347
                    const char *kernel_name)
348
{
91✔
349
    auto clCreateKernelF = get_clCreateKernel();
91✔
350
    if (clCreateKernelF == nullptr) {
91!
UNCOV
351
        return nullptr;
×
UNCOV
352
    }
×
353

354
    std::vector<cl_program> oclKB = get_native<cl_be>(kb);
91✔
355

356
    bool found = false;
91✔
357
    cl_kernel ocl_kernel_from_kb;
91✔
358
    for (auto &cl_pr : oclKB) {
91✔
359
        cl_int create_kernel_err_code = CL_SUCCESS;
91✔
360
        cl_kernel try_kern =
91✔
361
            clCreateKernelF(cl_pr, kernel_name, &create_kernel_err_code);
91✔
362
        if (create_kernel_err_code == CL_SUCCESS) {
91✔
363
            found = true;
90✔
364
            ocl_kernel_from_kb = try_kern;
90✔
365
            break;
90✔
366
        }
90✔
367
    }
91✔
368
    if (found) {
91✔
369
        try {
90✔
370
            context ctx = kb.get_context();
90✔
371

372
            const kernel &interop_kernel =
90✔
373
                make_kernel<cl_be>(ocl_kernel_from_kb, ctx);
90✔
374

375
            return wrap<kernel>(new kernel(interop_kernel));
90✔
376
        } catch (std::exception const &e) {
90✔
UNCOV
377
            error_handler(e, __FILE__, __func__, __LINE__);
×
UNCOV
378
            return nullptr;
×
UNCOV
379
        }
×
380
    }
90✔
381
    else {
1✔
382
        error_handler("Kernel " + std::string(kernel_name) + " not found.",
1✔
383
                      __FILE__, __func__, __LINE__);
1✔
384
        return nullptr;
1✔
385
    }
1✔
386
}
91✔
387

388
#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION
389

390
#ifdef __linux__
391
static const char *zeLoaderName = DPCTL_LIBZE_LOADER_FILENAME;
392
static const int zeLibLoadFlags = RTLD_NOLOAD | RTLD_NOW | RTLD_LOCAL;
393
#elif defined(_WIN64)
394
static const char *zeLoaderName = "ze_loader.dll";
395
static const int zeLibLoadFlags = 0;
396
#else
397
#error "Level Zero program compilation is unavailable for this platform"
398
#endif
399

400
static constexpr sycl::backend ze_be = sycl::backend::ext_oneapi_level_zero;
401

402
struct ze_loader
403
{
404
public:
405
    static ze_loader &get()
UNCOV
406
    {
×
UNCOV
407
        static ze_loader _loader;
×
UNCOV
408
        return _loader;
×
UNCOV
409
    }
×
410

411
    template <typename retTy> retTy getSymbol(const char *name)
UNCOV
412
    {
×
UNCOV
413
        if (!opened) {
×
UNCOV
414
            error_handler("The Level-Zero loader dynamic library could not "
×
UNCOV
415
                          "be opened.",
×
UNCOV
416
                          __FILE__, __func__, __LINE__);
×
417

UNCOV
418
            return nullptr;
×
UNCOV
419
        }
×
UNCOV
420
        return zeLib.getSymbol<retTy>(name);
×
UNCOV
421
    }
×
422

423
private:
424
    dpctl::DynamicLibHelper zeLib;
425
    bool opened;
426
    ze_loader() : zeLib(zeLoaderName, zeLibLoadFlags), opened(zeLib.opened()) {}
×
427
};
428

429
typedef ze_result_t (*zeModuleCreateFT)(ze_context_handle_t,
430
                                        ze_device_handle_t,
431
                                        const ze_module_desc_t *,
432
                                        ze_module_handle_t *,
433
                                        ze_module_build_log_handle_t *);
434
const char *zeModuleCreate_Name = "zeModuleCreate";
435
zeModuleCreateFT get_zeModuleCreate()
436
{
×
437
    static auto st_zeModuleCreateF =
×
438
        ze_loader::get().getSymbol<zeModuleCreateFT>(zeModuleCreate_Name);
×
439

UNCOV
440
    return st_zeModuleCreateF;
×
UNCOV
441
}
×
442

443
typedef ze_result_t (*zeModuleDestroyFT)(ze_module_handle_t);
444
const char *zeModuleDestroy_Name = "zeModuleDestroy";
445
zeModuleDestroyFT get_zeModuleDestroy()
UNCOV
446
{
×
UNCOV
447
    static auto st_zeModuleDestroyF =
×
UNCOV
448
        ze_loader::get().getSymbol<zeModuleDestroyFT>(zeModuleDestroy_Name);
×
449

UNCOV
450
    return st_zeModuleDestroyF;
×
UNCOV
451
}
×
452

453
typedef ze_result_t (*zeKernelCreateFT)(ze_module_handle_t,
454
                                        const ze_kernel_desc_t *,
455
                                        ze_kernel_handle_t *);
456
const char *zeKernelCreate_Name = "zeKernelCreate";
457
zeKernelCreateFT get_zeKernelCreate()
458
{
×
UNCOV
459
    static auto st_zeKernelCreateF =
×
UNCOV
460
        ze_loader::get().getSymbol<zeKernelCreateFT>(zeKernelCreate_Name);
×
461

UNCOV
462
    return st_zeKernelCreateF;
×
463
}
×
464

465
std::string _GetErrorCode_ze_impl(ze_result_t code)
UNCOV
466
{
×
467
    switch (code) {
×
468
        EnumCaseString(ZE_RESULT_ERROR_UNINITIALIZED);
×
UNCOV
469
        EnumCaseString(ZE_RESULT_ERROR_DEVICE_LOST);
×
UNCOV
470
        EnumCaseString(ZE_RESULT_ERROR_INVALID_NULL_HANDLE);
×
UNCOV
471
        EnumCaseString(ZE_RESULT_ERROR_INVALID_NULL_POINTER);
×
UNCOV
472
        EnumCaseString(ZE_RESULT_ERROR_INVALID_ENUMERATION);
×
UNCOV
473
        EnumCaseString(ZE_RESULT_ERROR_INVALID_NATIVE_BINARY);
×
UNCOV
474
        EnumCaseString(ZE_RESULT_ERROR_INVALID_SIZE);
×
475
        EnumCaseString(ZE_RESULT_ERROR_OUT_OF_HOST_MEMORY);
×
476
        EnumCaseString(ZE_RESULT_ERROR_OUT_OF_DEVICE_MEMORY);
×
477
        EnumCaseString(ZE_RESULT_ERROR_MODULE_BUILD_FAILURE);
×
UNCOV
478
        EnumCaseString(ZE_RESULT_ERROR_INVALID_MODULE_UNLINKED);
×
479
    default:
×
480
        return "<< UNRECOGNIZED ZE_RESULT_T CODE >> " + CodeStringSuffix(code);
×
UNCOV
481
    }
×
UNCOV
482
}
×
483

484
__dpctl_give DPCTLSyclKernelBundleRef
485
_CreateKernelBundleWithIL_ze_impl(const context &SyclCtx,
486
                                  const device &SyclDev,
487
                                  const void *IL,
488
                                  size_t il_length,
489
                                  const char *CompileOpts,
490
                                  size_t NumSpecConsts,
491
                                  const DPCTLSpecConst *SpecConsts)
492
{
×
493
    auto zeModuleCreateFn = get_zeModuleCreate();
×
494
    if (zeModuleCreateFn == nullptr) {
×
495
        error_handler("ZeModuleCreateFn is invalid.", __FILE__, __func__,
×
496
                      __LINE__);
×
497
        return nullptr;
×
498
    }
×
499

UNCOV
500
    backend_traits<ze_be>::return_type<context> ZeContext;
×
UNCOV
501
    ZeContext = get_native<ze_be>(SyclCtx);
×
502

UNCOV
503
    backend_traits<ze_be>::return_type<device> ZeDevice;
×
UNCOV
504
    ZeDevice = get_native<ze_be>(SyclDev);
×
505

NEW
506
    std::vector<std::uint32_t> spec_ids;
×
NEW
507
    std::vector<const void *> spec_values;
×
508

NEW
509
    if (SpecConsts != nullptr && NumSpecConsts > 0) {
×
NEW
510
        spec_ids.reserve(NumSpecConsts);
×
NEW
511
        spec_values.reserve(NumSpecConsts);
×
NEW
512
        for (size_t i = 0; i < NumSpecConsts; ++i) {
×
NEW
513
            spec_ids.push_back(SpecConsts[i].id);
×
NEW
514
            spec_values.push_back(SpecConsts[i].value);
×
NEW
515
        }
×
NEW
UNCOV
516
    }
×
517
    ze_module_constants_t ZeSpecConstants = {};
×
NEW
518
    ZeSpecConstants.numConstants = static_cast<std::uint32_t>(NumSpecConsts);
×
NEW
UNCOV
519
    ZeSpecConstants.pConstantIds = spec_ids.empty() ? nullptr : spec_ids.data();
×
NEW
520
    ZeSpecConstants.pConstantValues =
×
NEW
521
        spec_values.empty() ? nullptr : spec_values.data();
×
522

523
    // Populate the Level Zero module descriptions
524
    ze_module_desc_t ZeModuleDesc = {};
×
525
    ZeModuleDesc.stype = ZE_STRUCTURE_TYPE_MODULE_DESC;
×
526
    ZeModuleDesc.format = ZE_MODULE_FORMAT_IL_SPIRV;
×
527
    ZeModuleDesc.inputSize = il_length;
×
528
    ZeModuleDesc.pInputModule = (uint8_t *)IL;
×
529
    ZeModuleDesc.pBuildFlags = CompileOpts;
×
530
    ZeModuleDesc.pConstants = &ZeSpecConstants;
×
531

532
    ze_module_handle_t ZeModule;
×
533

534
    auto ret_code = zeModuleCreateFn(ZeContext, ZeDevice, &ZeModuleDesc,
×
535
                                     &ZeModule, nullptr);
×
536
    if (ret_code != ZE_RESULT_SUCCESS) {
×
537
        error_handler("Module creation failed " +
×
538
                          _GetErrorCode_ze_impl(ret_code),
×
UNCOV
539
                      __FILE__, __func__, __LINE__);
×
UNCOV
540
        return nullptr;
×
541
    }
×
542

543
    try {
×
544
        const auto &kb = make_kernel_bundle<ze_be, bundle_state::executable>(
×
545
            {ZeModule, ext::oneapi::level_zero::ownership::keep}, SyclCtx);
×
546

547
        return wrap<kernel_bundle<bundle_state::executable>>(
×
UNCOV
548
            new kernel_bundle<bundle_state::executable>(kb));
×
549
    } catch (std::exception const &e) {
×
UNCOV
550
        error_handler(e, __FILE__, __func__, __LINE__);
×
551
        auto zeModuleDestroyFn = get_zeModuleDestroy();
×
552
        if (zeModuleDestroyFn) {
×
553
            zeModuleDestroyFn(ZeModule);
×
554
        }
×
555
        return nullptr;
×
556
    }
×
557
}
×
558

559
__dpctl_give DPCTLSyclKernelRef
560
_GetKernel_ze_impl(const kernel_bundle<bundle_state::executable> &kb,
561
                   const char *kernel_name)
562
{
×
UNCOV
563
    auto zeKernelCreateFn = get_zeKernelCreate();
×
564
    if (zeKernelCreateFn == nullptr) {
×
565
        error_handler("Could not load zeKernelCreate function.", __FILE__,
×
566
                      __func__, __LINE__);
×
567
        return nullptr;
×
568
    }
×
569

570
    auto ZeKernelBundle = sycl::get_native<ze_be>(kb);
×
571
    bool found = false;
×
572

573
    // Populate the Level Zero kernel descriptions
574
    ze_kernel_desc_t ZeKernelDescr = {ZE_STRUCTURE_TYPE_KERNEL_DESC, nullptr,
×
UNCOV
575
                                      0, // flags
×
UNCOV
576
                                      kernel_name};
×
577

UNCOV
578
    std::unique_ptr<sycl::kernel> syclInteropKern_ptr;
×
579
    ze_kernel_handle_t ZeKern;
×
580
    for (auto &ZeM : ZeKernelBundle) {
×
581
        ze_result_t ze_status = zeKernelCreateFn(ZeM, &ZeKernelDescr, &ZeKern);
×
582

583
        if (ze_status == ZE_RESULT_SUCCESS) {
×
584
            found = true;
×
585
            const auto &ctx = kb.get_context();
×
UNCOV
586
            const auto &k = make_kernel<ze_be>(
×
587
                {kb, ZeKern, ext::oneapi::level_zero::ownership::keep}, ctx);
×
588
            syclInteropKern_ptr = std::unique_ptr<kernel>(new kernel(k));
×
UNCOV
589
            break;
×
UNCOV
590
        }
×
591
        else {
×
592
            if (ze_status != ZE_RESULT_ERROR_INVALID_KERNEL_NAME) {
×
593
                error_handler("zeKernelCreate failed: " +
×
UNCOV
594
                                  _GetErrorCode_ze_impl(ze_status),
×
595
                              __FILE__, __func__, __LINE__);
×
596
                return nullptr;
×
597
            }
×
598
        }
×
UNCOV
599
    }
×
600

601
    if (found) {
×
602
        return wrap<kernel>(new kernel(*syclInteropKern_ptr));
×
603
    }
×
604
    else {
×
605
        error_handler("Kernel named " + std::string(kernel_name) +
×
606
                          " could not be found.",
×
607
                      __FILE__, __func__, __LINE__, error_level::error);
×
608
        return nullptr;
×
609
    }
×
610
}
×
611

612
bool _HasKernel_ze_impl(const kernel_bundle<bundle_state::executable> &kb,
613
                        const char *kernel_name)
614
{
×
615
    auto zeKernelCreateFn = get_zeKernelCreate();
×
616
    if (zeKernelCreateFn == nullptr) {
×
UNCOV
617
        error_handler("Could not load zeKernelCreate function.", __FILE__,
×
618
                      __func__, __LINE__, error_level::error);
×
619
        return false;
×
620
    }
×
621

622
    auto ZeKernelBundle = sycl::get_native<ze_be>(kb);
×
623

624
    // Populate the Level Zero kernel descriptions
625
    ze_kernel_desc_t ZeKernelDescr = {ZE_STRUCTURE_TYPE_KERNEL_DESC, nullptr,
×
626
                                      0, // flags
×
627
                                      kernel_name};
×
628

UNCOV
629
    std::unique_ptr<sycl::kernel> syclInteropKern_ptr;
×
UNCOV
630
    ze_kernel_handle_t ZeKern;
×
631
    for (auto &ZeM : ZeKernelBundle) {
×
632
        ze_result_t ze_status = zeKernelCreateFn(ZeM, &ZeKernelDescr, &ZeKern);
×
633

634
        if (ze_status == ZE_RESULT_SUCCESS) {
×
635
            return true;
×
636
        }
×
637
        else {
×
UNCOV
638
            if (ze_status != ZE_RESULT_ERROR_INVALID_KERNEL_NAME) {
×
639
                error_handler("zeKernelCreate failed: " +
×
UNCOV
640
                                  _GetErrorCode_ze_impl(ze_status),
×
UNCOV
641
                              __FILE__, __func__, __LINE__, error_level::error);
×
642
                return false;
×
643
            }
×
644
        }
×
UNCOV
645
    }
×
646

647
    return false;
×
648
}
×
649

650
#endif /* #ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION */
651

652
} /* end of anonymous namespace */
653

654
__dpctl_give DPCTLSyclKernelBundleRef
655
DPCTLKernelBundle_CreateFromSpirv(__dpctl_keep const DPCTLSyclContextRef CtxRef,
656
                                  __dpctl_keep const DPCTLSyclDeviceRef DevRef,
657
                                  __dpctl_keep const void *IL,
658
                                  size_t length,
659
                                  const char *CompileOpts,
660
                                  size_t NumSpecConsts,
661
                                  const DPCTLSpecConst *SpecConsts)
662
{
75✔
663
    DPCTLSyclKernelBundleRef KBRef = nullptr;
75✔
664
    if (!CtxRef) {
75✔
665
        error_handler("Cannot create program from SPIR-V as the supplied SYCL "
3✔
666
                      "context is NULL.",
3✔
667
                      __FILE__, __func__, __LINE__);
3✔
668
        return KBRef;
3✔
669
    }
3✔
670
    if (!DevRef) {
72✔
671
        error_handler("Cannot create program from SPIR-V as the supplied SYCL "
3✔
672
                      "device is NULL.",
3✔
673
                      __FILE__, __func__, __LINE__);
3✔
674
        return KBRef;
3✔
675
    }
3✔
676
    if ((!IL) || (length == 0)) {
69!
677
        error_handler("Cannot create program from null SPIR-V buffer.",
3✔
678
                      __FILE__, __func__, __LINE__);
3✔
679
        return KBRef;
3✔
680
    }
3✔
681

682
    context *SyclCtx = unwrap<context>(CtxRef);
66✔
683
    device *SyclDev = unwrap<device>(DevRef);
66✔
684
    // get the backend type
685
    auto BE = SyclCtx->get_platform().get_backend();
66✔
686
    switch (BE) {
66✔
687
    case backend::opencl:
66!
688
        KBRef = _CreateKernelBundleWithIL_ocl_impl(*SyclCtx, *SyclDev, IL,
66✔
689
                                                   length, CompileOpts,
66✔
690
                                                   NumSpecConsts, SpecConsts);
66✔
691
        break;
66✔
UNCOV
692
    case backend::ext_oneapi_level_zero:
×
UNCOV
693
#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION
×
UNCOV
694
        KBRef = _CreateKernelBundleWithIL_ze_impl(*SyclCtx, *SyclDev, IL,
×
NEW
UNCOV
695
                                                  length, CompileOpts,
×
NEW
UNCOV
696
                                                  NumSpecConsts, SpecConsts);
×
UNCOV
697
        break;
×
UNCOV
698
#endif
×
UNCOV
699
    default:
×
UNCOV
700
        std::ostringstream os;
×
UNCOV
701
        os << "Backend " << BE << " is not supported";
×
UNCOV
702
        error_handler(os.str(), __FILE__, __func__, __LINE__);
×
UNCOV
703
        break;
×
704
    }
66✔
705
    return KBRef;
66✔
706
}
66✔
707

708
__dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromOCLSource(
709
    __dpctl_keep const DPCTLSyclContextRef Ctx,
710
    __dpctl_keep const DPCTLSyclDeviceRef Dev,
711
    __dpctl_keep const char *Source,
712
    __dpctl_keep const char *CompileOpts)
713
{
39✔
714
    context *SyclCtx = nullptr;
39✔
715
    device *SyclDev = nullptr;
39✔
716

717
    if (!Ctx) {
39✔
718
        error_handler("Input Ctx is nullptr.", __FILE__, __func__, __LINE__);
10✔
719
        return nullptr;
10✔
720
    }
10✔
721
    if (!Dev) {
29✔
722
        error_handler("Input Dev is nullptr.", __FILE__, __func__, __LINE__);
1✔
723
        return nullptr;
1✔
724
    }
1✔
725
    if (!Source) {
28!
UNCOV
726
        error_handler("Input Source is nullptr.", __FILE__, __func__, __LINE__);
×
UNCOV
727
        return nullptr;
×
UNCOV
728
    }
×
729

730
    SyclCtx = unwrap<context>(Ctx);
28✔
731
    SyclDev = unwrap<device>(Dev);
28✔
732

733
    // get the backend type
734
    auto BE = SyclCtx->get_platform().get_backend();
28✔
735
    switch (BE) {
28✔
736
    case backend::opencl:
28!
737
        try {
28✔
738
            return _CreateKernelBundleWithOCLSource_ocl_impl(
28✔
739
                *SyclCtx, *SyclDev, Source, CompileOpts);
28✔
740
        } catch (std::exception const &e) {
28✔
UNCOV
741
            error_handler(e, __FILE__, __func__, __LINE__);
×
UNCOV
742
            return nullptr;
×
743
        }
×
744
        break;
×
745
    case backend::ext_oneapi_level_zero:
×
UNCOV
746
        error_handler(
×
UNCOV
747
            "CreateFromSource is not supported for Level Zero backend.",
×
UNCOV
748
            __FILE__, __func__, __LINE__);
×
UNCOV
749
        return nullptr;
×
UNCOV
750
    default:
×
UNCOV
751
        error_handler("CreateFromSource is not supported in unknown backend.",
×
UNCOV
752
                      __FILE__, __func__, __LINE__);
×
UNCOV
753
        return nullptr;
×
754
    }
28✔
755
}
28✔
756

757
__dpctl_give DPCTLSyclKernelRef
758
DPCTLKernelBundle_GetKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef,
759
                            __dpctl_keep const char *KernelName)
760
{
115✔
761
    if (!KBRef) {
115✔
762
        error_handler("Input KBRef is nullptr", __FILE__, __func__, __LINE__);
21✔
763
        return nullptr;
21✔
764
    }
21✔
765
    if (!KernelName) {
94✔
766
        error_handler("Input KernelName is nullptr", __FILE__, __func__,
3✔
767
                      __LINE__);
3✔
768
        return nullptr;
3✔
769
    }
3✔
770
    auto SyclKB = unwrap<kernel_bundle<bundle_state::executable>>(KBRef);
91✔
771
    sycl::backend be = SyclKB->get_backend();
91✔
772
    switch (be) {
91✔
773
    case sycl::backend::opencl:
91!
774
        return _GetKernel_ocl_impl(*SyclKB, KernelName);
91✔
UNCOV
775
    case sycl::backend::ext_oneapi_level_zero:
×
UNCOV
776
#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION
×
UNCOV
777
        return _GetKernel_ze_impl(*SyclKB, KernelName);
×
UNCOV
778
#endif
×
UNCOV
779
    default:
×
UNCOV
780
        std::ostringstream os;
×
UNCOV
781
        os << "Backend " << be << " is not supported";
×
UNCOV
782
        error_handler(os.str(), __FILE__, __func__, __LINE__);
×
UNCOV
783
        return nullptr;
×
784
    }
91✔
785
}
91✔
786

787
bool DPCTLKernelBundle_HasKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef,
788
                                 __dpctl_keep const char *KernelName)
789
{
68✔
790
    if (!KBRef) {
68✔
791
        error_handler("Input KBRef is nullptr", __FILE__, __func__, __LINE__);
3✔
792
        return false;
3✔
793
    }
3✔
794
    if (!KernelName) {
65✔
795
        error_handler("Input KernelName is nullptr", __FILE__, __func__,
3✔
796
                      __LINE__);
3✔
797
        return false;
3✔
798
    }
3✔
799

800
    auto SyclKB = unwrap<kernel_bundle<bundle_state::executable>>(KBRef);
62✔
801
    sycl::backend be = SyclKB->get_backend();
62✔
802
    switch (be) {
62✔
803
    case sycl::backend::opencl:
62!
804
        return _HasKernel_ocl_impl(*SyclKB, KernelName);
62✔
UNCOV
805
    case sycl::backend::ext_oneapi_level_zero:
×
UNCOV
806
#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION
×
UNCOV
807
        return _HasKernel_ze_impl(*SyclKB, KernelName);
×
UNCOV
808
#endif
×
UNCOV
809
    default:
×
UNCOV
810
        std::ostringstream os;
×
UNCOV
811
        os << "Backend " << be << " is not supported";
×
UNCOV
812
        error_handler(os.str(), __FILE__, __func__, __LINE__);
×
UNCOV
813
        return false;
×
814
    }
62✔
815
}
62✔
816

817
void DPCTLKernelBundle_Delete(__dpctl_take DPCTLSyclKernelBundleRef KBRef)
818
{
106✔
819
    delete unwrap<kernel_bundle<bundle_state::executable>>(KBRef);
106✔
820
}
106✔
821

822
__dpctl_give DPCTLSyclKernelBundleRef
823
DPCTLKernelBundle_Copy(__dpctl_keep const DPCTLSyclKernelBundleRef KBRef)
824
{
8✔
825
    auto Bundle = unwrap<kernel_bundle<bundle_state::executable>>(KBRef);
8✔
826
    if (!Bundle) {
8✔
827
        error_handler(
3✔
828
            "Cannot copy DPCTLSyclKernelBundleRef as input is a nullptr",
3✔
829
            __FILE__, __func__, __LINE__);
3✔
830
        return nullptr;
3✔
831
    }
3✔
832
    try {
5✔
833
        auto CopiedBundle =
5✔
834
            new kernel_bundle<bundle_state::executable>(*Bundle);
5✔
835
        return wrap<kernel_bundle<bundle_state::executable>>(CopiedBundle);
5✔
836
    } catch (std::exception const &e) {
5✔
UNCOV
837
        error_handler(e, __FILE__, __func__, __LINE__);
×
UNCOV
838
        return nullptr;
×
UNCOV
839
    }
×
840
}
5✔
STATUS · Troubleshooting · Open an Issue · Sales · Support · CAREERS · ENTERPRISE · START FREE · SCHEDULE DEMO
ANNOUNCEMENTS · TWITTER · TOS & SLA · Supported CI Services · What's a CI service? · Automated Testing

© 2026 Coveralls, Inc