Skip to content

Commit 1aa9679

Browse files
authored
Device bitcode support for IPC - Use IPCContext explicitly to link correct device API implementation (#217)
* Device bitcode support for IPC - Use IPCContext explicitly * use base Context class for non-derived functions [ROCm/rocshmem commit: 2de8f45]
1 parent 7436abe commit 1aa9679

File tree

1 file changed

+11
-2
lines changed

1 file changed

+11
-2
lines changed

projects/rocshmem/src/rocshmem_gpu.cpp

Lines changed: 11 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -54,6 +54,9 @@
5454
#ifdef USE_RO
5555
#include "reverse_offload/context_ro_tmpl_device.hpp"
5656
#else
57+
#ifdef ENABLE_IPC_BITCODE
58+
#include "ipc/backend_ipc.hpp"
59+
#endif
5760
#include "ipc/context_ipc_tmpl_device.hpp"
5861
#endif
5962

@@ -67,6 +70,12 @@ __device__ rocshmem_ctx_t __attribute__((visibility("default"))) ROCSHMEM_CTX_D
6770

6871
__constant__ Backend *device_backend_proxy;
6972

73+
#ifdef ENABLE_IPC_BITCODE
74+
typedef IPCContext ContextTy;
75+
#else
76+
typedef Context ContextTy;
77+
#endif
78+
7079
__device__ void rocshmem_wg_init() {
7180
int provided;
7281

@@ -294,8 +303,8 @@ __host__ void set_internal_ctx(rocshmem_ctx_t *ctx) {
294303
hipMemcpyHostToDevice));
295304
}
296305

297-
__device__ Context *get_internal_ctx(rocshmem_ctx_t ctx) {
298-
return reinterpret_cast<Context *>(ctx.ctx_opaque);
306+
__device__ ContextTy *get_internal_ctx(rocshmem_ctx_t ctx) {
307+
return reinterpret_cast<ContextTy *>(ctx.ctx_opaque);
299308
}
300309

301310
__device__ int rocshmem_wg_ctx_create(long options, rocshmem_ctx_t *ctx) {

0 commit comments

Comments
 (0)