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/CHANGELOG.md b/CHANGELOG.md index f660fa239b..05a053bf6f 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -11,6 +11,9 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 * Added `add_event` and `add_cleanup_event` methods and `cleanup_events` and `num_cleanup_events` properties to the order manager, for tracking the events that gate the release of objects used by offloaded tasks [gh-2359](https://github.com/IntelPython/dpctl/pull/2359) * Added `SyclQueue_DeferRelease` C-API function, which defers a release callback until work submitted to a queue completes, and through which `dpctl::utils::keep_args_alive` keeps its arguments alive [gh-2359](https://github.com/IntelPython/dpctl/pull/2359) * Added `dpctl::utils::defer_release` to `dpctl4pybind11.hpp`, which calls a release callback once the given events have completed [gh-2359](https://github.com/IntelPython/dpctl/pull/2359) +* Added support for building `dpctl` with AdaptiveCpp as the SYCL implementation, selected with `--sycl-provider=AdaptiveCpp` in `scripts/build_locally.py` or `-DDPCTL_SYCL_PROVIDER=AdaptiveCpp` [gh-2384](https://github.com/IntelPython/dpctl/pull/2384) +* Added `dpctl.backend_type.openmp` for the AdaptiveCpp OpenMP host device [gh-2384](https://github.com/IntelPython/dpctl/pull/2384) +* Added `DPCTLService_GetSyclProvider` C-API function and `dpctl._diagnostics.sycl_provider`, which report the SYCL implementation `dpctl` was built with [gh-2384](https://github.com/IntelPython/dpctl/pull/2384) ### Changed * Object lifetime during offload is now managed by a list of releases deferred until the events that gate them have completed, rather than by a `host_task`. `dpctl` keeps one thread for looking over the list and running the releases that have come due, started by the first release deferred and again in the child of a `fork`, which takes the GIL to do so and lets a release call into the interpreter and free USM allocations [gh-2359](https://github.com/IntelPython/dpctl/pull/2359) @@ -25,6 +28,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 * Removed the deprecated `dpctl.program` submodule, use `dpctl.compiler` instead [gh-2407](https://github.com/IntelPython/dpctl/pull/2407) ### Fixed +* `dpctl.select_device_with_aspects` no longer raises `SyclDeviceCreationError` when every device with the required aspects has a selector score of zero [gh-2384](https://github.com/IntelPython/dpctl/pull/2384) * Removed stale `select_host_device` and `has_host_device` names from `dpctl.__all__`, which made `from dpctl import *` fail with `AttributeError` [gh-2410](https://github.com/IntelPython/dpctl/pull/2410) * A USM allocation that fails now drops the references held for offloaded tasks that have completed, and is attempted once more if that released anything, so that memory only waiting to be given up is not reported as unavailable [gh-2359](https://github.com/IntelPython/dpctl/pull/2359) diff --git a/CMakeLists.txt b/CMakeLists.txt index 594dd0bc24..b9c28e5390 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -26,9 +26,18 @@ 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) +if(NOT DPCTL_SYCL_PROVIDER MATCHES "^(Intel|AdaptiveCpp)$") + message(FATAL_ERROR + "DPCTL_SYCL_PROVIDER must be Intel or AdaptiveCpp, not '${DPCTL_SYCL_PROVIDER}'" + ) +endif() + 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." @@ -49,13 +58,10 @@ 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 +75,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 @@ -86,6 +102,11 @@ if ("x${DPCTL_SYCL_TARGETS}" STREQUAL "x") ) endif() endif() + + if(_acpp_targets_list) + # generic runs dpctl's kernels on the OpenMP and OpenCL devices + list(PREPEND _acpp_targets_list generic) + endif() else() if ("${DPCTL_SYCL_TARGETS}" MATCHES "amdgcn-amd-amdhsa") message(FATAL_ERROR @@ -94,12 +115,41 @@ 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 ACPP_TARGETS during find_package. The empty cache entry + # keeps the derived targets below out of the cache. + set(ACPP_TARGETS "" CACHE STRING + "AdaptiveCpp targets, derived from DPCTL_TARGET_CUDA, DPCTL_TARGET_HIP or DPCTL_SYCL_TARGETS if empty" + ) + if(_acpp_targets_list AND NOT ACPP_TARGETS AND "$ENV{ACPP_TARGETS}" STREQUAL "") + set(ACPP_TARGETS "${_acpp_targets_list}") + 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..9a6c05685c 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,91 @@ 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 +- peer access between devices, so ``enable_peer_access`` and + ``disable_peer_access`` raise ``ValueError`` +- 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..abeee9c940 100644 --- a/dpctl/_sycl_device.pyx +++ b/dpctl/_sycl_device.pyx @@ -157,6 +157,8 @@ import collections import functools import warnings +from ._diagnostics import sycl_provider + __all__ = [ "SyclDevice", "SyclDeviceCreationError", "SyclSubDeviceCreationError", ] @@ -247,6 +249,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" @@ -310,6 +314,8 @@ cdef inline void _raise_invalid_peer_access( """ Check peer access ahead of time and raise errors for invalid cases. """ + if sycl_provider() == "AdaptiveCpp": + raise ValueError("Peer access is not supported by AdaptiveCpp") cdef list _peer_access_backends = [ _backend_type._CUDA, _backend_type._HIP, @@ -560,6 +566,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..0236fb0934 100644 --- a/dpctl/tests/_device_attributes_checks.py +++ b/dpctl/tests/_device_attributes_checks.py @@ -18,6 +18,13 @@ import dpctl +from .helper import is_adaptivecpp + +# conftest imports this module before it registers unsupported_on_acpp +_unsupported_on_acpp = pytest.mark.skipif( + is_adaptivecpp(), reason="not supported by AdaptiveCpp" +) + list_of_standard_selectors = [ dpctl.select_accelerator_device, dpctl.select_cpu_device, @@ -866,7 +873,8 @@ 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=_unsupported_on_acpp), check_backend, check_device_type, check_global_mem_cache_type, @@ -883,17 +891,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=_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=_unsupported_on_acpp, + ), + pytest.param( + check_atomic_fence_order_capabilities, + marks=_unsupported_on_acpp, + ), + pytest.param( + check_atomic_memory_scope_capabilities, + marks=_unsupported_on_acpp, + ), + pytest.param( + check_atomic_fence_scope_capabilities, + marks=_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_device.py b/dpctl/tests/test_sycl_device.py index dd1a579272..d459363941 100644 --- a/dpctl/tests/test_sycl_device.py +++ b/dpctl/tests/test_sycl_device.py @@ -21,7 +21,7 @@ import dpctl from dpctl import SyclDeviceCreationError -from .helper import get_queue_or_skip +from .helper import get_queue_or_skip, is_adaptivecpp def test_standard_selectors(device_selector, check): @@ -389,6 +389,20 @@ def test_enable_disable_peer_access(platform_name): ) +@pytest.mark.skipif(not is_adaptivecpp(), reason="AdaptiveCpp only") +@pytest.mark.parametrize( + "method", ["enable_peer_access", "disable_peer_access"] +) +def test_peer_access_unsupported_on_acpp(method): + devices = dpctl.get_devices() + if len(devices) < 2: + pytest.skip("Need two devices") + dev0, dev1 = devices[0], devices[1] + assert not dev0.can_access_peer(dev1) + with pytest.raises(ValueError): + getattr(dev0, method)(dev1) + + @pytest.mark.parametrize( "method", [ 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..90e45eb625 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): @@ -169,6 +169,24 @@ def test_queue_submit_barrier(valid_filter): q.submit_barrier(range(3)) +def test_queue_submit_barrier_waits_for_prior_work(): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default queue could not be created") + from dpctl.memory import MemoryUSMDevice, USMAllocationError + + nbytes = 256 * 1024 * 1024 + try: + m1 = MemoryUSMDevice(nbytes, queue=q) + m2 = MemoryUSMDevice(nbytes, queue=q) + except USMAllocationError: + pytest.skip("Not enough memory") + ev = q.memcpy_async(m1, m2, nbytes) + q.submit_barrier().wait() + assert ev.execution_status == dpctl.event_status_type.complete + + def test_queue__repr__(): try: q1 = dpctl.SyclQueue(property=0) @@ -338,19 +356,32 @@ 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 + ) + acpp_skip = is_adaptivecpp() and (is_wsl_or_windows() or acpp_opencl_cpu) + if not acpp_skip: + 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..b872410dab 100644 --- a/dpctl/tests/test_sycl_usm.py +++ b/dpctl/tests/test_sycl_usm.py @@ -27,6 +27,8 @@ as_usm_memory, ) +from .helper import is_adaptivecpp + class Dummy(MemoryUSMShared): """ @@ -83,6 +85,22 @@ 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. + """ + if not is_adaptivecpp(): + return + try: + m = MemoryUSMDevice(8) if q is None else MemoryUSMDevice(8, queue=q) + except (dpctl.SyclDeviceCreationError, dpctl.SyclQueueCreationError): + pytest.skip("No SYCL device available") + 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 +109,7 @@ def _create_host_buf(nbytes): def test_memory_without_context(): + _skip_if_usm_types_not_reported() mobj = _create_memory() # Without context @@ -99,6 +118,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 +130,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 +139,7 @@ def test_memory_cpu_context(): def test_memory_gpu_context(): + _skip_if_usm_types_not_reported() mobj = _create_memory() # GPU context @@ -127,6 +149,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 +171,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 +184,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 +207,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 +248,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 +260,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 +268,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 +279,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 +340,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 +371,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 +401,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 +645,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 +674,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..ce8d8f12ed 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_type_casters.hpp +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_type_casters.hpp @@ -64,9 +64,12 @@ 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" 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..416ebea740 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,9 @@ DPCTLDevice_GetComponentDevices(__dpctl_keep const DPCTLSyclDeviceRef DRef) error_handler(e, __FILE__, __func__, __LINE__); return nullptr; } +#else + ComponentDevicesVectorPtr = new vecTy(); +#endif } return wrap(ComponentDevicesVectorPtr); } @@ -861,6 +867,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 +891,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 +917,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 +937,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 +960,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 +983,7 @@ void DPCTLDevice_DisablePeerAccess(__dpctl_keep const DPCTLSyclDeviceRef DRef, __func__, __LINE__); } } +#endif return; } @@ -1103,6 +1123,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 +1133,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 +1244,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 +1278,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) @@ -1270,6 +1343,7 @@ __dpctl_give int *DPCTLDevice_GetPartitionAffinityDomains( bool DPCTLDevice_CanCompileSPIRV(__dpctl_keep const DPCTLSyclDeviceRef DRef) { bool canCompile = false; +#ifndef __ADAPTIVECPP__ auto Dev = unwrap(DRef); if (Dev) { try { @@ -1280,12 +1354,14 @@ bool DPCTLDevice_CanCompileSPIRV(__dpctl_keep const DPCTLSyclDeviceRef DRef) error_handler(e, __FILE__, __func__, __LINE__); } } +#endif return canCompile; } bool DPCTLDevice_CanCompileOpenCL(__dpctl_keep const DPCTLSyclDeviceRef DRef) { bool canCompile = false; +#ifndef __ADAPTIVECPP__ auto Dev = unwrap(DRef); if (Dev) { try { @@ -1294,6 +1370,7 @@ bool DPCTLDevice_CanCompileOpenCL(__dpctl_keep const DPCTLSyclDeviceRef DRef) error_handler(e, __FILE__, __func__, __LINE__); } } +#endif return canCompile; } 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..8139cdeff7 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -25,20 +25,24 @@ /// //===----------------------------------------------------------------------===// -#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 +50,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 +687,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 +700,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 +749,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 +763,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 +805,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 +842,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 +878,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 +911,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 +1010,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 +1030,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..92271db84b 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,14 @@ using namespace sycl; 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 +137,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 +178,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 +275,8 @@ void set_kernel_args(handler &cgh, } } +#endif /* #ifndef __ADAPTIVECPP__ */ + std::unique_ptr create_property_list(int properties) { std::unique_ptr propList; @@ -347,7 +355,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 +524,7 @@ DPCTLQueue_SubmitRange(__dpctl_keep const DPCTLSyclKernelRef KRef, return nullptr; } +#ifndef __ADAPTIVECPP__ auto Kernel = unwrap(KRef); auto Queue = unwrap(QRef); event e; @@ -518,7 +534,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 +543,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 +552,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 +560,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 +572,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 +603,7 @@ DPCTLQueue_SubmitNDRange(__dpctl_keep const DPCTLSyclKernelRef KRef, return nullptr; } +#ifndef __ADAPTIVECPP__ auto Kernel = unwrap(KRef); auto Queue = unwrap(QRef); event e; @@ -594,7 +613,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 +623,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 +634,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 +643,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 +655,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) @@ -876,13 +897,24 @@ __dpctl_give DPCTLSyclEventRef DPCTLQueue_SubmitBarrierForEvents( event e; if (Q) { try { +#ifdef __ADAPTIVECPP__ + // A kernel waits only for its dependencies, so depend on all prior + // work + auto prior = Q->get_wait_list(); +#endif e = Q->submit([&](handler &cgh) { // Depend on any event that was specified by the caller. if (NDepEvents) for (auto i = 0ul; i < NDepEvents; ++i) cgh.depends_on(*unwrap(DepEvents[i])); +#ifndef __ADAPTIVECPP__ cgh.ext_oneapi_barrier(); +#else + cgh.depends_on(prior); + 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/libsyclinterface/tests/test_service.cpp b/libsyclinterface/tests/test_service.cpp index 95788fed83..60de3a0a8c 100644 --- a/libsyclinterface/tests/test_service.cpp +++ b/libsyclinterface/tests/test_service.cpp @@ -26,6 +26,7 @@ #include "Config/dpctl_config.h" #include "dpctl_service.h" +#include "dpctl_utils.h" #include #include @@ -46,3 +47,12 @@ TEST(TestServicesFns, ChkDPCPPVersion) // from compiler ASSERT_TRUE(found != std::string::npos); } + +TEST(TestServicesFns, ChkSyclProvider) +{ + auto c_prov = DPCTLService_GetSyclProvider(); + std::string prov = std::string(c_prov); + DPCTLCString_Delete(c_prov); + + ASSERT_TRUE(prov == "Intel" || prov == "AdaptiveCpp"); +} 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, )