AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge pull request #32 from Yiltan/ipc/sigop-bug

Browse Source 
[IPC] Fix ROCSHMEM_SIGNAL_ADD
This commit is contained in:
Yiltan
2025-02-03 16:48:05 -05:00
committed by GitHub
parent 7e5b533904 fd3eaa3f69
commit 2d9d09ea01
1 changed files with 4 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/ipc/context_ipc_device.cpp
+4 -4
View File
@@ -233,7 +233,7 @@ __device__ void IPCContext::putmem_signal(void *dest, const void *source, size_t
amo_set<uint64_t>(static_cast<void*>(sig_addr), signal, pe); amo_set<uint64_t>(static_cast<void*>(sig_addr), signal, pe);
break; break;
case ROCSHMEM_SIGNAL_ADD: case ROCSHMEM_SIGNAL_ADD:
amo_add<uint64_t>(static_cast<void*>(dest), signal, pe); amo_add<uint64_t>(static_cast<void*>(sig_addr), signal, pe);
break; break;
default: default:
DPRINTF("[%s] Invalid sig_op value (%d)\n", __func__, sig_op); 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<uint64_t>(static_cast<void*>(sig_addr), signal, pe); amo_set<uint64_t>(static_cast<void*>(sig_addr), signal, pe);
break; break;
case ROCSHMEM_SIGNAL_ADD: case ROCSHMEM_SIGNAL_ADD:
amo_add<uint64_t>(static_cast<void*>(dest), signal, pe); amo_add<uint64_t>(static_cast<void*>(sig_addr), signal, pe);
break; break;
default: default:
DPRINTF("[%s] Invalid sig_op value (%d)\n", __func__, sig_op); 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<uint64_t>(static_cast<void*>(sig_addr), signal, pe); amo_set<uint64_t>(static_cast<void*>(sig_addr), signal, pe);
break; break;
case ROCSHMEM_SIGNAL_ADD: case ROCSHMEM_SIGNAL_ADD:
amo_add<uint64_t>(static_cast<void*>(dest), signal, pe); amo_add<uint64_t>(static_cast<void*>(sig_addr), signal, pe);
break; break;
default: default:
DPRINTF("[%s] Invalid sig_op value (%d)\n", __func__, sig_op); 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, __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, uint64_t *sig_addr, uint64_t signal, int sig_op,
int pe) { 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) { __device__ uint64_t IPCContext::signal_fetch(const uint64_t *sig_addr) {