[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.
Este cometimento está contido em:
cometido por
GitHub
ascendente
5eaa152010
cometimento
cf6a53e81c
@@ -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;
|
||||
|
||||
}
|
||||
|
||||
|
||||
Criar uma nova questão referindo esta
Bloquear um utilizador