Skip to content

Commit e06d1b5

Browse files
TejaX-Alaghari@TejaX-Alaghari
TejaX-Alaghari
and
@TejaX-Alaghari
authored
[SYCL][PI][HIP] Enabling native HIP implementation methods of PI_HIP through interop (#5880)
Enabling native HIP implementation methods of PI_HIP through interop This minimal fix aims to let the HIP backend implementations of oneMKL libraries use get_native* methods of interop_handle to use PI_HIP backend to retrieve the native objects of HIP for the sycl runtime objects. The PR exposes the PI_HIP backend methods to the libs by adding the template specializations with HIP datatypes for "BackendInput", "BackendReturn" and "interop" classes. In addition to this, this PR also registers and implements "MemGetNativeHandle" for returning the device/host memory allocated with HIP API. This PR is tested for native objects of device, context, queue and memory allocation by successfully running a rocBLAS API through a host_task SYCL kernel with interop_handle on AMD GPU. Co-authored-by: @TejaX-Alaghari <[email protected]>
1 parent 7bc8447 commit e06d1b5

File tree

4 files changed

+155
-2
lines changed

4 files changed

+155
-2
lines changed

‎sycl/include/CL/sycl/backend.hpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -20,6 +20,9 @@
2020
#if SYCL_EXT_ONEAPI_BACKEND_CUDA
2121
#include <CL/sycl/detail/backend_traits_cuda.hpp>
2222
#endif
23+
#if SYCL_EXT_ONEAPI_BACKEND_HIP
24+
#include <CL/sycl/detail/backend_traits_hip.hpp>
25+
#endif
2326
#if SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO
2427
#include <CL/sycl/detail/backend_traits_level_zero.hpp>
2528
#endif

‎sycl/include/CL/sycl/detail/backend_traits_hip.hpp

Lines changed: 136 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,136 @@
1+
//===------- backend_traits_hip.hpp - Backend traits for HIP ---*-C++ -*-===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
//
9+
// This file defines the specializations of the sycl::detail::interop,
10+
// sycl::detail::BackendInput and sycl::detail::BackendReturn class templates
11+
// for the HIP backend but there is no sycl::detail::InteropFeatureSupportMap
12+
// specialization for the HIP backend.
13+
//===----------------------------------------------------------------------===//
14+
15+
#pragma once
16+
17+
#include <CL/sycl/accessor.hpp>
18+
#include <CL/sycl/context.hpp>
19+
#include <CL/sycl/detail/backend_traits.hpp>
20+
#include <CL/sycl/device.hpp>
21+
#include <CL/sycl/event.hpp>
22+
#include <CL/sycl/kernel_bundle.hpp>
23+
#include <CL/sycl/queue.hpp>
24+
25+
typedef int HIPdevice;
26+
typedef struct ihipCtx_t *HIPcontext;
27+
typedef struct ihipStream_t *HIPstream;
28+
typedef struct ihipEvent_t *HIPevent;
29+
typedef struct ihipModule_t *HIPmodule;
30+
typedef void *HIPdeviceptr;
31+
32+
__SYCL_INLINE_NAMESPACE(cl) {
33+
namespace sycl {
34+
namespace detail {
35+
36+
// TODO the interops for context, device, event, platform and program
37+
// may be removed after removing the deprecated 'get_native()' methods
38+
// from the corresponding classes. The interop<backend, queue> specialization
39+
// is also used in the get_queue() method of the deprecated class
40+
// interop_handler and also can be removed after API cleanup.
41+
template <> struct interop<backend::ext_oneapi_hip, context> {
42+
using type = HIPcontext;
43+
};
44+
45+
template <> struct interop<backend::ext_oneapi_hip, device> {
46+
using type = HIPdevice;
47+
};
48+
49+
template <> struct interop<backend::ext_oneapi_hip, event> {
50+
using type = HIPevent;
51+
};
52+
53+
template <> struct interop<backend::ext_oneapi_hip, queue> {
54+
using type = HIPstream;
55+
};
56+
57+
#ifdef __SYCL_INTERNAL_API
58+
template <> struct interop<backend::ext_oneapi_hip, program> {
59+
using type = HIPmodule;
60+
};
61+
#endif
62+
63+
// TODO the interops for accessor is used in the already deprecated class
64+
// interop_handler and can be removed after API cleanup.
65+
template <typename DataT, int Dimensions, access::mode AccessMode>
66+
struct interop<backend::ext_oneapi_hip,
67+
accessor<DataT, Dimensions, AccessMode, access::target::device,
68+
access::placeholder::false_t>> {
69+
using type = HIPdeviceptr;
70+
};
71+
72+
template <typename DataT, int Dimensions, access::mode AccessMode>
73+
struct interop<
74+
backend::ext_oneapi_hip,
75+
accessor<DataT, Dimensions, AccessMode, access::target::constant_buffer,
76+
access::placeholder::false_t>> {
77+
using type = HIPdeviceptr;
78+
};
79+
80+
template <typename DataT, int Dimensions, typename AllocatorT>
81+
struct BackendInput<backend::ext_oneapi_hip,
82+
buffer<DataT, Dimensions, AllocatorT>> {
83+
using type = HIPdeviceptr;
84+
};
85+
86+
template <typename DataT, int Dimensions, typename AllocatorT>
87+
struct BackendReturn<backend::ext_oneapi_hip,
88+
buffer<DataT, Dimensions, AllocatorT>> {
89+
using type = HIPdeviceptr;
90+
};
91+
92+
template <> struct BackendInput<backend::ext_oneapi_hip, context> {
93+
using type = HIPcontext;
94+
};
95+
96+
template <> struct BackendReturn<backend::ext_oneapi_hip, context> {
97+
using type = HIPcontext;
98+
};
99+
100+
template <> struct BackendInput<backend::ext_oneapi_hip, device> {
101+
using type = HIPdevice;
102+
};
103+
104+
template <> struct BackendReturn<backend::ext_oneapi_hip, device> {
105+
using type = HIPdevice;
106+
};
107+
108+
template <> struct BackendInput<backend::ext_oneapi_hip, event> {
109+
using type = HIPevent;
110+
};
111+
112+
template <> struct BackendReturn<backend::ext_oneapi_hip, event> {
113+
using type = HIPevent;
114+
};
115+
116+
template <> struct BackendInput<backend::ext_oneapi_hip, queue> {
117+
using type = HIPstream;
118+
};
119+
120+
template <> struct BackendReturn<backend::ext_oneapi_hip, queue> {
121+
using type = HIPstream;
122+
};
123+
124+
#ifdef __SYCL_INTERNAL_API
125+
template <> struct BackendInput<backend::ext_oneapi_hip, program> {
126+
using type = HIPmodule;
127+
};
128+
129+
template <> struct BackendReturn<backend::ext_oneapi_hip, program> {
130+
using type = HIPmodule;
131+
};
132+
#endif
133+
134+
} // namespace detail
135+
} // namespace sycl
136+
} // __SYCL_INLINE_NAMESPACE(cl)

‎sycl/plugins/hip/pi_hip.cpp

Lines changed: 7 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2101,6 +2101,12 @@ pi_result hip_piMemGetInfo(pi_mem memObj, pi_mem_info queriedInfo,
21012101
/// \param[out] nativeHandle Set to the native handle of the PI mem object.
21022102
///
21032103
/// \return PI_SUCCESS
2104+
pi_result hip_piextMemGetNativeHandle(pi_mem mem,
2105+
pi_native_handle *nativeHandle) {
2106+
*nativeHandle =
2107+
reinterpret_cast<pi_native_handle>(mem->mem_.buffer_mem_.get());
2108+
return PI_SUCCESS;
2109+
}
21042110

21052111
/// Created a PI mem object from a HIP mem handle.
21062112
/// TODO: Implement this.
@@ -4910,7 +4916,7 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
49104916
_PI_CL(piMemRetain, hip_piMemRetain)
49114917
_PI_CL(piMemRelease, hip_piMemRelease)
49124918
_PI_CL(piMemBufferPartition, hip_piMemBufferPartition)
4913-
//_PI_CL(piextMemGetNativeHandle, hip_piextMemGetNativeHandle)
4919+
_PI_CL(piextMemGetNativeHandle, hip_piextMemGetNativeHandle)
49144920
_PI_CL(piextMemCreateWithNativeHandle, hip_piextMemCreateWithNativeHandle)
49154921
// Program
49164922
_PI_CL(piProgramCreate, hip_piProgramCreate)

‎sycl/test/basic_tests/interop-backend-traits.cpp

Lines changed: 9 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,7 @@
11
// RUN: %clangxx -fsycl -DUSE_OPENCL %s
22
// RUN: %clangxx -fsycl -DUSE_L0 %s
33
// RUN: %clangxx -fsycl -DUSE_CUDA %s
4+
// RUN: %clangxx -fsycl -DUSE_HIP %s
45

56
#ifdef USE_OPENCL
67
#include <CL/cl.h>
@@ -24,6 +25,12 @@ constexpr auto Backend = sycl::backend::ext_oneapi_level_zero;
2425
constexpr auto Backend = sycl::backend::ext_oneapi_cuda;
2526
#endif
2627

28+
#ifdef USE_HIP
29+
#include <CL/sycl/detail/backend_traits_hip.hpp>
30+
31+
constexpr auto Backend = sycl::backend::ext_oneapi_hip;
32+
#endif
33+
2734
#include <sycl/sycl.hpp>
2835

2936
int main() {
@@ -43,7 +50,8 @@ int main() {
4350
#endif
4451

4552
// CUDA does not have a native type for platforms
46-
#ifndef USE_CUDA
53+
// HIP also should follow the same behavior - need confirmation
54+
#if !(defined(USE_CUDA) || defined(USE_HIP))
4755
static_assert(
4856
std::is_same_v<sycl::backend_traits<Backend>::input_type<sycl::platform>,
4957
sycl::detail::interop<Backend, sycl::platform>::type>);

0 commit comments

Comments
 (0)