Repository navigation
b11847
Pre-release
hip: fix bf16 WMMA operand vector type for gfx11/gfx12 (#30885)
The __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32 (and _gfx12) builtins
take their A/B operands as vectors of short (_ExtVector<16, short> /
_ExtVector<8, short>), not __bf16. Declaring the fragments as
ext_vector_type(N) __bf16 fails to compile:
error: cannot initialize a parameter of type 'short
attribute((ext_vector_type(16)))' with an lvalue of type 'const
bf16x16_t' (vector of 16 '__bf16' values)
short and __bf16 are both 16-bit and the value is only ever passed to
the builtin via reinterpret_cast over raw bytes, so changing the element
type is bit-identical. The operand type has been short since the builtin
was introduced, so this works across LLVM/ROCm versions without a
version gate. Fixes HIP builds for RDNA3 (gfx1150/1151) and RDNA4
(gfx1200/1201).
Signed-off-by: Pablo Cholaky waltercool@slash.cl
Website:
Attestations:
macOS/iOS:
- macOS Apple Silicon (arm64)
- macOS Apple Silicon (arm64, KleidiAI enabled) DISABLED
- macOS Intel (x64)
- iOS XCFramework
Linux:
- Ubuntu x64 (CPU)
- Ubuntu arm64 (CPU)
- Ubuntu s390x (CPU)
- Ubuntu x64 (Vulkan)
- Ubuntu arm64 (Vulkan)
- Ubuntu x64 (CUDA 12) - CUDA 12.8 libraries
- Ubuntu x64 (CUDA 13) - CUDA 13.4 libraries
- Ubuntu arm64 (CUDA 13) - CUDA 13.4 libraries
- Ubuntu x64 (ROCm 10.0)
- Ubuntu x64 (OpenVINO)
- Ubuntu x64 (SYCL FP32)
- Ubuntu x64 (SYCL FP16)
- Linux arm64 (Snapdragon: CPU, Adreno GPU, Hexagon NPU) - setup guide
Android:
Windows:
- Windows x64 (CPU)
- Windows arm64 (CPU)
- Windows arm64 (OpenCL Adreno)
- Windows x64 (CUDA 12) - CUDA 12.4 DLLs
- Windows x64 (CUDA 13) - CUDA 13.4 DLLs
- Windows arm64 (CUDA 13) - CUDA 13.4 DLLs
- Windows x64 (Vulkan)
- Windows arm64 (Vulkan)
- Windows x64 (OpenVINO)
- Windows x64 (SYCL)
- Windows x64 (ROCm 10.0)
openEuler:
- DISABLED
- openEuler x86 (310p)
- openEuler x86 (910b, ACL Graph)
- openEuler aarch64 (310p)
- openEuler aarch64 (910b, ACL Graph)
UI: