31424476294

Committed 10 Aug 2026 07:30PM UTC coverage: 74.344% (+0.3%) from 74.023%

Build # 31424476294

Build Type

Pull #2206

github

Committed by

web-flow

Commit Message

Merge cdcef0d86 into 6c5d2deda

Pull Request Pull Request #2206: Support SYCL source compilation

Coverage Stats

992 of 1404 branches covered (70.66%)

Branch coverage included in aggregate %.

184 of 207 new or added lines in 2 files covered. (88.89%)

14 existing lines in 1 file now uncovered.

3827 of 5078 relevant lines covered (75.36%)

277.21 hits per line

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

56.18

/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp

//===- dpctl_sycl_kernel_bundle_interface.cpp - Implements C API for
//    sycl::kernel_bundle<sycl::bundle_state::executable>  ---------------===//
//
//                      Data Parallel Control (dpctl)
//
// Copyright 2020 Intel Corporation
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
//    http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
//
//===----------------------------------------------------------------------===//
///
/// \file
/// This file implements the functions declared in
/// dpctl_sycl_kernel_bundle_interface.h.
///
//===----------------------------------------------------------------------===//

#include "dpctl_sycl_kernel_bundle_interface.h"
#include "Config/dpctl_config.h"
#include "dpctl_dynamic_lib_helper.h"
#include "dpctl_error_handlers.h"
#include "dpctl_sycl_type_casters.hpp"
#include <CL/cl.h> /* OpenCL headers     */
#include <cstdint>
#include <sstream>
#include <stddef.h>
#include <sycl/backend/opencl.hpp>
#include <sycl/sycl.hpp> /* Sycl headers       */
#include <utility>

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

using namespace sycl;

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

using namespace dpctl::syclinterface;

#ifdef __linux__
static const char *clLoaderName = DPCTL_LIBCL_LOADER_FILENAME;
static const int clLibLoadFlags = RTLD_NOLOAD | RTLD_NOW | RTLD_LOCAL;
#elif defined(_WIN64)
static const char *clLoaderName = "OpenCL.dll";
static const int clLibLoadFlags = 0;
#else
#error "OpenCL program compilation is unavailable for this platform"
#endif

#define CodeStringSuffix(code)                                                 \
    std::string(" (code=") + std::to_string(static_cast<int>(code)) + ")"

