diff --git a/.github/workflows/adaptivecpp-build.yml b/.github/workflows/adaptivecpp-build.yml new file mode 100644 index 0000000000..e627699040 --- /dev/null +++ b/.github/workflows/adaptivecpp-build.yml @@ -0,0 +1,172 @@ +name: Build with AdaptiveCpp + +on: + push: + branches: + - master + pull_request: + +permissions: + contents: read + +concurrency: + group: ${{ github.workflow }}-${{ github.head_ref || github.run_id }} + cancel-in-progress: true + +env: + CHANNELS: "-c conda-forge --override-channels" + # examples that AdaptiveCpp can build + SYCL_EXAMPLES: "cython/sycl_buffer cython/use_dpctl_sycl pybind11/external_usm_allocation pybind11/use_dpctl_sycl_queue" + C_EXAMPLES: "c/py_sycl_ls" + +jobs: + test_linux: + name: Build and test with AdaptiveCpp, Python ${{ matrix.python }} + runs-on: ubuntu-latest + timeout-minutes: 120 + defaults: + run: + shell: bash -el {0} + + strategy: + fail-fast: false + matrix: + include: + - python: '3.14' + python_spec: '3.14.*=*_cp314' + - python: '3.14t' + python_spec: '3.14.*=*_cp314t' + + steps: + - name: Checkout dpctl repo + uses: actions/checkout@3d3c42e5aac5ba805825da76410c181273ba90b1 # v7.0.1 + with: + persist-credentials: false + # the package version comes from git describe + fetch-depth: 0 + + - name: Cache conda packages + uses: actions/cache@55cc8345863c7cc4c66a329aec7e433d2d1c52a9 # v6.1.0 + env: + CACHE_NUMBER: 0 # Increase to reset cache + with: + # the package directory that setup-miniconda configures + path: ~/conda_pkgs_dir + key: + ${{ runner.os }}-conda-acpp-${{ env.CACHE_NUMBER }}-python-${{ matrix.python }}-${{ hashFiles('.github/workflows/adaptivecpp-build.yml') }} + restore-keys: | + ${{ runner.os }}-conda-acpp-${{ env.CACHE_NUMBER }}-python-${{ matrix.python }}- + ${{ runner.os }}-conda-acpp-${{ env.CACHE_NUMBER }}- + + - uses: conda-incubator/setup-miniconda@8ee1f361103df19b6f8c8655fd3967a8ecb162d5 # v4.0.1 + with: + auto-update-conda: true + miniforge-version: latest + activate-environment: acpp + channels: conda-forge + conda-remove-defaults: true + + - name: Install AdaptiveCpp and build dependencies + run: | + # the Python build string selects the GIL or the free-threaded variant, + # clang must come from the LLVM that adaptivecpp was built against, and + # intel-opencl-rt provides the OpenCL CPU device to run on; the HIP + # backend plugin links libamdhip64.so.6, which ROCm 7 does not provide + conda install -y ${{ env.CHANNELS }} \ + "python=${{ matrix.python_spec }}" \ + "adaptivecpp=*=cpu_llvm19*" "clang=19.*" "clangxx=19.*" \ + "hip-runtime-amd=6.*" \ + intel-opencl-rt \ + cmake cython ninja numpy pybind11 pytest scikit-build "setuptools<80" versioneer + + - name: Report the environment + run: | + conda list + python -VV + acpp-info -l + + - name: Build dpctl + run: | + python scripts/build_locally.py --sycl-provider=AdaptiveCpp + + - name: List platforms + run: | + python -m dpctl -f + python -c "from dpctl._diagnostics import sycl_provider; assert sycl_provider() == 'AdaptiveCpp'" + + - name: Run dpctl/tests on the OpenMP device + env: + ACPP_VISIBILITY_MASK: omp + run: | + python -m pytest -q --capture=sys dpctl/tests + + - name: Run dpctl/tests on the OpenCL CPU device + env: + ACPP_VISIBILITY_MASK: ocl + run: | + python -m pytest -q --capture=sys dpctl/tests + + - name: Build examples of extensions + env: + PYTHONPATH: ${{ github.workspace }} + run: | + for d in ${SYCL_EXAMPLES}; do + pushd "examples/${d}" > /dev/null + CC="${CONDA_PREFIX}/bin/clang" CXX="${CONDA_PREFIX}/bin/acpp" \ + python setup.py build_ext --inplace -G Ninja -- \ + -DDPCTL_SYCL_PROVIDER=AdaptiveCpp \ + -DDpctl_DIR="${GITHUB_WORKSPACE}/cmake" + popd > /dev/null + done + for d in ${C_EXAMPLES}; do + pushd "examples/${d}" > /dev/null + python setup.py build_ext --inplace + popd > /dev/null + done + + - name: Run examples of extensions on the OpenMP device + env: + ACPP_VISIBILITY_MASK: omp + PYTHONPATH: ${{ github.workspace }} + run: | + for d in ${SYCL_EXAMPLES} ${C_EXAMPLES}; do + # the OpenMP backend allocates plain host memory, so every USM kind + # reports usm_type "host" and this example's checks cannot hold + if [ "${d}" = "pybind11/external_usm_allocation" ]; then + continue + fi + pushd "examples/${d}" > /dev/null + python -m pytest -q tests + popd > /dev/null + done + + - name: Run examples of extensions on the OpenCL CPU device + env: + ACPP_VISIBILITY_MASK: ocl + PYTHONPATH: ${{ github.workspace }} + run: | + for d in ${SYCL_EXAMPLES} ${C_EXAMPLES}; do + pushd "examples/${d}" > /dev/null + python -m pytest -q tests + popd > /dev/null + done + + - name: Run Python examples on the OpenMP device + working-directory: examples/python + env: + ACPP_VISIBILITY_MASK: omp + run: | + while IFS= read -r script; do + echo "Executing ${script}" + python "${script}" --run all + done < <(find . \( -not -name "_*" -and -name "*.py" \)) + + - name: Run Python examples on the OpenCL CPU device + working-directory: examples/python + env: + ACPP_VISIBILITY_MASK: ocl + run: | + while IFS= read -r script; do + echo "Executing ${script}" + python "${script}" --run all + done < <(find . \( -not -name "_*" -and -name "*.py" \)) diff --git a/CMakeLists.txt b/CMakeLists.txt index 594dd0bc24..dc0c9e4bec 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/docs/doc_sources/contributor_guides/building.rst b/docs/doc_sources/contributor_guides/building.rst index e492a42b4f..b9bc4f4194 100644 --- a/docs/doc_sources/contributor_guides/building.rst +++ b/docs/doc_sources/contributor_guides/building.rst @@ -8,7 +8,8 @@ To run examples and test suite you would need GPU drivers and/or CPU OpenCL drivers. It is preferable to use the Intel(R) oneAPI DPC++ compiler available as part of oneAPI Base-Kit. However, it is possible to use a custom build of DPC++ to build :py:mod:`dpctl`, especially if you want to enable -CUDA support or try latest features. +CUDA support or try latest features. On Linux OS, AdaptiveCpp can be used +instead of DPC++, see `Building Using AdaptiveCpp`_. Building using oneAPI DPC++ --------------------------- @@ -210,6 +211,89 @@ Should a loader be installed under a different file name, override the name that ``LIBZE_SET_LOADER_FILENAME`` CMake variables. +Building Using AdaptiveCpp +-------------------------- + +:py:mod:`dpctl` can also be built with `AdaptiveCpp `_ +as the SYCL implementation, on Linux OS. The build was tested with AdaptiveCpp +25.10. + +Prerequisites for AdaptiveCpp +~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ + +Install AdaptiveCpp, the ``clang`` that it was built against, and a runtime for +the devices you intend to use. The conda packages, for example: + +.. code-block:: bash + + conda install -c conda-forge "adaptivecpp=*=cpu_llvm19*" "clang=19.*" \ + "clangxx=19.*" intel-opencl-rt tbb + +The ``clang`` version must match the LLVM version AdaptiveCpp was built with, +which the build string of the ``adaptivecpp`` package states, as in +``cpu_llvm19_h8255587_10``, hence the pins above. The ``intel-opencl-rt`` +package provides an OpenCL CPU device; AdaptiveCpp always provides its own +OpenMP host device, which needs no runtime. + +Run ``acpp-info -l`` to list the backends of the installed AdaptiveCpp, and set +the ``ACPP_VISIBILITY_MASK`` environment variable to select among them, for +example ``ACPP_VISIBILITY_MASK=omp`` for the OpenMP host device, or +``ACPP_VISIBILITY_MASK=ocl`` for the OpenCL devices. + +Build and Install Using AdaptiveCpp +~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ + +Use the driver script with the provider option: + +.. code-block:: bash + + python scripts/build_locally.py --sycl-provider=AdaptiveCpp + +The option defaults the compilers to ``clang`` and ``acpp``, which are taken +from the ``PATH`` unless ``--c-compiler`` and ``--cxx-compiler`` are given, and +it disables Level Zero kernel bundle creation, which only DPC++ implements. +Leaving it on would only check out the Level Zero headers for code that an +AdaptiveCpp build excludes; it does not affect whether a Level Zero device is +available, which depends on the backends AdaptiveCpp itself was built with. + +Building with ``setuptools`` directly takes the equivalent CMake variables: + +.. code-block:: bash + + python setup.py build_ext --inplace -G Ninja -- \ + -DCMAKE_C_COMPILER:PATH=clang -DCMAKE_CXX_COMPILER:PATH=acpp \ + -DDPCTL_SYCL_PROVIDER=AdaptiveCpp \ + -DDPCTL_ENABLE_L0_KERNEL_BUNDLE_CREATION=OFF + python -m pip install -e . + +Differences from a DPC++ build +~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ + +An AdaptiveCpp build reports its SYCL implementation through +``dpctl._diagnostics.sycl_provider()``. The following is unavailable in such a +build: + +- kernel bundle creation from OpenCL source and from SPIR-V, hence the + :py:mod:`dpctl.compiler` functions and kernel submission through them, both + through the OpenCL and the Level Zero backend +- the raw kernel argument extension +- the atomic memory order, fence order and scope capability descriptors of + devices and of contexts +- the DPC++ version, the backend version and the default selector score + +A few behaviors differ as well: a context can span platforms, and a platform is +reported for every backend AdaptiveCpp was built with, including backends with +no device. + +Tests covering the above carry the ``unsupported_on_acpp`` pytest marker and are +skipped automatically. + +The extension examples under ``examples/`` that build SYCL kernels take +``-DDPCTL_SYCL_PROVIDER=AdaptiveCpp`` as well, with the exception of +``pybind11/onemkl_gemv`` and ``pybind11/use_dpctl_sycl_kernel``, which are +DPC++ only. + + Building the libsyclinterface Library ======================================= diff --git a/dpctl/CMakeLists.txt b/dpctl/CMakeLists.txt index 9c9d4f9cad..0675214578 100644 --- a/dpctl/CMakeLists.txt +++ b/dpctl/CMakeLists.txt @@ -111,10 +111,13 @@ 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_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_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( @@ -133,7 +136,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/_backend.pxd b/dpctl/_backend.pxd index 3708ea9130..a73dc37738 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -47,6 +47,7 @@ cdef extern from "syclinterface/dpctl_sycl_enum_types.h": _HIP "DPCTL_HIP" _LEVEL_ZERO "DPCTL_LEVEL_ZERO" _OPENCL "DPCTL_OPENCL" + _OPENMP "DPCTL_OPENMP" _UNKNOWN_BACKEND "DPCTL_UNKNOWN_BACKEND" ctypedef enum _device_type "DPCTLSyclDeviceType": diff --git a/dpctl/_device_selection.py b/dpctl/_device_selection.py index 8aed10f194..893af6d93d 100644 --- a/dpctl/_device_selection.py +++ b/dpctl/_device_selection.py @@ -66,7 +66,9 @@ def select_device_with_aspects(required_aspects, excluded_aspects=None): if not hasattr(SyclDevice, "has_aspect_" + asp): raise AttributeError(f"The {asp} aspect is not supported in dpctl") devs = get_devices() - max_score = 0 + # get_devices only reports devices with a non-negative score, and a score + # of zero is a valid one + max_score = -1 selected_dev = None for dev in devs: diff --git a/dpctl/_diagnostics.pyx b/dpctl/_diagnostics.pyx index 299d959beb..b5aa11a604 100644 --- a/dpctl/_diagnostics.pyx +++ b/dpctl/_diagnostics.pyx @@ -24,10 +24,26 @@ import contextlib import os +from ._backend cimport DPCTLCString_Delete + cdef extern from "syclinterface/dpctl_service.h": cdef void DPCTLService_InitLogger(const char *, const char *) cdef void DPCTLService_ShutdownLogger() + cdef const char *DPCTLService_GetSyclProvider() + + +def sycl_provider(): + """Return the SYCL implementation `dpctl` was built with. + + Returns: + str: + Either ``"Intel"`` or ``"AdaptiveCpp"``. + """ + cdef const char *c_str = DPCTLService_GetSyclProvider() + cdef bytes byte_str = c_str + DPCTLCString_Delete(c_str) + return byte_str.decode("utf-8") def _init_logger(log_dir=None): diff --git a/dpctl/_sycl_device.pyx b/dpctl/_sycl_device.pyx index 624575a5d8..fd06c7fa76 100644 --- a/dpctl/_sycl_device.pyx +++ b/dpctl/_sycl_device.pyx @@ -247,6 +247,8 @@ cdef str _backend_type_to_filter_string_part(_backend_type BTy): return "level_zero" elif BTy == _backend_type._OPENCL: return "opencl" + elif BTy == _backend_type._OPENMP: + return "openmp" else: return "unknown" @@ -560,6 +562,8 @@ cdef class SyclDevice(_SyclDevice): return backend_type.level_zero elif BTy == _backend_type._OPENCL: return backend_type.opencl + elif BTy == _backend_type._OPENMP: + return backend_type.openmp else: raise ValueError("Unknown backend type.") diff --git a/dpctl/_sycl_device_factory.pyx b/dpctl/_sycl_device_factory.pyx index b94313d322..39d21998d1 100644 --- a/dpctl/_sycl_device_factory.pyx +++ b/dpctl/_sycl_device_factory.pyx @@ -81,6 +81,8 @@ cdef _backend_type _string_to_dpctl_sycl_backend_ty(str backend_str): return _backend_type._LEVEL_ZERO elif backend_str == "opencl": return _backend_type._OPENCL + elif backend_str == "openmp": + return _backend_type._OPENMP else: return _backend_type._UNKNOWN_BACKEND @@ -114,6 +116,8 @@ cdef _backend_type _enum_to_dpctl_sycl_backend_ty(BTy): return _backend_type._LEVEL_ZERO elif BTy == backend_type.opencl: return _backend_type._OPENCL + elif BTy == backend_type.openmp: + return _backend_type._OPENMP else: return _backend_type._UNKNOWN_BACKEND @@ -163,7 +167,7 @@ cpdef tuple get_devices( backend (str, :class:`dpctl.backend_type`, optional): A :class:`dpctl.backend_type` enum value or a string that specifies a SYCL backend. Currently, accepted values are: "cuda", - "hip", "opencl", "level_zero", or "all". + "hip", "opencl", "openmp", "level_zero", or "all". Default: ``dpctl.backend_type.all``. device_type (str, :class:`dpctl.device_type`, optional): A :class:`dpctl.device_type` enum value or a string that @@ -247,7 +251,7 @@ cpdef int get_num_devices( backend (str, :class:`dpctl.backend_type`, optional): A :class:`dpctl.backend_type` enum value or a string that specifies a SYCL backend. Currently, accepted values are: "cuda", - "hip", "opencl", "level_zero", or "all". + "hip", "opencl", "openmp", "level_zero", or "all". Default: ``dpctl.backend_type.all``. device_type (str, :class:`dpctl.device_type`, optional): A :class:`dpctl.device_type` enum value or a string that diff --git a/dpctl/_sycl_event.pyx b/dpctl/_sycl_event.pyx index 0eaa94c033..7075b6f2a8 100644 --- a/dpctl/_sycl_event.pyx +++ b/dpctl/_sycl_event.pyx @@ -336,6 +336,8 @@ cdef class SyclEvent(_SyclEvent): return backend_type.cuda elif BE == _backend_type._HIP: return backend_type.hip + elif BE == _backend_type._OPENMP: + return backend_type.openmp else: raise ValueError("Unknown backend type.") diff --git a/dpctl/_sycl_platform.pyx b/dpctl/_sycl_platform.pyx index d722e93b50..0600b715bf 100644 --- a/dpctl/_sycl_platform.pyx +++ b/dpctl/_sycl_platform.pyx @@ -310,6 +310,8 @@ cdef class SyclPlatform(_SyclPlatform): return backend_type.level_zero elif BTy == _backend_type._OPENCL: return backend_type.opencl + elif BTy == _backend_type._OPENMP: + return backend_type.openmp else: raise ValueError("Unknown backend type.") diff --git a/dpctl/_sycl_queue.pyx b/dpctl/_sycl_queue.pyx index 6dd7324691..5a928812f8 100644 --- a/dpctl/_sycl_queue.pyx +++ b/dpctl/_sycl_queue.pyx @@ -1306,6 +1306,8 @@ cdef class SyclQueue(_SyclQueue): return backend_type.cuda elif BE == _backend_type._HIP: return backend_type.hip + elif BE == _backend_type._OPENMP: + return backend_type.openmp else: raise ValueError("Unknown backend type.") diff --git a/dpctl/apis/include/dpctl4pybind11.hpp b/dpctl/apis/include/dpctl4pybind11.hpp index 8844ae3bdb..945ec6e929 100644 --- a/dpctl/apis/include/dpctl4pybind11.hpp +++ b/dpctl/apis/include/dpctl4pybind11.hpp @@ -538,6 +538,7 @@ template <> struct type_caster DPCTL_TYPE_CASTER(sycl::kernel, _("dpctl.compiler.SyclKernel")); }; +#ifndef __ADAPTIVECPP__ /* This type caster associates * ``sycl::kernel_bundle`` C++ class with * :class:`dpctl.compiler.SyclKernelBundle` for the purposes of generation of @@ -581,6 +582,7 @@ struct type_caster> DPCTL_TYPE_CASTER(sycl::kernel_bundle, _("dpctl.compiler.SyclKernelBundle")); }; +#endif /* This type caster associates * ``sycl::half`` C++ class with Python :class:`float` for the purposes diff --git a/dpctl/enum_types.py b/dpctl/enum_types.py index f445ebb7f6..c77c3aba21 100644 --- a/dpctl/enum_types.py +++ b/dpctl/enum_types.py @@ -79,6 +79,7 @@ class backend_type(Enum): | ``hip`` | ``level_zero`` | ``opencl`` + | ``openmp`` :Example: .. code-block:: python @@ -96,6 +97,7 @@ class backend_type(Enum): hip = auto() level_zero = auto() opencl = auto() + openmp = auto() class event_status_type(Enum): 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/_device_attributes_checks.py b/dpctl/tests/_device_attributes_checks.py index 7e00831fae..a6688b5521 100644 --- a/dpctl/tests/_device_attributes_checks.py +++ b/dpctl/tests/_device_attributes_checks.py @@ -866,7 +866,10 @@ def check_partition_affinity_domains(device): check_vendor, check_driver_version, check_name, - check_default_selector_score, + # AdaptiveCpp scores the OpenMP host device 0, not a positive value + pytest.param( + check_default_selector_score, marks=pytest.mark.unsupported_on_acpp + ), check_backend, check_device_type, check_global_mem_cache_type, @@ -883,17 +886,29 @@ def check_partition_affinity_domains(device): check_error_correction_support, check_is_available, check_version, - check_backend_version, + pytest.param(check_backend_version, marks=pytest.mark.unsupported_on_acpp), check_local_mem_type, check_partition_type_property, check_partition_type_affinity_domain, check_half_fp_config, check_single_fp_config, check_double_fp_config, - check_atomic_memory_order_capabilities, - check_atomic_fence_order_capabilities, - check_atomic_memory_scope_capabilities, - check_atomic_fence_scope_capabilities, + pytest.param( + check_atomic_memory_order_capabilities, + marks=pytest.mark.unsupported_on_acpp, + ), + pytest.param( + check_atomic_fence_order_capabilities, + marks=pytest.mark.unsupported_on_acpp, + ), + pytest.param( + check_atomic_memory_scope_capabilities, + marks=pytest.mark.unsupported_on_acpp, + ), + pytest.param( + check_atomic_fence_scope_capabilities, + marks=pytest.mark.unsupported_on_acpp, + ), check_partition_properties, check_partition_affinity_domains, ] diff --git a/dpctl/tests/conftest.py b/dpctl/tests/conftest.py index 69ae0a519b..aa76b19b8b 100644 --- a/dpctl/tests/conftest.py +++ b/dpctl/tests/conftest.py @@ -28,6 +28,7 @@ valid_filter, ) from ._numpy_warnings import suppress_invalid_numpy_warnings +from .helper import is_adaptivecpp sys.path.append(os.path.join(os.path.dirname(__file__), "helper")) @@ -46,6 +47,10 @@ def pytest_configure(config): "markers", "broken_complex: Specified again to remove warnings ", ) + config.addinivalue_line( + "markers", + "unsupported_on_acpp: Test exercises a feature AdaptiveCpp lacks", + ) def pytest_addoption(parser): @@ -64,6 +69,14 @@ def pytest_addoption(parser): def pytest_collection_modifyitems(config, items): + if is_adaptivecpp(): + skip_acpp = pytest.mark.skip( + reason="not supported by AdaptiveCpp", + ) + for item in items: + if "unsupported_on_acpp" in item.keywords: + item.add_marker(skip_acpp) + if config.getoption("--runcomplex"): return skip_complex = pytest.mark.skip( diff --git a/dpctl/tests/helper/__init__.py b/dpctl/tests/helper/__init__.py index be5f8d2e33..d9b40241fb 100644 --- a/dpctl/tests/helper/__init__.py +++ b/dpctl/tests/helper/__init__.py @@ -22,6 +22,8 @@ has_cpu, has_gpu, has_sycl_platforms, + is_adaptivecpp, + is_wsl_or_windows, ) __all__ = [ @@ -30,4 +32,6 @@ "has_gpu", "has_sycl_platforms", "get_queue_or_skip", + "is_adaptivecpp", + "is_wsl_or_windows", ] diff --git a/dpctl/tests/helper/_helper.py b/dpctl/tests/helper/_helper.py index 5857b23195..c016f9b1c4 100644 --- a/dpctl/tests/helper/_helper.py +++ b/dpctl/tests/helper/_helper.py @@ -14,9 +14,20 @@ # See the License for the specific language governing permissions and # limitations under the License. +import platform +import sys + import pytest import dpctl +from dpctl._diagnostics import sycl_provider + + +def is_adaptivecpp(): + """ + Return True if dpctl was built with AdaptiveCpp. + """ + return sycl_provider() == "AdaptiveCpp" def has_gpu(backend="opencl"): @@ -49,3 +60,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_raw_kernel_arg.py b/dpctl/tests/test_raw_kernel_arg.py index fd07930935..79cbc495e4 100644 --- a/dpctl/tests/test_raw_kernel_arg.py +++ b/dpctl/tests/test_raw_kernel_arg.py @@ -25,6 +25,11 @@ import dpctl import dpctl.memory as dpm +pytestmark = pytest.mark.skipif( + not dpctl.RawKernelArg.is_available(), + reason="Raw kernel arg extension not supported", +) + def get_spirv_abspath(fn): curr_dir = os.path.dirname(os.path.abspath(__file__)) @@ -59,9 +64,6 @@ class Params(ctypes.Structure): def launch_raw_arg_kernel(raw): - if not dpctl.RawKernelArg.is_available(): - pytest.skip("Raw kernel arg extension not supported") - try: q = dpctl.SyclQueue("level_zero") except dpctl.SyclQueueCreationError: diff --git a/dpctl/tests/test_service.py b/dpctl/tests/test_service.py index 224aed326e..7970cb0f49 100644 --- a/dpctl/tests/test_service.py +++ b/dpctl/tests/test_service.py @@ -28,6 +28,9 @@ import pytest import dpctl +from dpctl._diagnostics import sycl_provider + +from .helper import is_adaptivecpp def _get_mkl_version_if_present(): @@ -81,6 +84,14 @@ def test_get_include(): ) +def test_sycl_provider(): + provider = sycl_provider() + assert provider in ("Intel", "AdaptiveCpp") + assert is_adaptivecpp() == (provider == "AdaptiveCpp") + + +# AdaptiveCpp builds do not record a DPC++ version +@pytest.mark.unsupported_on_acpp def test_get_dpcppversion(): """Intent of this test is to verify that libraries from dpcpp_cpp_rt conda package used at run-time are not from an older oneAPI. Since these diff --git a/dpctl/tests/test_sycl_compiler.py b/dpctl/tests/test_sycl_compiler.py index a2d4be191a..99d257104f 100644 --- a/dpctl/tests/test_sycl_compiler.py +++ b/dpctl/tests/test_sycl_compiler.py @@ -216,6 +216,7 @@ def _check_multi_kernel_bundle(kb): assert kb2.has_sycl_kernel("axpy") +@pytest.mark.unsupported_on_acpp def test_create_kernel_bundle_from_source_ocl(): oclSrc = " \ kernel void add(global int* a, global int* b, global int* c) { \ @@ -231,6 +232,7 @@ def test_create_kernel_bundle_from_source_ocl(): _check_multi_kernel_bundle(kb) +@pytest.mark.unsupported_on_acpp def test_get_sycl_kernel_undefined_name(): try: q = dpctl.SyclQueue() @@ -246,6 +248,7 @@ def test_get_sycl_kernel_undefined_name(): kb.get_sycl_kernel("no_such_kernel") +@pytest.mark.unsupported_on_acpp def test_create_kernel_bundle_from_spirv_ocl(): q = _get_opencl_queue_or_skip() spirv_file = get_spirv_abspath("multi_kernel.spv") @@ -255,6 +258,7 @@ def test_create_kernel_bundle_from_spirv_ocl(): _check_multi_kernel_bundle(kb) +@pytest.mark.unsupported_on_acpp def test_create_kernel_bundle_from_spirv_l0(): q = _get_level_zero_queue_or_skip() spirv_file = get_spirv_abspath("multi_kernel.spv") @@ -291,6 +295,7 @@ def test_create_kernel_bundle_from_invalid_src_ocl(): dpc.create_kernel_bundle_from_source(q, invalid_oclSrc) +@pytest.mark.unsupported_on_acpp def test_create_kernel_bundle_with_spec_const(): try: q = dpctl.SyclQueue() @@ -325,6 +330,7 @@ def test_create_kernel_bundle_with_spec_const(): assert np.all(y == 43) +@pytest.mark.unsupported_on_acpp def test_create_kernel_bundle_with_composite_spec_const(): try: q = dpctl.SyclQueue() diff --git a/dpctl/tests/test_sycl_context.py b/dpctl/tests/test_sycl_context.py index 058bb7ad7f..18c1164747 100644 --- a/dpctl/tests/test_sycl_context.py +++ b/dpctl/tests/test_sycl_context.py @@ -162,7 +162,12 @@ def test_context_multi_device(): n2 = n - n1 if n1 == 0 or n2 == 0: pytest.skip() - d1, d2 = d.create_sub_devices(partition=(n1, n2)) + try: + d1, d2 = d.create_sub_devices(partition=(n1, n2)) + except dpctl.SyclSubDeviceCreationError: + pytest.skip( + "create_sub_devices can't create sub-devices on this device" + ) ctx = dpctl.SyclContext((d1, d2)) assert ctx.device_count == 2 assert type(repr(ctx)) is str @@ -266,6 +271,8 @@ def test_invalid_capsule(): dpctl.SyclContext(cap) +# AdaptiveCpp contexts span platforms +@pytest.mark.unsupported_on_acpp def test_multi_device_different_platforms(): devs = dpctl.get_devices() # all devices if len(devs) > 1 and len(set(d.sycl_platform for d in devs)) > 1: @@ -275,6 +282,8 @@ def test_multi_device_different_platforms(): pytest.skip("Insufficient amount of available devices for this test") +# AdaptiveCpp contexts span platforms +@pytest.mark.unsupported_on_acpp def test_context_sycl_platform(valid_filter): """ Test that :attr:`dpctl.SyclContext.sycl_platform` returns the @@ -290,6 +299,7 @@ def test_context_sycl_platform(valid_filter): assert d.sycl_platform == plat +@pytest.mark.unsupported_on_acpp def test_context_atomic_memory_order_capabilities(valid_filter): try: ctx = dpctl.SyclContext(valid_filter) @@ -317,6 +327,7 @@ def test_context_atomic_fence_order_capabilities(valid_filter): assert set(caps).issubset(set(d.atomic_fence_order_capabilities)) +@pytest.mark.unsupported_on_acpp def test_context_atomic_memory_scope_capabilities(valid_filter): try: ctx = dpctl.SyclContext(valid_filter) diff --git a/dpctl/tests/test_sycl_event.py b/dpctl/tests/test_sycl_event.py index 01c61ca7f6..228e3ec008 100644 --- a/dpctl/tests/test_sycl_event.py +++ b/dpctl/tests/test_sycl_event.py @@ -109,6 +109,7 @@ def test_execution_status(): assert event_status == esty.complete +@pytest.mark.unsupported_on_acpp def test_execution_status_nondefault_event(): try: event = produce_event() @@ -123,6 +124,7 @@ def test_execution_status_nondefault_event(): assert type(wl) is list +@pytest.mark.unsupported_on_acpp def test_event_backend(): if dpctl.get_num_devices() == 0: pytest.skip("No backends are available") @@ -140,6 +142,7 @@ def test_event_backend(): pytest.fail("Failed to get backend from event") +@pytest.mark.unsupported_on_acpp def test_get_wait_list(): try: q = dpctl.SyclQueue("opencl:cpu") @@ -184,6 +187,7 @@ def test_get_wait_list(): assert len(wait_list) >= 0 +@pytest.mark.unsupported_on_acpp def test_profiling_info(): try: event = produce_event(profiling=True) diff --git a/dpctl/tests/test_sycl_kernel_submit.py b/dpctl/tests/test_sycl_kernel_submit.py index 9c57186d81..103da46a72 100644 --- a/dpctl/tests/test_sycl_kernel_submit.py +++ b/dpctl/tests/test_sycl_kernel_submit.py @@ -40,6 +40,7 @@ ("double", np.dtype("f8"), ctypes.c_double), ], ) +@pytest.mark.unsupported_on_acpp def test_create_kernel_bundle_from_source(ctype_str, dtype, ctypes_ctor): try: q = dpctl.SyclQueue("opencl", property="enable_profiling") @@ -123,6 +124,7 @@ def test_create_kernel_bundle_from_source(ctype_str, dtype, ctypes_ctor): assert np.allclose(c, ref_c), f"Failed for {gr}, {lr}" +@pytest.mark.unsupported_on_acpp def test_submit_async(): try: q = dpctl.SyclQueue("opencl") @@ -313,6 +315,7 @@ def get_spirv_abspath(fn): # the process for generating the .spv files in this test is documented in # libsyclinterface/tests/test_sycl_queue_submit_local_accessor_arg.cpp # in a comment starting on line 123 +@pytest.mark.unsupported_on_acpp def test_submit_local_accessor_arg(): try: q = dpctl.SyclQueue("level_zero") @@ -344,6 +347,7 @@ def test_submit_local_accessor_arg(): assert np.all(res == expected) +@pytest.mark.unsupported_on_acpp def test_submit_wrong_number_of_args(): try: q = dpctl.SyclQueue() diff --git a/dpctl/tests/test_sycl_platform.py b/dpctl/tests/test_sycl_platform.py index 0660617752..471f42df4f 100644 --- a/dpctl/tests/test_sycl_platform.py +++ b/dpctl/tests/test_sycl_platform.py @@ -120,7 +120,9 @@ def check_hash_in_dict(platform): check_print_info, check_repr, check_default_context, - check_equal_and_hash, + # AdaptiveCpp contexts span platforms, so a context's devices need not + # belong to the platform the context was created for + pytest.param(check_equal_and_hash, marks=pytest.mark.unsupported_on_acpp), check_hash_in_dict, ] @@ -214,6 +216,8 @@ def test_get_platforms(): pytest.fail("Encountered an exception inside get_platforms().") +# AdaptiveCpp reports a platform per backend, even those with no device +@pytest.mark.unsupported_on_acpp def test_platform_get_devices(): platforms = dpctl.get_platforms() if platforms: diff --git a/dpctl/tests/test_sycl_queue.py b/dpctl/tests/test_sycl_queue.py index 0abe97b8bf..8ce9c7c439 100644 --- a/dpctl/tests/test_sycl_queue.py +++ b/dpctl/tests/test_sycl_queue.py @@ -24,7 +24,7 @@ import dpctl import dpctl.memory -from .helper import create_invalid_capsule +from .helper import create_invalid_capsule, is_adaptivecpp, is_wsl_or_windows def test_standard_selectors(device_selector, check): @@ -338,19 +338,31 @@ 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) + + # AdaptiveCpp prefetches by migrating the allocation to the device, which + # the OpenCL CPU runtime rejects with CL_INVALID_VALUE + acpp_opencl_cpu = ( + is_adaptivecpp() + and q.sycl_device.backend == dpctl.backend_type.opencl + and q.sycl_device.has_aspect_cpu + ) + if not is_wsl_or_windows() and not acpp_opencl_cpu: + q.prefetch(m1, 512) + with pytest.raises(TypeError): + q.prefetch([], 512) + + # AdaptiveCpp does not implement mem_advise + if not is_adaptivecpp(): + q.mem_advise(m1, 512, 0) with pytest.raises(TypeError): q.mem_advise([], 512, 0) diff --git a/dpctl/tests/test_sycl_timer.py b/dpctl/tests/test_sycl_timer.py index 6ce8c66f2e..3289fad1be 100644 --- a/dpctl/tests/test_sycl_timer.py +++ b/dpctl/tests/test_sycl_timer.py @@ -31,6 +31,10 @@ def profiling_queue(): pytest.skip( "Could not created profiling queue " "for default-selected device" ) + ev = q.submit_barrier() + ev.wait() + if ev.profiling_info_end == 0: + pytest.skip("Device does not report profiling timestamps") return q diff --git a/dpctl/tests/test_sycl_usm.py b/dpctl/tests/test_sycl_usm.py index 398a07b216..a11c0dddfb 100644 --- a/dpctl/tests/test_sycl_usm.py +++ b/dpctl/tests/test_sycl_usm.py @@ -83,6 +83,17 @@ def _create_memory(): return mobj +def _skip_if_usm_types_not_reported(q=None): + """Skip the test unless the device reports the type of a USM allocation. + + AdaptiveCpp's OpenMP backend allocates plain host memory, and reports + every allocation as USM-host. + """ + m = MemoryUSMDevice(8) if q is None else MemoryUSMDevice(8, queue=q) + if m.get_usm_type() != "device": + pytest.skip("Device does not report distinct USM allocation types") + + def _create_host_buf(nbytes): ba = bytearray(nbytes) for i in range(nbytes): @@ -91,6 +102,7 @@ def _create_host_buf(nbytes): def test_memory_without_context(): + _skip_if_usm_types_not_reported() mobj = _create_memory() # Without context @@ -99,6 +111,7 @@ def test_memory_without_context(): def test_memory_cpu_context(): + _skip_if_usm_types_not_reported() mobj = _create_memory() # USM type respective to the context in which @@ -110,6 +123,7 @@ def test_memory_cpu_context(): cpu_queue = dpctl.SyclQueue("cpu") except dpctl.SyclQueueCreationError: pytest.skip("SyclQueue('cpu') failed, skip further testing") + _skip_if_usm_types_not_reported(cpu_queue) # USM type as view from CPU queue usm_type = mobj.get_usm_type(cpu_queue) # USM type can be unknown if current queue is @@ -118,6 +132,7 @@ def test_memory_cpu_context(): def test_memory_gpu_context(): + _skip_if_usm_types_not_reported() mobj = _create_memory() # GPU context @@ -127,6 +142,7 @@ def test_memory_gpu_context(): gpu_queue = dpctl.SyclQueue("opencl:gpu") except dpctl.SyclQueueCreationError: pytest.skip("SyclQueue('opencl:gpu') failed, skipping") + _skip_if_usm_types_not_reported(gpu_queue) usm_type = mobj.get_usm_type(gpu_queue) assert usm_type in ["unknown", "shared"] @@ -148,6 +164,7 @@ def test_copy_host_roundtrip(): def test_zero_copy(): + _skip_if_usm_types_not_reported() mobj = _create_memory() mobj2 = type(mobj)(mobj) @@ -160,6 +177,7 @@ def test_zero_copy(): def test_pickling(memory_ctor): import pickle + _skip_if_usm_types_not_reported() try: mobj = memory_ctor(1024, alignment=64) except dpctl.SyclDeviceCreationError: @@ -182,6 +200,7 @@ def test_pickling(memory_ctor): def test_pickling_reconstructor_invalid_type(memory_ctor): import pickle + _skip_if_usm_types_not_reported() try: mobj = memory_ctor(1024, alignment=64) except dpctl.SyclDeviceCreationError: @@ -222,6 +241,7 @@ def test_create_with_size_and_alignment_and_queue(memory_ctor): q = dpctl.SyclQueue() except dpctl.SyclQueueCreationError: pytest.skip("SyclQueue() failed, skip further testing") + _skip_if_usm_types_not_reported(q) m = memory_ctor(1024, alignment=64, queue=q) assert m.nbytes == 1024 assert m.get_usm_type() == expected_usm_type_str(memory_ctor) @@ -233,6 +253,7 @@ def test_create_with_size_and_queue(memory_ctor): q = dpctl.SyclQueue() except dpctl.SyclQueueCreationError: pytest.skip("SyclQueue() failed, skip further testing") + _skip_if_usm_types_not_reported(q) m = memory_ctor(1024, queue=q) assert m.nbytes == 1024 assert m.get_usm_type() == expected_usm_type_str(memory_ctor) @@ -240,6 +261,7 @@ def test_create_with_size_and_queue(memory_ctor): def test_create_with_size_and_alignment(memory_ctor): + _skip_if_usm_types_not_reported() try: m = memory_ctor(1024, alignment=64) except dpctl.SyclDeviceCreationError: @@ -250,6 +272,7 @@ def test_create_with_size_and_alignment(memory_ctor): def test_usm_type_exceptions(): + _skip_if_usm_types_not_reported() ctor = MemoryUSMDevice try: m = ctor(1024) @@ -310,7 +333,7 @@ def test_suai_non_contig_1D(memory_ctor): Test of zero-copy using sycl_usm_array_interface with non-contiguous data. """ - + _skip_if_usm_types_not_reported() try: buf = memory_ctor(32) except Exception: @@ -341,6 +364,7 @@ def test_suai_non_contig_1D(memory_ctor): def test_suai_non_contig_2D(memory_ctor): + _skip_if_usm_types_not_reported() try: buf = memory_ctor(20) except Exception: @@ -370,6 +394,7 @@ def test_suai_non_contig_2D(memory_ctor): def test_suai_invalid_suai(): + _skip_if_usm_types_not_reported() n_bytes = 2 * 3 * 5 * 128 try: q = dpctl.SyclQueue() @@ -613,6 +638,7 @@ def test_cpython_api(memory_ctor): def test_memory_construction_from_other_memory_objects(): + _skip_if_usm_types_not_reported() try: q = dpctl.SyclQueue() except dpctl.SyclQueueCreationError: @@ -641,7 +667,12 @@ def test_memory_copy_between_contexts(): d = q.sycl_device n = d.max_compute_units n_half = n // 2 - d0, d1 = d.create_sub_devices(partition=[n_half, n - n_half]) + try: + d0, d1 = d.create_sub_devices(partition=[n_half, n - n_half]) + except dpctl.SyclSubDeviceCreationError: + pytest.skip( + "create_sub_devices can't create sub-devices on this device" + ) q0 = dpctl.SyclQueue(d0) q1 = dpctl.SyclQueue(d1) m0 = MemoryUSMDevice(256, queue=q0) diff --git a/dpctl/utils/CMakeLists.txt b/dpctl/utils/CMakeLists.txt index 95a673ffe7..fd7ae06871 100644 --- a/dpctl/utils/CMakeLists.txt +++ b/dpctl/utils/CMakeLists.txt @@ -27,9 +27,11 @@ list(APPEND _pybind11_targets ${python_module_name}) set(_linker_options "LINKER:${DPCTL_LDFLAGS}") foreach(python_module_name ${_pybind11_targets}) - 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_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) + endif() endif() target_link_options(${python_module_name} PRIVATE ${_linker_options}) 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/examples/README.md b/examples/README.md index 9bfb80282c..00a483b247 100644 --- a/examples/README.md +++ b/examples/README.md @@ -35,6 +35,21 @@ this functionality. Examples of Python extensions created using C are located in [c](c) folder. +## Building with AdaptiveCpp + +An extension must be built with the same SYCL implementation as the `dpctl` it is +built against, which `dpctl._diagnostics.sycl_provider()` reports. The `cython` +and `pybind11` examples accept `-DDPCTL_SYCL_PROVIDER=AdaptiveCpp` to build with +AdaptiveCpp rather than Intel DPC++: + +```bash +CC=clang CXX=acpp python setup.py build_ext --inplace -G Ninja -- \ + -DDPCTL_SYCL_PROVIDER=AdaptiveCpp +``` + +The `onemkl_gemv` and `use_dpctl_sycl_kernel` examples require Intel DPC++, see +their READMEs. + [platform]: https://intelpython.github.io/dpctl/latest/docfiles/user_guides/manual/dpctl/platforms.html [device_selection]: https://intelpython.github.io/dpctl/latest/docfiles/user_guides/manual/dpctl/device_selection.html diff --git a/examples/cython/sycl_buffer/CMakeLists.txt b/examples/cython/sycl_buffer/CMakeLists.txt index a30fbb2f35..fe06e30da9 100644 --- a/examples/cython/sycl_buffer/CMakeLists.txt +++ b/examples/cython/sycl_buffer/CMakeLists.txt @@ -5,7 +5,14 @@ project(example_cython_syclbuffer VERSION 0.1 LANGUAGES CXX set(DPCTL_CMAKE_MODULES_PATH "${CMAKE_SOURCE_DIR}/../../../cmake") set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${DPCTL_CMAKE_MODULES_PATH}) -find_package(IntelSYCL REQUIRED PATHS ${DPCTL_CMAKE_MODULES_PATH} NO_DEFAULT_PATH) +set(DPCTL_SYCL_PROVIDER "Intel" CACHE STRING "SYCL compiler provider (Intel or AdaptiveCpp)") +set_property(CACHE DPCTL_SYCL_PROVIDER PROPERTY STRINGS Intel AdaptiveCpp) + +if(DPCTL_SYCL_PROVIDER STREQUAL "AdaptiveCpp") + find_package(AdaptiveCpp CONFIG REQUIRED) +else() + find_package(IntelSYCL REQUIRED PATHS ${DPCTL_CMAKE_MODULES_PATH} NO_DEFAULT_PATH) +endif() set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} "${CMAKE_SOURCE_DIR}") @@ -41,6 +48,8 @@ foreach(_src_fn ${_sources}) COMPILE_OPTIONS "${_combined_options}" ) endforeach() -target_link_options(${py_module_name} PRIVATE -fsycl-device-code-split=per_kernel) +if(NOT DPCTL_SYCL_PROVIDER STREQUAL "AdaptiveCpp") + target_link_options(${py_module_name} PRIVATE -fsycl-device-code-split=per_kernel) +endif() set(ignoreMe "${SKBUILD}") diff --git a/examples/cython/sycl_buffer/src/custom_reduce.hpp b/examples/cython/sycl_buffer/src/custom_reduce.hpp new file mode 100644 index 0000000000..e008b84873 --- /dev/null +++ b/examples/cython/sycl_buffer/src/custom_reduce.hpp @@ -0,0 +1,56 @@ +//===--- custom_reduce.hpp - work-group reduction for AdaptiveCpp ---------===// +// +// 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. +// +//===----------------------------------------------------------------------===// +/// +/// \file +/// This file implements a reduction over a work-group through local memory, to +/// be used in place of `sycl::reduce_over_group`. +/// +//===----------------------------------------------------------------------===// + +#pragma once + +#include +#include + +template +T custom_reduce_over_group(const GroupTy &wg, + LocAccT local_mem_acc, + const T &local_val, + const OpT &op) +{ + const std::uint32_t lid = wg.get_local_linear_id(); + + local_mem_acc[lid] = local_val; + sycl::group_barrier(wg, sycl::memory_scope::work_group); + + // fold the upper half of the active range onto the lower one, until a + // single element is left + for (std::uint32_t n = wg.get_local_linear_range(); n > 1;) { + const std::uint32_t half = (n + 1) >> 1; + if (lid + half < n) { + local_mem_acc[lid] = + op(local_mem_acc[lid], local_mem_acc[lid + half]); + } + sycl::group_barrier(wg, sycl::memory_scope::work_group); + n = half; + } + + return local_mem_acc[0]; +} diff --git a/examples/cython/sycl_buffer/src/use_sycl_buffer.hpp b/examples/cython/sycl_buffer/src/use_sycl_buffer.hpp index 31b9bf35f0..6b26064b5b 100644 --- a/examples/cython/sycl_buffer/src/use_sycl_buffer.hpp +++ b/examples/cython/sycl_buffer/src/use_sycl_buffer.hpp @@ -23,7 +23,7 @@ /// provided as host C-contiguous allocation. SYCL kernels access this memory /// using `sycl::buffer`. Two routines are provided. One solves the task by /// calling BLAS function GEMV from Intel(R) Math Kernel Library, the other -/// performs the computation using DPC++ reduction group function and atomics. +/// performs the computation using a work-group reduction and atomics. /// //===----------------------------------------------------------------------===// @@ -33,6 +33,10 @@ #include #include +#ifdef __ADAPTIVECPP__ +#include "custom_reduce.hpp" +#endif + using std::size_t; inline size_t upper_multiple(size_t n, size_t wg) @@ -65,6 +69,9 @@ void columnwise_total(sycl::queue q, q.submit([&](sycl::handler &h) { sycl::accessor mat_acc{mat_buffer, h, sycl::read_only}; sycl::accessor ct_acc{ct_buffer, h}; +#ifdef __ADAPTIVECPP__ + sycl::local_accessor local_mem_acc{sycl::range<1>(wg), h}; +#endif sycl::range<2> global{upper_multiple(n, wg), m}; sycl::range<2> local{wg, 1}; @@ -72,10 +79,14 @@ void columnwise_total(sycl::queue q, h.parallel_for( sycl::nd_range<2>(global, local), [=](sycl::nd_item<2> it) { size_t i = it.get_global_id(0); - dataT group_sum = sycl::reduce_over_group( - it.get_group(), - (i < n) ? mat_acc[it.get_global_id()] : dataT(0), - std::plus()); + dataT val = (i < n) ? mat_acc[it.get_global_id()] : dataT(0); +#ifdef __ADAPTIVECPP__ + dataT group_sum = custom_reduce_over_group( + it.get_group(), local_mem_acc, val, std::plus()); +#else + dataT group_sum = sycl::reduce_over_group(it.get_group(), val, + std::plus()); +#endif if (it.get_group().leader()) { size_t j = it.get_global_id(1); sycl::atomic_ref(err_code)); + if (handler_) { + auto err_code = e.code().value(); + handler_(static_cast(err_code)); + } } } } diff --git a/libsyclinterface/helper/source/dpctl_utils_helper.cpp b/libsyclinterface/helper/source/dpctl_utils_helper.cpp index c03887a2b1..55b4630ed3 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,18 @@ 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_LEVEL_ZERO: + return backend::level_zero; + case DPCTLSyclBackendType::DPCTL_OPENCL: + return backend::ocl; + case DPCTLSyclBackendType::DPCTL_OPENMP: + return backend::omp; +#endif default: throw std::runtime_error("Unsupported backend type"); } @@ -102,6 +115,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 +124,18 @@ 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::level_zero: + return DPCTLSyclBackendType::DPCTL_LEVEL_ZERO; + case backend::ocl: + return DPCTLSyclBackendType::DPCTL_OPENCL; + case backend::omp: + return DPCTLSyclBackendType::DPCTL_OPENMP; +#endif default: return DPCTLSyclBackendType::DPCTL_UNKNOWN_BACKEND; } @@ -219,12 +245,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 +327,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 +388,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 +448,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 +603,7 @@ DPCTLPartitionPropertyType DPCTL_SyclPartitionPropertyToDPCTLType( } } +#ifndef __ADAPTIVECPP__ ext::oneapi::peer_access DPCTL_DPCTLPeerAccessTypeToSycl(DPCTLPeerAccessType PeerAccessTy) { @@ -590,6 +629,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 +657,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 +674,27 @@ 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::level_zero: + ss << "level_zero"; + break; + case backend::ocl: + ss << "opencl"; + break; + case backend::omp: + ss << "openmp"; + break; + default: + ss << "unknown"; + }; +#endif ss << filter_string_separator; ss << DPCTL_DeviceTypeToStr(Device.get_info()); diff --git a/libsyclinterface/include/syclinterface/Config/dpctl_config.h.in b/libsyclinterface/include/syclinterface/Config/dpctl_config.h.in index 055dd4ae8d..5e69fe80ac 100644 --- a/libsyclinterface/include/syclinterface/Config/dpctl_config.h.in +++ b/libsyclinterface/include/syclinterface/Config/dpctl_config.h.in @@ -34,6 +34,9 @@ /* The DPCPP version used to build dpctl */ #define DPCTL_DPCPP_VERSION "@IntelSyclCompiler_VERSION@" +/* The SYCL implementation used to build dpctl, "Intel" or "AdaptiveCpp" */ +#define DPCTL_SYCL_PROVIDER "@DPCTL_SYCL_PROVIDER@" + #define DPCTL_LIBZE_LOADER_FILENAME "@LIBZE_LOADER_FILENAME@" #define DPCTL_LIBCL_LOADER_FILENAME "@LIBCL_LOADER_FILENAME@" /* Defined when the SYCL runtime supports IPC memory (ipc::memory API). */ diff --git a/libsyclinterface/include/syclinterface/dpctl_device_selection.hpp b/libsyclinterface/include/syclinterface/dpctl_device_selection.hpp index 7af8ef5e86..f461510545 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 @@ -28,7 +27,9 @@ #pragma once #include "Support/DllExport.h" +#include #include +#include namespace dpctl { @@ -74,11 +75,17 @@ class DPCTL_API dpctl_cpu_selector : public dpctl_device_selector class DPCTL_API dpctl_filter_selector : public dpctl_device_selector { public: - dpctl_filter_selector(const std::string &fs) : _impl(fs) {} + dpctl_filter_selector(const std::string &fs); int operator()(const sycl::device &d) const override; private: +#ifndef __ADAPTIVECPP__ sycl::ext::oneapi::filter_selector _impl; +#else + // AdaptiveCpp has no filter_selector extension, so the filter string is + // resolved to devices when the selector is constructed. + std::vector _matches; +#endif }; } // namespace syclinterface diff --git a/libsyclinterface/include/syclinterface/dpctl_service.h b/libsyclinterface/include/syclinterface/dpctl_service.h index 1856f45702..1b13136755 100644 --- a/libsyclinterface/include/syclinterface/dpctl_service.h +++ b/libsyclinterface/include/syclinterface/dpctl_service.h @@ -43,6 +43,15 @@ DPCTL_C_EXTERN_C_BEGIN DPCTL_API __dpctl_give const char *DPCTLService_GetDPCPPVersion(void); +/*! + * @brief Get the SYCL implementation the library was compiled with. + * + * @return A C string, either "Intel" or "AdaptiveCpp". + * @ingroup Service + */ +DPCTL_API +__dpctl_give const char *DPCTLService_GetSyclProvider(void); + /*! * @brief Initialize logger if compiled to use logger, no-op otherwise. * diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_enum_types.h b/libsyclinterface/include/syclinterface/dpctl_sycl_enum_types.h index c66b8330f9..67ece16a69 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_enum_types.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_enum_types.h @@ -56,6 +56,7 @@ typedef enum DPCTL_UNKNOWN_BACKEND = 0, DPCTL_ALL_BACKENDS = ((1<<5)-1) << 16, DPCTL_HIP = 1 << 19, + DPCTL_OPENMP = 1 << 20, // clang-format on } DPCTLSyclBackendType; 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..24836fcdc2 100644 --- a/libsyclinterface/source/dpctl_device_selection.cpp +++ b/libsyclinterface/source/dpctl_device_selection.cpp @@ -27,14 +27,135 @@ #include "dpctl_device_selection.hpp" #include "Config/dpctl_config.h" +#include +#include +#include #include +#include namespace { +#ifndef __ADAPTIVECPP__ static_assert(__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_VERSION_REQUIRED, "The compiler does not meet minimum version requirement"); +#endif +} // namespace + +#ifdef __ADAPTIVECPP__ +namespace +{ + +std::vector split_string(const std::string &s, char sep) +{ + std::vector terms{}; + std::string::size_type pos = 0; + while (true) { + auto next = s.find(sep, pos); + if (next == std::string::npos) { + terms.emplace_back(s.substr(pos)); + return terms; + } + terms.emplace_back(s.substr(pos, next - pos)); + pos = next + 1; + } +} + +bool parse_backend(const std::string &term, sycl::backend &backend) +{ + if (term == "opencl") + backend = sycl::backend::ocl; + else if (term == "level_zero") + backend = sycl::backend::level_zero; + else if (term == "cuda") + backend = sycl::backend::cuda; + else if (term == "hip") + backend = sycl::backend::hip; + else if (term == "openmp") + backend = sycl::backend::omp; + else + return false; + return true; +} + +bool parse_device_type(const std::string &term, + sycl::info::device_type &device_type) +{ + if (term == "cpu") + device_type = sycl::info::device_type::cpu; + else if (term == "gpu") + device_type = sycl::info::device_type::gpu; + else if (term == "accelerator") + device_type = sycl::info::device_type::accelerator; + else if (term == "custom") + device_type = sycl::info::device_type::custom; + else + return false; + return true; +} + +bool parse_index(const std::string &term, long long &index) +{ + if (term.empty() || + term.find_first_not_of("0123456789") != std::string::npos) + return false; + index = std::stoll(term); + return true; } +/*! + * @brief Appends devices matched by a single "backend:device_type:index" + * filter, each term of which is optional, to the selected vector. + * + * Devices are enumerated the way DPCTLDeviceMgr_GetDevices enumerates them, so + * that an index term agrees with the relative id in a device's filter string. + */ +void resolve_filter(const std::string &filter, + std::vector &selected) +{ + auto terms = split_string(filter, ':'); + sycl::backend backend{}; + sycl::info::device_type device_type{}; + long long index = -1; + bool has_backend = false; + bool has_device_type = false; + + std::size_t pos = 0; + if (pos < terms.size() && parse_backend(terms[pos], backend)) { + has_backend = true; + ++pos; + } + if (pos < terms.size() && parse_device_type(terms[pos], device_type)) { + has_device_type = true; + ++pos; + } + if (pos < terms.size() && parse_index(terms[pos], index)) + ++pos; + if (pos != terms.size()) + throw std::invalid_argument("Could not parse filter string '" + filter + + "'"); + + dpctl::syclinterface::dpctl_default_selector ranker; + long long matched = 0; + for (const auto &d : sycl::device::get_devices()) { + if (ranker(d) < 0) + continue; + if (has_backend && d.get_platform().get_backend() != backend) + continue; + if (has_device_type && + d.get_info() != device_type) + continue; + if (index < 0) + selected.emplace_back(d); + else if (matched++ == index) { + selected.emplace_back(d); + break; + } + } +} + +} // namespace +#endif + namespace dpctl { namespace syclinterface @@ -47,7 +168,17 @@ int dpctl_device_selector::operator()(const sycl::device &) const int dpctl_accelerator_selector::operator()(const sycl::device &d) const { +#ifndef __ADAPTIVECPP__ return sycl::accelerator_selector_v(d); +#else + // AdaptiveCpp defines is_accelerator() as !is_cpu(), which makes its + // accelerator_selector_v accept GPUs as well + auto score = sycl::default_selector_v(d); + if (score < 0 || d.get_info() != + sycl::info::device_type::accelerator) + return REJECT_DEVICE; + return score + 1; +#endif } int dpctl_default_selector::operator()(const sycl::device &d) const @@ -66,9 +197,32 @@ int dpctl_cpu_selector::operator()(const sycl::device &d) const return sycl::cpu_selector_v(d); } +#ifndef __ADAPTIVECPP__ +dpctl_filter_selector::dpctl_filter_selector(const std::string &fs) : _impl(fs) +{ +} +#else +dpctl_filter_selector::dpctl_filter_selector(const std::string &fs) +{ + for (const auto &filter : split_string(fs, ',')) + resolve_filter(filter, _matches); +} +#endif + int dpctl_filter_selector::operator()(const sycl::device &d) const { +#ifndef __ADAPTIVECPP__ return _impl(d); +#else + for (const auto &m : _matches) + if (m == d) { + // scores of matched devices are shifted to be strictly positive, + // since AdaptiveCpp scores the OpenMP host device 0 + auto score = sycl::default_selector_v(d); + return (score < 0) ? REJECT_DEVICE : score + 1; + } + return REJECT_DEVICE; +#endif } } // namespace syclinterface diff --git a/libsyclinterface/source/dpctl_service.cpp b/libsyclinterface/source/dpctl_service.cpp index b0c1c1e300..699cba26f3 100644 --- a/libsyclinterface/source/dpctl_service.cpp +++ b/libsyclinterface/source/dpctl_service.cpp @@ -41,6 +41,12 @@ __dpctl_give const char *DPCTLService_GetDPCPPVersion(void) return dpctl::helper::cstring_from_string(version); } +__dpctl_give const char *DPCTLService_GetSyclProvider(void) +{ + std::string provider = DPCTL_SYCL_PROVIDER; + return dpctl::helper::cstring_from_string(provider); +} + #ifdef ENABLE_GLOG void DPCTLService_InitLogger(const char *app_name, const char *log_dir) diff --git a/libsyclinterface/source/dpctl_sycl_context_interface.cpp b/libsyclinterface/source/dpctl_sycl_context_interface.cpp index cffadad953..4097241377 100644 --- a/libsyclinterface/source/dpctl_sycl_context_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_context_interface.cpp @@ -38,8 +38,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 @@ -187,18 +189,7 @@ DPCTLContext_GetBackend(__dpctl_keep const DPCTLSyclContextRef CtxRef) auto BE = unwrap(CtxRef)->get_platform().get_backend(); - switch (BE) { - case backend::opencl: - return DPCTL_OPENCL; - case backend::ext_oneapi_level_zero: - return DPCTL_LEVEL_ZERO; - case backend::ext_oneapi_cuda: - return DPCTL_CUDA; - case backend::ext_oneapi_hip: - return DPCTL_HIP; - default: - return DPCTL_UNKNOWN_BACKEND; - } + return DPCTL_SyclBackendToDPCTLBackendType(BE); } size_t DPCTLContext_Hash(__dpctl_keep const DPCTLSyclContextRef CtxRef) @@ -230,6 +221,8 @@ DPCTLContext_GetPlatform(__dpctl_keep const DPCTLSyclContextRef CtxRef) return PRef; } +#ifndef __ADAPTIVECPP__ + namespace { @@ -298,3 +291,49 @@ __dpctl_give int *DPCTLContext_GetAtomicFenceScopeCapabilities( info::context::atomic_fence_scope_capabilities>( CtxRef, res_len, DPCTL_SyclMemoryScopeToDPCTLType); } + +#else + +namespace +{ + +int *unsupported_context_capabilities(size_t *res_len) +{ + error_handler("Atomic capabilities of a context are not queryable in " + "AdaptiveCpp", + __FILE__, __func__, __LINE__, error_level::error); + *res_len = 0; + return nullptr; +} + +} // end of anonymous namespace + +__dpctl_give int *DPCTLContext_GetAtomicMemoryOrderCapabilities( + __dpctl_keep const DPCTLSyclContextRef, + size_t *res_len) +{ + return unsupported_context_capabilities(res_len); +} + +__dpctl_give int *DPCTLContext_GetAtomicFenceOrderCapabilities( + __dpctl_keep const DPCTLSyclContextRef, + size_t *res_len) +{ + return unsupported_context_capabilities(res_len); +} + +__dpctl_give int *DPCTLContext_GetAtomicMemoryScopeCapabilities( + __dpctl_keep const DPCTLSyclContextRef, + size_t *res_len) +{ + return unsupported_context_capabilities(res_len); +} + +__dpctl_give int *DPCTLContext_GetAtomicFenceScopeCapabilities( + __dpctl_keep const DPCTLSyclContextRef, + size_t *res_len) +{ + return unsupported_context_capabilities(res_len); +} + +#endif /* #ifndef __ADAPTIVECPP__ */ diff --git a/libsyclinterface/source/dpctl_sycl_device_interface.cpp b/libsyclinterface/source/dpctl_sycl_device_interface.cpp index 9d4a066a02..14bfb96caa 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) { @@ -838,6 +840,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 = @@ -854,6 +857,7 @@ DPCTLDevice_GetComponentDevices(__dpctl_keep const DPCTLSyclDeviceRef DRef) error_handler(e, __FILE__, __func__, __LINE__); return nullptr; } +#endif } return wrap(ComponentDevicesVectorPtr); } @@ -861,6 +865,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; @@ -884,8 +889,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(); @@ -905,11 +915,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) { @@ -922,12 +935,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) { @@ -943,12 +958,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) { @@ -964,6 +981,7 @@ void DPCTLDevice_DisablePeerAccess(__dpctl_keep const DPCTLSyclDeviceRef DRef, __func__, __LINE__); } } +#endif return; } @@ -1103,6 +1121,7 @@ __dpctl_give const char * DPCTLDevice_GetBackendVersion(__dpctl_keep const DPCTLSyclDeviceRef DRef) { const char *cstr_version = nullptr; +#ifndef __ADAPTIVECPP__ auto D = unwrap(DRef); if (D) { try { @@ -1112,6 +1131,10 @@ DPCTLDevice_GetBackendVersion(__dpctl_keep const DPCTLSyclDeviceRef DRef) error_handler(e, __FILE__, __func__, __LINE__); } } +#else + error_handler("Backend version of a device is not queryable in AdaptiveCpp", + __FILE__, __func__, __LINE__, error_level::error); +#endif return cstr_version; } @@ -1219,6 +1242,8 @@ DPCTLDevice_GetDoubleFPConfig(__dpctl_keep const DPCTLSyclDeviceRef DRef, DRef, res_len, DPCTL_SyclFPConfigToDPCTLType); } +#ifndef __ADAPTIVECPP__ + __dpctl_give int *DPCTLDevice_GetAtomicMemoryOrderCapabilities( __dpctl_keep const DPCTLSyclDeviceRef DRef, size_t *res_len) @@ -1251,6 +1276,52 @@ __dpctl_give int *DPCTLDevice_GetAtomicFenceScopeCapabilities( DRef, res_len, DPCTL_SyclMemoryScopeToDPCTLType); } +#else + +namespace +{ + +int *unsupported_device_capabilities(size_t *res_len) +{ + error_handler("Atomic capabilities of a device are not queryable in " + "AdaptiveCpp", + __FILE__, __func__, __LINE__, error_level::error); + *res_len = 0; + return nullptr; +} + +} // end of anonymous namespace + +__dpctl_give int *DPCTLDevice_GetAtomicMemoryOrderCapabilities( + __dpctl_keep const DPCTLSyclDeviceRef, + size_t *res_len) +{ + return unsupported_device_capabilities(res_len); +} + +__dpctl_give int *DPCTLDevice_GetAtomicFenceOrderCapabilities( + __dpctl_keep const DPCTLSyclDeviceRef, + size_t *res_len) +{ + return unsupported_device_capabilities(res_len); +} + +__dpctl_give int *DPCTLDevice_GetAtomicMemoryScopeCapabilities( + __dpctl_keep const DPCTLSyclDeviceRef, + size_t *res_len) +{ + return unsupported_device_capabilities(res_len); +} + +__dpctl_give int *DPCTLDevice_GetAtomicFenceScopeCapabilities( + __dpctl_keep const DPCTLSyclDeviceRef, + size_t *res_len) +{ + return unsupported_device_capabilities(res_len); +} + +#endif /* #ifndef __ADAPTIVECPP__ */ + __dpctl_give int * DPCTLDevice_GetPartitionProperties(__dpctl_keep const DPCTLSyclDeviceRef DRef, size_t *res_len) @@ -1274,8 +1345,12 @@ bool DPCTLDevice_CanCompileSPIRV(__dpctl_keep const DPCTLSyclDeviceRef DRef) if (Dev) { try { auto Backend = Dev->get_platform().get_backend(); +#ifndef __ADAPTIVECPP__ canCompile = Backend == backend::opencl || Backend == backend::ext_oneapi_level_zero; +#else + canCompile = Backend == backend::ocl; +#endif } catch (std::exception const &e) { error_handler(e, __FILE__, __func__, __LINE__); } @@ -1289,7 +1364,11 @@ bool DPCTLDevice_CanCompileOpenCL(__dpctl_keep const DPCTLSyclDeviceRef DRef) auto Dev = unwrap(DRef); if (Dev) { try { +#ifndef __ADAPTIVECPP__ canCompile = Dev->get_platform().get_backend() == backend::opencl; +#else + canCompile = Dev->get_platform().get_backend() == backend::ocl; +#endif } catch (std::exception const &e) { error_handler(e, __FILE__, __func__, __LINE__); } diff --git a/libsyclinterface/source/dpctl_sycl_device_manager.cpp b/libsyclinterface/source/dpctl_sycl_device_manager.cpp index 2d647c33d0..d5b0c1d134 100644 --- a/libsyclinterface/source/dpctl_sycl_device_manager.cpp +++ b/libsyclinterface/source/dpctl_sycl_device_manager.cpp @@ -44,8 +44,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; @@ -376,12 +378,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_device_selector_interface.cpp b/libsyclinterface/source/dpctl_sycl_device_selector_interface.cpp index c79bc2c4a5..9d5bc516f6 100644 --- a/libsyclinterface/source/dpctl_sycl_device_selector_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_device_selector_interface.cpp @@ -34,8 +34,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 diff --git a/libsyclinterface/source/dpctl_sycl_event_interface.cpp b/libsyclinterface/source/dpctl_sycl_event_interface.cpp index 5e2b924e28..d75463dd64 100644 --- a/libsyclinterface/source/dpctl_sycl_event_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_event_interface.cpp @@ -36,8 +36,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 @@ -126,7 +128,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_ipc_memory_interface.cpp b/libsyclinterface/source/dpctl_sycl_ipc_memory_interface.cpp index bb70c57594..889a649fc4 100644 --- a/libsyclinterface/source/dpctl_sycl_ipc_memory_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_ipc_memory_interface.cpp @@ -50,8 +50,10 @@ using ipc_handle_data_t = ipc::handle_data_t; 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 diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp index 163f9efd8c..7ef0919d39 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_KERNEL_BUNDLE_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,16 +51,21 @@ #include // clang-format on #endif +#endif 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; +#ifndef __ADAPTIVECPP__ + #ifdef __linux__ static const char *clLoaderName = DPCTL_LIBCL_LOADER_FILENAME; static const int clLibLoadFlags = RTLD_NOLOAD | RTLD_NOW | RTLD_LOCAL; @@ -678,6 +688,8 @@ bool _HasKernel_ze_impl(const kernel_bundle &kb, #endif /* #ifdef DPCTL_ENABLE_L0_KERNEL_BUNDLE_CREATION */ +#endif /* #ifndef __ADAPTIVECPP__ */ + } /* end of anonymous namespace */ __dpctl_give DPCTLSyclKernelBundleRef @@ -689,6 +701,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 kernel bundle from SPIR-V as the supplied " @@ -737,6 +750,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 +764,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 +806,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 +843,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 +879,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,6 +912,11 @@ 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 } using build_option_list_t = std::vector; @@ -965,8 +1011,6 @@ const char *DPCTLKernelBuildLog_Get(__dpctl_keep DPCTLKernelBuildLogRef Ref) return reinterpret_cast(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. @@ -987,6 +1031,8 @@ bool DPCTLKernelBundle_CreateFromSYCLSource_Available() } #if (SUPPORTS_SYCL_COMPILATION > 0) +namespace syclex = sycl::ext::oneapi::experimental; + // 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 diff --git a/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp index e035ea7000..3659c88f32 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_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 @@ -89,7 +91,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 +116,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 +141,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 +167,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 +193,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 +219,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 +246,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..0998d04636 100644 --- a/libsyclinterface/source/dpctl_sycl_platform_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_platform_interface.cpp @@ -45,8 +45,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; @@ -342,12 +344,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..68a15b4181 100644 --- a/libsyclinterface/source/dpctl_sycl_platform_manager.cpp +++ b/libsyclinterface/source/dpctl_sycl_platform_manager.cpp @@ -42,8 +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; @@ -74,7 +76,24 @@ 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::level_zero) + ss << "level_zero"; + else if (be == sycl::backend::ocl) + ss << "opencl"; + else if (be == sycl::backend::omp) + ss << "openmp"; + 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 a7d2a49c8b..e81da06a48 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,11 +120,22 @@ 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 { +#ifndef __ADAPTIVECPP__ static_assert(__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_VERSION_REQUIRED, "The compiler does not meet minimum version requirement"); +#endif using namespace dpctl::syclinterface; @@ -133,6 +145,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 +186,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 +283,8 @@ void set_kernel_args(handler &cgh, } } +#endif /* #ifndef __ADAPTIVECPP__ */ + std::unique_ptr create_property_list(int properties) { std::unique_ptr propList; @@ -347,7 +363,14 @@ DPCTLQueue_Create(__dpctl_keep const DPCTLSyclContextRef CRef, } else { try { +#ifdef __ADAPTIVECPP__ + // a queue with no handler of its own inherits the default handler + // of the AdaptiveCpp context, which calls std::terminate + auto Queue = new queue(*ctx, *dev, DPCTL_AsyncErrorHandler(nullptr), + *propList); +#else auto Queue = new queue(*ctx, *dev, *propList); +#endif q = wrap(Queue); } catch (std::exception const &e) { error_handler(e, __FILE__, __func__, __LINE__); @@ -509,6 +532,7 @@ DPCTLQueue_SubmitRange(__dpctl_keep const DPCTLSyclKernelRef KRef, return nullptr; } +#ifndef __ADAPTIVECPP__ auto Kernel = unwrap(KRef); auto Queue = unwrap(QRef); event e; @@ -518,7 +542,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); @@ -528,7 +551,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); @@ -538,7 +560,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]}, @@ -547,8 +568,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; } @@ -560,6 +580,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 @@ -585,6 +611,7 @@ DPCTLQueue_SubmitNDRange(__dpctl_keep const DPCTLSyclKernelRef KRef, return nullptr; } +#ifndef __ADAPTIVECPP__ auto Kernel = unwrap(KRef); auto Queue = unwrap(QRef); event e; @@ -594,7 +621,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]}}, @@ -605,7 +631,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( @@ -617,7 +642,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]}, @@ -627,8 +651,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; } @@ -640,6 +663,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) @@ -882,7 +911,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__); diff --git a/libsyclinterface/source/dpctl_sycl_usm_interface.cpp b/libsyclinterface/source/dpctl_sycl_usm_interface.cpp index 82f566cf1e..4b6e1e4776 100644 --- a/libsyclinterface/source/dpctl_sycl_usm_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_usm_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 diff --git a/scripts/_build_helper.py b/scripts/_build_helper.py index dfbc86f46d..daaf3af827 100644 --- a/scripts/_build_helper.py +++ b/scripts/_build_helper.py @@ -25,13 +25,25 @@ def resolve_compilers( c_compiler: str, cxx_compiler: str, compiler_root: str, + sycl_provider: str = "Intel", ): is_linux = "linux" in sys.platform + if sycl_provider == "AdaptiveCpp": + if oneapi: + raise RuntimeError( + "--oneapi is specific to DPC++ and can not be used with the " + "AdaptiveCpp SYCL provider" + ) + default_c_compiler, default_cxx_compiler = "clang", "acpp" + else: + default_c_compiler = "icx" + default_cxx_compiler = "icpx" if is_linux else "icx" + if oneapi or ( c_compiler is None and cxx_compiler is None and compiler_root is None ): - return "icx", ("icpx" if is_linux else "icx") + return default_c_compiler, default_cxx_compiler if ( (c_compiler is None or not os.path.isabs(c_compiler)) @@ -45,9 +57,9 @@ def resolve_compilers( # default values if c_compiler is None: - c_compiler = "icx" + c_compiler = default_c_compiler if cxx_compiler is None: - cxx_compiler = "icpx" if is_linux else "icx" + cxx_compiler = default_cxx_compiler for name, opt_name in ( (c_compiler, "--c-compiler"), @@ -55,9 +67,12 @@ def resolve_compilers( ): if os.path.isabs(name): path = name - else: + elif compiler_root: path = os.path.join(compiler_root, name) - if not os.path.exists(path): + else: + # a default name with no compiler root is looked up in PATH + path = shutil.which(name) + if not path or not os.path.exists(path): raise RuntimeError(f"{opt_name} value {name} not found") return c_compiler, cxx_compiler @@ -92,6 +107,7 @@ def make_cmake_args( level_zero: bool = True, glog: bool = False, verbose: bool = False, + sycl_provider: str = "Intel", other_opts: str = None, ): args = [ @@ -100,6 +116,7 @@ def make_cmake_args( "-DDPCTL_ENABLE_L0_KERNEL_BUNDLE_CREATION=" f"{'ON' if level_zero else 'OFF'}", f"-DDPCTL_ENABLE_GLOG:BOOL={'ON' if glog else 'OFF'}", + f"-DDPCTL_SYCL_PROVIDER={sycl_provider}", ] if verbose: diff --git a/scripts/build_locally.py b/scripts/build_locally.py index 6918a429e5..e9fa09ef2c 100644 --- a/scripts/build_locally.py +++ b/scripts/build_locally.py @@ -51,6 +51,14 @@ def parse_args(): help="Path to compiler installation root", ) + p.add_argument( + "--sycl-provider", + type=str, + choices=["Intel", "AdaptiveCpp"], + default="Intel", + help="SYCL implementation to build with (defaults to Intel).", + ) + p.add_argument( "--oneapi", dest="oneapi", @@ -139,15 +147,23 @@ def main(): setup_dir = os.path.dirname(os.path.dirname(os.path.abspath(__file__))) c_compiler, cxx_compiler = resolve_compilers( - args.oneapi, args.c_compiler, args.cxx_compiler, args.compiler_root + args.oneapi, + args.c_compiler, + args.cxx_compiler, + args.compiler_root, + sycl_provider=args.sycl_provider, ) # clean build dir if --clean set if args.clean: clean_build_dir(setup_dir) - # Level Zero state (on unless explicitly disabled) - level_zero_enabled = False if args.no_level_zero else True + # Level Zero state (on unless explicitly disabled), the backend is only + # implemented for DPC++ + if args.sycl_provider == "AdaptiveCpp": + level_zero_enabled = False + else: + level_zero_enabled = False if args.no_level_zero else True cmake_args = make_cmake_args( c_compiler=c_compiler, @@ -155,6 +171,7 @@ def main(): level_zero=level_zero_enabled, glog=args.glog, verbose=args.verbose, + sycl_provider=args.sycl_provider, other_opts=args.cmake_opts, )