openpi / pi 0.5 multi-GPU NCCL + CUDA crashes

If your openpi pi 0.5 fine-tune dies with NCCL operation ncclGroupEnd() failed: unhandled cuda error, that message is only a wrapper. The real fault sits in the last NCCL warning line above it. That line names one of three things: a cuDNN convolution failure, a CUDA memory mapping error, or an illegal memory access. Each has its own fix, and only one of the three is really about VRAM.

Two crash families cause almost every openpi multi-GPU failure. In the log they look the same, and they are not the same problem. The first family is the NCCL group-end crash, which happens when a collective call fails while it is being launched. The second family is the plain CUDA fault, which stays inside one GPU and never touches the network. A third shape shows up too: it hangs, then dies on a timer, and that is not a crash. This guide covers all three, with the exact string, the cause behind it, and the change that fixes it.

The strings you will see, and what each one means

Why the group-end message is a wrapper

The NCCL documentation says this plainly. The codes ncclUnhandledCudaError and ncclSystemError mean that a call to an external library failed. The page tells you to read the NCCL warning message to resolve the problem (NCCL runtime and MPI issues).

So the NCCL message names the category but never the cause. The cause is one line up, in the last NCCL warning, and it is the line worth copying into a search box.

There is a second reason the message misleads. The group-end call is where NCCL launches the work for a whole group of collectives. The group-calls page says the operation starts as a whole during the last call to ncclGroupEnd (NCCL group calls). A group-end crash can therefore be the first visible sign of an earlier fault. The fault began several collectives earlier, on one rank, inside a kernel that had nothing to do with the network.

That single fact explains why the same line shows up in three different situations. One is a bad convolution on one GPU. Another is a memory mapping error tied to the memory layout of the host. A third is a genuinely broken link between two machines. Fixing the network settings changes nothing in two of those three cases.

Family one: the NCCL group-end crash

Cause A: a cuDNN convolution failure on a consumer Blackwell card

This is the openpi case that put the error on the map. The openpi issue thread carries the title when finetune 4*5090 pi0.5，NCCL operation ncclGroupEnd() failed: unhandled cuda error, and it was opened on 2026-10-10 against a fine-tune on four GeForce RTX 5090 GPUs.

The thread holds two failed runs with two different last lines. The first ends with:

Cuda failure 'an illegal memory access was encountered'

The second, at a lower batch size, ends with:

Cuda failure 2 'out of memory'

A third run reaches the cuDNN failure on its own, with no NCCL wrapper in front of it:

Execution of replica 0 failed: UNKNOWN: CUDNN_STATUS_EXECUTION_FAILED

The cause behind those lines is a version gap, and the fix lives in the version pins. The openpi lock file selects nvidia-cudnn-cu12==9.5.1.17 and nvidia-cuda-runtime-cu12==12.6.77 on Linux x86_64. The cuDNN 9.5.1 support matrix lists compute capability 9.0 down to 5.0 for CUDA 12.x and stops there, so compute capability 12.0 is absent from that list. The cuDNN 9.7.1 matrix does list 12.0 and names NVIDIA Blackwell, but it asks for CUDA 12.8 and a driver at 570.26 or newer.

An RTX 5090 is compute capability 12.0 on the NVIDIA compute-capability table. That gap is the whole distance between the card and the pinned library.

Two things keep this from being a one-line fix. The pinned torch==2.7.1 requires cuDNN 9.5.1.17 on this platform, so a cuDNN-only upgrade breaks the dependency tree. That run also used jaxlib 0.5.3, and the reporter traced the fault to the cuDNN build bundled inside it. The stack has to move together, and that is a slower fix than most people expect.

Cause B: the out-of-memory inside NCCL that is not about VRAM

Read Cuda failure 2 'out of memory' inside an NCCL log as the GPU is full and you will chase the wrong thing. NCCL maps host memory through the CUDA memory management API to hold its buffers. When that mapping fails, the driver reports an allocation failure. NCCL then prints an out-of-memory line, even though the GPU still has free memory.

The clearest record is a long PyTorch issue thread on the exact string (pytorch issue 152302). It carries the official error text:

ncclUnhandledCudaError: Call to CUDA function failed.

and then the last error:

Cuda failure 2 'out of memory'

In that thread the model is a toy network with two small layers. The fault shows up after PyTorch 2.7 with NCCL 2.26.2, and it needs more than one GPU before it appears. An NCCL maintainer traces it to host memory mapping calls and names the setting that works around it, NCCL_CUMEM_HOST_ENABLE=0. The same thread shows the fault on 4090s as well as 5090s, so this is not a consumer card defect. The fault also follows the memory layout of the machine: a pair of GPUs under one memory node works, and a pair split across nodes fails.

openpi pins NCCL 2.26.2 and PyTorch 2.7.1 on Linux x86_64. The thread names the same pair.

The fix has two parts. Set NCCL_CUMEM_HOST_ENABLE=0 before the run, and check the layout with nvidia-smi topo -m and numactl -H. If the pin is allowed to move, a later NCCL release carries a transparent check and a fallback for this mapping failure. That is the durable answer. One warning from the same thread: that setting on its own can trade the out-of-memory error for an illegal memory access during a checkpoint save. Test a save before you trust it.

