refactor(rocshmem/gda): refactor QueuePair class by omor1 · Pull Request #7217 · ROCm/rocm-systems

omor1 · 2026-06-13T00:10:39Z

Motivation

The QueuePair class [gda/queue_pair.{hpp,cpp}] functions as three classes in a trenchcoat that masquerade as a single class; it simultaneously has functionality for:

common device-side QP initialization and management, including host functions for registering additional local buffers; and
provider-specific send and completion queue management, including WQE posting and CQ polling; and
a multiplexing low-level interface that dispatches to the provider-specific functions based on which is at use at execution time; and
a high-level SHMEM-like communication interface that dispatches to the low-level interface.

Maintaining this complex class and its many interactions is complicated and has led to changes intended for one GDA provider causing unintended impacts on the others, as well making several types of optimizations far more difficult. The goal of this PR is to refactor this single class into several distinct classes that will help improve code readability and maintainability, as well as restoring the ability to compile directly for a specific GDA provider without runtime dispatch.

The end goal is to allow much of the QueuePair and GDA initialization code to be abstracted as a separate library that can be consumed by both rocSHMEM and other users of similar functionality, such as the RCCL-GIN GDA backend.

Technical Details

template <typename Provider> class QueuePairSHMEM
- gda/queue_pair.hpp
- High-level SHMEM-like communication interface
- Uses CRTP to dispatche to low-level functions implemented by Provider
- Contains no non-static data members
template <typename Provider> class QueuePairBase
- gda/queue_pair.hpp
- Common device-side QP initialization and management
- Inherits from QueuePairSHMEM<Provider>
class QueuePairIONIC
- gda/ionic/queue_pair_ionic.{hpp,cpp}
- Implements send and completion queue management for AMD Pensando 'ionic' RDMA NICs
- Inherits from QueuePairBase<QueuePairIONIC>
class QueuePairBNXT
- gda/bnxt/queue_pair_bnxt.{hpp,cpp}
- Implements send and completion queue management for Broadcom 'bnxt' RDMA NICs
- Inherits from QueuePairBase<QueuePairBNXT>
class QueuePairMLX5
- gda/mlx5/queue_pair_mlx5.{hpp,cpp}
- Implements send and completion queue management for NVIDIA/Mellanox 'mlx5' RDMA NICs
- Inherits from QueuePairBase<QueuePairMLX5>
class QueuePairMux
- gda/queue_pair_mux.{hpp,cpp}
- Multiplexing union-like class that implements the low-level interface by dispatching to an underlying GDA provider
- Holds one of QueuePairIONIC, QueuePairBNXT, or QueuePairMLX5
- Inherits from QueuePairSHMEM<QueuePairMux>
alias QueuePair
- gda/queue_pair_provider.hpp
- Alias for one of QueuePairIONIC, QueuePairBNXT, QueuePairMLX5, or QueuePairMux
- Simplifies generic use of any of the Queue Pair classes
- Selected by compile-time configuration

JIRA ID

JIRA ID: AIROCSHMEM-424

Test Plan

Use the rocSHMEM functional tests to validate correctness and performance regression analysis against current develop HEAD.

Test Result

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

omor1 · 2026-06-13T00:19:35Z

TODO:

Rebase on top of develop
Replace QueuePairMux::provider non-static data member with __constant__ memory, see [rocshmem] use constmem for backend and provider muxing #6880
- 795eb8d
Change post_wqe function interface to consume lkey and rkey directly
- c7a4fe6
Improve optimization template parameter interface
- b5c1293
Move QueuePairSHMEM functionality out of queue_pair.hpp and into GDAContext
Add QueuePairHostProxy for host-side management of Queue Pair classes, so that the device-side classes don't have host-only data members
Additional refactor of QueuePairIONIC code:
- Remove support for non-collapsed CQ (untested, dead code)
- Add support for optimization template parameters

omor1 · 2026-06-24T02:52:16Z

GIN will need to use QueuePair::to_provider_endianness<T>(T val) for passing explicit lkey/rkey to post_wqe_rma/amo, it expects them in the Provider's endianness. NCCL GIN GDAKI stores the lkeys and rkeys in BE, doing the byte swap once during memory registration.

