[Mlir-commits] [llvm] [mlir] [NFC][LLVM] Define range attribute for `llvm.prefetch` args (PR #207300)
Rahul Joshi
llvmlistbot at llvm.org
Fri Jul 3 00:46:31 PDT 2026
https://github.com/jurahul updated https://github.com/llvm/llvm-project/pull/207300
>From a8fa26814dfc0636e81a0f8bda0ad98e0f293eb7 Mon Sep 17 00:00:00 2001
From: Rahul Joshi <rjoshi at nvidia.com>
Date: Thu, 2 Jul 2026 17:47:20 -0700
Subject: [PATCH] [NFC][LLVM] Define range attribute for `llvm.prefetch` args
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`.
---
llvm/include/llvm/IR/Intrinsics.td | 6 ++--
llvm/lib/IR/Verifier.cpp | 14 ++-------
llvm/test/Assembler/invalid-immarg5.ll | 2 +-
.../AMDGPU/test-cvt-fp4f6f8-immarg-ranges.ll | 30 +++++++++----------
.../Verifier/NVPTX/fence-proxy.tensormap.ll | 8 ++---
llvm/test/Verifier/NVPTX/setmaxnreg.ll | 4 +--
llvm/test/Verifier/intrinsic-immarg.ll | 12 ++++++++
.../test/Target/LLVMIR/llvmir-intrinsics.mlir | 2 +-
8 files changed, 42 insertions(+), 36 deletions(-)
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/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