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

Update Barrier and Sync APIs (#73)

Browse Source 
* Add thread, wavefront, and workgroup-level `barrier` APIs in IPC and RO conduits; remove collectives on default context
 - Implemented `barrier` APIs for thread, wavefront, and workgroup scopes
 - Added support into both IPC and RO conduits
 - Added functional tests to cover all `barrier` APIs
 - Removed collective operations on default context

* Add thread, wavefront, and workgroup-level `sync` APIs in IPC and RO conduits.
  - Implemented `sync` APIs for thread, wavefront, and workgroup scopes
  - Added support into both IPC and RO conduits
  - Added functional tests to cover all `sync` APIs

* update naming convention for context-based `barrier` APIs

[ROCm/rocshmem commit: dc61bca066]
This commit is contained in:
Avinash Kethineedi
2025-04-08 11:25:31 -05:00
committed by GitHub
parent 426bbf525b
commit 9bd2b04899
16 changed files with 347 additions and 67 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocshmem/src/context_device.cpp
+26 -2
View File
@@ -161,11 +161,23 @@ __device__ void Context::barrier_all_wg() {
}
__device__ void Context::barrier(rocshmem_team_t team) {
ctxStats.incStat(NUM_BARRIER_ALL);
ctxStats.incStat(NUM_BARRIER);
DISPATCH(barrier(team));
}
__device__ void Context::barrier_wave(rocshmem_team_t team) {
ctxStats.incStat(NUM_BARRIER_WAVE);
DISPATCH(barrier_wave(team));
}
__device__ void Context::barrier_wg(rocshmem_team_t team) {
ctxStats.incStat(NUM_BARRIER_WG);
DISPATCH(barrier_wg(team));
}
__device__ void Context::sync_all() {
ctxStats.incStat(NUM_SYNC_ALL);
@@ -184,8 +196,20 @@ __device__ void Context::sync_all_wg() {
DISPATCH(sync_all_wg());
}
__device__ void Context::sync(rocshmem_team_t team) {
ctxStats.incStat(NUM_SYNC);
DISPATCH(sync(team));
}
__device__ void Context::sync_wave(rocshmem_team_t team) {
ctxStats.incStat(NUM_SYNC_WAVE);
DISPATCH(sync_wave(team));
}
__device__ void Context::sync_wg(rocshmem_team_t team) {
ctxStats.incStat(NUM_SYNC_ALL_WG);
ctxStats.incStat(NUM_SYNC_WG);
DISPATCH(sync_wg(team));
}