#define EnumCaseString(code)                                                   \
    case code:                                                                 \
        return std::string(#code) + CodeStringSuffix(code)

static constexpr backend cl_be = backend::opencl;

struct cl_loader
{
public:
    static cl_loader &get()
    {
        static cl_loader _loader;
        return _loader;
    }

    template <typename retTy> retTy getSymbol(const char *name)
    {
        if (!opened) {
            error_handler("The OpenCL loader dynamic library could not "
                          "be opened.",
                          __FILE__, __func__, __LINE__);

            return nullptr;
        }
        return clLib.getSymbol<retTy>(name);
    }

private:
    dpctl::DynamicLibHelper clLib;
    bool opened;
    cl_loader() : clLib(clLoaderName, clLibLoadFlags), opened(clLib.opened()) {}
};

typedef cl_program (*clCreateProgramWithSourceFT)(cl_context,
                                                  cl_uint,
                                                  const char **,
                                                  const size_t *,
                                                  cl_int *);
const char *clCreateProgramWithSource_Name = "clCreateProgramWithSource";
clCreateProgramWithSourceFT get_clCreateProgramWithSource()
{
    static auto st_clCreateProgramWithSourceF =
        cl_loader::get().getSymbol<clCreateProgramWithSourceFT>(
            clCreateProgramWithSource_Name);

    return st_clCreateProgramWithSourceF;
}

typedef cl_program (*clCreateProgramWithILFT)(cl_context,
                                              const void *,
                                              size_t,
                                              cl_int *);
const char *clCreateProgramWithIL_Name = "clCreateProgramWithIL";
clCreateProgramWithILFT get_clCreateProgramWithIL()
{
    static auto st_clCreateProgramWithILF =
        cl_loader::get().getSymbol<clCreateProgramWithILFT>(
            clCreateProgramWithIL_Name);

    return st_clCreateProgramWithILF;
}
typedef cl_int (*clBuildProgramFT)(cl_program,
                                   cl_uint,
                                   const cl_device_id *,
                                   const char *,
                                   void (*)(cl_program, void *),
                                   void *);
const char *clBuildProgram_Name = "clBuildProgram";
clBuildProgramFT get_clBuildProgram()
{
    static auto st_clBuildProgramF =
        cl_loader::get().getSymbol<clBuildProgramFT>(clBuildProgram_Name);

    return st_clBuildProgramF;
}

typedef cl_kernel (*clCreateKernelFT)(cl_program, const char *, cl_int *);
const char *clCreateKernel_Name = "clCreateKernel";
clCreateKernelFT get_clCreateKernel()
{
    static auto st_clCreateKernelF =
        cl_loader::get().getSymbol<clCreateKernelFT>(clCreateKernel_Name);

    return st_clCreateKernelF;
}

std::string _GetErrorCode_ocl_impl(cl_int code)
{
    switch (code) {
        EnumCaseString(CL_BUILD_PROGRAM_FAILURE);
        EnumCaseString(CL_INVALID_CONTEXT);
        EnumCaseString(CL_INVALID_DEVICE);
        EnumCaseString(CL_INVALID_VALUE);
        EnumCaseString(CL_OUT_OF_RESOURCES);
        EnumCaseString(CL_OUT_OF_HOST_MEMORY);
        EnumCaseString(CL_INVALID_OPERATION);
        EnumCaseString(CL_INVALID_BINARY);
    default:
        return "<< ERROR CODE UNRECOGNIZED >>" + CodeStringSuffix(code);
    }
}

typedef cl_int (*clSetProgramSpecializationConstantFT)(cl_program,
                                                       cl_uint,
                                                       size_t,
                                                       const void *);
const char *clSetProgramSpecializationConstant_Name =
    "clSetProgramSpecializationConstant";
clSetProgramSpecializationConstantFT get_clSetProgramSpecializationConstant()
{
    static auto st_clSetProgramSpecializationConstantF =
        cl_loader::get().getSymbol<clSetProgramSpecializationConstantFT>(
            clSetProgramSpecializationConstant_Name);

    return st_clSetProgramSpecializationConstantF;
}

typedef cl_int (*clReleaseProgramFT)(cl_program);
const char *clReleaseProgram_Name = "clReleaseProgram";
clReleaseProgramFT get_clReleaseProgram()
{
    static auto st_clReleaseProgramF =
        cl_loader::get().getSymbol<clReleaseProgramFT>(clReleaseProgram_Name);

    return st_clReleaseProgramF;
}

DPCTLSyclKernelBundleRef
_CreateKernelBundle_common_ocl_impl(cl_program clProgram,
                                    const context &ctx,
                                    const device &dev,
                                    const char *CompileOpts)
{
    backend_traits<cl_be>::return_type<device> clDevice;
    clDevice = get_native<cl_be>(dev);

    // Last two pointers are notification function pointer and user-data pointer
    // that can be passed to the notification function.
    auto clBuildProgramF = get_clBuildProgram();
    if (clBuildProgramF == nullptr) {
        auto clReleaseProgramF = get_clReleaseProgram();
        if (clReleaseProgramF) {
            clReleaseProgramF(clProgram);
        }
        return nullptr;
    }
    cl_int build_status =
        clBuildProgramF(clProgram, 1, &clDevice, CompileOpts, nullptr, nullptr);

    if (build_status != CL_SUCCESS) {
        error_handler("clBuildProgram failed: " +
                          _GetErrorCode_ocl_impl(build_status),
                      __FILE__, __func__, __LINE__);
        auto clReleaseProgramF = get_clReleaseProgram();
        if (clReleaseProgramF) {
            clReleaseProgramF(clProgram);
        }
        return nullptr;
    }

    using ekbTy = kernel_bundle<bundle_state::executable>;
    const ekbTy &kb =
        make_kernel_bundle<cl_be, bundle_state::executable>(clProgram, ctx);
    return wrap<ekbTy>(new ekbTy(kb));
}

DPCTLSyclKernelBundleRef
_CreateKernelBundleWithOCLSource_ocl_impl(const context &ctx,
                                          const device &dev,
                                          const char *oclSrc,
                                          const char *CompileOpts)
{
    auto clCreateProgramWithSourceF = get_clCreateProgramWithSource();
    if (clCreateProgramWithSourceF == nullptr) {
        return nullptr;
    }

    backend_traits<cl_be>::return_type<context> clContext;
    clContext = get_native<cl_be>(ctx);

    cl_int build_with_source_err_code = CL_SUCCESS;
    cl_program clProgram = clCreateProgramWithSourceF(
        clContext, 1, &oclSrc, nullptr, &build_with_source_err_code);

    if (build_with_source_err_code != CL_SUCCESS) {
        error_handler("clPCreateProgramWithSource failed with " +
                          _GetErrorCode_ocl_impl(build_with_source_err_code),
                      __FILE__, __func__, __LINE__);
        return nullptr;
    }

    return _CreateKernelBundle_common_ocl_impl(clProgram, ctx, dev,
                                               CompileOpts);
}

DPCTLSyclKernelBundleRef
_CreateKernelBundleWithIL_ocl_impl(const context &ctx,
                                   const device &dev,
                                   const void *IL,
                                   size_t il_length,
                                   const char *CompileOpts,
                                   size_t NumSpecConsts,
                                   const DPCTLSpecConst *SpecConsts)
{
    auto clCreateProgramWithILF = get_clCreateProgramWithIL();
    if (clCreateProgramWithILF == nullptr) {
        return nullptr;
    }

    backend_traits<cl_be>::return_type<context> clContext;
    clContext = get_native<cl_be>(ctx);

    if (NumSpecConsts > 0 && SpecConsts == nullptr) {
        error_handler("NumSpecConsts > 0 but SpecConsts is NULL.", __FILE__,
                      __func__, __LINE__);
        return nullptr;
    }

    cl_int create_err_code = CL_SUCCESS;
    cl_program clProgram =
        clCreateProgramWithILF(clContext, IL, il_length, &create_err_code);

    if (create_err_code != CL_SUCCESS) {
        error_handler("OpenCL program could not be created from the SPIR-V "
                      "binary. OpenCL Error " +
                          _GetErrorCode_ocl_impl(create_err_code),
                      __FILE__, __func__, __LINE__);
        return nullptr;
    }

    if (SpecConsts != nullptr && NumSpecConsts > 0) {
        auto clSetProgramSpecConstF = get_clSetProgramSpecializationConstant();
        if (clSetProgramSpecConstF) {
            for (size_t i = 0; i < NumSpecConsts; ++i) {
                cl_int spec_err = clSetProgramSpecConstF(
                    clProgram, SpecConsts[i].id, SpecConsts[i].size,
                    SpecConsts[i].value);
                if (spec_err != CL_SUCCESS) {
                    error_handler(
                        "clSetProgramSpecializationConstant failed for "
                        "spec constant id " +
                            std::to_string(SpecConsts[i].id) +
                            ". OpenCL Error " +
                            _GetErrorCode_ocl_impl(spec_err),
                        __FILE__, __func__, __LINE__);

                    auto clReleaseProgramF = get_clReleaseProgram();
                    if (clReleaseProgramF) {
                        clReleaseProgramF(clProgram);
                    }

                    return nullptr;
                }
            }
        }
        else {
            error_handler("clSetProgramSpecializationConstant is not available "
                          "in the OpenCL implementation.",
                          __FILE__, __func__, __LINE__);

            auto clReleaseProgramF = get_clReleaseProgram();
            if (clReleaseProgramF) {
                clReleaseProgramF(clProgram);
            }

            return nullptr;
        }
    }

    return _CreateKernelBundle_common_ocl_impl(clProgram, ctx, dev,
                                               CompileOpts);
}

bool _HasKernel_ocl_impl(const kernel_bundle<bundle_state::executable> &kb,
                         const char *kernel_name)
{
    auto clCreateKernelF = get_clCreateKernel();
    if (clCreateKernelF == nullptr) {
        return false;
    }

    std::vector<cl_program> oclKB = get_native<cl_be>(kb);

    bool found = false;
    for (auto &cl_pr : oclKB) {
        cl_int create_kernel_err_code = CL_SUCCESS;
        [[maybe_unused]] cl_kernel try_kern =
            clCreateKernelF(cl_pr, kernel_name, &create_kernel_err_code);
        if (create_kernel_err_code == CL_SUCCESS) {
            found = true;
            break;
        }
    }
    return found;
}

__dpctl_give DPCTLSyclKernelRef
_GetKernel_ocl_impl(const kernel_bundle<bundle_state::executable> &kb,
                    const char *kernel_name)
{
    auto clCreateKernelF = get_clCreateKernel();
    if (clCreateKernelF == nullptr) {
        return nullptr;
    }

    std::vector<cl_program> oclKB = get_native<cl_be>(kb);

    bool found = false;
    cl_kernel ocl_kernel_from_kb;
    for (auto &cl_pr : oclKB) {
        cl_int create_kernel_err_code = CL_SUCCESS;
        cl_kernel try_kern =
            clCreateKernelF(cl_pr, kernel_name, &create_kernel_err_code);
        if (create_kernel_err_code == CL_SUCCESS) {
            found = true;
            ocl_kernel_from_kb = try_kern;
            break;
        }
    }
    if (found) {
        try {
            context ctx = kb.get_context();

            const kernel &interop_kernel =
                make_kernel<cl_be>(ocl_kernel_from_kb, ctx);

            return wrap<kernel>(new kernel(interop_kernel));
        } catch (std::exception const &e) {
            error_handler(e, __FILE__, __func__, __LINE__);
            return nullptr;
        }
    }
    else {
        error_handler("Kernel " + std::string(kernel_name) + " not found.",
                      __FILE__, __func__, __LINE__);
        return nullptr;
    }
}

#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION

#ifdef __linux__
static const char *zeLoaderName = DPCTL_LIBZE_LOADER_FILENAME;
static const int zeLibLoadFlags = RTLD_NOLOAD | RTLD_NOW | RTLD_LOCAL;
#elif defined(_WIN64)
static const char *zeLoaderName = "ze_loader.dll";
static const int zeLibLoadFlags = 0;
#else
#error "Level Zero program compilation is unavailable for this platform"
#endif

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

struct ze_loader
{
public:
    static ze_loader &get()
    {
        static ze_loader _loader;
        return _loader;
    }

    template <typename retTy> retTy getSymbol(const char *name)
    {
        if (!opened) {
            error_handler("The Level-Zero loader dynamic library could not "
                          "be opened.",
                          __FILE__, __func__, __LINE__);

            return nullptr;
        }
        return zeLib.getSymbol<retTy>(name);
    }

private:
    dpctl::DynamicLibHelper zeLib;
    bool opened;
    ze_loader() : zeLib(zeLoaderName, zeLibLoadFlags), opened(zeLib.opened()) {}
};

typedef ze_result_t (*zeModuleCreateFT)(ze_context_handle_t,
                                        ze_device_handle_t,
                                        const ze_module_desc_t *,
                                        ze_module_handle_t *,
                                        ze_module_build_log_handle_t *);
const char *zeModuleCreate_Name = "zeModuleCreate";
zeModuleCreateFT get_zeModuleCreate()
{
    static auto st_zeModuleCreateF =
        ze_loader::get().getSymbol<zeModuleCreateFT>(zeModuleCreate_Name);

    return st_zeModuleCreateF;
}

typedef ze_result_t (*zeModuleDestroyFT)(ze_module_handle_t);
const char *zeModuleDestroy_Name = "zeModuleDestroy";
zeModuleDestroyFT get_zeModuleDestroy()
{
    static auto st_zeModuleDestroyF =
        ze_loader::get().getSymbol<zeModuleDestroyFT>(zeModuleDestroy_Name);

    return st_zeModuleDestroyF;
}

typedef ze_result_t (*zeKernelCreateFT)(ze_module_handle_t,
                                        const ze_kernel_desc_t *,
                                        ze_kernel_handle_t *);
const char *zeKernelCreate_Name = "zeKernelCreate";
zeKernelCreateFT get_zeKernelCreate()
{
    static auto st_zeKernelCreateF =
        ze_loader::get().getSymbol<zeKernelCreateFT>(zeKernelCreate_Name);

    return st_zeKernelCreateF;
}

std::string _GetErrorCode_ze_impl(ze_result_t code)
{
    switch (code) {
        EnumCaseString(ZE_RESULT_ERROR_UNINITIALIZED);
        EnumCaseString(ZE_RESULT_ERROR_DEVICE_LOST);
        EnumCaseString(ZE_RESULT_ERROR_INVALID_NULL_HANDLE);
        EnumCaseString(ZE_RESULT_ERROR_INVALID_NULL_POINTER);
        EnumCaseString(ZE_RESULT_ERROR_INVALID_ENUMERATION);
        EnumCaseString(ZE_RESULT_ERROR_INVALID_NATIVE_BINARY);
        EnumCaseString(ZE_RESULT_ERROR_INVALID_SIZE);
        EnumCaseString(ZE_RESULT_ERROR_OUT_OF_HOST_MEMORY);
        EnumCaseString(ZE_RESULT_ERROR_OUT_OF_DEVICE_MEMORY);
        EnumCaseString(ZE_RESULT_ERROR_MODULE_BUILD_FAILURE);
        EnumCaseString(ZE_RESULT_ERROR_INVALID_MODULE_UNLINKED);
    default:
        return "<< UNRECOGNIZED ZE_RESULT_T CODE >> " + CodeStringSuffix(code);
    }
}

__dpctl_give DPCTLSyclKernelBundleRef
_CreateKernelBundleWithIL_ze_impl(const context &SyclCtx,
                                  const device &SyclDev,
                                  const void *IL,
                                  size_t il_length,
                                  const char *CompileOpts,
                                  size_t NumSpecConsts,
                                  const DPCTLSpecConst *SpecConsts)
{
    auto zeModuleCreateFn = get_zeModuleCreate();
    if (zeModuleCreateFn == nullptr) {
        error_handler("ZeModuleCreateFn is invalid.", __FILE__, __func__,
                      __LINE__);
        return nullptr;
    }

    backend_traits<ze_be>::return_type<context> ZeContext;
    ZeContext = get_native<ze_be>(SyclCtx);

    backend_traits<ze_be>::return_type<device> ZeDevice;
    ZeDevice = get_native<ze_be>(SyclDev);

    if (NumSpecConsts > 0 && SpecConsts == nullptr) {
        error_handler("NumSpecConsts > 0 but SpecConsts is NULL.", __FILE__,
                      __func__, __LINE__);
        return nullptr;
    }

    std::vector<std::uint32_t> spec_ids;
    std::vector<const void *> spec_values;

    if (SpecConsts != nullptr && NumSpecConsts > 0) {
        spec_ids.reserve(NumSpecConsts);
        spec_values.reserve(NumSpecConsts);
        for (size_t i = 0; i < NumSpecConsts; ++i) {
            spec_ids.push_back(SpecConsts[i].id);
            spec_values.push_back(SpecConsts[i].value);
        }
    }
    ze_module_constants_t ZeSpecConstants = {};
    ZeSpecConstants.numConstants = static_cast<std::uint32_t>(spec_ids.size());
    ZeSpecConstants.pConstantIds = spec_ids.empty() ? nullptr : spec_ids.data();
    ZeSpecConstants.pConstantValues =
        spec_values.empty() ? nullptr : spec_values.data();

    // Populate the Level Zero module descriptions
    ze_module_desc_t ZeModuleDesc = {};
    ZeModuleDesc.stype = ZE_STRUCTURE_TYPE_MODULE_DESC;
    ZeModuleDesc.format = ZE_MODULE_FORMAT_IL_SPIRV;
    ZeModuleDesc.inputSize = il_length;
    ZeModuleDesc.pInputModule = (uint8_t *)IL;
    ZeModuleDesc.pBuildFlags = CompileOpts;
    ZeModuleDesc.pConstants = &ZeSpecConstants;

    ze_module_handle_t ZeModule;

    auto ret_code = zeModuleCreateFn(ZeContext, ZeDevice, &ZeModuleDesc,
                                     &ZeModule, nullptr);
    if (ret_code != ZE_RESULT_SUCCESS) {
        error_handler("Module creation failed " +
                          _GetErrorCode_ze_impl(ret_code),
                      __FILE__, __func__, __LINE__);
        return nullptr;
    }

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

        return wrap<kernel_bundle<bundle_state::executable>>(
            new kernel_bundle<bundle_state::executable>(kb));
    } catch (std::exception const &e) {
        error_handler(e, __FILE__, __func__, __LINE__);
        auto zeModuleDestroyFn = get_zeModuleDestroy();
        if (zeModuleDestroyFn) {
            zeModuleDestroyFn(ZeModule);
        }
        return nullptr;
    }
}

