[Mlir-commits] [llvm] [mlir] [NFC][LLVM] Define range attribute for `llvm.prefetch` args (PR #207300)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Fri Jul 3 06:17:35 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-mlir
Author: Rahul Joshi (jurahul)
<details>
<summary>Changes</summary>
Define valid range for `llvm.prefetch` immarg values and eliminate C++ verification code.
Additionally, change the range verification failure message to include the argument number and simplify formatting using `formatv`.
---
Patch is 28.20 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/207300.diff
13 Files Affected:
- (modified) llvm/include/llvm/IR/Intrinsics.td (+4-2)
- (modified) llvm/lib/IR/Verifier.cpp (+3-11)
- (modified) llvm/test/Assembler/invalid-immarg5.ll (+1-1)
- (modified) llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-g2s-invalid.ll (+3-5)
- (modified) llvm/test/CodeGen/NVPTX/fence-proxy-tensormap-invalid.ll (+2-2)
- (modified) llvm/test/CodeGen/NVPTX/tcgen05-mma-block-scale-invalid.ll (+8-11)
- (modified) llvm/test/CodeGen/NVPTX/tcgen05-mma-invalid.ll (+11-15)
- (modified) llvm/test/CodeGen/NVPTX/tcgen05-mma-scale-d-invalid.ll (+8-10)
- (modified) llvm/test/Verifier/AMDGPU/test-cvt-fp4f6f8-immarg-ranges.ll (+15-15)
- (modified) llvm/test/Verifier/NVPTX/fence-proxy.tensormap.ll (+4-4)
- (modified) llvm/test/Verifier/NVPTX/setmaxnreg.ll (+2-2)
- (modified) llvm/test/Verifier/intrinsic-immarg.ll (+12)
- (modified) mlir/test/Target/LLVMIR/llvmir-intrinsics.mlir (+1-1)
``````````diff
diff --git a/llvm/include/llvm/IR/Intrinsics.td b/llvm/include/llvm/IR/Intrinsics.td
index 0c6316ccf0a5c..f75335cd2c60a 100644
--- a/llvm/include/llvm/IR/Intrinsics.td
+++ b/llvm/include/llvm/IR/Intrinsics.td
@@ -1045,8 +1045,10 @@ def int_prefetch
[llvm_anyptr_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrInaccessibleMemOrArgMemOnly,
ReadOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>,
- ImmArg<ArgIndex<1>>, ImmArg<ArgIndex<2>>,
- ImmArg<ArgIndex<3>>]>;
+ ImmArg<ArgIndex<1>>, Range<ArgIndex<1>, 0, 2>,
+ ImmArg<ArgIndex<2>>, Range<ArgIndex<2>, 0, 4>,
+ ImmArg<ArgIndex<3>>, Range<ArgIndex<3>, 0, 2>
+ ]>;
def int_pcmarker : DefaultAttrsIntrinsic<[], [llvm_i32_ty]>;
def int_readcyclecounter : DefaultAttrsIntrinsic<[llvm_i64_ty]>;
diff --git a/llvm/lib/IR/Verifier.cpp b/llvm/lib/IR/Verifier.cpp
index 5f8b5e187bf14..f35e9066ab1f2 100644
--- a/llvm/lib/IR/Verifier.cpp
+++ b/llvm/lib/IR/Verifier.cpp
@@ -119,6 +119,7 @@
#include "llvm/Support/Casting.h"
#include "llvm/Support/CommandLine.h"
#include "llvm/Support/ErrorHandling.h"
+#include "llvm/Support/FormatVariadic.h"
#include "llvm/Support/MathExtras.h"
#include "llvm/Support/ModRef.h"
#include "llvm/Support/TimeProfiler.h"
@@ -3897,9 +3898,8 @@ void Verifier::visitCallBase(CallBase &Call) {
const ConstantRange &CR =
Call.getParamAttr(i, Attribute::Range).getValueAsConstantRange();
Check(CR.contains(CI->getValue()),
- "immarg value " + Twine(CI->getValue().getSExtValue()) +
- " out of range [" + Twine(CR.getLower().getSExtValue()) +
- ", " + Twine(CR.getUpper().getSExtValue()) + ")",
+ formatv("immarg value {} for arg {} out of range {}",
+ CI->getValue(), i, CR),
Call);
}
}
@@ -6289,14 +6289,6 @@ void Verifier::visitIntrinsicCall(Intrinsic::ID ID, CallBase &Call) {
"llvm.init_trampoline parameter #2 must resolve to a function.",
Call);
break;
- case Intrinsic::prefetch:
- Check(cast<ConstantInt>(Call.getArgOperand(1))->getZExtValue() < 2,
- "rw argument to llvm.prefetch must be 0-1", Call);
- Check(cast<ConstantInt>(Call.getArgOperand(2))->getZExtValue() < 4,
- "locality argument to llvm.prefetch must be 0-3", Call);
- Check(cast<ConstantInt>(Call.getArgOperand(3))->getZExtValue() < 2,
- "cache type argument to llvm.prefetch must be 0-1", Call);
- break;
case Intrinsic::reloc_none: {
Check(isa<MDString>(
cast<MetadataAsValue>(Call.getArgOperand(0))->getMetadata()),
diff --git a/llvm/test/Assembler/invalid-immarg5.ll b/llvm/test/Assembler/invalid-immarg5.ll
index deffbd3cad4cc..7c7bf11660e28 100644
--- a/llvm/test/Assembler/invalid-immarg5.ll
+++ b/llvm/test/Assembler/invalid-immarg5.ll
@@ -3,7 +3,7 @@
declare void @llvm.test.immarg.range.intrinsic.i32(i32 immarg range(i32 -3, 4))
define void @test_int_immarg_with_range() {
- ; CHECK: immarg value -4 out of range [-3, 4)
+ ; CHECK: immarg value -4 for arg 0 out of range [-3,4)
call void @llvm.test.immarg.range.intrinsic.i32(i32 -4)
ret void
}
diff --git a/llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-g2s-invalid.ll b/llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-g2s-invalid.ll
index 1c35fbead389e..0cef9c1d9405c 100644
--- a/llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-g2s-invalid.ll
+++ b/llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-g2s-invalid.ll
@@ -1,14 +1,12 @@
-; RUN: not llc < %s -mtriple=nvptx64 -mcpu=sm_100a -o /dev/null 2>&1 | FileCheck %s
-
-target triple = "nvptx64-nvidia-cuda"
+; RUN: not llvm-as -disable-output < %s 2>&1 | FileCheck %s
declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.1d(ptr addrspace(7) writeonly, ptr addrspace(3), ptr readonly, i32, i16, i64, i1 immarg, i1 immarg, i32 immarg range(i32 0, 3))
define void @test_cp_async_bulk_tensor_g2s_tile_1d(ptr addrspace(7) %d, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i16 %mc, i64 %ch) {
- ; CHECK: immarg value 3 out of range [0, 3)
+ ; CHECK: immarg value 3 for arg 8 out of range [0,3)
tail call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.1d(ptr addrspace(7) %d, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i16 %mc, i64 %ch, i1 0, i1 0, i32 3)
- ; CHECK: immarg value -1 out of range [0, 3)
+ ; CHECK: immarg value -1 for arg 8 out of range [0,3)
tail call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.1d(ptr addrspace(7) %d, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i16 %mc, i64 %ch, i1 0, i1 0, i32 -1)
ret void
diff --git a/llvm/test/CodeGen/NVPTX/fence-proxy-tensormap-invalid.ll b/llvm/test/CodeGen/NVPTX/fence-proxy-tensormap-invalid.ll
index ab35e4fb396d6..52f494ae590b6 100644
--- a/llvm/test/CodeGen/NVPTX/fence-proxy-tensormap-invalid.ll
+++ b/llvm/test/CodeGen/NVPTX/fence-proxy-tensormap-invalid.ll
@@ -1,7 +1,7 @@
-; RUN: not llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx83 -o /dev/null 2>&1 | FileCheck %s
+; RUN: not llvm-as -disable-output < %s 2>&1 | FileCheck %s
define void @test_fence_proxy_tensormap_generic_acquire(ptr addrspace(0) %addr) {
- ; CHECK: immarg value 130 out of range [128, 129)
+ ; CHECK: immarg value 130 for arg 1 out of range [128,129)
call void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.cta(ptr addrspace(0) %addr, i32 130);
ret void
diff --git a/llvm/test/CodeGen/NVPTX/tcgen05-mma-block-scale-invalid.ll b/llvm/test/CodeGen/NVPTX/tcgen05-mma-block-scale-invalid.ll
index c0f6f4c7c46bd..bd3ed2e447c0b 100644
--- a/llvm/test/CodeGen/NVPTX/tcgen05-mma-block-scale-invalid.ll
+++ b/llvm/test/CodeGen/NVPTX/tcgen05-mma-block-scale-invalid.ll
@@ -1,23 +1,20 @@
-; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
-; RUN: not llc < %s -o - -mcpu=sm_100a -march=nvptx64 -mattr=+ptx88 -o /dev/null 2>&1 | FileCheck %s
-; RUN: not llc < %s -o - -mcpu=sm_101a -march=nvptx64 -mattr=+ptx88 -o /dev/null 2>&1 | FileCheck %s
-; RUN: not llc < %s -o - -mcpu=sm_110a -march=nvptx64 -mattr=+ptx90 -o /dev/null 2>&1 | FileCheck %s
+; RUN: not llvm-as -disable-output < %s 2>&1 | FileCheck %s
define void @tcgen05_mma_block_scale_invalid_flags(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b) {
- ; CHECK: immarg value 0 out of range [1, 3)
+ ; CHECK: immarg value 0 for arg 7 out of range [1,3)
call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 0, i32 0)
- ; CHECK: immarg value 5 out of range [0, 4)
+ ; CHECK: immarg value 5 for arg 8 out of range [0,4)
call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 1, i32 5)
- ; CHECK: immarg value 0 out of range [1, 3)
+ ; CHECK: immarg value 0 for arg 7 out of range [1,3)
call void @llvm.nvvm.tcgen05.mma.shared.mxf4.block_scale.block32(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 0, i32 0)
- ; CHECK: immarg value 5 out of range [0, 4)
+ ; CHECK: immarg value 5 for arg 8 out of range [0,4)
call void @llvm.nvvm.tcgen05.mma.shared.mxf4.block_scale.block32(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 1, i32 5)
- ; CHECK: immarg value 0 out of range [1, 3)
+ ; CHECK: immarg value 0 for arg 7 out of range [1,3)
call void @llvm.nvvm.tcgen05.mma.shared.mxf4nvf4.block_scale.block32(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 0, i32 0)
- ; CHECK: immarg value 5 out of range [0, 4)
+ ; CHECK: immarg value 5 for arg 8 out of range [0,4)
call void @llvm.nvvm.tcgen05.mma.shared.mxf4nvf4.block_scale.block32(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 1, i32 5)
call void @llvm.nvvm.tcgen05.mma.shared.mxf4nvf4.block_scale.block16(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 0, i32 0)
- ; CHECK: immarg value 5 out of range [0, 4)
+ ; CHECK: immarg value 5 for arg 8 out of range [0,4)
call void @llvm.nvvm.tcgen05.mma.shared.mxf4nvf4.block_scale.block16(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 1, i32 5)
ret void
}
diff --git a/llvm/test/CodeGen/NVPTX/tcgen05-mma-invalid.ll b/llvm/test/CodeGen/NVPTX/tcgen05-mma-invalid.ll
index dff829ecf5321..c80755b6711d8 100644
--- a/llvm/test/CodeGen/NVPTX/tcgen05-mma-invalid.ll
+++ b/llvm/test/CodeGen/NVPTX/tcgen05-mma-invalid.ll
@@ -1,37 +1,33 @@
-; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
-; RUN: not llc < %s -mtriple=nvptx64 -mcpu=sm_100a -mattr=+ptx86 -o /dev/null 2>&1 | FileCheck %s
-; RUN: not llc < %s -mtriple=nvptx64 -mcpu=sm_101a -mattr=+ptx86 -o /dev/null 2>&1 | FileCheck %s
-; RUN: not llc < %s -mtriple=nvptx64 -mcpu=sm_110a -mattr=+ptx90 -o /dev/null 2>&1 | FileCheck %s
-target triple = "nvptx64-nvidia-cuda"
+; RUN: not llvm-as -disable-output < %s 2>&1 | FileCheck %s
define void @tcgen05_mma_invalid_flag_values(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d) {
- ; CHECK: immarg value 5 out of range [0, 4)
+ ; CHECK: immarg value 5 for arg 5 out of range [0,4)
call void @llvm.nvvm.tcgen05.mma.shared(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, i32 5, i32 1, i32 3)
- ; CHECK: immarg value 0 out of range [1, 3)
+ ; CHECK: immarg value 0 for arg 6 out of range [1,3)
call void @llvm.nvvm.tcgen05.mma.shared(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, i32 3, i32 0, i32 3)
- ; CHECK: immarg value 3 out of range [1, 3)
+ ; CHECK: immarg value 3 for arg 6 out of range [1,3)
call void @llvm.nvvm.tcgen05.mma.shared(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, i32 3, i32 3, i32 3)
- ; CHECK: immarg value 5 out of range [0, 4)
+ ; CHECK: immarg value 5 for arg 7 out of range [0,4)
call void @llvm.nvvm.tcgen05.mma.shared(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, i32 3, i32 2, i32 5)
- ; CHECK: immarg value 2 out of range [0, 2)
+ ; CHECK: immarg value 2 for arg 7 out of range [0,2)
call void @llvm.nvvm.tcgen05.mma.tensor.ashift(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, i32 3, i32 2, i32 2)
ret void
}
define void @tcgen05_mma_disable_output_lane_invalid_flag_values(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, <4 x i32> %disable_output_lanev4) {
- ; CHECK: immarg value 5 out of range [0, 4)
+ ; CHECK: immarg value 5 for arg 6 out of range [0,4)
call void @llvm.nvvm.tcgen05.mma.tensor.disable_output_lane.cg1(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, <4 x i32> %disable_output_lanev4, i32 5, i32 0)
- ; CHECK: immarg value 5 out of range [0, 4)
+ ; CHECK: immarg value 5 for arg 7 out of range [0,4)
call void @llvm.nvvm.tcgen05.mma.tensor.disable_output_lane.cg1(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, <4 x i32> %disable_output_lanev4, i32 0, i32 5)
ret void
}
define void @tcgen05_mma_ws_invalid_flag_values(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d) {
- ; CHECK: immarg value 5 out of range [0, 4)
+ ; CHECK: immarg value 5 for arg 5 out of range [0,4)
call void @llvm.nvvm.tcgen05.mma.ws.shared(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, i32 5, i32 0, i32 0)
- ; CHECK: immarg value 5 out of range [0, 4)
+ ; CHECK: immarg value 5 for arg 6 out of range [0,4)
call void @llvm.nvvm.tcgen05.mma.ws.shared(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, i32 0, i32 5, i32 0)
- ; CHECK: immarg value 5 out of range [0, 4)
+ ; CHECK: immarg value 5 for arg 7 out of range [0,4)
call void @llvm.nvvm.tcgen05.mma.ws.shared(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, i32 0, i32 0, i32 5)
ret void
}
diff --git a/llvm/test/CodeGen/NVPTX/tcgen05-mma-scale-d-invalid.ll b/llvm/test/CodeGen/NVPTX/tcgen05-mma-scale-d-invalid.ll
index 7c884a70b7530..5987668b10add 100644
--- a/llvm/test/CodeGen/NVPTX/tcgen05-mma-scale-d-invalid.ll
+++ b/llvm/test/CodeGen/NVPTX/tcgen05-mma-scale-d-invalid.ll
@@ -1,25 +1,23 @@
-; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
-; RUN: not llc < %s -mtriple=nvptx64 -mcpu=sm_100a -mattr=+ptx86 -o /dev/null 2>&1 | FileCheck %s
-target triple = "nvptx64-nvidia-cuda"
+; RUN: not llvm-as -disable-output < %s 2>&1 | FileCheck %s
define void @tcgen05_mma_scale_d_invalid_flag_values(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d) {
- ; CHECK: immarg value 16 out of range [0, 16)
+ ; CHECK: immarg value 16 for arg 5 out of range [0,16)
call void @llvm.nvvm.tcgen05.mma.shared.scale_d(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, i64 16, i32 0, i32 1, i32 0)
- ; CHECK: immarg value 3 out of range [0, 2)
+ ; CHECK: immarg value 3 for arg 6 out of range [0,2)
call void @llvm.nvvm.tcgen05.mma.shared.scale_d(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, i64 0, i32 3, i32 1, i32 0)
- ; CHECK: immarg value 0 out of range [1, 3)
+ ; CHECK: immarg value 0 for arg 7 out of range [1,3)
call void @llvm.nvvm.tcgen05.mma.shared.scale_d(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, i64 0, i32 0, i32 0, i32 0)
- ; CHECK: immarg value 5 out of range [0, 4)
+ ; CHECK: immarg value 5 for arg 8 out of range [0,4)
call void @llvm.nvvm.tcgen05.mma.shared.scale_d(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, i64 0, i32 0, i32 1, i32 5)
ret void
}
define void @tcgen05_mma_scale_d_disable_output_lane_invalid_flag_values(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, <4 x i32> %disable_output_lanev4) {
- ; CHECK: immarg value 16 out of range [0, 16)
+ ; CHECK: immarg value 16 for arg 5 out of range [0,16)
call void @llvm.nvvm.tcgen05.mma.tensor.scale_d.disable_output_lane.cg1(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, i64 16, <4 x i32> %disable_output_lanev4, i32 0, i32 0)
- ; CHECK: immarg value 3 out of range [0, 2)
+ ; CHECK: immarg value 3 for arg 7 out of range [0,2)
call void @llvm.nvvm.tcgen05.mma.tensor.scale_d.disable_output_lane.cg1(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, i64 0, <4 x i32> %disable_output_lanev4, i32 3, i32 0)
- ; CHECK: immarg value 5 out of range [0, 4)
+ ; CHECK: immarg value 5 for arg 8 out of range [0,4)
call void @llvm.nvvm.tcgen05.mma.tensor.scale_d.disable_output_lane.cg1(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, i64 0, <4 x i32> %disable_output_lanev4, i32 0, i32 5)
ret void
}
diff --git a/llvm/test/Verifier/AMDGPU/test-cvt-fp4f6f8-immarg-ranges.ll b/llvm/test/Verifier/AMDGPU/test-cvt-fp4f6f8-immarg-ranges.ll
index 4e8022711320d..98c73444ed62f 100644
--- a/llvm/test/Verifier/AMDGPU/test-cvt-fp4f6f8-immarg-ranges.ll
+++ b/llvm/test/Verifier/AMDGPU/test-cvt-fp4f6f8-immarg-ranges.ll
@@ -4,14 +4,14 @@
; llvm.amdgcn.cvt.sr.fp8.f16 - byte_sel out of range
; --------------------------------------------------------------------
-; CHECK: immarg value 4 out of range [0, 4)
+; CHECK: immarg value 4 for arg 3 out of range [0,4)
; CHECK-NEXT: %result = call i32 @llvm.amdgcn.cvt.sr.fp8.f16(half %src, i32 %seed, i32 %old, i32 4)
define i32 @test_cvt_sr_fp8_f16_byte_sel_out_of_range(half %src, i32 %seed, i32 %old) {
%result = call i32 @llvm.amdgcn.cvt.sr.fp8.f16(half %src, i32 %seed, i32 %old, i32 4)
ret i32 %result
}
-; CHECK: immarg value 10 out of range [0, 4)
+; CHECK: immarg value 10 for arg 3 out of range [0,4)
; CHECK-NEXT: %result = call i32 @llvm.amdgcn.cvt.sr.fp8.f16(half %src, i32 %seed, i32 %old, i32 10)
define i32 @test_cvt_sr_fp8_f16_byte_sel_way_out_of_range(half %src, i32 %seed, i32 %old) {
%result = call i32 @llvm.amdgcn.cvt.sr.fp8.f16(half %src, i32 %seed, i32 %old, i32 10)
@@ -22,7 +22,7 @@ define i32 @test_cvt_sr_fp8_f16_byte_sel_way_out_of_range(half %src, i32 %seed,
; llvm.amdgcn.cvt.sr.bf8.f16 - byte_sel out of range
; --------------------------------------------------------------------
-; CHECK: immarg value 4 out of range [0, 4)
+; CHECK: immarg value 4 for arg 3 out of range [0,4)
; CHECK-NEXT: %result = call i32 @llvm.amdgcn.cvt.sr.bf8.f16(half %src, i32 %seed, i32 %old, i32 4)
define i32 @test_cvt_sr_bf8_f16_byte_sel_out_of_range(half %src, i32 %seed, i32 %old) {
%result = call i32 @llvm.amdgcn.cvt.sr.bf8.f16(half %src, i32 %seed, i32 %old, i32 4)
@@ -33,14 +33,14 @@ define i32 @test_cvt_sr_bf8_f16_byte_sel_out_of_range(half %src, i32 %seed, i32
; llvm.amdgcn.cvt.scale.pk8.f16.fp8 - scale_sel out of range
; --------------------------------------------------------------------
-; CHECK: immarg value 16 out of range [0, 16)
+; CHECK: immarg value 16 for arg 2 out of range [0,16)
; CHECK-NEXT: %result = call <8 x half> @llvm.amdgcn.cvt.scale.pk8.f16.fp8(<2 x i32> %src, i32 0, i32 16)
define <8 x half> @test_cvt_scale_pk8_f16_fp8_scale_sel_out_of_range(<2 x i32> %src) {
%result = call <8 x half> @llvm.amdgcn.cvt.scale.pk8.f16.fp8(<2 x i32> %src, i32 0, i32 16)
ret <8 x half> %result
}
-; CHECK: immarg value 100 out of range [0, 16)
+; CHECK: immarg value 100 for arg 2 out of range [0,16)
; CHECK-NEXT: %result = call <8 x half> @llvm.amdgcn.cvt.scale.pk8.f16.fp8(<2 x i32> %src, i32 0, i32 100)
define <8 x half> @test_cvt_scale_pk8_f16_fp8_scale_sel_way_out_of_range(<2 x i32> %src) {
%result = call <8 x half> @llvm.amdgcn.cvt.scale.pk8.f16.fp8(<2 x i32> %src, i32 0, i32 100)
@@ -51,14 +51,14 @@ define <8 x half> @test_cvt_scale_pk8_f16_fp8_scale_sel_way_out_of_range(<2 x i3
; llvm.amdgcn.cvt.scalef32.f32.fp8 - src_sel out of range
; --------------------------------------------------------------------
-; CHECK: immarg value 4 out of range [0, 4)
+; CHECK: immarg value 4 for arg 2 out of range [0,4)
; CHECK-NEXT: %result = call float @llvm.amdgcn.cvt.scalef32.f32.fp8(i32 %src, float %scale, i32 4)
define float @test_cvt_scalef32_f32_fp8_src_sel_out_of_range(i32 %src, float %scale) {
%result = call float @llvm.amdgcn.cvt.scalef32.f32.fp8(i32 %src, float %scale, i32 4)
ret float %result
}
-; CHECK: immarg value 7 out of range [0, 4)
+; CHECK: immarg value 7 for arg 2 out of range [0,4)
; CHECK-NEXT: %result = call float @llvm.amdgcn.cvt.scalef32.f32.fp8(i32 %src, float %scale, i32 7)
define float @test_cvt_scalef32_f32_fp8_src_sel_way_out_of_range(i32 %src, float %scale) {
%result = call float @llvm.amdgcn.cvt.scalef32.f32.fp8(i32 %src, float %scale, i32 7)
@@ -69,14 +69,14 @@ define float @test_cvt_scalef32_f32_fp8_src_sel_way_out_of_range(i32 %src, float
; llvm.amdgcn.cvt.scalef32.f16.fp8 - src_sel_index out of range
; --------------------------------------------------------------------
-; CHECK: immarg value 4 out of range [0, 4)
+; CHECK: immarg value 4 for arg 3 out of range [0,4)
; CHECK-NEXT: %result = call <2 x half> @llvm.amdgcn.cvt.scalef32.f16.fp8(<2 x half> %old, i32 %src, float %scale, i32 4, i1 false)
define <2 x half> @test_cvt_scalef32_f16_fp8_src_sel_index_out_of_range(<2 x half> %old, i32 %src, float %scale) {
%result = call <2 x half> @llvm.amdgcn.cvt....
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/207300
More information about the Mlir-commits
mailing list