Summary
IREE's NVIDIA mma.sync lowering reshapes every lhs fragment to
vector<2x2x2>, which is valid for eight-element f16/bf16 fragments. FP8
m16n8k32 lhs fragments contain 16 elements, so the default CUDA heuristic
selects a supported intrinsic and then fails verification during lowering.
Reproducer
func.func @case(%lhs: tensor<256x256xf8E4M3FN>,
%rhs: tensor<256x256xf8E4M3FN>) -> tensor<256x256xf32> {
%z = arith.constant 0.0 : f32
%e = tensor.empty() : tensor<256x256xf32>
%i = linalg.fill ins(%z : f32) outs(%e : tensor<256x256xf32>) -> tensor<256x256xf32>
%r = linalg.matmul ins(%lhs,%rhs : tensor<256x256xf8E4M3FN>,tensor<256x256xf8E4M3FN>)
outs(%i : tensor<256x256xf32>) -> tensor<256x256xf32>
return %r : tensor<256x256xf32>
}Verified with iree-compile at a5ca3e2beaa8 (IREE compiler version
unknown, LLVM 24.0.0git, optimized with assertions):
$ iree-compile --iree-hal-target-device=cuda --iree-cuda-target=sm_89 repro.mlir -o /dev/null 2>&1 | sed -n 's/^.*error: /error: /p' | head -1
error: 'vector.shape_cast' op has different number of elements at source (16) and result (8)
Expected: the selected NV_MMA_SYNC_F32_16x8x32_F8E4M3FN intrinsic lowers
successfully.
Actual: compilation stops on a 16-to-8-element shape cast.
Agent Analysis
At a5ca3e2beaa8, compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/ IREEGPUAttrs.cpp:1262-1275 constructs VectorType::get({2,2,2}, elementType)
for every lhs. Derive the last extent from the 16-element FP8 fragment (or its
single-subgroup layout) instead of hard-coding two.
Related: issue #24625 is an elementwise FP8 LLVM-translation failure and
states that no tensor-core path was available before #24659; this is a later,
different failure inside the newly selected mma.sync lowering.
AI tool use disclosure
Adding this to comply with the project's
AI tool use policy
Assisted-by: Claude Code
Reproducer and root-cause analysis were prepared with LLM assistance and verified manually on the referenced commit.
Summary
IREE's NVIDIA
mma.synclowering reshapes every lhs fragment tovector<2x2x2>, which is valid for eight-element f16/bf16 fragments. FP8m16n8k32lhs fragments contain 16 elements, so the default CUDA heuristicselects a supported intrinsic and then fails verification during lowering.
Reproducer
Verified with
iree-compileata5ca3e2beaa8(IREE compiler versionunknown, LLVM 24.0.0git, optimized with assertions):Expected: the selected
NV_MMA_SYNC_F32_16x8x32_F8E4M3FNintrinsic lowerssuccessfully.
Actual: compilation stops on a 16-to-8-element shape cast.
Agent Analysis
At
a5ca3e2beaa8,compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/ IREEGPUAttrs.cpp:1262-1275constructsVectorType::get({2,2,2}, elementType)for every lhs. Derive the last extent from the 16-element FP8 fragment (or its
single-subgroup layout) instead of hard-coding two.
Related: issue #24625 is an elementwise FP8 LLVM-translation failure and
states that no tensor-core path was available before #24659; this is a later,
different failure inside the newly selected mma.sync lowering.
AI tool use disclosure
Adding this to comply with the project's
AI tool use policy
Assisted-by: Claude Code
Reproducer and root-cause analysis were prepared with LLM assistance and verified manually on the referenced commit.