__dpctl_give DPCTLSyclKernelRef
_GetKernel_ze_impl(const kernel_bundle<bundle_state::executable> &kb,
                   const char *kernel_name)
{
    auto zeKernelCreateFn = get_zeKernelCreate();
    if (zeKernelCreateFn == nullptr) {
        error_handler("Could not load zeKernelCreate function.", __FILE__,
                      __func__, __LINE__);
        return nullptr;
    }

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

    // Populate the Level Zero kernel descriptions
    ze_kernel_desc_t ZeKernelDescr = {ZE_STRUCTURE_TYPE_KERNEL_DESC, nullptr,
                                      0, // flags
                                      kernel_name};

    std::unique_ptr<sycl::kernel> syclInteropKern_ptr;
    ze_kernel_handle_t ZeKern;
    for (auto &ZeM : ZeKernelBundle) {
        ze_result_t ze_status = zeKernelCreateFn(ZeM, &ZeKernelDescr, &ZeKern);

        if (ze_status == ZE_RESULT_SUCCESS) {
            found = true;
            const auto &ctx = kb.get_context();
            const auto &k = make_kernel<ze_be>(
                {kb, ZeKern, ext::oneapi::level_zero::ownership::keep}, ctx);
            syclInteropKern_ptr = std::unique_ptr<kernel>(new kernel(k));
            break;
        }
        else {
            if (ze_status != ZE_RESULT_ERROR_INVALID_KERNEL_NAME) {
                error_handler("zeKernelCreate failed: " +
                                  _GetErrorCode_ze_impl(ze_status),
                              __FILE__, __func__, __LINE__);
                return nullptr;
            }
        }
    }

    if (found) {
        return wrap<kernel>(new kernel(*syclInteropKern_ptr));
    }
    else {
        error_handler("Kernel named " + std::string(kernel_name) +
                          " could not be found.",
                      __FILE__, __func__, __LINE__, error_level::error);
        return nullptr;
    }
}

