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

added notifier->sync() when we are not using cooperative groups

Browse Source 
updated scope bug
This commit is contained in:
Yiltan Hassan Temucin
2024-10-07 06:06:42 -07:00
parent 63667a3167
commit b1134e8633
2 changed files with 18 additions and 9 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/ipc/context_ipc_device_coll.cpp
+6 -9
View File
@@ -26,11 +26,6 @@
#include "../util.hpp"
#include "ipc_team.hpp"
#ifdef USE_COOPERATIVE_GROUPS
#include <hip/hip_cooperative_groups.h>
namespace cg = cooperative_groups;
#endif /* USE_COOPERATIVE_GROUPS */
namespace rocshmem {
__device__ void IPCContext::internal_direct_barrier(int pe, int PE_start,
@@ -92,11 +87,13 @@ __device__ void IPCContext::internal_sync(int pe, int PE_start, int stride,
#ifdef USE_COOPERATIVE_GROUPS
cg::grid_group grid = cg::this_grid();
grid.sync();
if (0 == grid.thread_rank()) {
if (0 == grid.thread_rank())
#else
__syncthreads();
if (is_thread_zero_in_block()) {
auto notifier{notifier_.get()};
notifier->sync();
if (is_thread_zero_in_block() && is_block_zero_in_grid())
#endif /* USE_COOPERATIVE_GROUPS */
{
if (PE_size < 64) {
internal_direct_barrier(pe, PE_start, stride, PE_size, pSync);
} else {
@@ -107,7 +104,7 @@ __device__ void IPCContext::internal_sync(int pe, int PE_start, int stride,
#ifdef USE_COOPERATIVE_GROUPS
grid.sync();
#else
__syncthreads();
notifier->sync();
#endif /* USE_COOPERATIVE_GROUPS */
}