Fix global dequantize buffer dtype mismatch when loading multiple 4-bit models with different dtypes in the same process. Adds dtype check alongside existing None check for WEIGHT_BUFFER in both CUDA/HIP and XPU paths.
Use 16 warps for RDNA in the chunked cross-entropy forward kernel
(large vocab > 65536), matching the existing CDNA optimization.
Benchmarked on W7900 (gfx1100) with actual unsloth kernels (5 trials, median):
- Chunked CE forward (BS=65536): 16 warps = 2.4-2.6x faster than 32
- All other kernels (LayerNorm, RoPE, SwiGLU): default heuristic is
already optimal for RDNA; no modification needed.
Depends on: #4109 (provides is_rdna() detection)
TMA (Tensor Memory Accelerator) is an NVIDIA Hopper+ feature that does
not exist on AMD GPUs. However, _check_tma_support() incorrectly
returns True on ROCm because:
1. torch.cuda.get_device_capability() returns (11, 0) for gfx1100,
satisfying the >= 9 check intended for Hopper (sm_90).
2. ROCm Triton exports tl.make_tensor_descriptor (the symbol exists
even though the hardware does not support TMA).
This would cause MoE grouped_gemm to attempt TMA operations on AMD
GPUs, leading to runtime failures.
Fix: early-return False for HIP devices, matching the existing XPU
guard.
* fix(Triton): ensure float32 eps in RMS LayerNorm rsqrt for HIP/ROCm
On HIP (AMD ROCm), Triton constexpr eps may not promote to float32
in rsqrt, causing numerical instability (NaN/Inf) on RDNA GPUs
(gfx1100, gfx1151 Strix Halo, etc.).
Use tl.full((), eps, tl.float32) to explicitly create a float32
scalar before adding to row_var in rsqrt. Applied to both standard
and Gemma RMS LayerNorm forward kernels.
Tested on W7900 (gfx1100): full test suite passed (dim 512-2048,
bf16/fp16, various seqlen).
Related: #3385, #3588
* Apply same float32 eps fix to layernorm.py for PR #4110
layernorm.py has the identical tl.constexpr eps pattern in
layernorm_forward that can misfire on HIP/ROCm. Apply the same
tl.full((), eps, tl.float32) fix for consistency.
Both testing_suite_layernorm (standard LayerNorm) and
testing_suite_layernorm (RMS LayerNorm) pass on NVIDIA after
this change.
---------
Co-authored-by: Daniel Han <danielhanchen@gmail.com>
* fix(ROCm): comprehensive RDNA GPU support - fix Gemma3 NaN & add is_rdna()
- Add is_rdna() detection for RDNA3/3.5/RDNA4 consumer GPUs (gfx11xx, gfx1151, gfx12xx)
- Disable torch.compile for Gemma3 on HIP to fix NaN loss (fixes#3385, #4029)
- Export is_cdna/is_rdna from kernels for downstream use
- Import is_rdna into cross_entropy_loss for future RDNA-specific tuning
Tested on AMD Radeon PRO W7900 (gfx1100) with ROCm 7.1:
✓ Gemma3-1B: loss 3.37→3.25 (no NaN)
✓ Llama-3.2-1B: loss 2.44→2.37 (no NaN)
✓ Qwen2.5-1.5B: loss 1.89→1.85 (no NaN)
✓ RMS LayerNorm Triton kernel: bf16/fp16 PASSED
✓ Cross Entropy Loss Triton kernel: 32K/256K vocab PASSED
* [pre-commit.ci] auto fixes from pre-commit.com hooks
for more information, see https://pre-commit.ci
* Address review: scope compile disable to RDNA only, use partial mode, remove unused import
Changes based on Daniel's review:
1. (HIGH) Replace DEVICE_TYPE=='hip' with is_rdna() to avoid disabling
torch.compile on CDNA GPUs (MI250X/MI300X/MI350) where it works fine
2. (MEDIUM) Use 'partial' instead of '1' for UNSLOTH_COMPILE_DISABLE to
only disable model forward compilation while keeping loss compilation,
matching the existing Sesame pattern
3. (LOW) Remove unused is_rdna import from cross_entropy_loss.py (F401)
* Remove redundant is_cdna/is_rdna exports from kernels/__init__.py
These functions are imported directly from .utils where needed
(e.g. cross_entropy_loss.py, loader.py). No external code imports
them from the unsloth.kernels namespace.
* [pre-commit.ci] auto fixes from pre-commit.com hooks
for more information, see https://pre-commit.ci
---------
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
The function (introduced in #3923) assumed that the absence of
`triton.runtime.triton_key` on ROCm means torch.compile will crash.
Investigation shows this is incorrect:
1. `triton.runtime.triton_key` was renamed/removed in the ROCm Triton
fork — it does not exist at that path. However,
`triton.compiler.compiler.triton_key` (the path torch._inductor
actually imports) EXISTS and works correctly on ROCm.
2. Both call-sites in torch._inductor (codecache.py and
async_compile.py) already wrap the import in try/except, so even a
genuinely missing triton_key would be handled gracefully.
3. Comprehensive testing on ROCm 7.1 + Triton 3.4.0 + gfx1100 confirms
torch.compile works correctly for matmul, cross-entropy, RMSNorm,
multi-layer transformer forward+backward, and LoRA — all without
triton.runtime.triton_key.
The original code was also ineffective (environment variables set after
torch import have no effect on torch._dynamo config), so removing it
has zero behavioral change on existing installations.
Supersedes the compile-disable portion of #3923.