bool _HasKernel_ze_impl(const kernel_bundle<bundle_state::executable> &kb,
                        const char *kernel_name)
{
    auto zeKernelCreateFn = get_zeKernelCreate();
    if (zeKernelCreateFn == nullptr) {
        error_handler("Could not load zeKernelCreate function.", __FILE__,
                      __func__, __LINE__, error_level::error);
        return false;
    }

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

    // Populate the Level Zero kernel descriptions
    ze_kernel_desc_t ZeKernelDescr = {ZE_STRUCTURE_TYPE_KERNEL_DESC, nullptr,
                                      0, // flags
                                      kernel_name};

    std::unique_ptr<sycl::kernel> syclInteropKern_ptr;
    ze_kernel_handle_t ZeKern;
    for (auto &ZeM : ZeKernelBundle) {
        ze_result_t ze_status = zeKernelCreateFn(ZeM, &ZeKernelDescr, &ZeKern);

        if (ze_status == ZE_RESULT_SUCCESS) {
            return true;
        }
        else {
            if (ze_status != ZE_RESULT_ERROR_INVALID_KERNEL_NAME) {
                error_handler("zeKernelCreate failed: " +
                                  _GetErrorCode_ze_impl(ze_status),
                              __FILE__, __func__, __LINE__, error_level::error);
                return false;
            }
        }
    }

    return false;
}

