From fd3eaa3f69543421dc488735cb7f0dce455c2b88 Mon Sep 17 00:00:00 2001 From: Yiltan Hassan Temucin Date: Wed, 29 Jan 2025 09:56:22 -0800 Subject: [PATCH] [IPC] Fix ROCSHMEM_SIGNAL_ADD --- src/ipc/context_ipc_device.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/src/ipc/context_ipc_device.cpp b/src/ipc/context_ipc_device.cpp index d214c663b8..de3386f581 100644 --- a/src/ipc/context_ipc_device.cpp +++ b/src/ipc/context_ipc_device.cpp @@ -233,7 +233,7 @@ __device__ void IPCContext::putmem_signal(void *dest, const void *source, size_t amo_set(static_cast(sig_addr), signal, pe); break; case ROCSHMEM_SIGNAL_ADD: - amo_add(static_cast(dest), signal, pe); + amo_add(static_cast(sig_addr), signal, pe); break; default: DPRINTF("[%s] Invalid sig_op value (%d)\n", __func__, sig_op); @@ -253,7 +253,7 @@ __device__ void IPCContext::putmem_signal_wg(void *dest, const void *source, siz amo_set(static_cast(sig_addr), signal, pe); break; case ROCSHMEM_SIGNAL_ADD: - amo_add(static_cast(dest), signal, pe); + amo_add(static_cast(sig_addr), signal, pe); break; default: DPRINTF("[%s] Invalid sig_op value (%d)\n", __func__, sig_op); @@ -274,7 +274,7 @@ __device__ void IPCContext::putmem_signal_wave(void *dest, const void *source, s amo_set(static_cast(sig_addr), signal, pe); break; case ROCSHMEM_SIGNAL_ADD: - amo_add(static_cast(dest), signal, pe); + amo_add(static_cast(sig_addr), signal, pe); break; default: DPRINTF("[%s] Invalid sig_op value (%d)\n", __func__, sig_op); @@ -298,7 +298,7 @@ __device__ void IPCContext::putmem_signal_nbi_wg(void *dest, const void *source, __device__ void IPCContext::putmem_signal_nbi_wave(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) { - putmem_signal_wave(dest, source, nelems, sig_addr, signal, sig_op, pe); + putmem_signal_wave(dest, source, nelems, sig_addr, signal, sig_op, pe); } __device__ uint64_t IPCContext::signal_fetch(const uint64_t *sig_addr) {