diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index 91d72d5ef9dfc..03488f8389aa2 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -653,19 +653,20 @@ def int_amdgcn_cvt_pk_bf8_f16 // byte_sel selects byte to write in vdst. def int_amdgcn_cvt_sr_fp8_f16 : DefaultAttrsIntrinsic< [llvm_i32_ty], [llvm_half_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], - [IntrNoMem, IntrSpeculatable, ImmArg>] + [IntrNoMem, IntrSpeculatable, ImmArg>, Range, 0, 4>] >, ClangBuiltin<"__builtin_amdgcn_cvt_sr_fp8_f16">; // llvm.amdgcn.cvt.sr.bf8.f16 i32 vdst, half src, i32 seed, i32 old, imm byte_sel [0..3] // byte_sel selects byte to write in vdst. def int_amdgcn_cvt_sr_bf8_f16 : DefaultAttrsIntrinsic< [llvm_i32_ty], [llvm_half_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], - [IntrNoMem, IntrSpeculatable, ImmArg>] + [IntrNoMem, IntrSpeculatable, ImmArg>, Range, 0, 4>] >, ClangBuiltin<"__builtin_amdgcn_cvt_sr_bf8_f16">; // llvm.amdgcn.cvt.scale.pk32.f16.bf6 v32f16 vdst, v6i32 src0, i32 scale_sel [0..15] class AMDGPUCvtScaleIntrinsic : DefaultAttrsIntrinsic< - [DstTy], [Src0Ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrSpeculatable, ImmArg>] + [DstTy], [Src0Ty, llvm_i32_ty, llvm_i32_ty], + [IntrNoMem, IntrSpeculatable, ImmArg>, Range, 0, 16>] >, ClangBuiltin<"__builtin_amdgcn_"#name>; class AMDGPUCvtScaleF32Intrinsic : DefaultAttrsIntrinsic< @@ -746,7 +747,8 @@ class AMDGPUCvtScaleFP4FP8BF8ToF1632Intrinsic : Def [llvm_i32_ty, // src llvm_float_ty, // scale llvm_i32_ty], // src_sel index [0..3] - [IntrNoMem, ImmArg>] + [IntrNoMem, + ImmArg>, Range, 0, 4>] >, ClangBuiltin<"__builtin_amdgcn_"#name>; class AMDGPUCvtScale_pk_FromFP8BF8Intrinsic : DefaultAttrsIntrinsic< @@ -783,7 +785,7 @@ class AMDGPUCvtScaleFP8BF8ToF16TiedInputIntrinsic : llvm_float_ty, // scale llvm_i32_ty, // src_sel_index[0..3] llvm_i1_ty], // dst_lo_hi_sel[true false] - [IntrNoMem, ImmArg>, ImmArg>] + [IntrNoMem, ImmArg>, Range, 0, 4>, ImmArg>] >, ClangBuiltin<"__builtin_amdgcn_"#name>; class AMDGPUCvtScaleF32ToFP4Intrinsic : DefaultAttrsIntrinsic< @@ -793,7 +795,7 @@ class AMDGPUCvtScaleF32ToFP4Intrinsic : DefaultAttrsIntrinsic< llvm_float_ty, // src1 llvm_float_ty, // scale llvm_i32_ty], // dst_sel_index[0..3] - [IntrNoMem, ImmArg>] + [IntrNoMem, ImmArg>, Range, 0, 4>] >, ClangBuiltin<"__builtin_amdgcn_"#name>; class AMDGPUCvtScaleF16ToFP4TiedInputIntrinsic : DefaultAttrsIntrinsic< @@ -802,7 +804,7 @@ class AMDGPUCvtScaleF16ToFP4TiedInputIntrinsic : De SrcTy, // src llvm_float_ty, // scale llvm_i32_ty], // dest_sel_index [0..3] - [IntrNoMem, ImmArg>] + [IntrNoMem, ImmArg>, Range, 0, 4>] >, ClangBuiltin<"__builtin_amdgcn_"#name>; class AMDGPUCvtScaleBF16F16F32SRToFP4BF8F8TiedInputIntrinsic : DefaultAttrsIntrinsic< @@ -812,7 +814,7 @@ class AMDGPUCvtScaleBF16F16F32SRToFP4BF8F8TiedInputIntrinsic>] + [IntrNoMem, ImmArg>, Range, 0, 4>] >, ClangBuiltin<"__builtin_amdgcn_"#name>; class AMDGPUCvtScaleSRF32ToBF16F16TiedInputIntrinsic : DefaultAttrsIntrinsic<