From 842addd7bbdb1aed24c2a0fc316c516e63109cab Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Wed, 19 Nov 2025 15:16:31 -0800 Subject: [PATCH 1/2] remove entire tensor submodule and modify cmake accordingly --- dpctl/program/utils/__init__.py | 26 -------------------------- 1 file changed, 26 deletions(-) delete mode 100644 dpctl/program/utils/__init__.py diff --git a/dpctl/program/utils/__init__.py b/dpctl/program/utils/__init__.py deleted file mode 100644 index 34c480838d..0000000000 --- a/dpctl/program/utils/__init__.py +++ /dev/null @@ -1,26 +0,0 @@ -# Data Parallel Control (dpctl) -# -# Copyright 2026 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. - -""" -A collection of utility functions for dpctl.program module. -""" - -from ._utils import SpecializationConstantInfo, parse_spirv_specializations - -__all__ = [ - "parse_spirv_specializations", - "SpecializationConstantInfo", -] From 6ae808c3340b27a8c956f08465e315ac118094b7 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Fri, 7 Aug 2026 22:59:04 -0700 Subject: [PATCH 2/2] Enable building with AdaptiveCpp --- CMakeLists.txt | 79 +++++++---- dpctl/CMakeLists.txt | 17 ++- dpctl/_host_task_util.hpp | 39 ++++- dpctl/apis/include/dpctl4pybind11.hpp | 49 ++++++- dpctl/apis/include/dpctl_acpp_host_task.hpp | 133 ++++++++++++++++++ dpctl/memory/CMakeLists.txt | 1 - dpctl/tests/helper/__init__.py | 2 + dpctl/tests/helper/_helper.py | 12 ++ dpctl/tests/test_sycl_queue.py | 19 +-- dpctl/utils/CMakeLists.txt | 11 +- dpctl/utils/src/device_queries.cpp | 22 +++ libsyclinterface/CMakeLists.txt | 29 ++-- .../helper/include/dpctl_utils_helper.h | 2 + .../helper/source/dpctl_utils_helper.cpp | 48 +++++++ .../syclinterface/dpctl_device_selection.hpp | 6 +- .../syclinterface/dpctl_sycl_type_casters.hpp | 5 + .../source/dpctl_device_selection.cpp | 10 +- .../source/dpctl_sycl_context_interface.cpp | 13 ++ .../source/dpctl_sycl_device_interface.cpp | 24 +++- .../source/dpctl_sycl_device_manager.cpp | 2 + .../source/dpctl_sycl_event_interface.cpp | 5 + .../dpctl_sycl_kernel_bundle_interface.cpp | 48 ++++++- .../source/dpctl_sycl_kernel_interface.cpp | 28 ++++ .../source/dpctl_sycl_platform_interface.cpp | 2 + .../source/dpctl_sycl_platform_manager.cpp | 13 ++ .../source/dpctl_sycl_queue_interface.cpp | 47 +++++-- 26 files changed, 591 insertions(+), 75 deletions(-) create mode 100644 dpctl/apis/include/dpctl_acpp_host_task.hpp diff --git a/CMakeLists.txt b/CMakeLists.txt index 0d27c9d6c3..c68ee6ebd2 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -26,9 +26,13 @@ option(DPCTL_GENERATE_COVERAGE_FOR_PYBIND11_EXTENSIONS OFF ) +# SYCL provider selection +set(DPCTL_SYCL_PROVIDER "Intel" CACHE STRING "SYCL compiler provider (Intel or AdaptiveCpp)") +set_property(CACHE DPCTL_SYCL_PROVIDER PROPERTY STRINGS Intel AdaptiveCpp) + set(DPCTL_TARGET_CUDA - "" - CACHE STRING + "" + CACHE STRING "Build DPCTL to target CUDA device. \ Set to a truthy value (e.g., ON, TRUE) to use default architecture (sm_50), \ or to a specific architecture like sm_80." @@ -38,24 +42,11 @@ set(DPCTL_TARGET_HIP CACHE STRING "Build DPCTL to target a HIP device architecture" ) -option( - DPCTL_WITH_REDIST - "Build DPCTL assuming DPC++ redistributable is installed into Python prefix" - OFF) -option( - DPCTL_OFFLOAD_COMPRESS - "Build using offload section compression feature of DPC++ to reduce \ -size of shared object with offloading sections" - OFF -) - -find_package(IntelSYCL REQUIRED PATHS ${CMAKE_SOURCE_DIR}/cmake NO_DEFAULT_PATH) - -set(_dpctl_sycl_target_compile_options) -set(_dpctl_sycl_target_link_options) +# architecture mapping set(_dpctl_sycl_targets) set(_dpctl_cuda_arch) +set(_acpp_targets_list) # Used strictly for AdaptiveCpp if ("x${DPCTL_SYCL_TARGETS}" STREQUAL "x") if (DPCTL_TARGET_CUDA) @@ -69,15 +60,25 @@ if ("x${DPCTL_SYCL_TARGETS}" STREQUAL "x") "Expected 'ON', 'TRUE', 'YES', 'Y', '1', or a CUDA architecture like 'sm_80'." ) endif() - set(_dpctl_sycl_targets "nvidia_gpu_${_dpctl_cuda_arch},spir64-unknown-unknown") + + # Map to Provider Formats + if(DPCTL_SYCL_PROVIDER STREQUAL "Intel") + set(_dpctl_sycl_targets "nvidia_gpu_${_dpctl_cuda_arch},spir64-unknown-unknown") + elseif(DPCTL_SYCL_PROVIDER STREQUAL "AdaptiveCpp") + list(APPEND _acpp_targets_list "cuda:${_dpctl_cuda_arch}") + endif() endif() if (DPCTL_TARGET_HIP) if(DPCTL_TARGET_HIP MATCHES "^gfx") - if(_dpctl_sycl_targets) - set(_dpctl_sycl_targets "amd_gpu_${DPCTL_TARGET_HIP},${_dpctl_sycl_targets}") - else() - set(_dpctl_sycl_targets "amd_gpu_${DPCTL_TARGET_HIP},spir64-unknown-unknown") + if(DPCTL_SYCL_PROVIDER STREQUAL "Intel") + if(_dpctl_sycl_targets) + set(_dpctl_sycl_targets "amd_gpu_${DPCTL_TARGET_HIP},${_dpctl_sycl_targets}") + else() + set(_dpctl_sycl_targets "amd_gpu_${DPCTL_TARGET_HIP},spir64-unknown-unknown") + endif() + elseif(DPCTL_SYCL_PROVIDER STREQUAL "AdaptiveCpp") + list(APPEND _acpp_targets_list "hip:${DPCTL_TARGET_HIP}") endif() else() message(FATAL_ERROR @@ -94,12 +95,38 @@ else() ) endif() set(_dpctl_sycl_targets ${DPCTL_SYCL_TARGETS}) + if(DPCTL_SYCL_PROVIDER STREQUAL "AdaptiveCpp") + set(_acpp_targets_list ${DPCTL_SYCL_TARGETS}) + endif() endif() -if (_dpctl_sycl_targets) - message(STATUS "Compiling for -fsycl-targets=${_dpctl_sycl_targets}") - list(APPEND _dpctl_sycl_target_compile_options -fsycl-targets=${_dpctl_sycl_targets}) - list(APPEND _dpctl_sycl_target_link_options -fsycl-targets=${_dpctl_sycl_targets}) +# Load the SYCL env +set(_dpctl_sycl_target_compile_options) +set(_dpctl_sycl_target_link_options) + +if(DPCTL_SYCL_PROVIDER STREQUAL "Intel") + message(STATUS "Using Intel DPC++ SYCL Provider") + find_package(IntelSYCL REQUIRED PATHS ${CMAKE_SOURCE_DIR}/cmake NO_DEFAULT_PATH) + + if (_dpctl_sycl_targets) + message(STATUS "Compiling for -fsycl-targets=${_dpctl_sycl_targets}") + list(APPEND _dpctl_sycl_target_compile_options -fsycl-targets=${_dpctl_sycl_targets}) + list(APPEND _dpctl_sycl_target_link_options -fsycl-targets=${_dpctl_sycl_targets}) + endif() + +elseif(DPCTL_SYCL_PROVIDER STREQUAL "AdaptiveCpp") + message(STATUS "Using AdaptiveCpp SYCL Provider") + + # AdaptiveCpp reads the ACPP_TARGETS variable directly during find_package + if(_acpp_targets_list) + string(REPLACE ";" ";" ACPP_TARGETS "${_acpp_targets_list}") + set(ACPP_TARGETS "${ACPP_TARGETS}" CACHE STRING "AdaptiveCpp Targets" FORCE) + message(STATUS "Compiling for ACPP_TARGETS=${ACPP_TARGETS}") + endif() + + # This automatically provides the add_sycl_to_target macro + find_package(AdaptiveCpp CONFIG REQUIRED) + endif() add_subdirectory(libsyclinterface) diff --git a/dpctl/CMakeLists.txt b/dpctl/CMakeLists.txt index a24c7443f9..f36b7fd56e 100644 --- a/dpctl/CMakeLists.txt +++ b/dpctl/CMakeLists.txt @@ -111,11 +111,14 @@ function(build_dpctl_ext _trgt _src _dest) Python_add_library(${_trgt} MODULE WITH_SOABI ${_generated_src}) if (BUILD_DPCTL_EXT_SYCL) add_sycl_to_target(TARGET ${_trgt} SOURCES ${_generated_src}) - target_compile_options(${_trgt} PRIVATE -fno-sycl-id-queries-fit-in-int) - target_link_options(${_trgt} PRIVATE -fsycl-device-code-split=per_kernel) - if (DPCTL_OFFLOAD_COMPRESS) - target_link_options(${_trgt} PRIVATE --offload-compress) - endif() + # Only apply DPC++ specific compiler optimizations + if(DPCTL_SYCL_PROVIDER STREQUAL "Intel") + target_compile_options(${_trgt} PRIVATE -fno-sycl-id-queries-fit-in-int) + target_link_options(${_trgt} PRIVATE -fsycl-device-code-split=per_kernel) + if (DPCTL_OFFLOAD_COMPRESS) + target_link_options(${_trgt} PRIVATE --offload-compress) + endif() + endif() if(_dpctl_sycl_targets) # make fat binary target_compile_options( @@ -134,7 +137,9 @@ function(build_dpctl_ext _trgt _src _dest) if (DPCTL_GENERATE_COVERAGE) target_compile_definitions(${_trgt} PRIVATE CYTHON_TRACE=1 CYTHON_TRACE_NOGIL=1) if (BUILD_DPCTL_EXT_SYCL) - target_compile_options(${_trgt} PRIVATE -fno-sycl-use-footer) + if(DPCTL_SYCL_PROVIDER STREQUAL "Intel") + target_compile_options(${_trgt} PRIVATE -fno-sycl-use-footer) + endif() endif() endif() target_link_libraries(${_trgt} PRIVATE DPCTLSyclInterface) diff --git a/dpctl/_host_task_util.hpp b/dpctl/_host_task_util.hpp index 6898893bdd..6ddd9bd7fd 100644 --- a/dpctl/_host_task_util.hpp +++ b/dpctl/_host_task_util.hpp @@ -39,6 +39,10 @@ #include "syclinterface/dpctl_data_types.h" #include "syclinterface/dpctl_sycl_type_casters.hpp" +#ifdef __ADAPTIVECPP__ +#include "apis/include/dpctl_acpp_host_task.hpp" +#endif + DPCTLSyclEventRef async_dec_ref(DPCTLSyclQueueRef QRef, PyObject **obj_array, size_t obj_array_size, @@ -49,12 +53,14 @@ DPCTLSyclEventRef async_dec_ref(DPCTLSyclQueueRef QRef, using dpctl::syclinterface::unwrap; using dpctl::syclinterface::wrap; - sycl::queue *q = unwrap(QRef); + [[maybe_unused]] sycl::queue *q = unwrap(QRef); std::vector obj_vec(obj_array, obj_array + obj_array_size); try { - sycl::event ht_ev = q->submit([&](sycl::handler &cgh) { + sycl::event ht_ev; +#ifndef __ADAPTIVECPP__ + ht_ev = q->submit([&](sycl::handler &cgh) { for (size_t ev_id = 0; ev_id < nDepERefs; ++ev_id) { cgh.depends_on(*(unwrap(depERefs[ev_id]))); } @@ -76,6 +82,35 @@ DPCTLSyclEventRef async_dec_ref(DPCTLSyclQueueRef QRef, } }); }); +#else + // Submit a dummy kernel to track dependencies + ht_ev = q->submit([&](sycl::handler &cgh) { + for (size_t ev_id = 0; ev_id < nDepERefs; ++ev_id) { + cgh.depends_on(*(unwrap(depERefs[ev_id]))); + } + class dpctl_async_decref_dummy; + cgh.single_task([=]() {}); + }); + + // Delegate to our custom thread pool + dpctl::detail::AcppHostTaskPool::get().submit( + ht_ev, [obj_array_size, obj_vec = std::move(obj_vec)]() { + const bool initialized = Py_IsInitialized(); +#if PY_VERSION_HEX < 0x30d0000 + const bool finalizing = _Py_IsFinalizing(); +#else + const bool finalizing = Py_IsFinalizing(); +#endif + if (initialized && !finalizing) { + PyGILState_STATE gstate; + gstate = PyGILState_Ensure(); + for (size_t i = 0; i < obj_array_size; ++i) { + Py_DECREF(obj_vec[i]); + } + PyGILState_Release(gstate); + } + }); +#endif static constexpr int result_ok = 0; diff --git a/dpctl/apis/include/dpctl4pybind11.hpp b/dpctl/apis/include/dpctl4pybind11.hpp index e8d48694f5..62aef0dc6c 100644 --- a/dpctl/apis/include/dpctl4pybind11.hpp +++ b/dpctl/apis/include/dpctl4pybind11.hpp @@ -39,6 +39,10 @@ #include #include +#ifdef __ADAPTIVECPP__ +#include "dpctl_acpp_host_task.hpp" +#endif + namespace py = pybind11; namespace dpctl @@ -517,6 +521,7 @@ template <> struct type_caster DPCTL_TYPE_CASTER(sycl::kernel, _("dpctl.program.SyclKernel")); }; +#ifndef __ADAPTIVECPP__ /* This type caster associates * ``sycl::kernel_bundle`` C++ class with * :class:`dpctl.program.SyclKernelBundle` for the purposes of generation of @@ -560,6 +565,7 @@ struct type_caster> DPCTL_TYPE_CASTER(sycl::kernel_bundle, _("dpctl.program.SyclKernelBundle")); }; +#endif /* This type caster associates * ``sycl::half`` C++ class with Python :class:`float` for the purposes @@ -643,7 +649,7 @@ class usm_memory : public py::object } /*! @brief Create usm_memory object from shared pointer that manages - * lifetime of the USM allocation. + * lifetime of the USM allocation. */ usm_memory(void *usm_ptr, std::size_t nbytes, @@ -827,6 +833,7 @@ sycl::event keep_args_alive(sycl::queue &q, sycl::event host_task_ev; if (n_usm_owners_held > 0) { +#ifndef __ADAPTIVECPP__ host_task_ev = q.submit([&](sycl::handler &cgh) { if (use_depends) { cgh.depends_on(depends); @@ -841,9 +848,28 @@ sycl::event keep_args_alive(sycl::queue &q, // kept alive }); }); +#else + host_task_ev = q.submit([&](sycl::handler &cgh) { + if (use_depends) { + cgh.depends_on(depends); + use_depends = false; + } + else { + cgh.depends_on(host_task_ev); + } + class dpctl_keep_alive_dummy_usm; + cgh.single_task([=]() {}); + }); + + dpctl::detail::AcppHostTaskPool::get().submit( + host_task_ev, [shp_usm = std::move(shp_usm)]() { + // shp_usm destructs naturally here + }); +#endif } if (n_objects_held > 0) { +#ifndef __ADAPTIVECPP__ host_task_ev = q.submit([&](sycl::handler &cgh) { if (use_depends) { cgh.depends_on(depends); @@ -860,6 +886,27 @@ sycl::event keep_args_alive(sycl::queue &q, } }); }); +#else + host_task_ev = q.submit([&](sycl::handler &cgh) { + if (use_depends) { + cgh.depends_on(depends); + use_depends = false; + } + else { + cgh.depends_on(host_task_ev); + } + class dpctl_keep_alive_dummy_arr; + cgh.single_task([=]() {}); + }); + + dpctl::detail::AcppHostTaskPool::get().submit( + host_task_ev, [n_objects_held, shp_arr = std::move(shp_arr)]() { + py::gil_scoped_acquire acquire; + for (std::size_t i = 0; i < n_objects_held; ++i) { + shp_arr[i]->dec_ref(); + } + }); +#endif } return host_task_ev; diff --git a/dpctl/apis/include/dpctl_acpp_host_task.hpp b/dpctl/apis/include/dpctl_acpp_host_task.hpp new file mode 100644 index 0000000000..a1ecf89a39 --- /dev/null +++ b/dpctl/apis/include/dpctl_acpp_host_task.hpp @@ -0,0 +1,133 @@ +//===--- dpctl_acpp_host_task.hpp - AdaptiveCpp host_task emulation -------===// +// +// Data Parallel Control (dpctl) +// +//===----------------------------------------------------------------------===// +/// +/// \file +/// This file implements a fixed-size thread pool to emulate SYCL host_task +/// functionality. It includes Windows-specific teardown logic mirroring +/// Intel LLVM's thread pool to prevent Loader Lock deadlocks during DLL unload. +/// +//===----------------------------------------------------------------------===// + +#pragma once + +#include + +#include +#include +#include +#include +#include +#include +#include +#include + +namespace dpctl +{ +namespace detail +{ + +class AcppHostTaskPool +{ +public: + static AcppHostTaskPool &get() + { + static AcppHostTaskPool instance(4); + return instance; + } + + void submit(sycl::event e, std::function task) + { + { + std::unique_lock lock(queue_mutex_); + tasks_.emplace(std::move(e), std::move(task)); + } + condition_.notify_one(); + } + + AcppHostTaskPool(const AcppHostTaskPool &) = delete; + AcppHostTaskPool &operator=(const AcppHostTaskPool &) = delete; + +private: + AcppHostTaskPool(std::size_t num_threads) : stop_(false) + { + for (std::size_t i = 0; i < num_threads; ++i) { + workers_.emplace_back([this] { + while (true) { + std::pair> item; + { + std::unique_lock lock(this->queue_mutex_); + this->condition_.wait(lock, [this] { + return this->stop_ || !this->tasks_.empty(); + }); + + if (this->stop_ && this->tasks_.empty()) { +#ifdef _WIN32 + // report to main thread we are about to die + { + std::lock_guard exit_lk( + this->win_exit_mutex_); + this->win_exit_count_++; + } + this->win_exit_cv_.notify_one(); +#endif + return; + } + + item = std::move(this->tasks_.front()); + this->tasks_.pop(); + } + + item.first.wait(); + item.second(); + } + }); + } + } + + ~AcppHostTaskPool() + { + { + std::unique_lock lock(queue_mutex_); + stop_ = true; + } + condition_.notify_all(); + +#ifdef _WIN32 + // wait for workers to check in before detach + { + std::unique_lock lock(win_exit_mutex_); + win_exit_cv_.wait( + lock, [this] { return win_exit_count_ == workers_.size(); }); + } + for (std::thread &worker : workers_) { + if (worker.joinable()) { + worker.detach(); + } + } +#else + for (std::thread &worker : workers_) { + if (worker.joinable()) { + worker.join(); + } + } +#endif + } + + std::vector workers_; + std::queue>> tasks_; + std::mutex queue_mutex_; + std::condition_variable condition_; + bool stop_; + +#ifdef _WIN32 + std::mutex win_exit_mutex_; + std::condition_variable win_exit_cv_; + std::size_t win_exit_count_ = 0; +#endif +}; + +} // namespace detail +} // namespace dpctl diff --git a/dpctl/memory/CMakeLists.txt b/dpctl/memory/CMakeLists.txt index 161f67ecba..7b0ab628f6 100644 --- a/dpctl/memory/CMakeLists.txt +++ b/dpctl/memory/CMakeLists.txt @@ -1,4 +1,3 @@ - set(_cy_file ${CMAKE_CURRENT_SOURCE_DIR}/_memory.pyx) get_filename_component(_trgt ${_cy_file} NAME_WLE) build_dpctl_ext(${_trgt} ${_cy_file} "dpctl/memory" SYCL RELATIVE_PATH "..") diff --git a/dpctl/tests/helper/__init__.py b/dpctl/tests/helper/__init__.py index be5f8d2e33..9141268541 100644 --- a/dpctl/tests/helper/__init__.py +++ b/dpctl/tests/helper/__init__.py @@ -22,6 +22,7 @@ has_cpu, has_gpu, has_sycl_platforms, + is_wsl_or_windows, ) __all__ = [ @@ -30,4 +31,5 @@ "has_gpu", "has_sycl_platforms", "get_queue_or_skip", + "is_wsl_or_windows", ] diff --git a/dpctl/tests/helper/_helper.py b/dpctl/tests/helper/_helper.py index 5857b23195..3fa6113062 100644 --- a/dpctl/tests/helper/_helper.py +++ b/dpctl/tests/helper/_helper.py @@ -14,6 +14,9 @@ # See the License for the specific language governing permissions and # limitations under the License. +import platform +import sys + import pytest import dpctl @@ -49,3 +52,12 @@ def get_queue_or_skip(args=()): except dpctl.SyclQueueCreationError: pytest.skip(f"Queue could not be created from {args}") return q + + +def is_wsl_or_windows(): + """ + Skip test on WSL or Windows. + Useful with AdaptiveCpp, which crashes more catastrophically for + unsupported CUDA features. + """ + return sys.platform == "win32" or "microsoft" in platform.release().lower() diff --git a/dpctl/tests/test_sycl_queue.py b/dpctl/tests/test_sycl_queue.py index af0fe71dcd..fc9071359e 100644 --- a/dpctl/tests/test_sycl_queue.py +++ b/dpctl/tests/test_sycl_queue.py @@ -23,7 +23,7 @@ import dpctl -from .helper import create_invalid_capsule +from .helper import create_invalid_capsule, is_wsl_or_windows def test_standard_selectors(device_selector, check): @@ -337,19 +337,22 @@ def test_queue_memops(): q = dpctl.SyclQueue() except dpctl.SyclQueueCreationError: pytest.skip("Failed to create device with supported filter") - from dpctl.memory import MemoryUSMDevice + from dpctl.memory import MemoryUSMShared - m1 = MemoryUSMDevice(512, queue=q) - m2 = MemoryUSMDevice(512, queue=q) + m1 = MemoryUSMShared(512, queue=q) + m2 = MemoryUSMShared(512, queue=q) q.memcpy(m1, m2, 512) - q.prefetch(m1, 512) - q.mem_advise(m1, 512, 0) with pytest.raises(TypeError): q.memcpy(m1, [], 512) with pytest.raises(TypeError): q.memcpy([], m2, 512) - with pytest.raises(TypeError): - q.prefetch([], 512) + + if not is_wsl_or_windows(): + q.prefetch(m1, 512) + with pytest.raises(TypeError): + q.prefetch([], 512) + + q.mem_advise(m1, 512, 0) with pytest.raises(TypeError): q.mem_advise([], 512, 0) diff --git a/dpctl/utils/CMakeLists.txt b/dpctl/utils/CMakeLists.txt index 8f27fb0566..6282ddf783 100644 --- a/dpctl/utils/CMakeLists.txt +++ b/dpctl/utils/CMakeLists.txt @@ -27,12 +27,13 @@ list(APPEND _pybind11_targets ${python_module_name}) set(_linker_options "LINKER:${DPCTL_LDFLAGS}") foreach(python_module_name ${_pybind11_targets}) - target_compile_options(${python_module_name} PRIVATE -fno-sycl-id-queries-fit-in-int) - target_link_options(${python_module_name} PRIVATE -fsycl-device-code-split=per_kernel) - if (DPCTL_OFFLOAD_COMPRESS) - target_link_options(${python_module_name} PRIVATE --offload-compress) + if(DPCTL_SYCL_PROVIDER STREQUAL "Intel") + target_compile_options(${python_module_name} PRIVATE -fsycl-device-code-split=per_kernel) + target_compile_options(${python_module_name} PRIVATE -fno-sycl-id-queries-fit-in-int) + if (DPCTL_OFFLOAD_COMPRESS) + target_link_options(${python_module_name} PRIVATE --offload-compress) + endif() endif() - target_link_options(${python_module_name} PRIVATE ${_linker_options}) if(DPCTL_GENERATE_COVERAGE) if(DPCTL_GENERATE_COVERAGE_FOR_PYBIND11_EXTENSIONS) diff --git a/dpctl/utils/src/device_queries.cpp b/dpctl/utils/src/device_queries.cpp index 05fea90e2a..e70c62218c 100644 --- a/dpctl/utils/src/device_queries.cpp +++ b/dpctl/utils/src/device_queries.cpp @@ -13,9 +13,11 @@ std::uint32_t py_intel_device_id(const sycl::device &d) { static constexpr std::uint32_t device_id_unavailable = 0; +#ifndef __ADAPTIVECPP__ if (d.has(sycl::aspect::ext_intel_device_id)) { return d.get_info(); } +#endif return device_id_unavailable; } @@ -24,9 +26,11 @@ std::uint32_t py_intel_gpu_eu_count(const sycl::device &d) { static constexpr std::uint32_t eu_count_unavailable = 0; +#ifndef __ADAPTIVECPP__ if (d.has(sycl::aspect::ext_intel_gpu_eu_count)) { return d.get_info(); } +#endif return eu_count_unavailable; } @@ -35,10 +39,12 @@ std::uint32_t py_intel_gpu_hw_threads_per_eu(const sycl::device &d) { static constexpr std::uint32_t thread_count_unavailable = 0; +#ifndef __ADAPTIVECPP__ if (d.has(sycl::aspect::ext_intel_gpu_hw_threads_per_eu)) { return d .get_info(); } +#endif return thread_count_unavailable; } @@ -47,9 +53,11 @@ std::uint32_t py_intel_gpu_eu_simd_width(const sycl::device &d) { static constexpr std::uint32_t width_unavailable = 0; +#ifndef __ADAPTIVECPP__ if (d.has(sycl::aspect::ext_intel_gpu_eu_simd_width)) { return d.get_info(); } +#endif return width_unavailable; } @@ -58,9 +66,11 @@ std::uint32_t py_intel_gpu_slices(const sycl::device &d) { static constexpr std::uint32_t count_unavailable = 0; +#ifndef __ADAPTIVECPP__ if (d.has(sycl::aspect::ext_intel_gpu_slices)) { return d.get_info(); } +#endif return count_unavailable; } @@ -69,10 +79,12 @@ std::uint32_t py_intel_gpu_subslices_per_slice(const sycl::device &d) { static constexpr std::uint32_t count_unavailable = 0; +#ifndef __ADAPTIVECPP__ if (d.has(sycl::aspect::ext_intel_gpu_subslices_per_slice)) { return d.get_info< sycl::ext::intel::info::device::gpu_subslices_per_slice>(); } +#endif return count_unavailable; } @@ -81,10 +93,12 @@ std::uint32_t py_intel_gpu_eu_count_per_subslice(const sycl::device &d) { static constexpr std::uint32_t count_unavailable = 0; +#ifndef __ADAPTIVECPP__ if (d.has(sycl::aspect::ext_intel_gpu_eu_count_per_subslice)) { return d.get_info< sycl::ext::intel::info::device::gpu_eu_count_per_subslice>(); } +#endif return count_unavailable; } @@ -93,9 +107,11 @@ std::uint64_t py_intel_max_mem_bandwidth(const sycl::device &d) { static constexpr std::uint64_t bandwidth_unavailable = 0; +#ifndef __ADAPTIVECPP__ if (d.has(sycl::aspect::ext_intel_max_mem_bandwidth)) { return d.get_info(); } +#endif return bandwidth_unavailable; } @@ -104,9 +120,11 @@ std::uint64_t py_intel_free_memory(const sycl::device &d) { static constexpr std::uint64_t free_memory_unavailable = 0; +#ifndef __ADAPTIVECPP__ if (d.has(sycl::aspect::ext_intel_free_memory)) { return d.get_info(); } +#endif return free_memory_unavailable; } @@ -114,9 +132,11 @@ std::uint32_t py_intel_memory_clock_rate(const sycl::device &d) { static constexpr std::uint32_t rate_unavailable = 0; +#ifndef __ADAPTIVECPP__ if (d.has(sycl::aspect::ext_intel_memory_clock_rate)) { return d.get_info(); } +#endif return rate_unavailable; } @@ -124,9 +144,11 @@ std::uint32_t py_intel_memory_bus_width(const sycl::device &d) { static constexpr std::uint32_t width_unavailable = 0; +#ifndef __ADAPTIVECPP__ if (d.has(sycl::aspect::ext_intel_memory_bus_width)) { return d.get_info(); } +#endif return width_unavailable; } diff --git a/libsyclinterface/CMakeLists.txt b/libsyclinterface/CMakeLists.txt index e71080602a..db0e682fc8 100644 --- a/libsyclinterface/CMakeLists.txt +++ b/libsyclinterface/CMakeLists.txt @@ -11,8 +11,14 @@ set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} "${CMAKE_CURRENT_SOURCE_DIR}/cmake/mo find_package(Git REQUIRED) -if(NOT DEFINED IntelSYCL_FOUND OR NOT IntelSYCL_FOUND) - find_package(IntelSYCL REQUIRED) +if(DPCTL_SYCL_PROVIDER STREQUAL "Intel") + if(NOT DEFINED IntelSYCL_FOUND OR NOT IntelSYCL_FOUND) + find_package(IntelSYCL REQUIRED) + endif() +elseif(DPCTL_SYCL_PROVIDER STREQUAL "AdaptiveCpp") + if(NOT DEFINED AdaptiveCpp_FOUND OR NOT AdaptiveCpp_FOUND) + find_package(AdaptiveCpp CONFIG REQUIRED) + endif() endif() # Option to turn on support for creating Level Zero interoperability programs @@ -52,12 +58,14 @@ set(LIBCL_SET_LOADER_FILENAME "" CACHE STRING "User-provided OpenCL ICD Loader f set(LIBCL_DEFAULT_LOADER_FILENAME "libOpenCL.so.1" CACHE STRING "Default OpenCL ICD Loader filename") -# Minimum version requirement only when oneAPI dpcpp is used. -if(DPCTL_DPCPP_FROM_ONEAPI) - find_package(IntelSyclCompiler 2021.3.0 REQUIRED) -else() - find_package(IntelSyclCompiler REQUIRED) -endif() +# SYCL RT driver discovery +if(DPCTL_SYCL_PROVIDER STREQUAL "Intel") + # Minimum version requirement only when oneAPI dpcpp is used. + if(DPCTL_DPCPP_FROM_ONEAPI) + find_package(IntelSyclCompiler 2021.3.0 REQUIRED) + else() + find_package(IntelSyclCompiler REQUIRED) + endif() if(DPCTL_ENABLE_L0_PROGRAM_CREATION) set(DPCTL_ENABLE_L0_PROGRAM_CREATION 1) @@ -325,6 +333,11 @@ if(_dpctl_sycl_targets) ) endif() +# Fake Intel's compiler version macro so dpctl's static_asserts pass +if(DPCTL_SYCL_PROVIDER STREQUAL "AdaptiveCpp") + target_compile_definitions(DPCTLSyclInterface PRIVATE __SYCL_COMPILER_VERSION=20240000) +endif() + if(DPCTL_GENERATE_COVERAGE) target_link_options(DPCTLSyclInterface PRIVATE -fprofile-instr-generate -fcoverage-mapping diff --git a/libsyclinterface/helper/include/dpctl_utils_helper.h b/libsyclinterface/helper/include/dpctl_utils_helper.h index c441a40aec..9a3b73cdad 100644 --- a/libsyclinterface/helper/include/dpctl_utils_helper.h +++ b/libsyclinterface/helper/include/dpctl_utils_helper.h @@ -179,6 +179,7 @@ DPCTL_API DPCTLPartitionAffinityDomainType DPCTL_SyclPartitionAffinityDomainToDPCTLType( sycl::info::partition_affinity_domain PartitionAffinityDomain); +#ifndef __ADAPTIVECPP__ /*! * @brief Converts a sycl::info::fp_config enum value to corresponding * DPCTLFPConfigType enum value. @@ -274,6 +275,7 @@ DPCTL_DPCTLPeerAccessTypeToSycl(DPCTLPeerAccessType PeerAccessTy); DPCTL_API DPCTLPeerAccessType DPCTL_SyclPeerAccessToDPCTLType(sycl::ext::oneapi::peer_access PeerAccess); +#endif /*! * @brief Gives the index of the given device with respective to all the other diff --git a/libsyclinterface/helper/source/dpctl_utils_helper.cpp b/libsyclinterface/helper/source/dpctl_utils_helper.cpp index c03887a2b1..b976fa78cc 100644 --- a/libsyclinterface/helper/source/dpctl_utils_helper.cpp +++ b/libsyclinterface/helper/source/dpctl_utils_helper.cpp @@ -84,6 +84,7 @@ info::device_type DPCTL_StrToDeviceType(const std::string &devTyStr) backend DPCTL_DPCTLBackendTypeToSyclBackend(DPCTLSyclBackendType BeTy) { switch (BeTy) { +#ifndef __ADAPTIVECPP__ case DPCTLSyclBackendType::DPCTL_CUDA: return backend::ext_oneapi_cuda; case DPCTLSyclBackendType::DPCTL_LEVEL_ZERO: @@ -94,6 +95,14 @@ backend DPCTL_DPCTLBackendTypeToSyclBackend(DPCTLSyclBackendType BeTy) return backend::all; case DPCTLSyclBackendType::DPCTL_HIP: return backend::ext_oneapi_hip; +#else + case DPCTLSyclBackendType::DPCTL_CUDA: + return backend::cuda; + case DPCTLSyclBackendType::DPCTL_HIP: + return backend::hip; + case DPCTLSyclBackendType::DPCTL_OPENCL: + return backend::ocl; +#endif default: throw std::runtime_error("Unsupported backend type"); } @@ -102,6 +111,7 @@ backend DPCTL_DPCTLBackendTypeToSyclBackend(DPCTLSyclBackendType BeTy) DPCTLSyclBackendType DPCTL_SyclBackendToDPCTLBackendType(backend B) { switch (B) { +#ifndef __ADAPTIVECPP__ case backend::ext_oneapi_cuda: return DPCTLSyclBackendType::DPCTL_CUDA; case backend::ext_oneapi_level_zero: @@ -110,6 +120,14 @@ DPCTLSyclBackendType DPCTL_SyclBackendToDPCTLBackendType(backend B) return DPCTLSyclBackendType::DPCTL_OPENCL; case backend::ext_oneapi_hip: return DPCTLSyclBackendType::DPCTL_HIP; +#else + case backend::cuda: + return DPCTLSyclBackendType::DPCTL_CUDA; + case backend::hip: + return DPCTLSyclBackendType::DPCTL_HIP; + case backend::ocl: + return DPCTLSyclBackendType::DPCTL_OPENCL; +#endif default: return DPCTLSyclBackendType::DPCTL_UNKNOWN_BACKEND; } @@ -219,12 +237,14 @@ std::string DPCTL_AspectToStr(aspect aspectTy) case aspect::emulated: ss << "emulated"; break; +#ifndef __ADAPTIVECPP__ case aspect::ext_oneapi_is_component: ss << "is_component"; break; case aspect::ext_oneapi_is_composite: ss << "is_composite"; break; +#endif #ifdef SYCL_EXT_ONEAPI_INTER_PROCESS_COMMUNICATION case aspect::ext_oneapi_ipc_memory: ss << "ext_oneapi_ipc_memory"; @@ -299,12 +319,14 @@ aspect DPCTL_StrToAspectType(const std::string &aspectTyStr) else if (aspectTyStr == "emulated") { aspectTy = aspect::emulated; } +#ifndef __ADAPTIVECPP__ else if (aspectTyStr == "is_component") { aspectTy = aspect::ext_oneapi_is_component; } else if (aspectTyStr == "is_composite") { aspectTy = aspect::ext_oneapi_is_composite; } +#endif #ifdef SYCL_EXT_ONEAPI_INTER_PROCESS_COMMUNICATION else if (aspectTyStr == "ext_oneapi_ipc_memory") { aspectTy = aspect::ext_oneapi_ipc_memory; @@ -358,10 +380,16 @@ aspect DPCTL_DPCTLAspectTypeToSyclAspect(DPCTLSyclAspectType AspectTy) return aspect::host_debuggable; case DPCTLSyclAspectType::emulated: return aspect::emulated; +#ifndef __ADAPTIVECPP__ case DPCTLSyclAspectType::is_component: return aspect::ext_oneapi_is_component; case DPCTLSyclAspectType::is_composite: return aspect::ext_oneapi_is_composite; +#else + case DPCTLSyclAspectType::is_component: + case DPCTLSyclAspectType::is_composite: + throw std::runtime_error("Aspect type unsupported in AdaptiveCpp"); +#endif #ifdef SYCL_EXT_ONEAPI_INTER_PROCESS_COMMUNICATION case DPCTLSyclAspectType::ext_oneapi_ipc_memory: return aspect::ext_oneapi_ipc_memory; @@ -412,10 +440,12 @@ DPCTLSyclAspectType DPCTL_SyclAspectToDPCTLAspectType(aspect Aspect) return DPCTLSyclAspectType::host_debuggable; case aspect::emulated: return DPCTLSyclAspectType::emulated; +#ifndef __ADAPTIVECPP__ case aspect::ext_oneapi_is_composite: return DPCTLSyclAspectType::is_composite; case aspect::ext_oneapi_is_component: return DPCTLSyclAspectType::is_component; +#endif #ifdef SYCL_EXT_ONEAPI_INTER_PROCESS_COMMUNICATION case aspect::ext_oneapi_ipc_memory: return DPCTLSyclAspectType::ext_oneapi_ipc_memory; @@ -565,6 +595,7 @@ DPCTLPartitionPropertyType DPCTL_SyclPartitionPropertyToDPCTLType( } } +#ifndef __ADAPTIVECPP__ ext::oneapi::peer_access DPCTL_DPCTLPeerAccessTypeToSycl(DPCTLPeerAccessType PeerAccessTy) { @@ -590,6 +621,7 @@ DPCTL_SyclPeerAccessToDPCTLType(ext::oneapi::peer_access PeerAccess) throw std::runtime_error("Unsupported peer_access type"); } } +#endif int64_t DPCTL_GetRelativeDeviceId(const device &Device) { @@ -617,6 +649,7 @@ std::string DPCTL_GetDeviceFilterString(const device &Device) auto be = Device.get_platform().get_backend(); +#ifndef __ADAPTIVECPP__ switch (be) { case backend::ext_oneapi_level_zero: ss << "level_zero"; @@ -633,6 +666,21 @@ std::string DPCTL_GetDeviceFilterString(const device &Device) default: ss << "unknown"; }; +#else + switch (be) { + case backend::cuda: + ss << "cuda"; + break; + case backend::hip: + ss << "hip"; + break; + case backend::ocl: + ss << "opencl"; + break; + default: + ss << "unknown"; + }; +#endif ss << filter_string_separator; ss << DPCTL_DeviceTypeToStr(Device.get_info()); diff --git a/libsyclinterface/include/syclinterface/dpctl_device_selection.hpp b/libsyclinterface/include/syclinterface/dpctl_device_selection.hpp index 7af8ef5e86..c95cddb0d6 100644 --- a/libsyclinterface/include/syclinterface/dpctl_device_selection.hpp +++ b/libsyclinterface/include/syclinterface/dpctl_device_selection.hpp @@ -1,7 +1,6 @@ //===-- dpctl_device_selection.h - // Device selector class declaration --*-C++-*- =// // -// // Data Parallel Control (dpctl) // // Copyright 2022 Intel Corporation @@ -29,6 +28,7 @@ #include "Support/DllExport.h" #include +#include namespace dpctl { @@ -78,7 +78,11 @@ class DPCTL_API dpctl_filter_selector : public dpctl_device_selector int operator()(const sycl::device &d) const override; private: +#ifndef __ADAPTIVECPP__ sycl::ext::oneapi::filter_selector _impl; +#else + std::string _impl; // Standard string fallback for AdaptiveCpp +#endif }; } // namespace syclinterface diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_type_casters.hpp b/libsyclinterface/include/syclinterface/dpctl_sycl_type_casters.hpp index 4d4156d92a..6c0e2207ed 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_type_casters.hpp +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_type_casters.hpp @@ -64,9 +64,14 @@ DEFINE_SIMPLE_CONVERSION_FUNCTIONS(void, DPCTLSyclUSMRef) DEFINE_SIMPLE_CONVERSION_FUNCTIONS(sycl::platform, DPCTLSyclPlatformRef) DEFINE_SIMPLE_CONVERSION_FUNCTIONS(sycl::event, DPCTLSyclEventRef) DEFINE_SIMPLE_CONVERSION_FUNCTIONS(sycl::kernel, DPCTLSyclKernelRef) + +#ifndef __ADAPTIVECPP__ DEFINE_SIMPLE_CONVERSION_FUNCTIONS( sycl::kernel_bundle, DPCTLSyclKernelBundleRef) +#endif + +#include "dpctl_sycl_device_manager.h" #include "dpctl_sycl_device_manager.h" DEFINE_SIMPLE_CONVERSION_FUNCTIONS(std::vector, diff --git a/libsyclinterface/source/dpctl_device_selection.cpp b/libsyclinterface/source/dpctl_device_selection.cpp index 2e216162f4..b99ac7779e 100644 --- a/libsyclinterface/source/dpctl_device_selection.cpp +++ b/libsyclinterface/source/dpctl_device_selection.cpp @@ -31,9 +31,11 @@ namespace { +#ifndef __ADAPTIVECPP__ static_assert(__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_VERSION_REQUIRED, "The compiler does not meet minimum version requirement"); -} +#endif +} // namespace namespace dpctl { @@ -68,7 +70,13 @@ int dpctl_cpu_selector::operator()(const sycl::device &d) const int dpctl_filter_selector::operator()(const sycl::device &d) const { +#ifndef __ADAPTIVECPP__ return _impl(d); +#else + // AdaptiveCpp does not natively support ext::oneapi::filter_selector + // Fall back to default behavior + return sycl::default_selector_v(d); +#endif } } // namespace syclinterface diff --git a/libsyclinterface/source/dpctl_sycl_context_interface.cpp b/libsyclinterface/source/dpctl_sycl_context_interface.cpp index c27cd9a343..fc9210c122 100644 --- a/libsyclinterface/source/dpctl_sycl_context_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_context_interface.cpp @@ -37,8 +37,10 @@ using namespace sycl; namespace { +#ifndef __ADAPTIVECPP__ static_assert(__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_VERSION_REQUIRED, "The compiler does not meet minimum version requirement"); +#endif using namespace dpctl::syclinterface; } // end of anonymous namespace @@ -186,6 +188,7 @@ DPCTLContext_GetBackend(__dpctl_keep const DPCTLSyclContextRef CtxRef) auto BE = unwrap(CtxRef)->get_platform().get_backend(); +#ifndef __ADAPTIVECPP__ switch (BE) { case backend::opencl: return DPCTL_OPENCL; @@ -198,6 +201,16 @@ DPCTLContext_GetBackend(__dpctl_keep const DPCTLSyclContextRef CtxRef) default: return DPCTL_UNKNOWN_BACKEND; } +#else + switch (BE) { + case backend::cuda: + return DPCTL_CUDA; + case backend::hip: + return DPCTL_HIP; + default: + return DPCTL_UNKNOWN_BACKEND; + } +#endif } size_t DPCTLContext_Hash(__dpctl_keep const DPCTLSyclContextRef CtxRef) diff --git a/libsyclinterface/source/dpctl_sycl_device_interface.cpp b/libsyclinterface/source/dpctl_sycl_device_interface.cpp index d7e1b092ef..c65a9fdef6 100644 --- a/libsyclinterface/source/dpctl_sycl_device_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_device_interface.cpp @@ -42,9 +42,10 @@ using namespace sycl; namespace { - +#ifndef __ADAPTIVECPP__ static_assert(__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_VERSION_REQUIRED, "The compiler does not meet minimum version requirement"); +#endif using namespace dpctl::syclinterface; @@ -62,11 +63,12 @@ DPCTLDevice__GetMaxWorkItemSizes(__dpctl_keep const DPCTLSyclDeviceRef DRef) auto D = unwrap(DRef); if (D) { try { -#if __SYCL_COMPILER_VERSION >= __SYCL_COMPILER_MAX_WORK_ITEM_SIZE_THRESHOLD +#if defined(__ADAPTIVECPP__) || \ + (__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_MAX_WORK_ITEM_SIZE_THRESHOLD) auto id_sizes = D->get_info>(); #else - auto id_sizes = D->get_info(); + auto id_sizes = D->get_info>(); #endif sizes = new size_t[dim]; for (auto i = 0ul; i < dim; ++i) { @@ -854,6 +856,7 @@ DPCTLDevice_GetComponentDevices(__dpctl_keep const DPCTLSyclDeviceRef DRef) using vecTy = std::vector; vecTy *ComponentDevicesVectorPtr = nullptr; if (DRef) { +#ifndef __ADAPTIVECPP__ auto D = unwrap(DRef); try { auto componentDevices = @@ -870,6 +873,7 @@ DPCTLDevice_GetComponentDevices(__dpctl_keep const DPCTLSyclDeviceRef DRef) error_handler(e, __FILE__, __func__, __LINE__); return nullptr; } +#endif } return wrap(ComponentDevicesVectorPtr); } @@ -877,6 +881,7 @@ DPCTLDevice_GetComponentDevices(__dpctl_keep const DPCTLSyclDeviceRef DRef) __dpctl_give DPCTLSyclDeviceRef DPCTLDevice_GetCompositeDevice(__dpctl_keep const DPCTLSyclDeviceRef DRef) { +#ifndef __ADAPTIVECPP__ auto D = unwrap(DRef); if (D) { bool is_component = false; @@ -900,8 +905,13 @@ DPCTLDevice_GetCompositeDevice(__dpctl_keep const DPCTLSyclDeviceRef DRef) } else return nullptr; +#else + return nullptr; +#endif } +#ifndef __ADAPTIVECPP__ + static inline bool _CallPeerAccess(device dev, device peer) { auto BE1 = dev.get_backend(); @@ -921,11 +931,14 @@ static inline bool _CallPeerAccess(device dev, device peer) return false; } +#endif /* #ifndef __ADAPTIVECPP__ */ + bool DPCTLDevice_CanAccessPeer(__dpctl_keep const DPCTLSyclDeviceRef DRef, __dpctl_keep const DPCTLSyclDeviceRef PDRef, DPCTLPeerAccessType PT) { bool canAccess = false; +#ifndef __ADAPTIVECPP__ auto D = unwrap(DRef); auto PD = unwrap(PDRef); if (D && PD) { @@ -938,12 +951,14 @@ bool DPCTLDevice_CanAccessPeer(__dpctl_keep const DPCTLSyclDeviceRef DRef, } } } +#endif return canAccess; } void DPCTLDevice_EnablePeerAccess(__dpctl_keep const DPCTLSyclDeviceRef DRef, __dpctl_keep const DPCTLSyclDeviceRef PDRef) { +#ifndef __ADAPTIVECPP__ auto D = unwrap(DRef); auto PD = unwrap(PDRef); if (D && PD) { @@ -959,12 +974,14 @@ void DPCTLDevice_EnablePeerAccess(__dpctl_keep const DPCTLSyclDeviceRef DRef, __func__, __LINE__); } } +#endif return; } void DPCTLDevice_DisablePeerAccess(__dpctl_keep const DPCTLSyclDeviceRef DRef, __dpctl_keep const DPCTLSyclDeviceRef PDRef) { +#ifndef __ADAPTIVECPP__ auto D = unwrap(DRef); auto PD = unwrap(PDRef); if (D && PD) { @@ -980,6 +997,7 @@ void DPCTLDevice_DisablePeerAccess(__dpctl_keep const DPCTLSyclDeviceRef DRef, __func__, __LINE__); } } +#endif return; } diff --git a/libsyclinterface/source/dpctl_sycl_device_manager.cpp b/libsyclinterface/source/dpctl_sycl_device_manager.cpp index 2d647c33d0..aadabaacac 100644 --- a/libsyclinterface/source/dpctl_sycl_device_manager.cpp +++ b/libsyclinterface/source/dpctl_sycl_device_manager.cpp @@ -376,12 +376,14 @@ __dpctl_give DPCTLDeviceVectorRef DPCTLDeviceMgr_GetCompositeDevices() } try { +#ifndef __ADAPTIVECPP__ auto composite_devices = ext::oneapi::experimental::get_composite_devices(); Devices->reserve(composite_devices.size()); for (const auto &CDev : composite_devices) { Devices->emplace_back(wrap(new device(std::move(CDev)))); } +#endif return wrap(Devices); } catch (std::exception const &e) { delete Devices; diff --git a/libsyclinterface/source/dpctl_sycl_event_interface.cpp b/libsyclinterface/source/dpctl_sycl_event_interface.cpp index 5e2b924e28..641d687971 100644 --- a/libsyclinterface/source/dpctl_sycl_event_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_event_interface.cpp @@ -126,7 +126,12 @@ DPCTLSyclBackendType DPCTLEvent_GetBackend(__dpctl_keep DPCTLSyclEventRef ERef) DPCTLSyclBackendType BTy = DPCTLSyclBackendType::DPCTL_UNKNOWN_BACKEND; auto E = unwrap(ERef); if (E) { +#ifndef __ADAPTIVECPP__ BTy = DPCTL_SyclBackendToDPCTLBackendType(E->get_backend()); +#else + // AdaptiveCpp does not natively expose get_backend() on sycl::event + BTy = DPCTLSyclBackendType::DPCTL_UNKNOWN_BACKEND; +#endif } else { error_handler("Backend cannot be looked up for a NULL event.", __FILE__, diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp index ec94e5a077..c987a33895 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -25,20 +25,25 @@ /// //===----------------------------------------------------------------------===// -#include "dpctl_sycl_kernel_bundle_interface.h" +#ifndef __ADAPTIVECPP__ +#include /* OpenCL headers */ +#include +#endif + #include "Config/dpctl_config.h" #include "dpctl_dynamic_lib_helper.h" #include "dpctl_error_handlers.h" +#include "dpctl_sycl_kernel_bundle_interface.h" #include "dpctl_sycl_type_casters.hpp" #include /* OpenCL headers */ #include #include #include -#include #include /* Sycl headers */ #include #ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION +#ifndef __ADAPTIVECPP__ // Note: include ze_api.h before level_zero.hpp. Make sure clang-format does // not reorder the includes. // clang-format off @@ -46,6 +51,7 @@ #include // clang-format on #endif +#endif using namespace sycl; @@ -56,6 +62,8 @@ static_assert(__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_VERSION_REQUIRED, using namespace dpctl::syclinterface; +#ifndef __ADAPTIVECPP__ + #ifdef __linux__ static const char *clLoaderName = DPCTL_LIBCL_LOADER_FILENAME; static const int clLibLoadFlags = RTLD_NOLOAD | RTLD_NOW | RTLD_LOCAL; @@ -678,6 +686,8 @@ bool _HasKernel_ze_impl(const kernel_bundle &kb, #endif /* #ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION */ +#endif /* #ifndef __ADAPTIVECPP__ */ + } /* end of anonymous namespace */ __dpctl_give DPCTLSyclKernelBundleRef @@ -689,6 +699,7 @@ DPCTLKernelBundle_CreateFromSpirv(__dpctl_keep const DPCTLSyclContextRef CtxRef, size_t NumSpecConsts, const DPCTLSpecConst *SpecConsts) { +#ifndef __ADAPTIVECPP__ DPCTLSyclKernelBundleRef KBRef = nullptr; if (!CtxRef) { error_handler("Cannot create program from SPIR-V as the supplied SYCL " @@ -737,6 +748,12 @@ DPCTLKernelBundle_CreateFromSpirv(__dpctl_keep const DPCTLSyclContextRef CtxRef, return nullptr; } return KBRef; +#else + error_handler( + "Dynamic kernel bundle creation is not supported in AdaptiveCpp", + __FILE__, __func__, __LINE__, error_level::error); + return nullptr; +#endif } __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromOCLSource( @@ -745,6 +762,7 @@ __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromOCLSource( __dpctl_keep const char *Source, __dpctl_keep const char *CompileOpts) { +#ifndef __ADAPTIVECPP__ context *SyclCtx = nullptr; device *SyclDev = nullptr; @@ -786,12 +804,19 @@ __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromOCLSource( __FILE__, __func__, __LINE__); return nullptr; } +#else + error_handler( + "Dynamic kernel bundle creation is not supported in AdaptiveCpp", + __FILE__, __func__, __LINE__, error_level::error); + return nullptr; +#endif } __dpctl_give DPCTLSyclKernelRef DPCTLKernelBundle_GetKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, __dpctl_keep const char *KernelName) { +#ifndef __ADAPTIVECPP__ if (!KBRef) { error_handler("Input KBRef is nullptr", __FILE__, __func__, __LINE__); return nullptr; @@ -816,11 +841,17 @@ DPCTLKernelBundle_GetKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, error_handler(os.str(), __FILE__, __func__, __LINE__); return nullptr; } +#else + error_handler("Dynamic kernel querying is not supported in AdaptiveCpp", + __FILE__, __func__, __LINE__, error_level::error); + return nullptr; +#endif } bool DPCTLKernelBundle_HasKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, __dpctl_keep const char *KernelName) { +#ifndef __ADAPTIVECPP__ if (!KBRef) { error_handler("Input KBRef is nullptr", __FILE__, __func__, __LINE__); return false; @@ -846,16 +877,24 @@ bool DPCTLKernelBundle_HasKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, error_handler(os.str(), __FILE__, __func__, __LINE__); return false; } +#else + error_handler("Dynamic kernel querying is not supported in AdaptiveCpp", + __FILE__, __func__, __LINE__, error_level::error); + return false; +#endif } void DPCTLKernelBundle_Delete(__dpctl_take DPCTLSyclKernelBundleRef KBRef) { +#ifndef __ADAPTIVECPP__ delete unwrap>(KBRef); +#endif } __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_Copy(__dpctl_keep const DPCTLSyclKernelBundleRef KBRef) { +#ifndef __ADAPTIVECPP__ auto Bundle = unwrap>(KBRef); if (!Bundle) { error_handler( @@ -871,4 +910,9 @@ DPCTLKernelBundle_Copy(__dpctl_keep const DPCTLSyclKernelBundleRef KBRef) error_handler(e, __FILE__, __func__, __LINE__); return nullptr; } +#else + error_handler("Kernel bundle copies are not supported in AdaptiveCpp", + __FILE__, __func__, __LINE__, error_level::error); + return nullptr; +#endif } diff --git a/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp index e035ea7000..b98a328868 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp @@ -89,7 +89,11 @@ size_t DPCTLKernel_GetWorkGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef) } auto sycl_kern = unwrap(KRef); +#ifndef __ADAPTIVECPP__ auto devs = sycl_kern->get_kernel_bundle().get_devices(); +#else + auto devs = sycl_kern->get_context().get_devices(); +#endif if (devs.empty()) { error_handler("Input DPCTKSyclKernelRef has no associated device.", __FILE__, __func__, __LINE__); @@ -110,7 +114,11 @@ size_t DPCTLKernel_GetPreferredWorkGroupSizeMultiple( } auto sycl_kern = unwrap(KRef); +#ifndef __ADAPTIVECPP__ auto devs = sycl_kern->get_kernel_bundle().get_devices(); +#else + auto devs = sycl_kern->get_context().get_devices(); +#endif if (devs.empty()) { error_handler("Input DPCTKSyclKernelRef has no associated device.", __FILE__, __func__, __LINE__); @@ -131,7 +139,11 @@ size_t DPCTLKernel_GetPrivateMemSize(__dpctl_keep const DPCTLSyclKernelRef KRef) } auto sycl_kern = unwrap(KRef); +#ifndef __ADAPTIVECPP__ auto devs = sycl_kern->get_kernel_bundle().get_devices(); +#else + auto devs = sycl_kern->get_context().get_devices(); +#endif if (devs.empty()) { error_handler("Input DPCTKSyclKernelRef has no associated device.", __FILE__, __func__, __LINE__); @@ -153,7 +165,11 @@ DPCTLKernel_GetMaxNumSubGroups(__dpctl_keep const DPCTLSyclKernelRef KRef) } auto sycl_kern = unwrap(KRef); +#ifndef __ADAPTIVECPP__ auto devs = sycl_kern->get_kernel_bundle().get_devices(); +#else + auto devs = sycl_kern->get_context().get_devices(); +#endif if (devs.empty()) { error_handler("Input DPCTKSyclKernelRef has no associated device.", __FILE__, __func__, __LINE__); @@ -175,7 +191,11 @@ DPCTLKernel_GetMaxSubGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef) } auto sycl_kern = unwrap(KRef); +#ifndef __ADAPTIVECPP__ auto devs = sycl_kern->get_kernel_bundle().get_devices(); +#else + auto devs = sycl_kern->get_context().get_devices(); +#endif if (devs.empty()) { error_handler("Input DPCTKSyclKernelRef has no associated device.", __FILE__, __func__, __LINE__); @@ -197,7 +217,11 @@ DPCTLKernel_GetCompileNumSubGroups(__dpctl_keep const DPCTLSyclKernelRef KRef) } auto sycl_kern = unwrap(KRef); +#ifndef __ADAPTIVECPP__ auto devs = sycl_kern->get_kernel_bundle().get_devices(); +#else + auto devs = sycl_kern->get_context().get_devices(); +#endif if (devs.empty()) { error_handler("Input DPCTKSyclKernelRef has no associated device.", __FILE__, __func__, __LINE__); @@ -220,7 +244,11 @@ DPCTLKernel_GetCompileSubGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef) } auto sycl_kern = unwrap(KRef); +#ifndef __ADAPTIVECPP__ auto devs = sycl_kern->get_kernel_bundle().get_devices(); +#else + auto devs = sycl_kern->get_context().get_devices(); +#endif if (devs.empty()) { error_handler("Input DPCTKSyclKernelRef has no associated device.", __FILE__, __func__, __LINE__); diff --git a/libsyclinterface/source/dpctl_sycl_platform_interface.cpp b/libsyclinterface/source/dpctl_sycl_platform_interface.cpp index d19d9844bb..c0382d4a9f 100644 --- a/libsyclinterface/source/dpctl_sycl_platform_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_platform_interface.cpp @@ -342,12 +342,14 @@ DPCTLPlatform_GetCompositeDevices(__dpctl_keep const DPCTLSyclPlatformRef PRef) } try { +#ifndef __ADAPTIVECPP__ auto composite_devices = P->ext_oneapi_get_composite_devices(); DevicesVectorPtr->reserve(composite_devices.size()); for (const auto &Dev : composite_devices) { DevicesVectorPtr->emplace_back( wrap(new device(std::move(Dev)))); } +#endif return wrap(DevicesVectorPtr); } catch (std::exception const &e) { delete DevicesVectorPtr; diff --git a/libsyclinterface/source/dpctl_sycl_platform_manager.cpp b/libsyclinterface/source/dpctl_sycl_platform_manager.cpp index e490bfef9d..57a440561f 100644 --- a/libsyclinterface/source/dpctl_sycl_platform_manager.cpp +++ b/libsyclinterface/source/dpctl_sycl_platform_manager.cpp @@ -74,7 +74,20 @@ std::string platform_print_info_impl(const platform &p, size_t verbosity) << p.get_info() << _endl << std::setw(4) << " " << std::left << std::setw(12) << "Vendor" << vendor << _endl << std::setw(4) << " " << std::left << std::setw(12) << "Backend"; + +#ifndef __ADAPTIVECPP__ ss << p.get_backend(); +#else + auto be = p.get_backend(); + if (be == sycl::backend::cuda) + ss << "cuda"; + else if (be == sycl::backend::hip) + ss << "hip"; + else if (be == sycl::backend::ocl) + ss << "opencl"; + else + ss << static_cast(be); +#endif ss << _endl; // Get number of devices on the platform diff --git a/libsyclinterface/source/dpctl_sycl_queue_interface.cpp b/libsyclinterface/source/dpctl_sycl_queue_interface.cpp index 575a3c13fa..7e7d2d4eca 100644 --- a/libsyclinterface/source/dpctl_sycl_queue_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_queue_interface.cpp @@ -50,6 +50,7 @@ using namespace sycl; +#ifndef __ADAPTIVECPP__ #define SET_LOCAL_ACCESSOR_ARG(CGH, NDIM, ARGTY, R, IDX) \ do { \ switch ((ARGTY)) { \ @@ -119,6 +120,15 @@ using namespace sycl; return false; \ } \ } while (0); +#else +#define SET_LOCAL_ACCESSOR_ARG(CGH, NDIM, ARGTY, R, IDX) \ + do { \ + error_handler("Local accessors and dynamic kernel args are not " \ + "supported in AdaptiveCpp.", \ + __FILE__, __func__, __LINE__, error_level::error); \ + return false; \ + } while (0); +#endif namespace { @@ -133,6 +143,8 @@ typedef struct complex std::uint64_t imag; } complexNumber; +#ifndef __ADAPTIVECPP__ + void set_dependent_events(handler &cgh, __dpctl_keep const DPCTLSyclEventRef *DepEvents, size_t NDepEvents) @@ -172,7 +184,7 @@ bool set_local_accessor_arg(handler &cgh, * @brief Set the kernel arg object * * @param cgh SYCL command group handler using which a kernel is going to - * be submitted. + * be submitted. * @param idx The position of the argument in the list of arguments passed * to a kernel. * @param Arg A void* representing a kernel argument. @@ -269,6 +281,8 @@ void set_kernel_args(handler &cgh, } } +#endif /* #ifndef __ADAPTIVECPP__ */ + std::unique_ptr create_property_list(int properties) { std::unique_ptr propList; @@ -498,6 +512,7 @@ DPCTLQueue_SubmitRange(__dpctl_keep const DPCTLSyclKernelRef KRef, __dpctl_keep const DPCTLSyclEventRef *DepEvents, size_t NDepEvents) { +#ifndef __ADAPTIVECPP__ auto Kernel = unwrap(KRef); auto Queue = unwrap(QRef); event e; @@ -507,7 +522,6 @@ DPCTLQueue_SubmitRange(__dpctl_keep const DPCTLSyclKernelRef KRef, case 1: { e = Queue->submit([&](handler &cgh) { - // Depend on any event that was specified by the caller. set_dependent_events(cgh, DepEvents, NDepEvents); set_kernel_args(cgh, Args, ArgTypes, NArgs); cgh.parallel_for(range<1>{Range[0]}, *Kernel); @@ -517,7 +531,6 @@ DPCTLQueue_SubmitRange(__dpctl_keep const DPCTLSyclKernelRef KRef, case 2: { e = Queue->submit([&](handler &cgh) { - // Depend on any event that was specified by the caller. set_dependent_events(cgh, DepEvents, NDepEvents); set_kernel_args(cgh, Args, ArgTypes, NArgs); cgh.parallel_for(range<2>{Range[0], Range[1]}, *Kernel); @@ -527,7 +540,6 @@ DPCTLQueue_SubmitRange(__dpctl_keep const DPCTLSyclKernelRef KRef, case 3: { e = Queue->submit([&](handler &cgh) { - // Depend on any event that was specified by the caller. set_dependent_events(cgh, DepEvents, NDepEvents); set_kernel_args(cgh, Args, ArgTypes, NArgs); cgh.parallel_for(range<3>{Range[0], Range[1], Range[2]}, @@ -536,8 +548,7 @@ DPCTLQueue_SubmitRange(__dpctl_keep const DPCTLSyclKernelRef KRef, return wrap(new event(std::move(e))); } default: - error_handler("Range cannot be greater than three " - "dimensions.", + error_handler("Range cannot be greater than three dimensions.", __FILE__, __func__, __LINE__, error_level::error); return nullptr; } @@ -549,6 +560,12 @@ DPCTLQueue_SubmitRange(__dpctl_keep const DPCTLSyclKernelRef KRef, __LINE__, error_level::error); return nullptr; } +#else + error_handler("Dynamic OpenCL-style kernel execution is not supported in " + "AdaptiveCpp.", + __FILE__, __func__, __LINE__, error_level::error); + return nullptr; +#endif } __dpctl_give DPCTLSyclEventRef @@ -563,6 +580,7 @@ DPCTLQueue_SubmitNDRange(__dpctl_keep const DPCTLSyclKernelRef KRef, __dpctl_keep const DPCTLSyclEventRef *DepEvents, size_t NDepEvents) { +#ifndef __ADAPTIVECPP__ auto Kernel = unwrap(KRef); auto Queue = unwrap(QRef); event e; @@ -572,7 +590,6 @@ DPCTLQueue_SubmitNDRange(__dpctl_keep const DPCTLSyclKernelRef KRef, case 1: { e = Queue->submit([&](handler &cgh) { - // Depend on any event that was specified by the caller. set_dependent_events(cgh, DepEvents, NDepEvents); set_kernel_args(cgh, Args, ArgTypes, NArgs); cgh.parallel_for(nd_range<1>{{gRange[0]}, {lRange[0]}}, @@ -583,7 +600,6 @@ DPCTLQueue_SubmitNDRange(__dpctl_keep const DPCTLSyclKernelRef KRef, case 2: { e = Queue->submit([&](handler &cgh) { - // Depend on any event that was specified by the caller. set_dependent_events(cgh, DepEvents, NDepEvents); set_kernel_args(cgh, Args, ArgTypes, NArgs); cgh.parallel_for( @@ -595,7 +611,6 @@ DPCTLQueue_SubmitNDRange(__dpctl_keep const DPCTLSyclKernelRef KRef, case 3: { e = Queue->submit([&](handler &cgh) { - // Depend on any event that was specified by the caller. set_dependent_events(cgh, DepEvents, NDepEvents); set_kernel_args(cgh, Args, ArgTypes, NArgs); cgh.parallel_for(nd_range<3>{{gRange[0], gRange[1], gRange[2]}, @@ -605,8 +620,7 @@ DPCTLQueue_SubmitNDRange(__dpctl_keep const DPCTLSyclKernelRef KRef, return wrap(new event(std::move(e))); } default: - error_handler("Range cannot be greater than three " - "dimensions.", + error_handler("Range cannot be greater than three dimensions.", __FILE__, __func__, __LINE__, error_level::error); return nullptr; } @@ -618,6 +632,12 @@ DPCTLQueue_SubmitNDRange(__dpctl_keep const DPCTLSyclKernelRef KRef, __LINE__, error_level::error); return nullptr; } +#else + error_handler("Dynamic OpenCL-style kernel execution is not supported in " + "AdaptiveCpp.", + __FILE__, __func__, __LINE__, error_level::error); + return nullptr; +#endif } void DPCTLQueue_Wait(__dpctl_keep DPCTLSyclQueueRef QRef) @@ -860,7 +880,12 @@ __dpctl_give DPCTLSyclEventRef DPCTLQueue_SubmitBarrierForEvents( for (auto i = 0ul; i < NDepEvents; ++i) cgh.depends_on(*unwrap(DepEvents[i])); +#ifndef __ADAPTIVECPP__ cgh.ext_oneapi_barrier(); +#else + class dpctl_barrier_task; + cgh.single_task([=]() {}); +#endif }); } catch (std::exception const &e) { error_handler(e, __FILE__, __func__, __LINE__);