[Tracking] AMD: training on RDNA 1 cards (RX 5700 XT, gfx1010) #11614

Tracking issue for training on AMD RDNA1 (RX 5000 series, gfx1010) with Unsloth. Today it does not work at all; with the three changes linked below it does, on Windows, for the models that do not need Triton dot products. On Linux (WSL2) the same stack is confirmed on RDNA2 but the RDNA1 card is stopped by the WSL GPU driver before ROCm even starts, details under Still open.

Short version

• Unsloth Studio cannot install for an RDNA1 card. Its ROCm 7.13 wheels come per GPU family and there is no gfx1010 family, so the installer either picks another card's family (if one is present) or falls back to CPU torch. Unsloth Studio (AMD/ROCm): a second AMD GPU the installed PyTorch cannot run is still offered for training, and the installer never warns #11569 / Unsloth Studio (AMD/ROCm): warn about, and refuse, a GPU the installed PyTorch has no kernels for #11571 cover the "another card is present" half.
• With AMD's multi-arch index (torch[device-gfx1010]; the stable one at repo.amd.com/rocm/whl-multi-arch, torch 2.11.0+rocm7.14.1, and the 10.2 nightly both measured) torch itself runs fine on the card. Then three things break in turn, each with a fix:

What trains on an RX 5700 XT once those are in

Windows 11, torch 2.12.0+rocm10.2 (AMD multi-arch nightly; the stable 2.11.0+rocm7.14.1 gives the same numbers, see the installer PR), triton-windows 3.8.0.post28, unsloth 2026.9.7, seq 1024, batch 2, 16-bit LoRA, adamw_torch (no bitsandbytes kernels on that stack).

Background for anyone new to this

On AMD, PyTorch and Triton compile GPU code per card generation. RDNA1 (2019: RX 5700 XT, 5600 XT, 5500 XT) is the oldest generation Triton still targets and it is missing two things newer cards have: bf16 support and dot-product instructions. Everything above follows from that. RDNA2 (RX 6000) has the same bf16 gap, which is what unslothai/unsloth-zoo#1316 was found on, but does have dot instructions.

Still open

• Linux (WSL2). Tried. The RDNA1 card is blocked one layer lower than anything Unsloth touches: the WSL ROCm backend (librocdxg) has to reserve one GPU virtual-address range that every adapter accepts, and the 5700 XT's Windows driver refuses every reservation size tried, from the default 16 GB down to 1 GB (dxgkio_reserve_gpu_va: EOVERFLOW in dmesg). With both cards present AMD's shipped library crashes at init for both; a locally rebuilt one with a per-process adapter filter brings the 6500 XT back, and on it the same Linux multi-arch stack (torch 2.12.0+rocm10.2, AMD's own triton 3.8.0+rocm10.2) trains with GPUs without bf16 (AMD RX 5000 / 6000): keep offloaded activations in their own dtype so Qwen3.5 training doesn't crash unsloth-zoo#1316: Qwen3.5-0.8B 0.0366620 0.0366620 0.0281449, Qwen3-1.7B 0.0207567 0.0207567 0.0191435, both 3/3 steps. So Linux-on-WSL is confirmed for RDNA2 and not reachable for RDNA1 until the driver accepts a reservation; a bare-metal Linux box would not go through this path at all and is untested.
• Installer. PR: Unsloth Studio (AMD RDNA 1, Windows): install ROCm torch for the RX 5700 XT from AMD's multi-arch index #11755 (draft). Windows installers route gfx1010 / gfx1011 / gfx1012 to AMD's stable multi-arch index (repo.amd.com/rocm/whl-multi-arch), pinned to 2.11.0+rocm7.14.1 with torchvision and torchaudio on the same tag (an earlier revision used the 10.2 nightly), no torchaudio (none published). A fresh install from that branch on this box with only the 5700 XT visible produced the multi-arch torch and trained Qwen3.5-0.8B (LoRA) and Qwen3-1.7B (QLoRA) with the same losses as the hand-built stack. Linux is left as is until someone runs it on bare metal.
• gemma compiled path. With unsloth's compiled modules, gemma-3-270m gives a different loss curve on every run on RDNA1, and on the RX 6500 XT under ROCm 10.2 it went NaN from the first backward in 2 of 3 runs. UNSLOTH_COMPILE_DISABLE=1 makes both cards bit-identical and correct, and plain transformers+peft in fp32 is bit-identical across both cards. Separate issue: to follow
• Triton upstream. The buffer-ops silence and the FDOT2 abort are Triton AMD backend bugs for gfx1010. Upstream report: to follow
• QLoRA. Works. The AMD Windows bitsandbytes wheel Unsloth Studio already installs (libbitsandbytes_rocm715.dll) runs NF4 and 8-bit AdamW correctly on gfx1010 under the ROCm 10.2 torch: Qwen3.5-0.8B QLoRA 0.0463648662 0.0463648662 0.0361770093, Qwen3-1.7B QLoRA 0.0209256057 0.0209256057 0.0192451496, matching the RX 6500 XT. (Earlier note here said no build existed; that was true on WSL only.)

Reproduce the Triton buffer-ops problem in ten lines

import torch, triton, triton.language as tl
@triton.jit
def k(x_ptr, y_ptr, n, BLOCK: tl.constexpr):
    offs = tl.program_id(0) * BLOCK + tl.arange(0, BLOCK); m = offs < n
    tl.store(y_ptr + offs, tl.load(x_ptr + offs, mask = m) * 2, mask = m)
x = torch.arange(1024, device = "cuda", dtype = torch.float32); y = torch.full_like(x, -1)
k[(4,)](x, y, 1024, BLOCK = 256); torch.cuda.synchronize()
print((y == x * 2).sum().item(), "of 1024 correct")   # gfx1010: 0. with AMDGCN_USE_BUFFER_OPS=0: 1024

The generated ISA uses buffer_load_dwordx4 / buffer_store_dwordx4; gfx10.1 reads the buffer resource descriptor differently from gfx10.3, and the accesses go nowhere.