[Codegen][LLVMGPU] Lower FP8 elementwise extensions to NVVM

[weimin023]

Summary

Adds a CUDA/NVVM lowering for general FP8-to-FP32 elementwise conversion.

This complements the existing FP8 Tensor Core MMA path.
The earlier work lowers supported FP8 matrix contractions through nvgpu.mma.sync / nvvm.mma.sync;
this change handles non-MMA conversion operations by decoding FP8 storage bits with regular LLVM integer operations and bitcasting the reconstructed IEEE-754 FP32 bits.

Fixes #24625.

Implementation

The failing tensor-level stablehlo.convert is lowered to scalar/vector arith.extf before the CUDA LLVM/NVVM conversion boundary. The custom pattern therefore matches only:

arith.extf : f8E4M3FN/f8E5M2 -> f32

At this boundary, LLVMTypeConverter represents the FP8 source operand as i8 storage (or vector). The lowering:

1. extracts the sign, exponent, and mantissa fields from the i8 bits;
2. constructs IEEE-754 FP32 bits with LLVM integer operations;
3. handles zero, subnormal, normal, NaN, and E5M2 infinity encodings; and
4. uses llvm.bitcast to reinterpret the resulting i32 bits as f32.

The specialized pattern has PatternBenefit(2) and is registered before the generic Arith-to-LLVM patterns so that FP8 extensions use this software decode path. The implementation supports both scalar and vector operands.

This does not add arith.truncf (f32 -> f8) lowering or general FP8 arithmetic.

• CUDA FP8 support status

Path	Status	Target requirements
FP8 Tensor Core matrix multiply	Existing mma.sync lowering path	FP8 Tensor Core capable targets, e.g. sm_89 / sm_120
General FP8 elementwise f8E4M3FN/f8E5M2 -> f32	Added by this PR	CUDA Core software decode; does not require FP8 Tensor Cores
f32 -> FP8 conversion	Not included	Requires separate arith.truncf lowering with defined rounding, overflow, subnormal, and NaN semantics

• Explanation of createFP8ToF32Bits

The decode rules below reflect the two formats handled by this PR.
In particular, E4M3FN is finite-number-only and does not encode infinity.

Exponent	Mantissa	E4M3FN	E5M2
0	0	signed zero	signed zero
0	nonzero	subnormal	subnormal
1 .. max-1	any	normal	normal
max	0	normal finite value	infinity
max	1 .. max-1	normal finite value	NaN
max	max	NaN	NaN

Previous failure

A StableHLO conversion (from #24625):

func.func @c(%a: tensor<1024xf8E4M3FN>) -> tensor<1024xf32> {
  %0 = stablehlo.convert %a : (tensor<1024xf8E4M3FN>) -> tensor<1024xf32>
  return %0 : tensor<1024xf32>
}

iree-compile convert_f8_to_f32.mlir \
  --iree-hal-target-backends=cuda --iree-cuda-target=sm_86 -o /dev/null

eventually reached CUDA final lowering as:

%14 = "builtin.unrealized_conversion_cast"(%13)
    : (vector<1xi8>) -> vector<1xf8E4M3FN>
%15 = "arith.extf"(%14)
    : (vector<1xf8E4M3FN>) -> vector<1xf32>

and failed during LLVM translation:

error: LLVM Translation failed for operation:
builtin.unrealized_conversion_cast

Testing

./tools/iree-compile convert_f8_to_f32.mlir \
  --iree-hal-target-backends=cuda \
  --iree-cuda-target=sm_86 \
  -o /dev/null

./tools/iree-compile convert_f8_to_f32.mlir \
  --iree-hal-target-backends=cuda \
  --iree-cuda-target=sm_89 \
  -o /dev/null

./tools/iree-compile convert_f8_to_f32.mlir \
  --iree-hal-target-backends=cuda \
  --iree-cuda-target=sm_120 \
  --iree-cuda-target-features=+ptx87 \
  -o /dev/null

[krzysz00]

This already exists - it's one of the clone() methods. Or setElementType or the like, I'm pretty sure

[weimin023]

@krzysz00
Oh thanks! I inlined the scalar/vector type handling at the only use site and use VectorType::cloneWith for vectors, so the local helper is no longer needed.