#endif /* #ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION */

} /* end of anonymous namespace */

__dpctl_give DPCTLSyclKernelBundleRef
DPCTLKernelBundle_CreateFromSpirv(__dpctl_keep const DPCTLSyclContextRef CtxRef,
                                  __dpctl_keep const DPCTLSyclDeviceRef DevRef,
                                  __dpctl_keep const void *IL,
                                  size_t length,
                                  const char *CompileOpts,
                                  size_t NumSpecConsts,
                                  const DPCTLSpecConst *SpecConsts)
{
    DPCTLSyclKernelBundleRef KBRef = nullptr;
    if (!CtxRef) {
        error_handler("Cannot create program from SPIR-V as the supplied SYCL "
                      "context is NULL.",
                      __FILE__, __func__, __LINE__);
        return KBRef;
    }
    if (!DevRef) {
        error_handler("Cannot create program from SPIR-V as the supplied SYCL "
                      "device is NULL.",
                      __FILE__, __func__, __LINE__);
        return KBRef;
    }
    if ((!IL) || (length == 0)) {
        error_handler("Cannot create program from null SPIR-V buffer.",
                      __FILE__, __func__, __LINE__);
        return KBRef;
    }

    context *SyclCtx = unwrap<context>(CtxRef);
    device *SyclDev = unwrap<device>(DevRef);
    // get the backend type
    try {
        auto BE = SyclCtx->get_platform().get_backend();
        switch (BE) {
        case backend::opencl:
            KBRef = _CreateKernelBundleWithIL_ocl_impl(
                *SyclCtx, *SyclDev, IL, length, CompileOpts, NumSpecConsts,
                SpecConsts);
            break;
        case backend::ext_oneapi_level_zero:
#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION
            KBRef = _CreateKernelBundleWithIL_ze_impl(
                *SyclCtx, *SyclDev, IL, length, CompileOpts, NumSpecConsts,
                SpecConsts);
            break;
#endif
        default:
            std::ostringstream os;
            os << "Backend " << BE << " is not supported";
            error_handler(os.str(), __FILE__, __func__, __LINE__);
            break;
        }
    } catch (std::exception const &e) {
        error_handler(e, __FILE__, __func__, __LINE__);
        return nullptr;
    }
    return KBRef;
}

__dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromOCLSource(
    __dpctl_keep const DPCTLSyclContextRef Ctx,
    __dpctl_keep const DPCTLSyclDeviceRef Dev,
    __dpctl_keep const char *Source,
    __dpctl_keep const char *CompileOpts)
{
    context *SyclCtx = nullptr;
    device *SyclDev = nullptr;

    if (!Ctx) {
        error_handler("Input Ctx is nullptr.", __FILE__, __func__, __LINE__);
        return nullptr;
    }
    if (!Dev) {
        error_handler("Input Dev is nullptr.", __FILE__, __func__, __LINE__);
        return nullptr;
    }
    if (!Source) {
        error_handler("Input Source is nullptr.", __FILE__, __func__, __LINE__);
        return nullptr;
    }

    SyclCtx = unwrap<context>(Ctx);
    SyclDev = unwrap<device>(Dev);

    // get the backend type
    auto BE = SyclCtx->get_platform().get_backend();
    switch (BE) {
    case backend::opencl:
        try {
            return _CreateKernelBundleWithOCLSource_ocl_impl(
                *SyclCtx, *SyclDev, Source, CompileOpts);
        } catch (std::exception const &e) {
            error_handler(e, __FILE__, __func__, __LINE__);
            return nullptr;
        }
        break;
    case backend::ext_oneapi_level_zero:
        error_handler(
            "CreateFromSource is not supported for Level Zero backend.",
            __FILE__, __func__, __LINE__);
        return nullptr;
    default:
        error_handler("CreateFromSource is not supported in unknown backend.",
                      __FILE__, __func__, __LINE__);
        return nullptr;
    }
}

__dpctl_give DPCTLSyclKernelRef
DPCTLKernelBundle_GetKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef,
                            __dpctl_keep const char *KernelName)
{
    if (!KBRef) {
        error_handler("Input KBRef is nullptr", __FILE__, __func__, __LINE__);
        return nullptr;
    }
    if (!KernelName) {
        error_handler("Input KernelName is nullptr", __FILE__, __func__,
                      __LINE__);
        return nullptr;
    }
    auto SyclKB = unwrap<kernel_bundle<bundle_state::executable>>(KBRef);
    sycl::backend be = SyclKB->get_backend();
    switch (be) {
    case sycl::backend::opencl:
        return _GetKernel_ocl_impl(*SyclKB, KernelName);
    case sycl::backend::ext_oneapi_level_zero:
#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION
        return _GetKernel_ze_impl(*SyclKB, KernelName);
#endif
    default:
        std::ostringstream os;
        os << "Backend " << be << " is not supported";
        error_handler(os.str(), __FILE__, __func__, __LINE__);
        return nullptr;
    }
}

bool DPCTLKernelBundle_HasKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef,
                                 __dpctl_keep const char *KernelName)
{
    if (!KBRef) {
        error_handler("Input KBRef is nullptr", __FILE__, __func__, __LINE__);
        return false;
    }
    if (!KernelName) {
        error_handler("Input KernelName is nullptr", __FILE__, __func__,
                      __LINE__);
        return false;
    }

    auto SyclKB = unwrap<kernel_bundle<bundle_state::executable>>(KBRef);
    sycl::backend be = SyclKB->get_backend();
    switch (be) {
    case sycl::backend::opencl:
        return _HasKernel_ocl_impl(*SyclKB, KernelName);
    case sycl::backend::ext_oneapi_level_zero:
#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION
        return _HasKernel_ze_impl(*SyclKB, KernelName);
#endif
    default:
        std::ostringstream os;
        os << "Backend " << be << " is not supported";
        error_handler(os.str(), __FILE__, __func__, __LINE__);
        return false;
    }
}

void DPCTLKernelBundle_Delete(__dpctl_take DPCTLSyclKernelBundleRef KBRef)
{
    delete unwrap<kernel_bundle<bundle_state::executable>>(KBRef);
}

__dpctl_give DPCTLSyclKernelBundleRef
DPCTLKernelBundle_Copy(__dpctl_keep const DPCTLSyclKernelBundleRef KBRef)
{
    auto Bundle = unwrap<kernel_bundle<bundle_state::executable>>(KBRef);
    if (!Bundle) {
        error_handler(
            "Cannot copy DPCTLSyclKernelBundleRef as input is a nullptr",
            __FILE__, __func__, __LINE__);
        return nullptr;
    }
    try {
        auto CopiedBundle =
            new kernel_bundle<bundle_state::executable>(*Bundle);
        return wrap<kernel_bundle<bundle_state::executable>>(CopiedBundle);
    } catch (std::exception const &e) {
        error_handler(e, __FILE__, __func__, __LINE__);
        return nullptr;
    }
}

using build_option_list_t = std::vector<std::string>;

__dpctl_give DPCTLBuildOptionListRef DPCTLBuildOptionList_Create()
{
    auto BuildOptionList =
        std::unique_ptr<build_option_list_t>(new build_option_list_t());
    auto *RetVal =
        reinterpret_cast<DPCTLBuildOptionListRef>(BuildOptionList.get());
    BuildOptionList.release();
    return RetVal;
}

void DPCTLBuildOptionList_Delete(__dpctl_take DPCTLBuildOptionListRef Ref)
{
    delete reinterpret_cast<build_option_list_t *>(Ref);
}

void DPCTLBuildOptionList_Append(__dpctl_keep DPCTLBuildOptionListRef Ref,
                                 __dpctl_keep const char *Option)
{
    reinterpret_cast<build_option_list_t *>(Ref)->emplace_back(Option);
}