therock-pr-bot · 2026-06-30T00:37:13Z

❌ PR Check — Action Required

Check	Status	Details
🌿 Branch Name	✅ Pass	—
📝 PR Title/Description	❌ Fail	Error: Title does not follow Conventional Commits style. Expected: start with a valid type (feat, fix, docs, …). Desired format: `type(optional-scope): short description`
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	❌ Fail	Error: Source/code files changed without an accompanying unit test. Expected: add at least one test file named like `test_<name>.py` / `test_<name>.cpp` (or `<name>_test.`). Current:* code file(s) changed: `projects/rocshmem/src/constmem.cpp`, `projects/rocshmem/src/gda/backend_gda.cpp`, `projects/rocshmem/src/gda/backend_gda.hpp`, `projects/rocshmem/src/gda/bnxt/backend_gda_bnxt.cpp`, `projects/rocshmem/src/gda/bnxt/provider_gda_bnxt.hpp` (+18 more); no test file found
🔎 pre-commit	⏳ Pending	⏳ Still running…
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—

⚠️ 2 policy check(s) failed. Please address the issues above before this PR can be Reviewed.

🚫 Please fix the failed policies

❌ PR Title/Description

❌ Unit Test

The Not ready to Review label was added to this PR. Once all policies pass, the label is removed automatically.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

therock-pr-bot · 2026-06-30T00:37:14Z

🚫 Please fix the failed policies before requesting reviews.

The following policy checks failed:

❌ PR Title/Description
❌ Unit Test

The Not ready to Review label has been added to this PR.
Once all policies pass, the label will be removed automatically.

Copilot

Pull request overview

Refactors rocSHMEM’s GDA QueuePair implementation by splitting the prior monolithic class into provider-specific queue-pair classes (IONIC/BNXT/MLX5) plus an optional runtime-dispatching mux, enabling cleaner separation of responsibilities and compile-time provider selection.

Changes:

Introduces provider-specific QueuePair* implementations and a QueuePairMux runtime-dispatch wrapper, with QueuePair selected via a new queue_pair_provider.hpp alias.
Updates GDA backend initialization and device-side context code to use the new QueuePair API (including new posting options usage).
Updates build system and bitcode compilation to support GDA_MUX and provider-specific compilation/selection.

Reviewed changes

Copilot reviewed 30 out of 30 changed files in this pull request and generated 8 comments.

Show a summary per file