Cause C: an illegal address from a multi-step action horizon

A different openpi issue reports a crash tied to action_horizon above one (openpi issue 539). The strings are:

CUDA error: : CUDA_ERROR_ILLEGAL_ADDRESS: an illegal memory access was encountered

jaxlib.xla_extension.XlaRuntimeError: INTERNAL: CUDA error: Could not synchronize CUDA stream

The reporter notes that the same setting runs on CPU, and the default horizon also runs on eight H100s. It appears when the horizon is two or three. Returning to action_horizon = 1 is the state the reporter confirmed working. The thread then closed elsewhere: the maintainer considered the data-transform path unlikely and suspected a conflict between the system-installed CUDA and the pip-installed JAX stack, and the reporter confirmed that making the JAX and CUDA versions compatible fixed it. The data-transform hunt was the reporter's original hypothesis, and the thread moved past it.

Cause D: the shared memory transport

The same wrapper also appears with unhandled system error instead of unhandled cuda error. The error-class table calls this one ncclSystemError, which covers a system call or an external library call that failed. In the PyTorch thread the system error variant shows the shared memory attach failing, with a line of the form Error while attaching to shared memory segment /dev/shm/nccl-.... The NCCL documentation on runtime and MPI issues names this failure mode and gives the container fix: increase the shared memory size, for example with a setting like --shm-size=1g. The layout check still applies first.

Family two: CUDA faults that never reach NCCL

The cuDNN failure on its own

CUDNN_STATUS_EXECUTION_FAILED shows up by itself in another openpi issue about multi-GPU training (openpi issue 725), where the JAX traceback reads jaxlib.xla_extension.XlaRuntimeError: UNKNOWN: CUDNN_STATUS_EXECUTION_FAILED. That fault has the same root as cause A. The convolution engine that the compiler picks does not run on that card with that cuDNN build. Treat this string and the group-end wrapper as the same fault seen from two angles.

The PyTorch out of memory

A single RTX 4090 and a Pi0 base fine-tune produce this in the openpi PyTorch path (openpi issue 677):

torch.OutOfMemoryError: CUDA out of memory. Tried to allocate 64.00 MiB. GPU 0 has a total capacity of 23.63 GiB of which 226.56 MiB is free. Including non-PyTorch memory, this process has 22.77 GiB memory in use. Of the allocated memory 21.98 GiB is allocated by PyTorch, and 320.27 MiB is reserved by PyTorch but unallocated. If reserved but unallocated memory is large try setting PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True to avoid fragmentation.

That one is a real memory shortage, and the sentence structure tells you which kind. A small tried to allocate next to a large reserved figure means the card is full. A large reserved-but-unallocated figure next to free memory means the allocator is fragmented. The openpi README troubleshooting table gives the openpi answers: set XLA_PYTHON_CLIENT_MEM_FRACTION=0.9 before training, use --fsdp-devices <n> to shard, and consider disabling EMA. The same table notes that the current script does not support multi-node training, so a multi-node plan is not a way out.

The dtype failure

RuntimeError: Expected Float, got BFloat16 appears in an openpi PyTorch multi-GPU report where one card through seven cards worked and eight cards failed (openpi issue 758). The error is raised during the backward pass. The README says the PyTorch path trains in full bfloat16 by default and does not support mixed precision yet. The fix is to pick one dtype for the whole run. Set it in the config rather than mixing.

The Triton shared memory warning

No valid triton configs. OutOfResources: out of resource: shared memory, Required: 131072, Hardware limit: 101376. Reducing block sizes or `num_stages` may help.. shows up while the PyTorch compiler autotunes a pi 0.5 run (openpi issue 654). On its own this is not a crash: the compiler ignores the failing candidate and tries the next one. It matters when every candidate fails, because then the slow fallback path runs for the whole job.

The third shape: a hang is not a crash

A run that sits still and then dies on a timer is a different failure with different tools. NCCL ships a watchdog that monitors collectives, and it aborts the operation when one runs past the timeout. The timeout is set by NCCL_TIMEOUT, and the default is 1800 seconds (NCCL errors in the OSMO documentation). The same page lists the symptoms. A run hangs before the timeout fires, and progress across GPUs turns uneven. It names four usual causes: uneven work per rank, a rank that skipped a collective, resource contention, and network or device faults.

The fix path is the one the NCCL troubleshooting guide describes. Start with NCCL_DEBUG=INFO. Pin the network interface with NCCL_SOCKET_IFNAME when a box has several. Turn off the peer-to-peer path with NCCL_P2P_DISABLE=1 to test whether it is the link. None of these will help if the crash class is what you have, because there is no hang to fix.

A diagnosis order that works

