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

[GDA] Implement fetching atomics for BNXT (#253)

Browse Source 
* Indent driver script
* Implemented fetching atomics BNXT
This commit is contained in:
Yiltan
2025-09-18 09:50:42 -04:00
committed by GitHub
parent 758c4a43f6
commit f5aefd15f3
3 changed files with 50 additions and 38 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/gda/bnxt/queue_pair_bnxt.cpp
+13 -2
View File
@@ -336,6 +336,7 @@ __device__ uint64_t QueuePair::post_wqe_amo(int pe, int32_t length, uintptr_t *r
uint64_t active_lane_mask;
uint8_t active_lane_count;
uint8_t active_lane_id;
uint32_t atomic_idx = 0;
active_lane_mask = get_active_lane_mask();
active_lane_count = get_active_lane_count(active_lane_mask);
@@ -378,8 +379,14 @@ __device__ uint64_t QueuePair::post_wqe_amo(int pe, int32_t length, uintptr_t *r
amo.swp_dt = atomic_data;
/* Populate SG Segment - (Return address of atomic) */
sge.pa = (uint64_t) nonfetching_atomic;
sge.lkey = nonfetching_atomic_lkey;
if (fetching) {
atomic_idx = fetching_atomic_idx++ % FETCHING_ATOMIC_CNT;
sge.pa = (uint64_t) &fetching_atomic[atomic_idx];
sge.lkey = fetching_atomic_lkey;
} else {
sge.pa = (uint64_t) nonfetching_atomic;
sge.lkey = nonfetching_atomic_lkey;
}
sge.length = length;
/* Write WQE to SQ */
@@ -406,6 +413,10 @@ __device__ uint64_t QueuePair::post_wqe_amo(int pe, int32_t length, uintptr_t *r
release_lock(&sq.lock);
}
if (fetching) {
return fetching_atomic[atomic_idx];
}
return 0;
}
src/gda/queue_pair.hpp
+3 -2
View File
@@ -311,9 +311,12 @@ class QueuePair {
uint64_t* nonfetching_atomic{nullptr};
uint32_t nonfetching_atomic_lkey{0};
struct ibv_mr *mr_nonfetching_atomic;
uint64_t* fetching_atomic{nullptr};
uint32_t fetching_atomic_lkey{0};
uint32_t fetching_atomic_idx{0};
struct ibv_mr *mr_fetching_atomic;
static const uint32_t FETCHING_ATOMIC_CNT{1024};
static_assert(FETCHING_ATOMIC_CNT % WF_SIZE == 0);
@@ -322,8 +325,6 @@ class QueuePair {
HIPAllocator allocator{};
struct ibv_mr *mr_nonfetching_atomic;
struct ibv_mr *mr_fetching_atomic;
};
} // namespace rocshmem