File	Description
projects/rocshmem/src/gda/queue_pair.cpp	Removes the legacy monolithic `QueuePair` implementation (replaced by refactored classes).
projects/rocshmem/src/gda/queue_pair_provider.hpp	Adds `QueuePair` alias selection for provider/mux at compile time.
projects/rocshmem/src/gda/queue_pair_mux.hpp	Adds mux `QueuePair` type and device-side dispatch helpers.
projects/rocshmem/src/gda/queue_pair_mux.cpp	Implements mux union lifecycle and host-side dispatch helpers.
projects/rocshmem/src/gda/mlx5/queue_pair_mlx5.hpp	Adds MLX5 provider-specific queue pair class interface + templated posting.
projects/rocshmem/src/gda/mlx5/queue_pair_mlx5.cpp	Moves MLX5 queue-pair implementation out of legacy code into provider class.
projects/rocshmem/src/gda/mlx5/provider_gda_mlx5.hpp	Refactors MLX5 inline decision helper to a template form.
projects/rocshmem/src/gda/mlx5/CMakeLists.txt	Adjusts MLX5 sources inclusion based on provider build option.
projects/rocshmem/src/gda/mlx5/backend_gda_mlx5.cpp	Updates MLX5 GPU QP initialization to construct new provider QP object.
projects/rocshmem/src/gda/ionic/queue_pair_ionic.hpp	Adds IONIC provider-specific queue pair class interface + templated posting.
projects/rocshmem/src/gda/ionic/queue_pair_ionic.cpp	Moves IONIC queue-pair implementation out of legacy code into provider class.
projects/rocshmem/src/gda/ionic/provider_gda_ionic.hpp	Introduces device-queue structs and namespaces for IONIC provider internals.
projects/rocshmem/src/gda/ionic/CMakeLists.txt	Adjusts IONIC sources inclusion based on provider build option.
projects/rocshmem/src/gda/ionic/backend_gda_ionic.cpp	Updates IONIC GPU QP initialization to construct new provider QP object.
projects/rocshmem/src/gda/bnxt/queue_pair_bnxt.hpp	Adds BNXT provider-specific queue pair class interface + templated posting.
projects/rocshmem/src/gda/bnxt/queue_pair_bnxt.cpp	Moves BNXT queue-pair implementation out of legacy code into provider class.
projects/rocshmem/src/gda/bnxt/provider_gda_bnxt.hpp	Adds missing includes and constructors for BNXT device queue structs.
projects/rocshmem/src/gda/bnxt/CMakeLists.txt	Adjusts BNXT sources inclusion based on provider build option.
projects/rocshmem/src/gda/bnxt/backend_gda_bnxt.cpp	Updates BNXT GPU QP initialization to construct new provider QP object.
projects/rocshmem/src/gda/gda_enums.hpp	Changes provider enum to `enum class` and keeps `GDAProvider` alias.
projects/rocshmem/src/gda/debug_gda.hpp	Minor annotation update (`[[maybe_unused]]`).
projects/rocshmem/src/gda/context_gda_tmpl_device.hpp	Updates device-side calls to new QP AMO/RMA APIs and posting options.
projects/rocshmem/src/gda/context_gda_device.hpp	Switches include to `queue_pair_provider.hpp`.
projects/rocshmem/src/gda/backend_gda.hpp	Updates includes and provider-guarded declarations for backend init paths.
projects/rocshmem/src/gda/backend_gda.cpp	Updates buffer registration paths and GPU QP setup/teardown for refactor.
projects/rocshmem/src/gda/CMakeLists.txt	Removes legacy queue_pair.cpp and conditionally adds mux sources.
projects/rocshmem/CMakeLists.txt	Adds `GDA_MUX` option, provider-count checks, and backend string updates.
projects/rocshmem/cmake/rocshmem_config.h.in	Adds `GDA_MUX` config define.
projects/rocshmem/cmake/DeviceBitcode.cmake	Updates bitcode sources to compile mux and provider-specific QP sources.

Comments suppressed due to low confidence (1)

projects/rocshmem/src/gda/backend_gda.cpp:587

Host-side buffer_register/buffer_unregister calls QueuePair methods through gpu_qps, but gpu_qps is allocated with hipMalloc (device memory). Under standard HIP semantics the host cannot dereference/call methods on device allocations, so registration/destruction will be invalid (and was previously done via a host mirror + hipMemcpy). Consider restoring a host-resident QP array for host-only operations (MR map + buffer_info updates) and explicitly copying the device-visible portions to gpu_qps, or switching gpu_qps to a host-accessible allocation (e.g., hipMallocManaged/hipHostMalloc with the required coherency guarantees).

  /* Register with QPs */
  for (size_t i = 0; i < num_qps; i++) {
    err = gpu_qps[i].buffer_register(addr, length);
    if (ROCSHMEM_SUCCESS != err) {
      qp_registration_failed = true;
    }
  }

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

abouteiller

Ongoing review

abouteiller · 2026-07-06T16:03:50Z


