From 1aa9679fa768e1eb6ff367981ebc7ab6a5f0523a Mon Sep 17 00:00:00 2001 From: Dimple Prajapati Date: Mon, 25 Aug 2025 11:11:38 -0700 Subject: [PATCH] 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: 2de8f4579f0160a52866f12351f31170994c64d2] --- projects/rocshmem/src/rocshmem_gpu.cpp | 13 +++++++++++-- 1 file changed, 11 insertions(+), 2 deletions(-) diff --git a/projects/rocshmem/src/rocshmem_gpu.cpp b/projects/rocshmem/src/rocshmem_gpu.cpp index d612df4abb..2e3fc0821d 100644 --- a/projects/rocshmem/src/rocshmem_gpu.cpp +++ b/projects/rocshmem/src/rocshmem_gpu.cpp @@ -54,6 +54,9 @@ #ifdef USE_RO #include "reverse_offload/context_ro_tmpl_device.hpp" #else +#ifdef ENABLE_IPC_BITCODE + #include "ipc/backend_ipc.hpp" +#endif #include "ipc/context_ipc_tmpl_device.hpp" #endif @@ -67,6 +70,12 @@ __device__ rocshmem_ctx_t __attribute__((visibility("default"))) ROCSHMEM_CTX_D __constant__ Backend *device_backend_proxy; +#ifdef ENABLE_IPC_BITCODE + typedef IPCContext ContextTy; +#else + typedef Context ContextTy; +#endif + __device__ void rocshmem_wg_init() { int provided; @@ -294,8 +303,8 @@ __host__ void set_internal_ctx(rocshmem_ctx_t *ctx) { hipMemcpyHostToDevice)); } -__device__ Context *get_internal_ctx(rocshmem_ctx_t ctx) { - return reinterpret_cast(ctx.ctx_opaque); +__device__ ContextTy *get_internal_ctx(rocshmem_ctx_t ctx) { + return reinterpret_cast(ctx.ctx_opaque); } __device__ int rocshmem_wg_ctx_create(long options, rocshmem_ctx_t *ctx) {