From ebcfdb2f75c6e94c0dcc114c9c2de33aedcf8b52 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Tue, 30 Jul 2024 12:52:06 -0500 Subject: [PATCH 1/2] Remove `dev_mtx_` that is no longer in the backend --- src/gpu_ib/backend_ib.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/src/gpu_ib/backend_ib.cpp b/src/gpu_ib/backend_ib.cpp index 9ae3acd482..1357aea78d 100644 --- a/src/gpu_ib/backend_ib.cpp +++ b/src/gpu_ib/backend_ib.cpp @@ -133,8 +133,6 @@ __device__ bool GPUIBBackend::create_ctx(int64_t options, return false; } ctx_ = pop_result.value; - - ctx_->dev_mtx_.shareable_ = static_cast(options >> 3); ctx->ctx_opaque = ctx_; return true; } From 3162d49b5610be4604f591648d7ac4c66d9dcdab Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Tue, 17 Sep 2024 20:34:18 -0500 Subject: [PATCH 2/2] Vectorize WQe segments writes Signed-off-by: Muhammad Awad --- src/gpu_ib/segment_builder.cpp | 52 ++++++++++++++++++++-------------- 1 file changed, 30 insertions(+), 22 deletions(-) diff --git a/src/gpu_ib/segment_builder.cpp b/src/gpu_ib/segment_builder.cpp index 205c890861..6f2a77a23e 100644 --- a/src/gpu_ib/segment_builder.cpp +++ b/src/gpu_ib/segment_builder.cpp @@ -22,6 +22,7 @@ #include "segment_builder.hpp" +#include "../util.hpp" #include "endian.hpp" namespace rocshmem { @@ -35,9 +36,9 @@ __device__ SegmentBuilder::SegmentBuilder(uint64_t wqe_idx, void *base) { __device__ void SegmentBuilder::update_cntrl_seg( uint8_t opcode, uint16_t wqe_idx, uint32_t ctrl_qp_sq, uint64_t ctrl_sig, ConnectionImpl *connection_policy, bool zero_byte_rd) { - mlx5_wqe_ctrl_seg *ctrl_seg = &seg_ptr->ctrl_seg; + mlx5_wqe_ctrl_seg ctrl_seg; - ctrl_seg->opmod_idx_opcode = (opcode << 24) | (wqe_idx << 8); + ctrl_seg.opmod_idx_opcode = (opcode << 24) | (wqe_idx << 8); uint32_t DS = 2; if (zero_byte_rd == false) { @@ -48,39 +49,43 @@ __device__ void SegmentBuilder::update_cntrl_seg( DS += connection_policy->wqeCntrlOffset(); - ctrl_seg->qpn_ds = (DS << 24) | ctrl_qp_sq; + ctrl_seg.qpn_ds = (DS << 24) | ctrl_qp_sq; - ctrl_seg->signature = ctrl_sig; + ctrl_seg.signature = ctrl_sig; - ctrl_seg->fm_ce_se = ctrl_sig >> 24; + ctrl_seg.fm_ce_se = ctrl_sig >> 24; - ctrl_seg->imm = ctrl_sig >> 32; + ctrl_seg.imm = ctrl_sig >> 32; + + memcpy(&seg_ptr->ctrl_seg, &ctrl_seg, sizeof(mlx5_wqe_ctrl_seg)); seg_ptr++; } __device__ void SegmentBuilder::update_atomic_data_seg(uint64_t atomic_data, uint64_t atomic_cmp) { - mlx5_wqe_atomic_seg *atomic_seg = &seg_ptr->atomic_seg; + mlx5_wqe_atomic_seg atomic_seg; - swap_endian_store(reinterpret_cast(&atomic_seg->swap_add), + swap_endian_store(reinterpret_cast(&atomic_seg.swap_add), atomic_data); - swap_endian_store(reinterpret_cast(&atomic_seg->compare), + swap_endian_store(reinterpret_cast(&atomic_seg.compare), atomic_cmp); + memcpy(&seg_ptr->atomic_seg, &atomic_seg, sizeof(mlx5_wqe_atomic_seg)); seg_ptr++; } __device__ void SegmentBuilder::update_rdma_seg(uintptr_t *raddr, uint32_t rkey) { - mlx5_wqe_raddr_seg *raddr_seg = &seg_ptr->raddr_seg; + mlx5_wqe_raddr_seg raddr_seg; - raddr_seg->rkey = rkey; + raddr_seg.rkey = rkey; - swap_endian_store(reinterpret_cast(&raddr_seg->raddr), + swap_endian_store(reinterpret_cast(&raddr_seg.raddr), reinterpret_cast(raddr)); + memcpy(&seg_ptr->raddr_seg, &raddr_seg, sizeof(mlx5_wqe_raddr_seg)); seg_ptr++; } @@ -90,33 +95,36 @@ __device__ void SegmentBuilder::update_data_seg(uintptr_t *laddr, int32_t size, return; } - mlx5_wqe_data_seg *data_seg = &seg_ptr->data_seg; + mlx5_wqe_data_seg data_seg; + data_seg.lkey = lkey; - data_seg->lkey = lkey; - - swap_endian_store(&data_seg->byte_count, size & 0x7FFFFFFFU); - - swap_endian_store(reinterpret_cast(&data_seg->addr), + swap_endian_store(&data_seg.byte_count, size & 0x7FFFFFFFU); + swap_endian_store(reinterpret_cast(&data_seg.addr), reinterpret_cast(laddr)); + memcpy(&seg_ptr->data_seg, &data_seg, sizeof(mlx5_wqe_data_seg)); seg_ptr++; } __device__ void SegmentBuilder::update_inl_data_seg(uintptr_t *laddr, int32_t size) { - mlx5_wqe_inl_data_seg *inl_data_seg = &seg_ptr->inl_data_seg; + mlx5_wqe_inl_data_seg inl_data_seg; - swap_endian_store(&inl_data_seg->byte_count, (size & 0x3FF) | 0x80000000); + swap_endian_store(&inl_data_seg.byte_count, (size & 0x3FF) | 0x80000000); // Assume fence HDP flush // TODO(khamidou): Rework fence interface to avoid this + size_t field_size{sizeof(mlx5_wqe_inl_data_seg)}; if (!laddr) { uint8_t flush_val = 1; - memcpy(inl_data_seg + 1, &flush_val, sizeof(flush_val)); + memcpy(&inl_data_seg + 1, &flush_val, sizeof(flush_val)); + field_size += sizeof(flush_val); } else { - memcpy(inl_data_seg + 1, laddr, size); + memcpy(&inl_data_seg + 1, laddr, size); + field_size += size; } + memcpy(&seg_ptr->inl_data_seg, &inl_data_seg, field_size); seg_ptr++; }