diff --git a/sycl/include/sycl/backend.hpp b/sycl/include/sycl/backend.hpp index b12b0310206f2..e9b89dbc3fcfd 100644 --- a/sycl/include/sycl/backend.hpp +++ b/sycl/include/sycl/backend.hpp @@ -25,8 +25,12 @@ #endif #endif #if SYCL_EXT_ONEAPI_BACKEND_HIP +#ifdef SYCL_EXT_ONEAPI_BACKEND_HIP_EXPERIMENTAL +#include +#else #include #endif +#endif #if SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO #include #endif @@ -110,10 +114,8 @@ auto get_native_buffer(const buffer &Obj) // No check for backend mismatch because buffer can be allocated on different // backends if (BackendName == backend::ext_oneapi_level_zero) - throw sycl::runtime_error( - errc::feature_not_supported, - "Buffer interop is not supported by level zero yet", - PI_ERROR_INVALID_OPERATION); + throw sycl::exception(sycl::errc::feature_not_supported, + "Buffer interop is not supported by level zero yet"); return Obj.template getNative(); } #endif @@ -122,10 +124,8 @@ auto get_native_buffer(const buffer &Obj) template auto get_native(const SyclObjectT &Obj) -> backend_return_t { - // TODO use SYCL 2020 exception when implemented if (Obj.get_backend() != BackendName) { - throw sycl::runtime_error(errc::backend_mismatch, "Backends mismatch", - PI_ERROR_INVALID_OPERATION); + throw sycl::exception(sycl::errc::backend_mismatch, "Backends mismatch"); } return reinterpret_cast>( Obj.getNative()); @@ -134,10 +134,8 @@ auto get_native(const SyclObjectT &Obj) template auto get_native(const kernel_bundle &Obj) -> backend_return_t> { - // TODO use SYCL 2020 exception when implemented if (Obj.get_backend() != BackendName) { - throw sycl::runtime_error(errc::backend_mismatch, "Backends mismatch", - PI_ERROR_INVALID_OPERATION); + throw sycl::exception(sycl::errc::backend_mismatch, "Backends mismatch"); } return Obj.template getNative(); } @@ -153,10 +151,8 @@ auto get_native(const buffer &Obj) template <> inline backend_return_t get_native(const event &Obj) { - // TODO use SYCL 2020 exception when implemented if (Obj.get_backend() != backend::opencl) { - throw sycl::runtime_error(errc::backend_mismatch, "Backends mismatch", - PI_ERROR_INVALID_OPERATION); + throw sycl::exception(sycl::errc::backend_mismatch, "Backends mismatch"); } backend_return_t ReturnValue; for (auto const &element : Obj.getNativeVector()) { @@ -173,10 +169,8 @@ get_native(const event &Obj) { template <> inline backend_return_t get_native(const device &Obj) { - // TODO use SYCL 2020 exception when implemented if (Obj.get_backend() != backend::ext_oneapi_cuda) { - throw sycl::runtime_error(errc::backend_mismatch, "Backends mismatch", - PI_ERROR_INVALID_OPERATION); + throw sycl::exception(sycl::errc::backend_mismatch, "Backends mismatch"); } // CUDA uses a 32-bit int instead of an opaque pointer like other backends, // so we need a specialization with static_cast instead of reinterpret_cast. @@ -185,6 +179,20 @@ get_native(const device &Obj) { } #endif +#if SYCL_EXT_ONEAPI_BACKEND_HIP +template <> +inline backend_return_t +get_native(const device &Obj) { + if (Obj.get_backend() != backend::ext_oneapi_hip) { + throw sycl::exception(sycl::errc::backend_mismatch, "Backends mismatch"); + } + // HIP uses a 32-bit int instead of an opaque pointer like other backends, + // so we need a specialization with static_cast instead of reinterpret_cast. + return static_cast>( + Obj.getNative()); +} +#endif + // Native handle of an accessor should be accessed through interop_handler template +#include +#include +#include +#include +#include +#include + +#include + +typedef int HIPdevice; +typedef struct HIPctx_st *HIPcontext; +typedef struct HIPstream_st *HIPstream; +typedef struct HIPevent_st *HIPevent; +typedef struct HIPmod_st *HIPmodule; + +namespace sycl { +__SYCL_INLINE_VER_NAMESPACE(_V1) { +namespace detail { + +// TODO the interops for context, device, event, platform and program +// may be removed after removing the deprecated 'get_native()' methods +// from the corresponding classes. The interop specialization +// is also used in the get_queue() method of the deprecated class +// interop_handler and also can be removed after API cleanup. +template <> struct interop { + using type = HIPcontext; +}; + +template <> struct interop { + using type = HIPdevice; +}; + +template <> struct interop { + using type = HIPevent; +}; + +template <> struct interop { + using type = HIPstream; +}; + +template <> struct interop { + using type = std::vector; +}; + +template +struct BackendInput> { + using type = DataT *; +}; + +template +struct BackendReturn> { + using type = DataT *; +}; + +template <> struct BackendInput { + using type = HIPcontext; +}; + +template <> struct BackendReturn { + using type = std::vector; +}; + +template <> struct BackendInput { + using type = HIPdevice; +}; + +template <> struct BackendReturn { + using type = HIPdevice; +}; + +template <> struct BackendInput { + using type = HIPevent; +}; + +template <> struct BackendReturn { + using type = HIPevent; +}; + +template <> struct BackendInput { + using type = HIPstream; +}; + +template <> struct BackendReturn { + using type = HIPstream; +}; + +template <> struct BackendInput { + using type = std::vector; +}; + +template <> struct BackendReturn { + using type = std::vector; +}; + +template <> struct InteropFeatureSupportMap { + static constexpr bool MakePlatform = false; + static constexpr bool MakeDevice = true; + static constexpr bool MakeContext = true; + static constexpr bool MakeQueue = true; + static constexpr bool MakeEvent = true; + static constexpr bool MakeBuffer = false; + static constexpr bool MakeKernel = false; + static constexpr bool MakeKernelBundle = false; +}; + +} // namespace detail +} // __SYCL_INLINE_VER_NAMESPACE(_V1) +} // namespace sycl diff --git a/sycl/include/sycl/ext/oneapi/experimental/backend/hip.cpp b/sycl/include/sycl/ext/oneapi/experimental/backend/hip.cpp new file mode 100644 index 0000000000000..2eadf82847deb --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/experimental/backend/hip.cpp @@ -0,0 +1,97 @@ +//==--------- hip.hpp - SYCL HIP backend ---------------------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include +#include +#include + +#include + +namespace sycl { +__SYCL_INLINE_VER_NAMESPACE(_V1) { +namespace ext { +namespace oneapi { +namespace hip { + +// Implementation of ext_oneapi_hip::make +inline __SYCL_EXPORT device make_device(pi_native_handle NativeHandle) { + return sycl::detail::make_device(NativeHandle, backend::ext_oneapi_hip); +} + +// Implementation of hip::has_native_event +inline __SYCL_EXPORT bool has_native_event(event sycl_event) { + if (sycl_event.get_backend() == backend::ext_oneapi_hip) + return get_native(sycl_event) != nullptr; + + return false; +} + +} // namespace hip +} // namespace oneapi +} // namespace ext + +// HIP context specialization +template <> +inline auto get_native(const context &C) + -> backend_return_t { + // create a vector to be returned + backend_return_t ret; + + // get the native HIP context from the SYCL object + auto native = reinterpret_cast< + backend_return_t::value_type>( + C.getNative()); + ret.push_back(native); + + return ret; +} + +// Specialisation of interop_handles get_native_context +template <> +inline backend_return_t +interop_handle::get_native_context() const { +#ifndef __SYCL_DEVICE_ONLY__ + return std::vector{reinterpret_cast(getNativeContext())}; +#else + // we believe this won't be ever called on device side + return {}; +#endif +} + +// HIP device specialization +template <> +inline device make_device( + const backend_input_t &BackendObject) { + pi_native_handle NativeHandle = static_cast(BackendObject); + return ext::oneapi::hip::make_device(NativeHandle); +} + +// HIP event specialization +template <> +inline event make_event( + const backend_input_t &BackendObject, + const context &TargetContext) { + return detail::make_event(detail::pi::cast(BackendObject), + TargetContext, true, + /*Backend*/ backend::ext_oneapi_hip); +} + +// HIP queue specialization +template <> +inline queue make_queue( + const backend_input_t &BackendObject, + const context &TargetContext, const async_handler Handler) { + return detail::make_queue(detail::pi::cast(BackendObject), + TargetContext, nullptr, true, Handler, + /*Backend*/ backend::ext_oneapi_hip); +} + +} // __SYCL_INLINE_VER_NAMESPACE(_V1) +} // namespace sycl diff --git a/sycl/plugins/hip/pi_hip.hpp b/sycl/plugins/hip/pi_hip.hpp index f0f46567ab585..8e0e20ddc37b8 100644 --- a/sycl/plugins/hip/pi_hip.hpp +++ b/sycl/plugins/hip/pi_hip.hpp @@ -77,15 +77,15 @@ struct _pi_device { private: using native_type = hipDevice_t; - native_type cuDevice_; + native_type hipDevice_; std::atomic_uint32_t refCount_; pi_platform platform_; public: - _pi_device(native_type cuDevice, pi_platform platform) - : cuDevice_(cuDevice), refCount_{1}, platform_(platform) {} + _pi_device(native_type hipDevice, pi_platform platform) + : hipDevice_(hipDevice), refCount_{1}, platform_(platform) {} - native_type get() const noexcept { return cuDevice_; }; + native_type get() const noexcept { return hipDevice_; }; pi_uint32 get_reference_count() const noexcept { return refCount_; } diff --git a/sycl/test/basic_tests/interop-hip.cpp b/sycl/test/basic_tests/interop-hip.cpp new file mode 100644 index 0000000000000..412759d3d4ced --- /dev/null +++ b/sycl/test/basic_tests/interop-hip.cpp @@ -0,0 +1,98 @@ +// REQUIRES: hip +// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s -o %t.out +// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note -D__SYCL_INTERNAL_API %s -o %t.out +// +/// Also test the experimental HIP interop interface +// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note -DSYCL_EXT_ONEAPI_BACKEND_HIP_EXPERIMENTAL %s -o %t.out +// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note -D__SYCL_INTERNAL_API -DSYCL_EXT_ONEAPI_BACKEND_HIP_EXPERIMENTAL %s -o %t.out +// expected-no-diagnostics + +// Test for legacy and experimental HIP interop API + +#ifdef SYCL_EXT_ONEAPI_BACKEND_HIP_EXPERIMENTAL +#include +#endif + +#include + +using namespace sycl; + +// +// 4.5.1 SYCL application interoperability may be provided for +// platform, +// device, +// context, +// queue, +// event, +// buffer, +// device_image, +// sampled_image, +// unsampled_image. + +int main() { + + // Create SYCL objects + device Device; + context Context(Device); + queue Queue(Device); + event Event; + + // 4.5.1.1 For each SYCL runtime class T which supports SYCL application + // interoperability with the SYCL backend, a specialization of return_type + // must be defined as the type of SYCL application interoperability native + // backend object associated with T for the SYCL backend, specified in the + // SYCL backend specification. + // + // return_type is used when retrieving the backend specific native object from + // a SYCL object. See the relevant backend specification for details. + + backend_traits::return_type hip_device; + backend_traits::return_type hip_context; + backend_traits::return_type hip_event; + backend_traits::return_type hip_queue; + + // 4.5.1.2 For each SYCL runtime class T which supports SYCL application + // interoperability, a specialization of get_native must be defined, which + // takes an instance of T and returns a SYCL application interoperability + // native backend object associated with syclObject which can be used for SYCL + // application interoperability. The lifetime of the object returned are + // backend-defined and specified in the backend specification. + + hip_device = get_native(Device); + hip_context = get_native(Context); + hip_event = get_native(Event); + hip_queue = get_native(Queue); + + // 4.5.1.1 For each SYCL runtime class T which supports SYCL application + // interoperability with the SYCL backend, a specialization of input_type must + // be defined as the type of SYCL application interoperability native backend + // object associated with T for the SYCL backend, specified in the SYCL + // backend specification. input_type is used when constructing SYCL objects + // from backend specific native objects. See the relevant backend + // specification for details. + + // 4.5.1.3 For each SYCL runtime class T which supports SYCL application + // interoperability, a specialization of the appropriate template function + // make_{sycl_class} where {sycl_class} is the class name of T, must be + // defined, which takes a SYCL application interoperability native backend + // object and constructs and returns an instance of T. The availability and + // behavior of these template functions is defined by the SYCL backend + // specification document. + +#ifdef SYCL_EXT_ONEAPI_BACKEND_HIP_EXPERIMENTAL + backend_input_t InteropDeviceInput{ + hip_device}; + device InteropDevice = + make_device(InteropDeviceInput); + + backend_input_t InteropContextInput{ + hip_context[0]}; + context InteropContext = + make_context(InteropContextInput); + event InteropEvent = make_event(hip_event, Context); + + queue InteropQueue = make_queue(hip_queue, Context); +#endif + + return 0; +}