using kernel_name_list_t = std::vector<std::string>;

__dpctl_give DPCTLKernelNameListRef DPCTLKernelNameList_Create()
{
    auto KernelNameList =
        std::unique_ptr<kernel_name_list_t>(new kernel_name_list_t());
    auto *RetVal =
        reinterpret_cast<DPCTLKernelNameListRef>(KernelNameList.get());
    KernelNameList.release();
    return RetVal;
}

void DPCTLKernelNameList_Delete(__dpctl_take DPCTLKernelNameListRef Ref)
{
    delete reinterpret_cast<kernel_name_list_t *>(Ref);
}

void DPCTLKernelNameList_Append(__dpctl_keep DPCTLKernelNameListRef Ref,
                                __dpctl_keep const char *Option)
{
    reinterpret_cast<kernel_name_list_t *>(Ref)->emplace_back(Option);
}

using virtual_header_list_t = std::vector<std::pair<std::string, std::string>>;

__dpctl_give DPCTLVirtualHeaderListRef DPCTLVirtualHeaderList_Create()
{
    auto HeaderList =
        std::unique_ptr<virtual_header_list_t>(new virtual_header_list_t());
    auto *RetVal =
        reinterpret_cast<DPCTLVirtualHeaderListRef>(HeaderList.get());
    HeaderList.release();
    return RetVal;
}

void DPCTLVirtualHeaderList_Delete(__dpctl_take DPCTLVirtualHeaderListRef Ref)
{
    delete reinterpret_cast<virtual_header_list_t *>(Ref);
}

void DPCTLVirtualHeaderList_Append(__dpctl_keep DPCTLVirtualHeaderListRef Ref,
                                   __dpctl_keep const char *Name,
                                   __dpctl_keep const char *Content)
{
    auto Header = std::make_pair<std::string, std::string>(Name, Content);
    reinterpret_cast<virtual_header_list_t *>(Ref)->push_back(Header);
}

using kernel_build_log_t = std::string;

__dpctl_give DPCTLKernelBuildLogRef DPCTLKernelBuildLog_Create()
{
    auto BuildLog =
        std::unique_ptr<kernel_build_log_t>(new kernel_build_log_t(""));
    auto *RetVal = reinterpret_cast<DPCTLKernelBuildLogRef>(BuildLog.get());
    BuildLog.release();
    return RetVal;
}

void DPCTLKernelBuildLog_Delete(__dpctl_take DPCTLKernelBuildLogRef Ref)
{
    delete reinterpret_cast<kernel_build_log_t *>(Ref);
}

const char *DPCTLKernelBuildLog_Get(__dpctl_keep DPCTLKernelBuildLogRef Ref)
{
    return reinterpret_cast<kernel_build_log_t *>(Ref)->data();
}

namespace syclex = sycl::ext::oneapi::experimental;

#if defined(SYCL_EXT_ONEAPI_KERNEL_COMPILER) &&                                \
    defined(__SYCL_COMPILER_VERSION) && !defined(SUPPORTS_SYCL_COMPILATION)
// SYCL source code compilation is supported from 2025.1 onwards.
#if __SYCL_COMPILER_VERSION >= 20250317u
#define SUPPORTS_SYCL_COMPILATION 1
#else
#define SUPPORTS_SYCL_COMPILATION 0
#endif
#endif

bool DPCTLKernelBundle_CreateFromSYCLSource_Available()
{
#if (SUPPORTS_SYCL_COMPILATION > 0)
    return true;
#else
    return false;
#endif
}

#if (SUPPORTS_SYCL_COMPILATION > 0)
// The property for registering names was renamed between DPC++ versions 2025.1
// and 2025.2. The original name was `registered_kernel_names`, the new name is
// `registered_names`. To select the correct name without being overly reliant
// on the SYCL compiler version definition, we forward declare both names and
// then select the new name if it is defined (i.e., not only declared).
namespace sycl::ext::oneapi::experimental
{
struct registered_names;
struct registered_kernel_names;
} // namespace sycl::ext::oneapi::experimental

template <typename NewT, typename FallbackT, typename = void>
struct new_type_if_defined
{
    using type = FallbackT;
};

template <typename NewT, typename FallbackT>
struct new_type_if_defined<NewT, FallbackT, std::void_t<decltype(sizeof(NewT))>>
{
    using type = NewT;
};

using registered_names_property_t =
    new_type_if_defined<syclex::registered_names,
                        syclex::registered_kernel_names>::type;
#endif

__dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource(
    __dpctl_keep const DPCTLSyclContextRef Ctx,
    __dpctl_keep const DPCTLSyclDeviceRef Dev,
    __dpctl_keep const char *Source,
    __dpctl_keep DPCTLVirtualHeaderListRef Headers,
    __dpctl_keep DPCTLKernelNameListRef Names,
    __dpctl_keep DPCTLBuildOptionListRef BuildOptions,
    __dpctl_keep DPCTLKernelBuildLogRef BuildLog)
{
#if (SUPPORTS_SYCL_COMPILATION > 0)
    // optional arguments and BuildLog may be NULL, which is equivalent to
    // passing an empty list or discarding the log
    // Writing to the log is routed through a helper
    auto *RawBuildLog = reinterpret_cast<kernel_build_log_t *>(BuildLog);
    auto set_build_log = [RawBuildLog](const std::string &Msg) {
        if (RawBuildLog)
            *RawBuildLog = Msg;
    };
    auto fail = [&set_build_log](const char *Msg, int Line) {
        set_build_log(Msg);
        error_handler(Msg, __FILE__, "DPCTLKernelBundle_CreateFromSYCLSource",
                      Line);
        return nullptr;
    };

    if (!Ctx)
        return fail("Input Ctx is nullptr.", __LINE__);
    if (!Dev)
        return fail("Input Dev is nullptr.", __LINE__);
    if (!Source)
        return fail("Input Source is nullptr.", __LINE__);

    static const virtual_header_list_t EmptyHeaders{};
    static const kernel_name_list_t EmptyNames{};
    static const build_option_list_t EmptyOptions{};

    const auto &IncludeFiles =
        Headers ? *reinterpret_cast<virtual_header_list_t *>(Headers)
                : EmptyHeaders;
    const auto &KernelNames =
        Names ? *reinterpret_cast<kernel_name_list_t *>(Names) : EmptyNames;
    const auto &Options =
        BuildOptions ? *reinterpret_cast<build_option_list_t *>(BuildOptions)
                     : EmptyOptions;

    context *SyclCtx = unwrap<context>(Ctx);
    device *SyclDev = unwrap<device>(Dev);
    if (!SyclDev->ext_oneapi_can_compile(syclex::source_language::sycl)) {
        set_build_log("Device does not support compilation of SYCL source.");
        return nullptr;
    }
    try {
        std::unique_ptr<kernel_bundle<bundle_state::ext_oneapi_source>>
            SrcBundle;
        std::string Src(Source);
        // The following logic is to work around a bug in DPC++ version 2025.1.
        // This version declares a constructor with no parameters for the
        // `include_files` property, but does not implement it. Therefore, the
        // only way to create `include_files` is with the name and content of
        // the first virtual header, if any.
        if (!IncludeFiles.empty()) {
            auto IncludeFileIt = IncludeFiles.begin();
            syclex::include_files IncludeFilesProp{IncludeFileIt->first,
                                                   IncludeFileIt->second};
            for (std::advance(IncludeFileIt, 1);
                 IncludeFileIt != IncludeFiles.end(); ++IncludeFileIt)
            {
                IncludeFilesProp.add(IncludeFileIt->first,
                                     IncludeFileIt->second);
            }
            SrcBundle = std::make_unique<
                kernel_bundle<bundle_state::ext_oneapi_source>>(
                syclex::create_kernel_bundle_from_source(
                    *SyclCtx, syclex::source_language::sycl, Src,
                    syclex::properties{IncludeFilesProp}));
        }
        else {
            SrcBundle = std::make_unique<
                kernel_bundle<bundle_state::ext_oneapi_source>>(
                syclex::create_kernel_bundle_from_source(
                    *SyclCtx, syclex::source_language::sycl, Src));
        }

        registered_names_property_t RegisteredNames;
        for (const std::string &Name : KernelNames) {
            RegisteredNames.add(Name);
        }

        syclex::build_options Opts{Options};

        std::vector<sycl::device> Devices({*SyclDev});

        auto ExeBundle = syclex::build(
            *SrcBundle, Devices, syclex::properties{RegisteredNames, Opts});
        auto ResultBundle =
            std::make_unique<sycl::kernel_bundle<bundle_state::executable>>(
                ExeBundle);
        return wrap<kernel_bundle<bundle_state::executable>>(
            ResultBundle.release());
    } catch (const std::exception &e) {
        set_build_log(e.what());
        return nullptr;
    }
#else
    return nullptr;
#endif
}

__dpctl_give DPCTLSyclKernelRef
DPCTLKernelBundle_GetSyclKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef,
                                __dpctl_keep const char *KernelName)
{
#if (SUPPORTS_SYCL_COMPILATION > 0)
    if (!KBRef) {
        error_handler("Input KBRef is nullptr", __FILE__, __func__, __LINE__);
        return nullptr;
    }
    if (!KernelName) {
        error_handler("Input KernelName is nullptr", __FILE__, __func__,
                      __LINE__);
        return nullptr;
    }
    try {
        auto KernelBundle =
            unwrap<sycl::kernel_bundle<bundle_state::executable>>(KBRef);
        auto Kernel = KernelBundle->ext_oneapi_get_kernel(KernelName);
        return wrap<sycl::kernel>(new sycl::kernel(Kernel));
    } catch (const std::exception &e) {
        error_handler(e, __FILE__, __func__, __LINE__);
        return nullptr;
    }
#else
    return nullptr;
#endif
}

bool DPCTLKernelBundle_HasSyclKernel(__dpctl_keep DPCTLSyclKernelBundleRef
                                         KBRef,
                                     __dpctl_keep const char *KernelName)
{
#if (SUPPORTS_SYCL_COMPILATION > 0)
    if (!KBRef) {
        error_handler("Input KBRef is nullptr", __FILE__, __func__, __LINE__);
        return false;
    }
    if (!KernelName) {
        error_handler("Input KernelName is nullptr", __FILE__, __func__,
                      __LINE__);
        return false;
    }
    try {
        auto KernelBundle =
            unwrap<sycl::kernel_bundle<bundle_state::executable>>(KBRef);
        return KernelBundle->ext_oneapi_has_kernel(KernelName);
    } catch (const std::exception &e) {
        error_handler(e, __FILE__, __func__, __LINE__);
        return false;
    }
#else
    return false;
#endif
}

IntelPython / dpctl / 31424476294

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

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