[BugFix] Fix rocshmem_get_device_ctx to return ctx_opaque pointer (#359)

Changed rocshmem_get_device_ctx() to properly copy the full rocshmem_ctx_t
structure and return only the ctx_opaque pointer instead of trying to
copy directly to a void pointer.

Prior implementation would cause undefined behavior or memory corruption
as it was copying 16 bytes of data to 8 bytes. It worked so far beucase ctx_opaque
field is at proper offsest, but incorrectly memcpy would overwrite some other allocations
and cause issues.
This fixes the context memory handling when passing device context from
host to device kernels.

[ROCm/rocshmem commit: cf6a53e81c]
Tento commit je obsažen v:
Dimple Prajapati
2025-12-19 07:01:02 -08:00
odevzdal GitHub
rodič dde4902844
revize e21c087f2a
+2 -2
Zobrazit soubor
@@ -119,11 +119,11 @@ __device__ void rocshmem_wg_finalize() {}
******************************************************************************/
__host__ void * rocshmem_get_device_ctx() {
void *ctx = nullptr;
rocshmem_ctx_t ctx;
CHECK_HIP(hipMemcpyFromSymbol(&ctx, HIP_SYMBOL(ROCSHMEM_CTX_DEFAULT),
sizeof(rocshmem_ctx_t)));
return ctx;
return ctx.ctx_opaque;
}