-__device__ static inline void bnxt_re_fill_psns_for_msntbl(struct bnxt_device_sq *sq,
-                                                           uint32_t msg_len) {
+__device__ void QueuePairBNXT::fill_psns_for_msntbl(bnxt_device_sq& sq, uint32_t msg_len) {


this was static inline, and is not anymore?

fill_psns_for_msntbl() is called from within template <OpCode Op, bool CheckSQ> write_rma_wqe() (which must be defined in the header since it's a template), so one of the following must be true:

It is declared as a static inline free function and defined in the header, which also "pollutes" the namespace of anything that includes the header.

It is declared as an inline static member function in QueuePairBNXT and defined in the header; it is declared as a private member, so not accessible by anything outside the class, but the definition is still provided in the header, which might make it more difficult to read and might mean we need to add a [[maybe_unused]] annotation to quiet warnings.

It is declared as a regular (not-inline) static member function in QueuePairBNXT and defined in the source code file and the compiler likely inlines it when linking.

I generally chose to only have template definitions in the header and have all other functions defined in the corresponding cpp file. As another option, we could declare it as a non-static member function (whether defined inline or not) and remove the sq argument.

Note for clarification:

In C++, when declaring a static member function, you only use the keyword static in the declaration inside the class body. The out-of-line definition does not use the keyword.

In C++, when a function is declared inline, it has the following properties:

It must have a definition available in all "translation units" (source code files compiled separately) where that function is accessed.

It can have more than one definition (which is usually disallowed, the so-called "One Definition Rule") as long as each translation unit only has a single definition, it is declared inline in those translation units, and all definitions are identical. The linker will merge these definitions such that all definitions have the same address.

A static object defined in an inline function all refer to the same object across all translation units.

Tl;dr an inline function (and variable, since C++17) really just means that it can have multiple definitions (e.g. be defined in a header file) and all those definitions will get merged by the linker. If a member function is defined within the class body, it is implicitly inline.

abouteiller · 2026-07-06T16:05:21Z

 }

-__device__ static inline void unlock(uint32_t *lock) {
+__device__ void QueuePairBNXT::release_lock(uint32_t* lock) {


I think we prefer type *var

missing static inline

I think we prefer type *var

I think C++ generally prefers type* name and type& name (and type&& name), while C generally prefers type *name. I've usually been using the C style for variables coming from C code e.g. from libibverbs, but using the C++ style otherwise.

I have no strong preference either way, but we should be consistent throughout the code; this is part of the larger style issues we need to come to a consensus on. (I do have a slight preference for type& name over type &name for references, but definitely think that we should be consistent across references and pointers, regardless of which style we use.)

missing static inline

See above comment.

abouteiller · 2026-07-06T16:28:55Z

+}
+
+__device__ uint32_t QueuePairMux::get_lkey(uintptr_t addr) {
+  switch (constmem.gda_provider) {


Its unfortunate that we have to duplicate the switch here, but that comes from upleveling the get_lkey, rkey. Are the implementations actually different for each provider though?

They all use the implementation from QueuePairBase<Provider>::get_lkey, but that refers to a member variable that might be located in different offsets in each provider. (They probably aren't, but we can't guarantee that!)

We could do some really ugly hacks to check whether the offsets are the same at compile time and then access it directly at runtime without checking which one is "correct". Hopefully though the compiler is smart enough to be able to merge the duplicate switches and then lift the access out of the conditional since it's invariant.

Refactor monolithic QueuePair class into: * template class QueuePairSHMEM<typename Provider> a CRTP mixin class implementing a high-level SHMEM-like interface * template class QueuePairBase<typename Provider> : public QueuePairSHMEM<Provider> a CRTP base class implementing data members and intialization common to all QueuePair provider implementations * QueuePairMLX5 : public QueuePairBase<QueuePairMLX5> a concrete QueuePair implementation for mlx5 * QueuePairGeneric : public QueuePairSHMEM<QueuePairGeneric> a variant (tagged union) class that dispatches to an underlying concrete implementation at runtime * QueuePairBNXT and QueuePairIONIC implementations will follow shortly * extern "C" ::rocshmem_gin_init_constmem definition has been moved to constmem.cpp Signed-off-by: Omri Mor <Omri.Mor@amd.com>

Signed-off-by: Omri Mor <Omri.Mor@amd.com>

First pass for QueuePairIONIC refactor. Needs some subsequent code cleanup. Signed-off-by: Omri Mor <Omri.Mor@amd.com>

Signed-off-by: Omri Mor <Omri.Mor@amd.com>

GIN GDA needs to pass lkey/rkey explicitly to the QP classes. (Why?) Signed-off-by: Omri Mor <Omri.Mor@amd.com>

Compilers sometimes find it easier to optimize `static constexpr` block-scope variables over `constexpr` block-scope variables. Signed-off-by: Omri Mor <Omri.Mor@amd.com>

Signed-off-by: Omri Mor <Omri.Mor@amd.com>

Posting options and optimization parameters can be passed to QueuePair::post_wqe_rma/amo via an additional function argument of type PostOpt<typename... Option>. For example: qp.post_wqe_rma(laddr, lkey, raddr, rkey, size, PostOpt{RingDB<true>, ThreadSafe<false>, CheckSQ<false>, UpdateCQ<UpdateThread::Last>}); These options can also be passed from the higher-level `SHMEM' QP interface. Signed-off-by: Omri Mor <Omri.Mor@amd.com>

Signed-off-by: Omri Mor <Omri.Mor@amd.com>

Copilot suggests that QueuePairMux::operator=(QueuePairMux&&) can cause undefined behavior when this == &other i.e. self move-assignment. This shouldn't occur in normal usage, but better to be safe and fix this. The underlying QueuePair provider's move assignment operator should handle self move-assignment correctly, so defer to the operator when this->provider == other.provider. Add QueuePairMux::QueuePairUnion::assign(QueuePairUnion&&, GDAProvider) that dispatches to the correct assignment operator depending on which provider is active. Signed-off-by: Omri Mor <Omri.Mor@amd.com>

…::create_cqs path" This reverts commit 0869052.

…atomic_add Align with OpenSHMEM naming. Signed-off-by: Omri Mor <Omri.Mor@amd.com>

…fetch Align with OpenSHMEM naming. Signed-off-by: Omri Mor <Omri.Mor@amd.com>

…ic_compare_swap Align with OpenSHMEM naming. Signed-off-by: Omri Mor <Omri.Mor@amd.com>

… compare) Align with naming conventions used by GDA providers and in InfiniBand spec. Signed-off-by: Omri Mor <Omri.Mor@amd.com>

Signed-off-by: Omri Mor <Omri.Mor@amd.com>

…educe diff Revert order of function defintions in the following files, so as to minimize the differences as compared to develop HEAD: * ionic/queue_pair_ionic.cpp * bnxt/queue_pair_ionic.cpp * mlx5/queue_pair_mlx5.cpp Signed-off-by: Omri Mor <Omri.Mor@amd.com>

QueuePairXYZ::post_wqe* and similar functions require a large number of VGPRs; if these functions are inlined by the compiler, the resulting kernel function can require too many VGPRs and cannot execute correctly in some configurations. Conversely, QueuePairMux simply dispatches to some underlying provider-specific function (that itself should be __noinline__), so the multiplexing parent function *should* be inlined into its caller. Signed-off-by: Omri Mor <Omri.Mor@amd.com>

Mirror other providers by calling the mlx5_devx_qp::dump function from the GDABackend::initialize_gpu_qp parent call instead of within GDABackend::mlx5_initialize_gpu_qp. Signed-off-by: Omri Mor <Omri.Mor@amd.com>

… host_qps) Remove implicit dependency on unified memory in QueuePair creation: * Allocate array of QueuePair on host, GDABackend::host_qps * Construct QueuePair objects in host memory * hipMemcpy(gda_qps, host_qps, sizeof(QueuePair) * num_qps) * Host and Device copies point internally to same locations in device memory Signed-off-by: Omri Mor <Omri.Mor@amd.com>

std::unordered_map uses a hash table rather than the tree used in std::map. Since we know the maximum number of buffers users can register, we can reserve enough space in the hash table for those buffers. std::unordered_map operations are O(1) on average versus O(log N) for std::map. Additionally, remove some more uses of implicit unified memory. Signed-off-by: Omri Mor <Omri.Mor@amd.com>

Instead of directly searching device memory for the next free buffer_info entry, track this on the host using a freelist. The freelist presents an interface similar to std::priority_queue, but rather than using a heap the data structure makes use of the assumption that buffers are usually unregistered in the reverse order that they were registered: by using a sorted std::vector, freelist::top and freelist::pop can be implemented using vector::back and vector::pop_back, respectively. Most freelist::push operations are done using vector::push_back, which has O(1) complexity since it will never require a reallocation. Only when buffers are not unregistered in reverse order of registration does an insertion in the middle of the vector occur, which requires an O(log N) binary search using std::upper_bound to find the insertion position and then a call to vector::insert, for worst-case performance O(N). Signed-off-by: Omri Mor <Omri.Mor@amd.com>

Signed-off-by: Omri Mor <Omri.Mor@amd.com>

github-actions Bot added the project: rocshmem label Jun 13, 2026

systems-assistant Bot added the organization: ROCm label Jun 13, 2026

omor1 force-pushed the users/omor1/rocshmem/refactor_gda_qp branch from 5172b65 to 795eb8d Compare June 13, 2026 03:48

omor1 mentioned this pull request Jun 16, 2026

[rocshmem] gda/mlx5: skip get_lkey() for inline RMA sends #7294

Closed

1 task

omor1 force-pushed the users/omor1/rocshmem/refactor_gda_qp branch 2 times, most recently from 62b6284 to b5c1293 Compare June 30, 2026 00:36

therock-pr-bot Bot added the Not ready to Review label Jun 30, 2026

omor1 removed the Not ready to Review label Jun 30, 2026

omor1 marked this pull request as ready for review June 30, 2026 00:38

omor1 requested a review from a team as a code owner June 30, 2026 00:38

Copilot AI review requested due to automatic review settings June 30, 2026 00:38

Copilot started reviewing on behalf of omor1 June 30, 2026 00:39 View session

therock-pr-bot Bot added the Not ready to Review label Jun 30, 2026

Copilot AI reviewed Jun 30, 2026

View reviewed changes

omor1 changed the title ~~[rocshmem] gda: refactor QueuePair class~~ refactor(rocshmem/gda): refactor QueuePair class Jun 30, 2026

abouteiller reviewed Jul 6, 2026

View reviewed changes

omor1 mentioned this pull request Jul 7, 2026

Add SQTT marker support to rocSHMEM to improve device-level performance tracing #8176

Open

1 task

omor1 force-pushed the users/omor1/rocshmem/refactor_gda_qp branch 3 times, most recently from 4ca9fab to 8ef949a Compare July 13, 2026 20:26

omor1 added 6 commits July 14, 2026 00:40

[rocshmem] gda: refactor QueuePairBNXT

962f65e

Signed-off-by: Omri Mor <Omri.Mor@amd.com>

[rocshmem] gda: fix building w/multiple providers & w/o generic provider

6d2b8e1

Signed-off-by: Omri Mor <Omri.Mor@amd.com>

[rocshmem] gda: fix invalid memory accesses from QueuePairBNXT refactor

148b526

Signed-off-by: Omri Mor <Omri.Mor@amd.com>

[rocshmem] gda: QueuePairGeneric -> QueuePairMux, GDA_GENERIC -> GDA_MUX

3f8a04c

Signed-off-by: Omri Mor <Omri.Mor@amd.com>

[rocshmem] gda: consistent enum order for ATOMIC_CS & ATOMIC_FA opcodes

c9bc963

Signed-off-by: Omri Mor <Omri.Mor@amd.com>

omor1 added 21 commits July 14, 2026 00:40

[rocshmem] gda: refactor QueuePairIONIC

f6a4c5d

First pass for QueuePairIONIC refactor. Needs some subsequent code cleanup. Signed-off-by: Omri Mor <Omri.Mor@amd.com>

[rocshmem] bitcode: add QueuePairMux (GDA_MUX) sources

5933898

Signed-off-by: Omri Mor <Omri.Mor@amd.com>

[rocshmem] gda/ionic: fix QueuePairIONIC build issues

f81430f

Signed-off-by: Omri Mor <Omri.Mor@amd.com>

[rocshmem] gda/ionic: missing <new> header for placement-new

922998b

Signed-off-by: Omri Mor <Omri.Mor@amd.com>

[rocshmem] gda: rebase fixes

28cab96

Signed-off-by: Omri Mor <Omri.Mor@amd.com>

[rocshmem] gda: use constmem.gda_provider in QueuePairMux

8cf74b7

Signed-off-by: Omri Mor <Omri.Mor@amd.com>

[rocshmem] gda: explicit lkey/rkey parameters in QP post_wqe_rma/amo

b57df63

GIN GDA needs to pass lkey/rkey explicitly to the QP classes. (Why?) Signed-off-by: Omri Mor <Omri.Mor@amd.com>

[rocsmem] gda: make all local constexpr variables static

c50c350

Compilers sometimes find it easier to optimize `static constexpr` block-scope variables over `constexpr` block-scope variables. Signed-off-by: Omri Mor <Omri.Mor@amd.com>

[rocshmem] gda: minor doxygen formatting

05ec71d

Signed-off-by: Omri Mor <Omri.Mor@amd.com>

[rocshmem] gda: minor function signature reformatting

e6c6083

Signed-off-by: Omri Mor <Omri.Mor@amd.com>

docs(rocshmem): [gda] fix typo in comments

c648e68

Signed-off-by: Omri Mor <Omri.Mor@amd.com>

fix(rocshmem): [gda] guard against invalid opcodes in mux provider

34f1616

Signed-off-by: Omri Mor <Omri.Mor@amd.com>

Revert "[rocshmem] gda/mlx5: remove mlx5 code from generic GDABackend…

fadd862

…::create_cqs path" This reverts commit 0869052.

style(rocshmem): [gda] rename QueuePairSHMEM<XYZ>::atomic_add_nbi to …

06c325d

…atomic_add Align with OpenSHMEM naming. Signed-off-by: Omri Mor <Omri.Mor@amd.com>

style(rocshmem): [gda] rename atomic_cas_nbi_nofetch to atomic_cas_no…

815a9d0

…fetch Align with OpenSHMEM naming. Signed-off-by: Omri Mor <Omri.Mor@amd.com>

style(rocshmem): [gda] rename QueuePairSHMEM<XYZ>::atomic_cas to atom…

c454162

…ic_compare_swap Align with OpenSHMEM naming. Signed-off-by: Omri Mor <Omri.Mor@amd.com>

style(rocshmem): [gda] rename post_wqe_amo(value, cond) to (swap_add,…

21f0503

… compare) Align with naming conventions used by GDA providers and in InfiniBand spec. Signed-off-by: Omri Mor <Omri.Mor@amd.com>

fix(rocshmem): [gda] remove unused QueuePairIONIC::print_cqe_error

b646152

Signed-off-by: Omri Mor <Omri.Mor@amd.com>

omor1 force-pushed the users/omor1/rocshmem/refactor_gda_qp branch from 8ef949a to 6da6c27 Compare July 14, 2026 00:40

This was referenced Jul 14, 2026

feat(rocshmem): add _av targeted-ordering variants #8503

Open

feat(rocshmem): [scripts] add --variant-dir option to run_perf_compare.sh #8579

Open

omor1 added 5 commits July 20, 2026 20:38

chore(rocshmem): [gda] call mlx5_devx_qp::dump from initialize_gpu_qp

6b62b2a

Mirror other providers by calling the mlx5_devx_qp::dump function from the GDABackend::initialize_gpu_qp parent call instead of within GDABackend::mlx5_initialize_gpu_qp. Signed-off-by: Omri Mor <Omri.Mor@amd.com>

perf(rocshmem): [gda] annotate some expected-true branches as [[likely]]

693ac68

Signed-off-by: Omri Mor <Omri.Mor@amd.com>

Uh oh!

Conversation

omor1 commented Jun 13, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical Details

JIRA ID

Test Plan

Test Result

Submission Checklist

Uh oh!

omor1 commented Jun 13, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

omor1 commented Jun 24, 2026

Uh oh!

therock-pr-bot Bot commented Jun 30, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

❌ PR Check — Action Required

Uh oh!

therock-pr-bot Bot commented Jun 30, 2026

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

abouteiller left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

omor1 commented Jun 13, 2026 •

edited

Loading

omor1 commented Jun 13, 2026 •

edited

Loading

therock-pr-bot Bot commented Jun 30, 2026 •

edited

Loading