From 82d296db73d8fcb28b38da9fc0fcd685e2250e9b Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Tue, 22 Oct 2024 16:15:26 +0000 Subject: [PATCH] Fix quiet and fence of default context * Update tinfo of default context [ROCm/rocshmem commit: d5ea5868e32b2ab4c6f3f3ffd867e8416ee6bb10] --- projects/rocshmem/src/ipc/ipc_context_proxy.hpp | 1 + projects/rocshmem/src/roc_shmem_gpu.cpp | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/projects/rocshmem/src/ipc/ipc_context_proxy.hpp b/projects/rocshmem/src/ipc/ipc_context_proxy.hpp index 87ca4371e0..29261b62c6 100644 --- a/projects/rocshmem/src/ipc/ipc_context_proxy.hpp +++ b/projects/rocshmem/src/ipc/ipc_context_proxy.hpp @@ -45,6 +45,7 @@ class IPCDefaultContextProxy { : constructed_{true} { auto ctx{proxy_.get()}; new (ctx) IPCContext(reinterpret_cast(backend)); + ctx->tinfo = tinfo; roc_shmem_ctx_t local{ctx, tinfo}; set_internal_ctx(&local); } diff --git a/projects/rocshmem/src/roc_shmem_gpu.cpp b/projects/rocshmem/src/roc_shmem_gpu.cpp index 4ff19b20a2..582249eec9 100644 --- a/projects/rocshmem/src/roc_shmem_gpu.cpp +++ b/projects/rocshmem/src/roc_shmem_gpu.cpp @@ -263,7 +263,7 @@ __device__ int roc_shmem_wg_ctx_create(long option, roc_shmem_ctx_t *ctx) { bool result{true}; if (get_flat_block_id() == 0) { ctx->team_opaque = reinterpret_cast(ROC_SHMEM_CTX_DEFAULT.team_opaque); - device_backend_proxy->create_ctx(option, ctx); + result = device_backend_proxy->create_ctx(option, ctx); reinterpret_cast(ctx->ctx_opaque)->setFence(option); } __syncthreads();