1. Find the last NCCL warning line. Not the exception, the warning: it is the only line that names a cause. If NCCL_DEBUG was not set, set it and run again with a short step count.
2. Match that line against the table at the top of this page. The last line decides the family, and the family decides the tool.
3. A cuDNN failure or an illegal memory access as the last line means you should check compute capability against the cuDNN version your stack actually loaded. Print the version from the installed package, not from the README.
4. An out-of-memory inside NCCL as the last line needs a layout check. Run nvidia-smi topo -m and numactl -H before you free any memory, and if the layout is split, set NCCL_CUMEM_HOST_ENABLE=0.
5. A PyTorch allocator message as the last line gives you five fields to read. Then apply the README answers for memory and sharding.
6. Reproduce small. A crash at step 20 with a batch of 2 costs a minute. At step 5000 the same crash costs an hour of waiting.
7. Check the pre-flight guards last. openpi raises Batch size {n} must be divisible by the number of devices {d}. and Number of devices {d} must be divisible by the number of FSDP devices {n}. before it starts. Those are configuration errors, not crashes, and they are worth fixing before the next run.

Prevention: set the environment before the first step

export NCCL_DEBUG=INFO
export NCCL_CUMEM_HOST_ENABLE=0
export XLA_PYTHON_CLIENT_MEM_FRACTION=0.9

Three rules carry most of the value. Turn on NCCL logging before the first run, so the cause line sits in the file when the crash comes. Set the JAX memory fraction to 0.9, as the openpi README asks, so the allocator sees the whole card. Read the layout once with nvidia-smi topo -m, so you know which GPUs sit under which memory node before you trust a multi-GPU run.

On a consumer Blackwell card, check the cuDNN version against the support matrix before you spend an evening on the crash. If the matrix does not list your card, no flag will fix it.

When you would rather not own the stack

Repairing a GPU stack is real work, and it is not always the right use of a week. If you would rather not own that, the same policy can run somewhere else. The robot model page states that family's current status, and the local versus hosted break-even guide works through the trade with your own numbers. For what any metered route costs, the answer lives on the pricing guide, never in prose. You can sign up and make a first call when you are ready.

If you do keep the stack, use the evaluation guide to confirm that a repaired run learned something. Stopping the crash is not the same as learning. The neighbouring failure family, where missing statistics produce wrong motion with no error at all, has its own guide.

Common questions

Does an ncclGroupEnd error mean my network is broken?

Usually not. It means a CUDA call failed inside the collective launch. The NCCL docs point you at the warning line above it for the cause. Read that line first. In the openpi case that started this guide, the fault was a cuDNN convolution on one GPU, so no network change would have helped.

Is out of memory always about VRAM?

No. Two different failures print that line. The PyTorch allocator sentence is about VRAM, and it names the card, the free amount and the fragmentation. A bare NCCL warn line with Cuda failure 2 'out of memory' is usually the CUDA memory mapping path failing on the host side. That is a layout problem rather than a capacity problem.

Can I fine-tune pi 0.5 on an RTX 5090?

The card is compute capability 12.0. The cuDNN version that openpi pins for Linux x86_64 does not list 12.0 in its support matrix. The PyTorch pin ties that cuDNN version in place. A newer cuDNN build lists 12.0 but wants CUDA 12.8 and a driver at 570.26 or newer. Moving the whole stack together works; moving one piece does not.

Do I need more than one GPU for pi 0.5?

For a LoRA fine-tune, no. The openpi README puts a LoRA fine-tune above 22.5 GB and a full fine-tune above 70 GB. The JAX path shards a full fine-tune across devices with fsdp_devices. Add a second GPU when you need the memory or the throughput, not to avoid a crash.

Why does the same command work on one GPU and fail on four?

The cuDNN and illegal-address faults in this family fire inside one device. The group-end wrapper only shows up once a collective is in play. The mapping failure needs at least two devices before it appears, so the multi-GPU run surfaces a fault that was always there.

What is the difference between a crash and a hang?

A crash throws and stops. A hang waits for a collective that never completes, and the watchdog ends it when the timeout fires. The NCCL timeout default is 1800 seconds. If your log ends with a timeout message, you have a hang and you need the transport and topology tools. If it ends with a CUDA or cuDNN message, you have a crash and you need the version pins.

Where do I look if the error is still not clear?

Start with the issue search in the openpi repository. The crash families are indexed there by their exact strings, so a new report is worth filing with the last NCCL warning line attached.

The short version

The group-end message is a wrapper, and the cause sits in the last NCCL warning line above it. Three faults live behind it. The first is a cuDNN convolution that the pinned cuDNN does not support on a compute capability 12.0 card. The second is a CUDA memory mapping failure that reads as out of memory but is about host memory layout. The third is an illegal address from a multi-step action horizon. Around them sit the plain CUDA faults, the PyTorch allocator running out of room, and a dtype mismatch, while a Triton autotune warning is not a crash at all. A hang that fires on a timer is a fourth shape, and it needs the transport tools instead of the version pins. Set NCCL_DEBUG=INFO, set the memory fraction to 0.9, read the memory layout once, and match your last warning line against the table at the top. Costs for any hosted route live on the pricing guide. When you are ready to make a first call, sign up.

Last verified: 2026-10-26