[Mlir-commits] [mlir] 780298c - [NFC][LLVM] Define range attribute for `llvm.prefetch` args (#207300)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Fri Jul 3 08:25:49 PDT 2026
Author: Rahul Joshi
Date: 2026-07-03T08:25:42-07:00
New Revision: 780298c4424514f0205df5e82c2a135717e6a723
URL: https://github.com/llvm/llvm-project/commit/780298c4424514f0205df5e82c2a135717e6a723
DIFF: https://github.com/llvm/llvm-project/commit/780298c4424514f0205df5e82c2a135717e6a723.diff
LOG: [NFC][LLVM] Define range attribute for `llvm.prefetch` args (#207300)
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`.
Added:
Modified:
llvm/include/llvm/IR/Intrinsics.td
llvm/lib/IR/Verifier.cpp
llvm/test/Assembler/invalid-immarg5.ll
llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-g2s-invalid.ll
llvm/test/CodeGen/NVPTX/fence-proxy-tensormap-invalid.ll
llvm/test/CodeGen/NVPTX/tcgen05-mma-block-scale-invalid.ll
llvm/test/CodeGen/NVPTX/tcgen05-mma-invalid.ll
llvm/test/CodeGen/NVPTX/tcgen05-mma-scale-d-invalid.ll
llvm/test/Verifier/AMDGPU/test-cvt-fp4f6f8-immarg-ranges.ll
llvm/test/Verifier/NVPTX/fence-proxy.tensormap.ll
llvm/test/Verifier/NVPTX/setmaxnreg.ll
llvm/test/Verifier/intrinsic-immarg.ll
mlir/test/Target/LLVMIR/llvmir-intrinsics.mlir
Removed:
################################################################################
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.scalef32.f16.fp8(<2 x half> %old, i32 %src, float %scale, i32 4, i1 false)
ret <2 x half> %result
}
-; CHECK: immarg value 15 out of range [0, 4)
+; CHECK: immarg value 15 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 15, i1 true)
define <2 x half> @test_cvt_scalef32_f16_fp8_src_sel_index_way_out_of_range(<2 x half> %old, i32 %src, float %scale) {
%result = call <2 x half> @llvm.amdgcn.cvt.scalef32.f16.fp8(<2 x half> %old, i32 %src, float %scale, i32 15, i1 true)
@@ -87,14 +87,14 @@ define <2 x half> @test_cvt_scalef32_f16_fp8_src_sel_index_way_out_of_range(<2 x
; llvm.amdgcn.cvt.scalef32.pk.fp4.f32 - dst_sel_index out of range
; --------------------------------------------------------------------
-; CHECK: immarg value 4 out of range [0, 4)
+; CHECK: immarg value 4 for arg 4 out of range [0,4)
; CHECK-NEXT: %result = call i32 @llvm.amdgcn.cvt.scalef32.pk.fp4.f32(i32 %old, float %src0, float %src1, float %scale, i32 4)
define i32 @test_cvt_scalef32_pk_fp4_f32_dst_sel_index_out_of_range(i32 %old, float %src0, float %src1, float %scale) {
%result = call i32 @llvm.amdgcn.cvt.scalef32.pk.fp4.f32(i32 %old, float %src0, float %src1, float %scale, i32 4)
ret i32 %result
}
-; CHECK: immarg value 8 out of range [0, 4)
+; CHECK: immarg value 8 for arg 4 out of range [0,4)
; CHECK-NEXT: %result = call i32 @llvm.amdgcn.cvt.scalef32.pk.fp4.f32(i32 %old, float %src0, float %src1, float %scale, i32 8)
define i32 @test_cvt_scalef32_pk_fp4_f32_dst_sel_index_way_out_of_range(i32 %old, float %src0, float %src1, float %scale) {
%result = call i32 @llvm.amdgcn.cvt.scalef32.pk.fp4.f32(i32 %old, float %src0, float %src1, float %scale, i32 8)
@@ -105,14 +105,14 @@ define i32 @test_cvt_scalef32_pk_fp4_f32_dst_sel_index_way_out_of_range(i32 %old
; llvm.amdgcn.cvt.scalef32.pk.fp4.f16 - dest_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 i32 @llvm.amdgcn.cvt.scalef32.pk.fp4.f16(i32 %old, <2 x half> %src, float %scale, i32 4)
define i32 @test_cvt_scalef32_pk_fp4_f16_dest_sel_index_out_of_range(i32 %old, <2 x half> %src, float %scale) {
%result = call i32 @llvm.amdgcn.cvt.scalef32.pk.fp4.f16(i32 %old, <2 x half> %src, float %scale, i32 4)
ret i32 %result
}
-; CHECK: immarg value 12 out of range [0, 4)
+; CHECK: immarg value 12 for arg 3 out of range [0,4)
; CHECK-NEXT: %result = call i32 @llvm.amdgcn.cvt.scalef32.pk.fp4.f16(i32 %old, <2 x half> %src, float %scale, i32 12)
define i32 @test_cvt_scalef32_pk_fp4_f16_dest_sel_index_way_out_of_range(i32 %old, <2 x half> %src, float %scale) {
%result = call i32 @llvm.amdgcn.cvt.scalef32.pk.fp4.f16(i32 %old, <2 x half> %src, float %scale, i32 12)
@@ -123,14 +123,14 @@ define i32 @test_cvt_scalef32_pk_fp4_f16_dest_sel_index_way_out_of_range(i32 %ol
; llvm.amdgcn.cvt.scalef32.sr.pk.fp4.f16 - dst_sel_index out of range
; --------------------------------------------------------------------
-; CHECK: immarg value 4 out of range [0, 4)
+; CHECK: immarg value 4 for arg 4 out of range [0,4)
; CHECK-NEXT: %result = call i32 @llvm.amdgcn.cvt.scalef32.sr.pk.fp4.f16(i32 %old, <2 x half> %src, i32 %seed, float %scale, i32 4)
define i32 @test_cvt_scalef32_sr_pk_fp4_f16_dst_sel_index_out_of_range(i32 %old, <2 x half> %src, i32 %seed, float %scale) {
%result = call i32 @llvm.amdgcn.cvt.scalef32.sr.pk.fp4.f16(i32 %old, <2 x half> %src, i32 %seed, float %scale, i32 4)
ret i32 %result
}
-; CHECK: immarg value 9 out of range [0, 4)
+; CHECK: immarg value 9 for arg 4 out of range [0,4)
; CHECK-NEXT: %result = call i32 @llvm.amdgcn.cvt.scalef32.sr.pk.fp4.f16(i32 %old, <2 x half> %src, i32 %seed, float %scale, i32 9)
define i32 @test_cvt_scalef32_sr_pk_fp4_f16_dst_sel_index_way_out_of_range(i32 %old, <2 x half> %src, i32 %seed, float %scale) {
%result = call i32 @llvm.amdgcn.cvt.scalef32.sr.pk.fp4.f16(i32 %old, <2 x half> %src, i32 %seed, float %scale, i32 9)
diff --git a/llvm/test/Verifier/NVPTX/fence-proxy.tensormap.ll b/llvm/test/Verifier/NVPTX/fence-proxy.tensormap.ll
index 4fa7a7ae71001..652990a700a6e 100644
--- a/llvm/test/Verifier/NVPTX/fence-proxy.tensormap.ll
+++ b/llvm/test/Verifier/NVPTX/fence-proxy.tensormap.ll
@@ -1,16 +1,16 @@
; RUN: not llvm-as %s -o /dev/null 2>&1 | FileCheck %s
define void @test_fence_proxy_tensormap_generic_acquire(ptr addrspace(0) %addr) {
- ; CHECK: immarg value 127 out of range [128, 129)
+ ; CHECK: immarg value 127 for arg 1 out of range [128,129)
call void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.cta(ptr addrspace(0) %addr, i32 127);
- ; CHECK: immarg value 129 out of range [128, 129)
+ ; CHECK: immarg value 129 for arg 1 out of range [128,129)
call void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.cluster(ptr addrspace(0) %addr, i32 129);
- ; CHECK: immarg value 127 out of range [128, 129)
+ ; CHECK: immarg value 127 for arg 1 out of range [128,129)
call void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.gpu(ptr addrspace(0) %addr, i32 127);
- ; CHECK: immarg value 129 out of range [128, 129)
+ ; CHECK: immarg value 129 for arg 1 out of range [128,129)
call void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.sys(ptr addrspace(0) %addr, i32 129);
ret void
diff --git a/llvm/test/Verifier/NVPTX/setmaxnreg.ll b/llvm/test/Verifier/NVPTX/setmaxnreg.ll
index 1afebeab4742c..79714a0521c36 100644
--- a/llvm/test/Verifier/NVPTX/setmaxnreg.ll
+++ b/llvm/test/Verifier/NVPTX/setmaxnreg.ll
@@ -7,10 +7,10 @@ define void @test_set_maxn_reg() {
; CHECK: reg_count argument to nvvm.setmaxnreg must be in multiples of 8
call void @llvm.nvvm.setmaxnreg.inc.sync.aligned.u32(i32 95)
- ; CHECK: immarg value 16 out of range [24, 257)
+ ; CHECK: immarg value 16 for arg 0 out of range [24,257)
call void @llvm.nvvm.setmaxnreg.dec.sync.aligned.u32(i32 16)
- ; CHECK: immarg value 264 out of range [24, 257)
+ ; CHECK: immarg value 264 for arg 0 out of range [24,257)
call void @llvm.nvvm.setmaxnreg.dec.sync.aligned.u32(i32 264)
ret void
}
diff --git a/llvm/test/Verifier/intrinsic-immarg.ll b/llvm/test/Verifier/intrinsic-immarg.ll
index c95c7214fde71..39ba3cb480b4d 100644
--- a/llvm/test/Verifier/intrinsic-immarg.ll
+++ b/llvm/test/Verifier/intrinsic-immarg.ll
@@ -154,6 +154,18 @@ define void @test_prefetch(ptr %ptr, i32 %arg0, i32 %arg1) {
; CHECK-NEXT: i32 %arg1
call void @llvm.prefetch(ptr %ptr, i32 %arg0, i32 0, i32 0)
call void @llvm.prefetch(ptr %ptr, i32 0, i32 %arg1, i32 0)
+
+ ; CHECK: immarg value 2 for arg 1 out of range [0,2)
+ ; CHECK-NEXT: call void @llvm.prefetch.p0(ptr %ptr, i32 2, i32 0, i32 0)
+ call void @llvm.prefetch(ptr %ptr, i32 2, i32 0, i32 0)
+
+ ; CHECK: immarg value 8 for arg 2 out of range [0,4)
+ ; CHECK-NEXT: call void @llvm.prefetch.p0(ptr %ptr, i32 0, i32 8, i32 0)
+ call void @llvm.prefetch(ptr %ptr, i32 0, i32 8, i32 0)
+
+ ; CHECK: immarg value 4 for arg 3 out of range [0,2)
+ ; CHECK-NEXT: call void @llvm.prefetch.p0(ptr %ptr, i32 0, i32 0, i32 4)
+ call void @llvm.prefetch(ptr %ptr, i32 0, i32 0, i32 4)
ret void
}
diff --git a/mlir/test/Target/LLVMIR/llvmir-intrinsics.mlir b/mlir/test/Target/LLVMIR/llvmir-intrinsics.mlir
index ea393dd445eda..5266470891b4d 100644
--- a/mlir/test/Target/LLVMIR/llvmir-intrinsics.mlir
+++ b/mlir/test/Target/LLVMIR/llvmir-intrinsics.mlir
@@ -1420,7 +1420,7 @@ llvm.func @vector_scmp(%a: vector<4 x i32>, %b: vector<4 x i32>) -> vector<4 x i
// CHECK-DAG: declare <8 x float> @llvm.fma.v8f32(<8 x float>, <8 x float>, <8 x float>) #0
// CHECK-DAG: declare float @llvm.fmuladd.f32(float, float, float)
// CHECK-DAG: declare <8 x float> @llvm.fmuladd.v8f32(<8 x float>, <8 x float>, <8 x float>) #0
-// CHECK-DAG: declare void @llvm.prefetch.p0(ptr readonly captures(none), i32 immarg, i32 immarg, i32 immarg)
+// CHECK-DAG: declare void @llvm.prefetch.p0(ptr readonly captures(none), i32 immarg range(i32 0, 2), i32 immarg range(i32 0, 4), i32 immarg range(i32 0, 2))
// CHECK-DAG: declare i1 @llvm.is.fpclass.f32(float, i32 immarg)
// CHECK-DAG: declare float @llvm.exp.f32(float)
// CHECK-DAG: declare <8 x float> @llvm.exp.v8f32(<8 x float>) #0
More information about the Mlir-commits
mailing list