Does the LRLink PCIe Bridge Fix RCCL's Broken P2P on Dual R9700? (Update: Yes it does!)

High-Performance Computing Machine Learning, LLMs, & AI
,
1

@wendell ,

I caught the recent video on the LRLink PCIe bridge with the dual R9700 setup and it got my brain spinning. I’m running 2x R9700s on a Threadripper TRX40 box for local LLM inference and I’ve been digging into whether local training is viable yet on these cards. Short answer: it’s getting there, but multi-GPU training on R9700s has a nasty blocker.

RCCL (AMD’s NCCL equivalent) has a known bug where hipIpcGetMemHandle crashes when two R9700s try to set up peer-to-peer communication through the CPU’s root complex. It’s been sitting in AMD’s triage for about five months now (rccl-tests issue #162). The workaround is NCCL_P2P_DISABLE=1 which forces everything through host-staged copies and roughly halves your inter-GPU bandwidth. It works, but you’re leaving a lot on the table.

Here’s where it gets interesting… The crash happens specifically when RCCL tries to export IPC memory handles through the CPU root complex. But with a LRLink Broadcom bridge, both GPUs sit behind a shared PCIe switch fabric. RCCL does topology detection and differentiates between “same switch” and “separate root complex endpoints.” So the question is: does putting both R9700s behind the bridge change the topology enough that RCCL either takes a different P2P path that doesn’t crash, or succeeds on the IPC path that currently fails through the root complex?

Nobody has tested this. that I could see. I’ve dug through GitHub issues, forums, and Reddit with zero results for anyone trying a PCIe switch as a workaround for this specific bug on RDNA 4.

There’s also a bonus data point that makes this even more interesting. A brand new vLLM issue (#40980, filed a few days ago) shows that TP=2 on dual R9700 actually worked on vLLM v0.14.0 with RCCL 2.27.3 back in December, but deadlocks on v0.19.x with RCCL 2.27.7. So it’s a regression, not a fundamental hardware limitation. Which means there might be multiple angles here.

As far as I can tell, the test is pretty straightforward:

  • Install rccl-tests
  • Run all_reduce_perf -g 2 in three configs:
    1. Both GPUs through CPU root complex (normal setup) with NCCL_P2P_DISABLE=1
    2. Both GPUs behind the LRLink bridge without NCCL_P2P_DISABLE
    3. Both GPUs behind the LRLink bridge with NCCL_P2P_DISABLE=1
  • Compare bandwidth numbers

If the bridge topology lets RCCL’s P2P path succeed where the root complex path crashes, that’s a hardware workaround for a paywall software bug AMD hasn’t fixed in half a year. If it doesn’t work, we’ve at least eliminated the idea and saved everyone else the experiment.

Either way it seems like it’d make for a great video. You’ve literally got all the hardware sitting right there already.

For context, I’m building toward local LoRA training on the R9700s and this P2P bug is the main thing standing between single-GPU training (which works) and being able to use both cards for larger models.

Thanks for the consistently great content. The PCIe lanes are fungible line lives rent-free in my head now.

6 Likes
Understanding PCI-e P2P? Requirements, Use Cases and Possible Pitfalls
Dual R9700's AMD VLLM-Container woes
Someone needs to figure out how to adapt MI250 GPUs to PCIe
Can someone recommend a reasonably priced PCIe 5 Xeon platform for me?
Project Deep Thoughts (Or: Watch me light money on fire and call it learning)
2

I mean, while we are wishing for videos, maybe one on whether using one off the LRLink PCIe Bridge on a PCIe Gen 3 MB will still allow the GPUs to connect at Gen 5 speeds, which makes using them on older boards a much more viable option.

And are there any of these that go beyond just 32 lanes? I could see one card for 8 GPUs being extremely useful for long running tasks.

And a really crazy setup might be to have a High Speed Network Adapter on the card, and one on a second card, routing directly from the GPUs memory to others.

3

In theory, it would give somewhat of a boost with two cards talking to each other, but likely not as much as native Gen 4-5 lanes to the CPU would.

I decided to buy one of these https://www.ebay.com/itm/136730781261?var=435322480730 that should allow three 16x PCIe Gen 4 connections. I have a TRX40 board so I only get 2 slots that can do 16x. This is cheaper than moving up to a TRX80 (with current market prices). So even if it doesn’t fix my training issue it will at least give me more PCIe lanes to expand from 2 R9700s to 3-4.

It may be interesting to test how much lowering the PCI Gen on the Mobo down to 3-2 from 4 changes tok/s, then again with the cards connected to the LRLink.

My card is coming on the slow boat, but if Wendell doesn’t have time to test it, I can post numbers.

4

That would be great. There is that 6 slot package bundle, so I could see using 4 for a GPU, one for 100/200/400GbE and one port for NVMe storage on each of the expanders, then a second set on the other, if the networking worked out, kinda doing what Nvidia did with their setup, but much cheaper, and more jank.

Gotta love the jank, right? :smile:

5

There are indeed such as:

He is working on a 160 lane version for around 2900 euros if I remember correctly.

1 Like
6

Very cool. He has some nice PCIe 4 boards as well.

Wonder if they allow CXL?

I just got to thinking about it, but having 2 GPUs and 2 CXL large memory dedicated to them might clear some bottlenecks while massively expanding the total amount of memory they have available.

This would take engine level work in vLLM or llama.cpp, but could make using cheaper GPUs viable, just page the needed parts into and out of the GPU at runtime.

7

Got the card installed today, mixed results.

The short version is using the PCIe switch DOES solve the p2p crashing issue with AMD cards. However, there are a lot of gotchas with it, At least with my hardware and an eBay card.

I’m only getting 0.67-1.35BG/s with p2p. While it works, that’s not viable for actual training. I did try a few other settings to try and tune it, but nothing seemed to help much.

Test command: HIP_VISIBLE_DEVICES=0,1 ./build/all_reduce_perf -b 8 -e 128M -f 2 -g 2

| Variant                          | Saturation busbw | Avg busbw |
|----------------------------------|-----------------|-----------|
| P2P enabled (through switch)     | 0.67-1.35 GB/s  | 0.37-0.70 |
| P2P disabled (host-staged)       | 0.55-1.07 GB/s  | 0.28-0.52 |
| HSA_ENABLE_SDMA=0                | no change       | no change |
| HSA_ENABLE_PEER_SDMA=0           | no change       | no change |
| NCCL_DMABUF_ENABLE=1             | no change       | no change |

P2P advantage: 22-45% depending on physical cable configuration

This is a quick compare to what NV cards would get:

| Setup                              | all_reduce busbw  | vs R9700 best |
|------------------------------------|-------------------|---------------|
| R9700 through PEX88096 (this test) | 0.67-1.35 GB/s    | 1x            |
| RTX 4090 PCIe P2P (no NVLink)      | ~12-24 GB/s       | 9-18x faster  |
| RTX 3090 NVLink pair               | ~56 GB/s          | 41x faster    |
| V100 NVLink pair                   | ~50 GB/s          | 37x faster    |
| A100 NVLink                        | ~300 GB/s         | 222x faster   |

Debugging didn’t seem to give me any clues either.

NCCL INFO Check P2P Type isAllDirectP2p 1 directMode 1
NCCL INFO Channel 00/0 : 0[7000] -> 1[d000] via P2P/direct pointer
NCCL INFO Channel 01/0 : 0[7000] -> 1[d000] via P2P/direct pointer
NCCL INFO Connected all rings, use ring PXN 0 GDR 1

There is also, in-fact, a loss of overall speed for general Vulkan inference which using the switch I bought. Native gen4 16x on two slots for a TRX40 board:

| Model  | Layout | Test   | pre (no switch) | post (switch) | Delta  |
|--------|--------|--------|-----------------|---------------|--------|
| 4B     | single | tg128  | 168.3 t/s       | 162.3 t/s     | -3.6%  |
| 30B    | single | tg128  | 173.9 t/s       | 166.7 t/s     | -4.1%  |
| 30B    | split  | tg128  | 130.6 t/s       | 126.8 t/s     | -3.0%  |
| 80B    | split  | tg128  | 67.1 t/s        | 65.5 t/s      | -2.5%  |

So, one step forward, two steps back. I’m going to play around with this a bit more, but open to ideas.

8

Doesn’t AMD have their own version of P2P? Its not NCCL, but keeps the CCL part?

9

Correct. There were TWO things breaking that was killing performance:

  1. The cables. I would lose speed if the cables were too close to power or crossing over each other. I would NOT bother with a kit. Get just the card and source your own cables and adapters. I have some higher quality cables on order that should be here Friday.

  2. ReBAR was poorly configured. It was enabled in the BIOS, but it wasn’t correctly adjusting in the OS. Once I turned that, HUGE speed boost!

ReBAR results:

# nThread 1 nGpus 2 minBytes 1048576 maxBytes 134217728 step: 2(factor)

     Size         busbw (out-of-place)    busbw (in-place)
     1048576      4.99 GB/s               5.01 GB/s
     2097152      5.19                    5.20
     4194304      5.30                    5.31
     8388608      5.36                    5.36
    16777216      5.39                    5.39
    33554432      5.39                    5.39
    67108864      5.40                    5.40
   134217728      5.41                    5.41

Commands I used:

# Full GRUB line:
GRUB_CMDLINE_LINUX_DEFAULT="quiet splash iommu=pt amdgpu.runpm=0 pci=realloc"

# Apply:
sudo update-grub && sudo reboot

# Verify after reboot:
cat /proc/cmdline
# Should contain: iommu=pt amdgpu.runpm=0 pci=realloc

# Verify ReBAR:
lspci -vvs <GPU_BDF> | grep "Region 0"
# Should show [size=32G] not [size=256M]

# Verify IOMMU passthrough:
journalctl -k -b 0 | grep -i 'iommu.*passthrough'
# Should show: Default domain type: Passthrough

This is HUGE! This enables real training on AMD cards. I’m hoping with better cables I can get to 10-15GB/s. Still slower than Nvidia cards, but better $ value than anything else.

10

Yes, rccl. There have been regressions noted in ROCm issues/discussions.

11

ROCm 7.2.0 results 2xR9700

Bytes Per Direction 268435456
Unidirectional copy peak bandwidth GB/s [Local read / Remote write] (GPU-Executor: GFX)
 SRC+EXE\DST    CPU 00       GPU 00    GPU 01
  CPU 00  ->     45.37        21.52     21.53

  GPU 00  ->     28.10       294.36     28.12
  GPU 01  ->     28.18        28.11    299.76
                           CPU->CPU  CPU->GPU  GPU->CPU  GPU->GPU
Averages (During UniDir):       N/A     21.53     28.14     28.11

Bidirectional copy peak bandwidth GB/s [Local read / Remote write] (GPU-Executor: GFX)
     SRC\DST    CPU 00       GPU 00    GPU 01
  CPU 00  ->       N/A        20.87     20.72
  CPU 00 <-        N/A        27.91     27.91
  CPU 00 <->       N/A        48.77     48.63


  GPU 00  ->     27.90          N/A     27.59
  GPU 00 <-      20.86          N/A     27.59
  GPU 00 <->     48.76          N/A     55.18

  GPU 01  ->     27.93        27.59       N/A
  GPU 01 <-      20.64        27.59       N/A
  GPU 01 <->     48.57        55.19       N/A
                           CPU->CPU  CPU->GPU  GPU->CPU  GPU->GPU
Averages (During  BiDir):       N/A     24.35     24.33     27.59
12

Several months back I went down this rabbit hole on Asus SAGE WRX80 Wifi-II w/ 5995WX and 2xR9700, bios:

Preferred IO
Infinity Fabric Performance State = P0
DF C-States = Disabled
PCIe Ten-Bit Tag Support = Enabled

Kernel cmdline:

GRUB_CMDLINE_LINUX_DEFAULT="quiet amd_iommu=on iommu=pt pcie_aspm=off amdgpu.ras_enable=0 pci=realloc=on amdgpu.runpm=0 amdgpu.gpu_recovery=1 pcie_acs_override=downstream,multifunction"

Tried various pci slot combos. Hardware doesn’t have a switch, infinity traffic limits.

        size    busbw (out-of-place)    busbw (in-place)
     1048576    14.07 GB/s				14.07 GB/s
     2097152    16.39    				16.43
     4194304    17.84    				17.79
     8388608    18.64    				18.63
    16777216    18.98    				18.99
    33554432    19.18    				19.17
    67108864    19.29    				19.16
   134217728    19.23    				19.22
1 Like
Dual R9700's AMD VLLM-Container woes
13

I am assuming the vulkan inference speeds go up now, right?

Glad you got that figured out. Where your Nvidia numbers on that same system with the PCIe risers? Or just a different system you are using as a baseline?

I wonder if you are in some mode that is just using one of the cables per slot?

That looks like about half of the Gen 3 bandwidth of an x16 slot, right? Maybe there is some switches or software to tell the card to pair 2 of them together? Or on the PCIe boards?

It might double your speed if I am right.

Next, is there a way to artificially shield the cables? Maybe wrap them in aluminum and ground it?

It might make do. Don’t let it be a temporary solution that never changes.

14

This was the part I needed. Huge boost in speed!

Component State
Cmdline iommu=pt amdgpu.runpm=0 pci=realloc
ReBAR 32 GB on both GPUs
ACS Cleared on 7 PEX880xx DSPs; persistent via clear-pex-acs.service (enabled)
RCCL allreduce 19.60 GB/s at 128M (was 5.41)
Raw P2P (TransferBench) 28.07 GB/s unidir (was 7.00)
infinity_emb Running, /health 200, model on both cards

The whole testing arc, compressed

Stage RCCL busbw Factor
Original (pre-PCIe-switch, Vulkan host-staged) ~0.4 1.00×
Post-switch, ReBAR 256M (yesterday morning) 0.67 1.7×
Post-cable-reseat, ReBAR 256M 1.85 4.6×
Post-pci=realloc (ReBAR 32G) 5.41 13.5×
Post-ACS-clear (today) 19.60 49×

So, now this is a clear path forward for people that want to use multiple AMD GPUs for testing. The PCIe Switch removes the crashing issue and these system settings get the bandwidth upto some Nvidia cards.

2 Likes
15

The switch is buying you crash avoidance, which is absolutely a performance win. I’m hoping AMD tracks down and addresses rccl bug. :crossed_fingers:

1 Like
16

Holy, that’s a large jump. So, it wasn’t the cables or the PCIe risers?

Do you still recommend getting better quality ones?

1 Like
Case suggestions for a Multi-GPU pcie backplane
17

Pure tuning.

That said, I still know for a fact the cables that came with the eBay kit are trash and still waiting on the new ones. That may get me closer to saturating the Gen4 bus, but already this has gone from “It kinda works…” to “It works and competes with lower end Nvidia cards!”

Next thing is getting better cooling on this card. The included fan is only enough to keep it from melting, but that’s with a VERY well vented case. I have a generic water block on order that I’m going to add to the loop.

18

Do you have a thread on this forum with these details or more?

19

I do not, but summary of tuned BIOS settings and kernel flags are posted above. You can skip “Preferred IO” BIOS settings, those had no measurable impact.

The adventure started much like this thread. Poor P2P bandwidth number from rocm_bandwidth_test. I had noted topological 2 hops, but was not able to reduce it on the Asus Pro WS WRX80E Sage Wifi II board with any pcie slot (wish Asus published a schematic in their manual). Physically everything goes through CPU infinity fabric. That’s my limiting factor currently.

I do have ACS disabled in bios. But the kernel parameters were for better iommu grouping. ReBar defaults to auto in my BIOS. The kernel flags were added ensuring maximum 32GB is allocated.

All other settings were minor tweaks governing power, mostly an attempt to address / mitigate a separate, but common issue with this board, AER. My system occasionally emits AER on primary m.2, a Solidigm P44.

1 Like
20

So I can say that a PCIe switch card does solve the P2P bug for AMD cards, but there is a deeper issue with training on dual cards after you make it past the P2P block. DDP training crashes with hipErrorIllegalAddress when tensor shapes vary across steps (dynamic padding). The PyTorch caching allocator fragments VRAM under varying allocation sizes, and on gfx1201 that fragmentation eventually produces an invalid address.

Thankfully, the fix seems simple:

# CRASHES:
tokenizer(texts, padding=True, truncation=True, max_length=256)

# WORKS:
tokenizer(texts, padding="max_length", truncation=True, max_length=256)

I was able to test this on Qwen3-8B with LoRA r=16, 1000 Alpaca samples, BF16, padding=“max_length” at max_length=256. 62 optimizer steps, 125 micro-steps, 14.5 samples/sec, final loss 1.372, 29.55 GB peak VRAM per rank.

So far, it’s stable!

There is a tradeoff though: With median sequence length 62 and max_length 256, ~75% of tokens are padding. You waste some compute but the allocator never fragments. At 14.5 samples/sec this is still fast enough for all practical LoRA workloads. Reducing max_length to 128 would cut waste if your dataset allows it (check your sequence length distribution first).

I also wanted to compare this setup to a budget RunPod option that is close to my setup (as close as I could reasonably get…).

Here are the AI summarized results:

Model: Qwen3-8B, LoRA r=16, 1000 Alpaca samples, BF16, 1 epoch

Config RunPod 2x RTX PRO 4500 Local 2x R9700 PRO
1-GPU samples/sec 9.81 8.79
2-GPU DDP samples/sec 17.93 14.50
1→2 GPU scaling 1.83x 1.65x
Final loss 1.4369 (1-GPU) 1.372
Peak VRAM — 29.55 GB/rank
Cost/hr $1.28 $0

Per-GPU: R9700 is 11.6% behind RTX PRO 4500 Blackwell (1-GPU comparison).
2-GPU: R9700 pair is 19% behind (14.50 vs 17.93), wider gap due to
fixed-padding overhead reducing scaling efficiency (1.65x vs 1.83x).
Losses match to 0.06% at 1-GPU — numerical parity across backends.
RTX PRO 4500 retails ~$4,000/card. R9700 PRO ~$1,300/card.

So even with the cost of a PCIe Gen4 switch, I’m nudging at a RTX 4500 x2 RunPod. I’m happy with that!

Root cause was fun to track down. Had to eliminate multiple variables until the issue was tracked down:

To find the root cause we built 17 standalone reproducers, each
adding one ingredient. Every stage either passed or crashed,
isolating the trigger by elimination.

HIP-level (no PyTorch):
  A1  Single-GPU hipblasSgemm × 100                    PASS
  A2  2-thread, 2-GPU, matmul + P2P + barrier          PASS
  A3  + watchdog thread (hipSetDevice every 10ms)       PASS
  B   + ncclAllReduce of 16MB grad buffer               PASS
  B2  10×1.6MB bucketed allreduces (250 collectives)    PASS
  B3  Two streams, compute/comm overlap, no sync        PASS
  B4  15 GB VRAM pressure per GPU                       PASS

PyTorch-level (no HuggingFace):
  P1  Bare DDP, 8M params, 20 steps                     PASS
  P2  + gradient checkpointing                           PASS
  P3  1.34B model, full DDP gradient flow                PASS
  P4  + PEFT LoRA                                        PASS
  P5  1.34B + per-layer checkpointing                    PASS
  P6  + gradient_accumulation_steps=2 + no_sync()        PASS

HuggingFace-level:
  P7  Qwen3-8B + PEFT + ckpt + grad_accum + DDP         PASS (20 steps)
  P7  Same but 200 micro-steps, fixed shape              PASS
  P7  Same but real data + dynamic padding               CRASH (step 40)

Every layer below PyTorch's allocator is clean. The HIP runtime,
RCCL, the kernel module, P2P transport — all exonerated by the
MRE suite. The bug needs PyTorch's allocator + varying tensor
shapes + gfx1201's memory management interacting together.

These are the likely upstream issues:

1 Like
Dual R9700's AMD VLLM-Container woes