From 614d982e1cc226c1a7f996fb8f5ce652f5a6323a Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Thu, 10 Dec 2020 22:01:45 -0600 Subject: [PATCH 1/9] WIP level zero support. --- .../source/dpctl_sycl_program_interface.cpp | 57 ++++++++++++++++++- 1 file changed, 54 insertions(+), 3 deletions(-) diff --git a/dpctl-capi/source/dpctl_sycl_program_interface.cpp b/dpctl-capi/source/dpctl_sycl_program_interface.cpp index 7f8b75cf63..b4d0afbe7b 100644 --- a/dpctl-capi/source/dpctl_sycl_program_interface.cpp +++ b/dpctl-capi/source/dpctl_sycl_program_interface.cpp @@ -27,9 +27,11 @@ #include "dpctl_sycl_program_interface.h" #include "Support/CBindingWrapping.h" -#include /* Sycl headers */ -#include /* OpenCL headers */ - +#include /* Sycl headers */ +#include /* OpenCL headers */ +#if 0 +#include /* Level Zero headers */ +#endif using namespace cl::sycl; namespace @@ -82,7 +84,56 @@ createOpenCLInterOpProgram (const context &SyclCtx, return nullptr; } } +#if 0 +template +__dpctl_give DPCTLSyclProgramRef +createLevelZeroInterOpProgram (const context &SyclCtx, + const void *IL, + size_t length, + const char* BuildOptions) +{ + auto ZeCtx = SyclCtx.get_native(); + auto SyclDevices = SyclCtx.get_devices(); + if(SyclDevices.size() > 1) { + // We only support build to one device with Level Zero now. + // TODO: log error + return nullptr; + } + + // Specialization constants are not yet supported. + // Refer https://bit.ly/33UEDYN for details on specialization constants. + ze_module_constants_t ZeSpecConstants = {}; + ZeSpecConstants.numConstants = 0; + // Populate the Level Zero module descriptions + ze_module_desc_t ZeModuleDesc = {}; + ZeModuleDesc.format = ZE_MODULE_FORMAT_IL_SPIRV; + ZeModuleDesc.inputSize = length; + ZeModuleDesc.pInputModule = (uint8_t*)IL; + ZeModuleDesc.pBuildFlags = BuildOptions; + ZeModuleDesc.pConstants = &ZeSpecConstants; + + auto ZeDevice = SyclDevices[0].get_native(); + ze_module_handle_t ZeModule; + auto ret = zeModuleCreate(ZeContext, ZeDevice, &ZeModuleDesc, &ZeModule, + nullptr); + if(ret != ZE_RESULT_SUCCESS) { + // TODO: handle error + return nullptr; + } +#if 0 + // Create the Sycl program from the ZeModule + try { + auto SyclProgram = new program(SyclCtx, ZeModule); + return wrap(SyclProgram); + } catch (invalid_object_error &e) { + // \todo record error + std::cerr << e.what() << '\n'; + return nullptr; + } +#endif +} +#endif } /* end of anonymous namespace */ __dpctl_give DPCTLSyclProgramRef From a10ef0af78f6c385ec8edc3c1d8db8cb2d74868d Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Fri, 15 Jan 2021 19:19:16 -0600 Subject: [PATCH 2/9] Add support to build a SYCL program from SPIR-V. --- dpctl-capi/dbg_build.sh | 19 ++++---- dpctl-capi/include/Config/dpctl_config.h | 32 +++++++++++++ .../include/dpctl_sycl_program_interface.h | 9 ++-- .../source/dpctl_sycl_program_interface.cpp | 46 +++++++++++-------- .../tests/test_sycl_program_interface.cpp | 46 +++++++++++++++---- dpctl/_backend.pxd | 9 ++-- dpctl/program/_program.pxd | 3 +- dpctl/program/_program.pyx | 9 +++- dpctl/tests/test_sycl_program.py | 1 - scripts/build_backend.py | 1 + 10 files changed, 126 insertions(+), 49 deletions(-) create mode 100644 dpctl-capi/include/Config/dpctl_config.h diff --git a/dpctl-capi/dbg_build.sh b/dpctl-capi/dbg_build.sh index 60c8674a42..5b42f82273 100755 --- a/dpctl-capi/dbg_build.sh +++ b/dpctl-capi/dbg_build.sh @@ -9,15 +9,16 @@ rm -rf ${INSTALL_PREFIX} export ONEAPI_ROOT=/opt/intel/oneapi DPCPP_ROOT=${ONEAPI_ROOT}/compiler/latest/linux -cmake \ - -DCMAKE_BUILD_TYPE=Debug \ - -DCMAKE_INSTALL_PREFIX=${INSTALL_PREFIX} \ - -DCMAKE_PREFIX_PATH=${INSTALL_PREFIX} \ - -DDPCPP_INSTALL_DIR=${DPCPP_ROOT} \ - -DCMAKE_C_COMPILER:PATH=${DPCPP_ROOT}/bin/clang \ - -DCMAKE_CXX_COMPILER:PATH=${DPCPP_ROOT}/bin/dpcpp \ - -DDPCTL_BUILD_CAPI_TESTS=ON \ - -DDPCTL_GENERATE_COVERAGE=ON \ +cmake \ + -DCMAKE_BUILD_TYPE=Debug \ + -DCMAKE_INSTALL_PREFIX=${INSTALL_PREFIX} \ + -DCMAKE_PREFIX_PATH=${INSTALL_PREFIX} \ + -DDPCPP_INSTALL_DIR=${DPCPP_ROOT} \ + -DCMAKE_C_COMPILER:PATH=${DPCPP_ROOT}/bin/clang \ + -DCMAKE_CXX_COMPILER:PATH=${DPCPP_ROOT}/bin/dpcpp \ + -DDPCTL_ENABLE_LO_PROGRAM_CREATION=ON \ + -DDPCTL_BUILD_CAPI_TESTS=ON \ + -DDPCTL_GENERATE_COVERAGE=ON \ .. make V=1 -n -j 4 && make check && make install diff --git a/dpctl-capi/include/Config/dpctl_config.h b/dpctl-capi/include/Config/dpctl_config.h new file mode 100644 index 0000000000..63eaeae6e0 --- /dev/null +++ b/dpctl-capi/include/Config/dpctl_config.h @@ -0,0 +1,32 @@ +//===---- dpctl-capi/Config/dpCtl-config.h - dpctl-C API -------*- C++ -*-===// +// +// Data Parallel Control Library (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 exports a set of dpCtl C API configurations. +/// +//===----------------------------------------------------------------------===// + +#pragma once + +/* Defined when dpCtl was built with level zero program creation enabled. */ +#define DPCTL_ENABLE_LO_PROGRAM_CREATION ON + +/* The DPCPP version used to build dpCtl */ +#define DPCTL_DPCPP_VERSION "Intel(R) oneAPI DPC++ Compiler 2021.1.2 (2020.10.0.1214)" diff --git a/dpctl-capi/include/dpctl_sycl_program_interface.h b/dpctl-capi/include/dpctl_sycl_program_interface.h index c931b4b166..8e25d1df40 100644 --- a/dpctl-capi/include/dpctl_sycl_program_interface.h +++ b/dpctl-capi/include/dpctl_sycl_program_interface.h @@ -52,14 +52,17 @@ DPCTL_C_EXTERN_C_BEGIN * @param Ctx An opaque pointer to a sycl::context * @param IL SPIR-V binary * @param Length The size of the IL binary in bytes. + * @param CompileOpts Optional compiler flags used when compiling th + * SPIR-V binary. * @return A new SyclProgramRef pointer if the program creation succeeded, * else returns NULL. */ DPCTL_API __dpctl_give DPCTLSyclProgramRef -DPCTLProgram_CreateFromOCLSpirv (__dpctl_keep const DPCTLSyclContextRef Ctx, - __dpctl_keep const void *IL, - size_t Length); +DPCTLProgram_CreateFromSpirv (__dpctl_keep const DPCTLSyclContextRef Ctx, + __dpctl_keep const void *IL, + size_t Length, + const char *CompileOpts); /*! * @brief Create a Sycl program from an OpenCL kernel source string. diff --git a/dpctl-capi/source/dpctl_sycl_program_interface.cpp b/dpctl-capi/source/dpctl_sycl_program_interface.cpp index b4d0afbe7b..ae735abbe2 100644 --- a/dpctl-capi/source/dpctl_sycl_program_interface.cpp +++ b/dpctl-capi/source/dpctl_sycl_program_interface.cpp @@ -25,13 +25,16 @@ //===----------------------------------------------------------------------===// #include "dpctl_sycl_program_interface.h" +#include "Config/dpctl_config.h" #include "Support/CBindingWrapping.h" #include /* Sycl headers */ #include /* OpenCL headers */ -#if 0 +#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION #include /* Level Zero headers */ +#include #endif + using namespace cl::sycl; namespace @@ -43,7 +46,8 @@ DEFINE_SIMPLE_CONVERSION_FUNCTIONS(kernel, DPCTLSyclKernelRef) __dpctl_give DPCTLSyclProgramRef createOpenCLInterOpProgram (const context &SyclCtx, __dpctl_keep const void *IL, - size_t length) + size_t length, + const char * /* */) { cl_int err; auto CLCtx = SyclCtx.get(); @@ -84,15 +88,15 @@ createOpenCLInterOpProgram (const context &SyclCtx, return nullptr; } } -#if 0 -template + +#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION __dpctl_give DPCTLSyclProgramRef createLevelZeroInterOpProgram (const context &SyclCtx, const void *IL, size_t length, - const char* BuildOptions) + const char *CompileOpts) { - auto ZeCtx = SyclCtx.get_native(); + auto ZeCtx = SyclCtx.get_native(); auto SyclDevices = SyclCtx.get_devices(); if(SyclDevices.size() > 1) { // We only support build to one device with Level Zero now. @@ -110,36 +114,38 @@ createLevelZeroInterOpProgram (const context &SyclCtx, ZeModuleDesc.format = ZE_MODULE_FORMAT_IL_SPIRV; ZeModuleDesc.inputSize = length; ZeModuleDesc.pInputModule = (uint8_t*)IL; - ZeModuleDesc.pBuildFlags = BuildOptions; + ZeModuleDesc.pBuildFlags = CompileOpts; ZeModuleDesc.pConstants = &ZeSpecConstants; - auto ZeDevice = SyclDevices[0].get_native(); + auto ZeDevice = SyclDevices[0].get_native(); ze_module_handle_t ZeModule; - auto ret = zeModuleCreate(ZeContext, ZeDevice, &ZeModuleDesc, &ZeModule, + auto ret = zeModuleCreate(ZeCtx, ZeDevice, &ZeModuleDesc, &ZeModule, nullptr); if(ret != ZE_RESULT_SUCCESS) { // TODO: handle error return nullptr; } -#if 0 + // Create the Sycl program from the ZeModule try { - auto SyclProgram = new program(SyclCtx, ZeModule); - return wrap(SyclProgram); + auto ZeProgram = new program(sycl::level_zero::make_program( + SyclCtx, reinterpret_cast(ZeModule)) + ); + return wrap(ZeProgram); } catch (invalid_object_error &e) { // \todo record error std::cerr << e.what() << '\n'; return nullptr; } -#endif } #endif } /* end of anonymous namespace */ __dpctl_give DPCTLSyclProgramRef -DPCTLProgram_CreateFromOCLSpirv (__dpctl_keep const DPCTLSyclContextRef CtxRef, - __dpctl_keep const void *IL, - size_t length) +DPCTLProgram_CreateFromSpirv (__dpctl_keep const DPCTLSyclContextRef CtxRef, + __dpctl_keep const void *IL, + size_t length, + const char *CompileOpts) { DPCTLSyclProgramRef Pref = nullptr; context *SyclCtx = nullptr; @@ -147,21 +153,23 @@ DPCTLProgram_CreateFromOCLSpirv (__dpctl_keep const DPCTLSyclContextRef CtxRef, // \todo handle error return Pref; } - SyclCtx = unwrap(CtxRef); // get the backend type auto BE = SyclCtx->get_platform().get_backend(); switch (BE) { case backend::opencl: - Pref = createOpenCLInterOpProgram(*SyclCtx, IL, length); + Pref = createOpenCLInterOpProgram(*SyclCtx, IL, length, CompileOpts); break; case backend::level_zero: +#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION + Pref = createLevelZeroInterOpProgram(*SyclCtx, IL, length, + CompileOpts); +#endif break; default: break; } - return Pref; } diff --git a/dpctl-capi/tests/test_sycl_program_interface.cpp b/dpctl-capi/tests/test_sycl_program_interface.cpp index 0e0061b674..791907c1c7 100644 --- a/dpctl-capi/tests/test_sycl_program_interface.cpp +++ b/dpctl-capi/tests/test_sycl_program_interface.cpp @@ -1,4 +1,4 @@ -//===---------- test_sycl_program_interface.cpp - dpctl-C_API --*-- C++ -*-===// +//===---------- test_sycl_program_interface.cpp - dpctl-C_API ---*- C++ -*-===// // // Data Parallel Control Library (dpCtl) // @@ -29,7 +29,7 @@ #include "dpctl_sycl_program_interface.h" #include "dpctl_sycl_queue_interface.h" #include "dpctl_sycl_queue_manager.h" - +#include "Config/dpctl_config.h" #include #include #include @@ -127,12 +127,16 @@ struct TestDPCTLSyclProgramInterface : public ::testing::Test size_t spirvFileSize = 0; std::vector spirvBuffer; size_t nOpenCLGpuQ = 0; +#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION + size_t nL0GpuQ = 0; +#endif TestDPCTLSyclProgramInterface () : spirvFile{"./multi_kernel.spv", std::ios::binary | std::ios::ate}, spirvFileSize(std::filesystem::file_size("./multi_kernel.spv")), spirvBuffer(spirvFileSize), - nOpenCLGpuQ(DPCTLQueueMgr_GetNumQueues(DPCTL_OPENCL, DPCTL_GPU)) + nOpenCLGpuQ(DPCTLQueueMgr_GetNumQueues(DPCTL_OPENCL, DPCTL_GPU)), + nL0GpuQ(DPCTLQueueMgr_GetNumQueues(DPCTL_LEVEL_ZERO, DPCTL_GPU)) { spirvFile.seekg(0, std::ios::beg); spirvFile.read(spirvBuffer.data(), spirvFileSize); @@ -152,7 +156,7 @@ TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromOCLSource) auto QueueRef = DPCTLQueueMgr_GetQueue(DPCTL_OPENCL, DPCTL_GPU, 0); auto CtxRef = DPCTLQueue_GetContext(QueueRef); auto PRef = DPCTLProgram_CreateFromOCLSource(CtxRef, CLProgramStr, - CompileOpts); + CompileOpts); ASSERT_TRUE(PRef != nullptr); ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "add")); ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "axpy")); @@ -162,15 +166,36 @@ TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromOCLSource) DPCTLProgram_Delete(PRef); } -TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromOCLSpirv) +TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromSpirvOCL) { if(!nOpenCLGpuQ) GTEST_SKIP_("Skipping as no OpenCL GPU device found.\n"); auto QueueRef = DPCTLQueueMgr_GetQueue(DPCTL_OPENCL, DPCTL_GPU, 0); auto CtxRef = DPCTLQueue_GetContext(QueueRef); - auto PRef = DPCTLProgram_CreateFromOCLSpirv(CtxRef, spirvBuffer.data(), - spirvFileSize); + auto PRef = DPCTLProgram_CreateFromSpirv(CtxRef, spirvBuffer.data(), + spirvFileSize, + nullptr); + ASSERT_TRUE(PRef != nullptr); + ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "add")); + ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "axpy")); + + DPCTLQueue_Delete(QueueRef); + DPCTLContext_Delete(CtxRef); + DPCTLProgram_Delete(PRef); +} + +#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION +TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromSpirvL0) +{ + if(!nL0GpuQ) + GTEST_SKIP_("Skipping as no OpenCL GPU device found.\n"); + + auto QueueRef = DPCTLQueueMgr_GetQueue(DPCTL_LEVEL_ZERO, DPCTL_GPU, 0); + auto CtxRef = DPCTLQueue_GetContext(QueueRef); + auto PRef = DPCTLProgram_CreateFromSpirv(CtxRef, spirvBuffer.data(), + spirvFileSize, + nullptr); ASSERT_TRUE(PRef != nullptr); ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "add")); ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "axpy")); @@ -179,6 +204,7 @@ TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromOCLSpirv) DPCTLContext_Delete(CtxRef); DPCTLProgram_Delete(PRef); } +#endif TEST_F (TestDPCTLSyclProgramInterface, CheckGetKernelOCLSource) { @@ -207,15 +233,15 @@ TEST_F (TestDPCTLSyclProgramInterface, CheckGetKernelOCLSource) DPCTLProgram_Delete(PRef); } -TEST_F (TestDPCTLSyclProgramInterface, CheckGetKernelOCLSpirv) +TEST_F (TestDPCTLSyclProgramInterface, CheckGetKernelSpirv) { if(!nOpenCLGpuQ) GTEST_SKIP_("Skipping as no OpenCL GPU device found.\n"); auto QueueRef = DPCTLQueueMgr_GetQueue(DPCTL_OPENCL, DPCTL_GPU, 0); auto CtxRef = DPCTLQueue_GetContext(QueueRef); - auto PRef = DPCTLProgram_CreateFromOCLSpirv(CtxRef, spirvBuffer.data(), - spirvFileSize); + auto PRef = DPCTLProgram_CreateFromSpirv(CtxRef, spirvBuffer.data(), + spirvFileSize, nullptr); auto AddKernel = DPCTLProgram_GetKernel(PRef, "add"); auto AxpyKernel = DPCTLProgram_GetKernel(PRef, "axpy"); diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index ff6d0c536f..9e9dbca948 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -140,14 +140,15 @@ cdef extern from "dpctl_sycl_context_interface.h": cdef extern from "dpctl_sycl_program_interface.h": - cdef DPCTLSyclProgramRef DPCTLProgram_CreateFromOCLSpirv ( + cdef DPCTLSyclProgramRef DPCTLProgram_CreateFromSpirv ( const DPCTLSyclContextRef Ctx, const void *IL, - size_t Length) + size_t Length, + const char *CompileOpts) cdef DPCTLSyclProgramRef DPCTLProgram_CreateFromOCLSource ( const DPCTLSyclContextRef Ctx, - const char* Source, - const char* CompileOpts) + const char *Source, + const char *CompileOpts) cdef DPCTLSyclKernelRef DPCTLProgram_GetKernel (DPCTLSyclProgramRef PRef, const char *KernelName) cdef bool DPCTLProgram_HasKernel (DPCTLSyclProgramRef PRef, diff --git a/dpctl/program/_program.pxd b/dpctl/program/_program.pxd index 27f0716791..1190d5a12c 100644 --- a/dpctl/program/_program.pxd +++ b/dpctl/program/_program.pxd @@ -60,4 +60,5 @@ cdef class SyclProgram: cpdef create_program_from_source (SyclQueue q, unicode source, unicode copts=*) -cpdef create_program_from_spirv (SyclQueue q, const unsigned char[:] IL) +cpdef create_program_from_spirv (SyclQueue q, const unsigned char[:] IL, + unicode copts=*) diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index efa47f7649..2ea8ef6ea8 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -162,7 +162,8 @@ cpdef create_program_from_source(SyclQueue q, unicode src, unicode copts=""): cimport cython.array -cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL): +cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL, + unicode copts=""): """ Creates a Sycl interoperability program from an SPIR-V binary. @@ -173,6 +174,8 @@ cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL): q (SyclQueue): The :class:`SyclQueue` for which the :class:`SyclProgram` is going to be built. IL (const char[:]) : SPIR-V binary IL file for an OpenCL program. + copts (unicode) : Optional compilation flags that will be used + when compiling the program. Returns: 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): cdef const unsigned char *dIL = &IL[0] cdef DPCTLSyclContextRef CRef = q.get_sycl_context().get_context_ref() cdef size_t length = IL.shape[0] - Pref = DPCTLProgram_CreateFromOCLSpirv(CRef, dIL, length) + cdef bytes bCOpts = copts.encode('utf8') + cdef const char *COpts = bCOpts + Pref = DPCTLProgram_CreateFromSpirv(CRef, dIL, length, COpts) if Pref is NULL: raise SyclProgramCompilationError() diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index 5eb1c2b66b..cd31c6e059 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -87,7 +87,6 @@ def test_create_program_from_spirv(self): "No Level0 GPU queues available", ) class TestProgramForLevel0GPU(unittest.TestCase): - @unittest.expectedFailure def test_create_program_from_spirv(self): CURR_DIR = os.path.dirname(os.path.abspath(__file__)) diff --git a/scripts/build_backend.py b/scripts/build_backend.py index 2708094dd6..c68dd37b40 100644 --- a/scripts/build_backend.py +++ b/scripts/build_backend.py @@ -43,6 +43,7 @@ "-DDPCPP_INSTALL_DIR=" + DPCPP_ROOT, "-DCMAKE_C_COMPILER:PATH=" + os.path.join(DPCPP_ROOT, "bin", "clang"), "-DCMAKE_CXX_COMPILER:PATH=" + os.path.join(DPCPP_ROOT, "bin", "clang++"), + "-DDPCTL_ENABLE_LO_PROGRAM_CREATION=ON", backends, ] subprocess.check_call(cmake_args, stderr=subprocess.STDOUT, shell=False) From 05c4b3db97883be7eb40b62f64b029a2c01a6f00 Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Fri, 15 Jan 2021 19:24:13 -0600 Subject: [PATCH 3/9] Ignore the generated config file. --- dpctl-capi/include/Config/dpctl_config.h | 32 ------------------------ 1 file changed, 32 deletions(-) delete mode 100644 dpctl-capi/include/Config/dpctl_config.h diff --git a/dpctl-capi/include/Config/dpctl_config.h b/dpctl-capi/include/Config/dpctl_config.h deleted file mode 100644 index 63eaeae6e0..0000000000 --- a/dpctl-capi/include/Config/dpctl_config.h +++ /dev/null @@ -1,32 +0,0 @@ -//===---- dpctl-capi/Config/dpCtl-config.h - dpctl-C API -------*- C++ -*-===// -// -// Data Parallel Control Library (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 exports a set of dpCtl C API configurations. -/// -//===----------------------------------------------------------------------===// - -#pragma once - -/* Defined when dpCtl was built with level zero program creation enabled. */ -#define DPCTL_ENABLE_LO_PROGRAM_CREATION ON - -/* The DPCPP version used to build dpCtl */ -#define DPCTL_DPCPP_VERSION "Intel(R) oneAPI DPC++ Compiler 2021.1.2 (2020.10.0.1214)" From 015249abd8c8786b9327578e1d1d6faf3e91c40c Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Fri, 15 Jan 2021 22:28:39 -0600 Subject: [PATCH 4/9] Skip test on Windows. --- dpctl/tests/test_sycl_program.py | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index cd31c6e059..ed889e17fa 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -87,8 +87,10 @@ def test_create_program_from_spirv(self): "No Level0 GPU queues available", ) class TestProgramForLevel0GPU(unittest.TestCase): + # Level zero proram creation from SPIR-V is only supported on Windows. + import sys + @unittest.skipIf(sys.platform in ["win32", "cygwin"]) def test_create_program_from_spirv(self): - CURR_DIR = os.path.dirname(os.path.abspath(__file__)) spirv_file = os.path.join(CURR_DIR, "input_files/multi_kernel.spv") with open(spirv_file, "rb") as fin: From f36a915c012b8901afb11677b86c56a8dbbe5c19 Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Fri, 15 Jan 2021 22:32:38 -0600 Subject: [PATCH 5/9] Black formatting. --- dpctl/tests/test_sycl_program.py | 1 + 1 file changed, 1 insertion(+) diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index ed889e17fa..2ff2634a39 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -89,6 +89,7 @@ def test_create_program_from_spirv(self): class TestProgramForLevel0GPU(unittest.TestCase): # Level zero proram creation from SPIR-V is only supported on Windows. import sys + @unittest.skipIf(sys.platform in ["win32", "cygwin"]) def test_create_program_from_spirv(self): CURR_DIR = os.path.dirname(os.path.abspath(__file__)) From 9c4eb764b7ddeeaa53ce84a797db4c08feb5b63c Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Fri, 15 Jan 2021 23:26:16 -0600 Subject: [PATCH 6/9] Skip test properly. --- dpctl/tests/test_sycl_program.py | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index 2ff2634a39..e28e37bfe1 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -90,7 +90,10 @@ class TestProgramForLevel0GPU(unittest.TestCase): # Level zero proram creation from SPIR-V is only supported on Windows. import sys - @unittest.skipIf(sys.platform in ["win32", "cygwin"]) + @unittest.skipIf( + sys.platform in ["win32", "cygwin"], + "Level Zero module creation unsupported on Windows.", + ) def test_create_program_from_spirv(self): CURR_DIR = os.path.dirname(os.path.abspath(__file__)) spirv_file = os.path.join(CURR_DIR, "input_files/multi_kernel.spv") From b3ebc8c29fba0944dc02d4ac7439bd5f7de74b6d Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Wed, 20 Jan 2021 17:11:51 -0600 Subject: [PATCH 7/9] Fix comment. --- dpctl/tests/test_sycl_program.py | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index e28e37bfe1..1b72250a1e 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -87,9 +87,10 @@ def test_create_program_from_spirv(self): "No Level0 GPU queues available", ) class TestProgramForLevel0GPU(unittest.TestCase): - # Level zero proram creation from SPIR-V is only supported on Windows. - import sys + import sys + # Level zero program creation from a SPIR-V binary is not supported + # on Windows. @unittest.skipIf( sys.platform in ["win32", "cygwin"], "Level Zero module creation unsupported on Windows.", From 954e07bd6101f7fe97f95192bc22e492246bc41f Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Wed, 20 Jan 2021 17:21:32 -0600 Subject: [PATCH 8/9] Fix black error. --- dpctl/tests/test_sycl_program.py | 1 + 1 file changed, 1 insertion(+) diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index 1b72250a1e..ce7f6051f2 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -89,6 +89,7 @@ def test_create_program_from_spirv(self): class TestProgramForLevel0GPU(unittest.TestCase): import sys + # Level zero program creation from a SPIR-V binary is not supported # on Windows. @unittest.skipIf( From 834172ec1d4d42183bc52d64f4191948e50eb14d Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Wed, 20 Jan 2021 18:33:02 -0600 Subject: [PATCH 9/9] Fix comments. --- dpctl/_sycl_core.pyx | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/dpctl/_sycl_core.pyx b/dpctl/_sycl_core.pyx index 25d01330b2..835df44db1 100644 --- a/dpctl/_sycl_core.pyx +++ b/dpctl/_sycl_core.pyx @@ -678,7 +678,7 @@ cdef class _SyclRTManager: raise UnsupportedDeviceError("Device can only be gpu or cpu") except KeyError: raise UnsupportedBackendError("Backend can only be opencl or " - "level-0") + "level0") def _remove_current_queue(self): DPCTLQueueMgr_PopQueue() @@ -970,7 +970,7 @@ cdef class _SyclRTManager: raise UnsupportedDeviceError("Device can only be gpu or cpu") except KeyError: raise UnsupportedBackendError("Backend can only be opencl or " - "level-0") + "level0") # This private instance of the _SyclQueueManager should not be directly