[llvm] [LLVM][Intrinsics] Print note if manual name matches default name (PR #164716)
Rahul Joshi via llvm-commits
llvm-commits at lists.llvm.org
Thu Oct 23 10:07:43 PDT 2025
https://github.com/jurahul updated https://github.com/llvm/llvm-project/pull/164716
>From 4af5513f3b80f9f895cb3b97a07f1bd78636cbd1 Mon Sep 17 00:00:00 2001
From: Rahul Joshi <rjoshi at nvidia.com>
Date: Wed, 22 Oct 2025 12:51:38 -0700
Subject: [PATCH 1/5] [LLVM][Intrinsics] Print note if manual name matches
default name
Print a note when the manually specified name in an intrinsic
matches the default name it would have been assigned based on the
record name, in which case the manual specification is redundant
and can be eliminated.
Also remove existing redundant manual names.
---
llvm/include/llvm/IR/Intrinsics.td | 12 ++---
llvm/include/llvm/IR/IntrinsicsNVVM.td | 49 ++++++++++++-------
.../TableGen/Basic/CodeGenIntrinsics.cpp | 10 +++-
3 files changed, 44 insertions(+), 27 deletions(-)
diff --git a/llvm/include/llvm/IR/Intrinsics.td b/llvm/include/llvm/IR/Intrinsics.td
index e6cce9a4eea1d..4d59ee8676b9e 100644
--- a/llvm/include/llvm/IR/Intrinsics.td
+++ b/llvm/include/llvm/IR/Intrinsics.td
@@ -1487,24 +1487,23 @@ def int_eh_sjlj_setup_dispatch : Intrinsic<[], []>;
//
def int_var_annotation : DefaultAttrsIntrinsic<
[], [llvm_anyptr_ty, llvm_anyptr_ty, LLVMMatchType<1>, llvm_i32_ty, LLVMMatchType<1>],
- [IntrInaccessibleMemOnly], "llvm.var.annotation">;
+ [IntrInaccessibleMemOnly]>;
def int_ptr_annotation : DefaultAttrsIntrinsic<
[llvm_anyptr_ty],
[LLVMMatchType<0>, llvm_anyptr_ty, LLVMMatchType<1>, llvm_i32_ty, LLVMMatchType<1>],
- [IntrInaccessibleMemOnly], "llvm.ptr.annotation">;
+ [IntrInaccessibleMemOnly]>;
def int_annotation : DefaultAttrsIntrinsic<
[llvm_anyint_ty],
[LLVMMatchType<0>, llvm_anyptr_ty, LLVMMatchType<1>, llvm_i32_ty],
- [IntrInaccessibleMemOnly], "llvm.annotation">;
+ [IntrInaccessibleMemOnly]>;
// Annotates the current program point with metadata strings which are emitted
// as CodeView debug info records. This is expensive, as it disables inlining
// and is modelled as having side effects.
def int_codeview_annotation : DefaultAttrsIntrinsic<[], [llvm_metadata_ty],
- [IntrInaccessibleMemOnly, IntrNoDuplicate],
- "llvm.codeview.annotation">;
+ [IntrInaccessibleMemOnly, IntrNoDuplicate]>;
//===------------------------ Trampoline Intrinsics -----------------------===//
//
@@ -1881,8 +1880,7 @@ def int_clear_cache : Intrinsic<[], [llvm_ptr_ty, llvm_ptr_ty],
// Intrinsic to detect whether its argument is a constant.
def int_is_constant : DefaultAttrsIntrinsic<[llvm_i1_ty], [llvm_any_ty],
- [IntrNoMem, IntrConvergent],
- "llvm.is.constant">;
+ [IntrNoMem, IntrConvergent]>;
// Introduce a use of the argument without generating any code.
def int_fake_use : DefaultAttrsIntrinsic<[], [llvm_vararg_ty],
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 3af1750ffcf3f..7bf835d2a170f 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -533,26 +533,34 @@ class MMA_SP_NAME<string Metadata, string Kind, int Satfinite,
# signature;
}
+class IntrinsicName<string name> {
+ string record = !subst(".", "_",
+ !subst("llvm.", "int_", name));
+ // Use explicit intrinsic name if it has an _ in it, else rely on LLVM
+ // assigned default name.
+ string intr = !if(!ne(!find(name, "_"), -1), name, "");
+}
+
class LDMATRIX_NAME<WMMA_REGS Frag, int Trans> {
- string intr = "llvm.nvvm.ldmatrix.sync.aligned"
+ defvar name = "llvm.nvvm.ldmatrix.sync.aligned"
# "." # Frag.geom
# "." # Frag.frag
# !if(Trans, ".trans", "")
# "." # Frag.ptx_elt_type
;
- string record = !subst(".", "_",
- !subst("llvm.", "int_", intr));
+ string intr = IntrinsicName<name>.intr;
+ string record = IntrinsicName<name>.record;
}
class STMATRIX_NAME<WMMA_REGS Frag, int Trans> {
- string intr = "llvm.nvvm.stmatrix.sync.aligned"
+ defvar name = "llvm.nvvm.stmatrix.sync.aligned"
# "." # Frag.geom
# "." # Frag.frag
# !if(Trans, ".trans", "")
# "." # Frag.ptx_elt_type
;
- string record = !subst(".", "_",
- !subst("llvm.", "int_", intr));
+ string intr = IntrinsicName<name>.intr;
+ string record = IntrinsicName<name>.record;
}
// Generates list of 4-tuples of WMMA_REGS representing a valid MMA op.
@@ -1042,45 +1050,49 @@ class NVVM_TCGEN05_MMA_BASE<string Space, bit Sp> {
class NVVM_TCGEN05_MMA<bit Sp, string Space,
bit AShift, bit ScaleInputD>:
NVVM_TCGEN05_MMA_BASE<Space, Sp> {
- string intr = "llvm.nvvm.tcgen05.mma"
+ string name = "llvm.nvvm.tcgen05.mma"
# !if(!eq(Sp, 1), ".sp", "")
# "." # Space
# !if(!eq(ScaleInputD, 1), ".scale_d", "")
# !if(!eq(AShift, 1), ".ashift", "");
- string record = !subst(".", "_", !subst("llvm.", "int_", intr));
+ string intr = IntrinsicName<name>.intr;
+ string record = IntrinsicName<name>.record;
}
class NVVM_TCGEN05_MMA_BLOCKSCALE<bit Sp, string Space,
string Kind, string ScaleVecSize>:
NVVM_TCGEN05_MMA_BASE<Space, Sp> {
- string intr = "llvm.nvvm.tcgen05.mma"
+ string name = "llvm.nvvm.tcgen05.mma"
# !if(!eq(Sp, 1), ".sp", "")
# "." # Space
# "." # Kind
# ".block_scale" # ScaleVecSize;
- string record = !subst(".", "_", !subst("llvm.", "int_", intr));
+ string intr = IntrinsicName<name>.intr;
+ string record = IntrinsicName<name>.record;
}
class NVVM_TCGEN05_MMA_WS<bit Sp, string Space, bit ZeroColMask>:
NVVM_TCGEN05_MMA_BASE<Space, Sp> {
- string intr = "llvm.nvvm.tcgen05.mma.ws"
+ string name = "llvm.nvvm.tcgen05.mma.ws"
# !if(!eq(Sp, 1), ".sp", "")
# "." # Space
# !if(!eq(ZeroColMask, 1), ".zero_col_mask", "");
- string record = !subst(".", "_", !subst("llvm.", "int_", intr));
+ string intr = IntrinsicName<name>.intr;
+ string record = IntrinsicName<name>.record;
}
class NVVM_TCGEN05_MMA_DISABLE_OUTPUT_LANE<bit Sp, string Space,
int CtaGroup, bit AShift,
bit ScaleInputD>:
NVVM_TCGEN05_MMA_BASE<Space, Sp> {
- string intr = "llvm.nvvm.tcgen05.mma"
+ string name = "llvm.nvvm.tcgen05.mma"
# !if(!eq(Sp, 1), ".sp", "")
# "." # Space
# !if(!eq(ScaleInputD, 1), ".scale_d", "")
# ".disable_output_lane.cg" # CtaGroup
# !if(!eq(AShift, 1), ".ashift", "");
- string record = !subst(".", "_", !subst("llvm.", "int_", intr));
+ string intr = IntrinsicName<name>.intr;
+ string record = IntrinsicName<name>.record;
}
class NVVM_TCGEN05_MMA_BLOCKSCALE_SUPPORTED<string Kind, string ScaleVecSize> {
@@ -2767,14 +2779,15 @@ foreach cta_group = ["cg1", "cg2"] in {
"64x128b_warpx2_02_13",
"64x128b_warpx2_01_23",
"32x128b_warpx4"] in {
- defvar intr_suffix = StrJoin<"_", [shape, src_fmt, cta_group]>.ret;
- defvar name_suffix = StrJoin<".", [shape, src_fmt, cta_group]>.ret;
+ defvar name = "llvm.nvvm.tcgen05.cp." #
+ StrJoin<".", [shape, src_fmt, cta_group]>.ret;
- def int_nvvm_tcgen05_cp_ # intr_suffix : Intrinsic<[],
+ defvar intname = IntrinsicName<name>;
+ def intname.record : Intrinsic<[],
[llvm_tmem_ptr_ty, // tmem_addr
llvm_i64_ty], // smem descriptor
[IntrConvergent, IntrInaccessibleMemOrArgMemOnly, NoCapture<ArgIndex<0>>],
- "llvm.nvvm.tcgen05.cp." # name_suffix>;
+ intname.intr>;
}
}
}
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
index be7537c83da3a..d0993e8d7cd91 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
@@ -278,15 +278,21 @@ CodeGenIntrinsic::CodeGenIntrinsic(const Record *R,
TargetPrefix = R->getValueAsString("TargetPrefix");
Name = R->getValueAsString("LLVMName").str();
+ std::string DerivedName = "llvm." + EnumName.str();
+ llvm::replace(DerivedName, '_', '.');
+
if (Name == "") {
// If an explicit name isn't specified, derive one from the DefName.
- Name = "llvm." + EnumName.str();
- llvm::replace(Name, '_', '.');
+ Name = std::move(DerivedName);
} else {
// Verify it starts with "llvm.".
if (!StringRef(Name).starts_with("llvm."))
PrintFatalError(DefLoc, "Intrinsic '" + DefName +
"'s name does not start with 'llvm.'!");
+
+ if (Name == DerivedName)
+ PrintNote(DefLoc, "Explicitly specified name matches default name, "
+ "consider dropping it.");
}
// If TargetPrefix is specified, make sure that Name starts with
>From f4de61c21c770fbfa23d3b5afc76b3a01b57587b Mon Sep 17 00:00:00 2001
From: Rahul Joshi <rjoshi at nvidia.com>
Date: Thu, 23 Oct 2025 09:33:53 -0700
Subject: [PATCH 2/5] Review feedback
---
llvm/include/llvm/IR/IntrinsicsNVVM.td | 61 ++++++++++---------
llvm/test/TableGen/intrinsic-manual-name.td | 6 ++
.../TableGen/Basic/CodeGenIntrinsics.cpp | 2 +-
3 files changed, 39 insertions(+), 30 deletions(-)
create mode 100644 llvm/test/TableGen/intrinsic-manual-name.td
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 7bf835d2a170f..a6511a567cfbf 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -533,12 +533,15 @@ class MMA_SP_NAME<string Metadata, string Kind, int Satfinite,
# signature;
}
+// Helper class that takes an intrinsic name and construct a record name.
+// Additionally, sets `intr_name` to be non-empty if the default name assigned
+// to this intrinsic will not match the name given.
class IntrinsicName<string name> {
- string record = !subst(".", "_",
+ string record_name = !subst(".", "_",
!subst("llvm.", "int_", name));
// Use explicit intrinsic name if it has an _ in it, else rely on LLVM
// assigned default name.
- string intr = !if(!ne(!find(name, "_"), -1), name, "");
+ string intr_name = !if(!ne(!find(name, "_"), -1), name, "");
}
class LDMATRIX_NAME<WMMA_REGS Frag, int Trans> {
@@ -548,8 +551,8 @@ class LDMATRIX_NAME<WMMA_REGS Frag, int Trans> {
# !if(Trans, ".trans", "")
# "." # Frag.ptx_elt_type
;
- string intr = IntrinsicName<name>.intr;
- string record = IntrinsicName<name>.record;
+ string intr_name = IntrinsicName<name>.intr_name;
+ string record_name = IntrinsicName<name>.record_name;
}
class STMATRIX_NAME<WMMA_REGS Frag, int Trans> {
@@ -559,8 +562,8 @@ class STMATRIX_NAME<WMMA_REGS Frag, int Trans> {
# !if(Trans, ".trans", "")
# "." # Frag.ptx_elt_type
;
- string intr = IntrinsicName<name>.intr;
- string record = IntrinsicName<name>.record;
+ string intr_name = IntrinsicName<name>.intr_name;
+ string record_name = IntrinsicName<name>.record_name;
}
// Generates list of 4-tuples of WMMA_REGS representing a valid MMA op.
@@ -1055,8 +1058,8 @@ class NVVM_TCGEN05_MMA<bit Sp, string Space,
# "." # Space
# !if(!eq(ScaleInputD, 1), ".scale_d", "")
# !if(!eq(AShift, 1), ".ashift", "");
- string intr = IntrinsicName<name>.intr;
- string record = IntrinsicName<name>.record;
+ string intr_name = IntrinsicName<name>.intr_name;
+ string record_name = IntrinsicName<name>.record_name;
}
class NVVM_TCGEN05_MMA_BLOCKSCALE<bit Sp, string Space,
@@ -1067,8 +1070,8 @@ class NVVM_TCGEN05_MMA_BLOCKSCALE<bit Sp, string Space,
# "." # Space
# "." # Kind
# ".block_scale" # ScaleVecSize;
- string intr = IntrinsicName<name>.intr;
- string record = IntrinsicName<name>.record;
+ string intr_name = IntrinsicName<name>.intr_name;
+ string record_name = IntrinsicName<name>.record_name;
}
class NVVM_TCGEN05_MMA_WS<bit Sp, string Space, bit ZeroColMask>:
@@ -1077,8 +1080,8 @@ class NVVM_TCGEN05_MMA_WS<bit Sp, string Space, bit ZeroColMask>:
# !if(!eq(Sp, 1), ".sp", "")
# "." # Space
# !if(!eq(ZeroColMask, 1), ".zero_col_mask", "");
- string intr = IntrinsicName<name>.intr;
- string record = IntrinsicName<name>.record;
+ string intr_name = IntrinsicName<name>.intr_name;
+ string record_name = IntrinsicName<name>.record_name;
}
class NVVM_TCGEN05_MMA_DISABLE_OUTPUT_LANE<bit Sp, string Space,
@@ -1091,8 +1094,8 @@ class NVVM_TCGEN05_MMA_DISABLE_OUTPUT_LANE<bit Sp, string Space,
# !if(!eq(ScaleInputD, 1), ".scale_d", "")
# ".disable_output_lane.cg" # CtaGroup
# !if(!eq(AShift, 1), ".ashift", "");
- string intr = IntrinsicName<name>.intr;
- string record = IntrinsicName<name>.record;
+ string intr_name = IntrinsicName<name>.intr_name;
+ string record_name = IntrinsicName<name>.record_name;
}
class NVVM_TCGEN05_MMA_BLOCKSCALE_SUPPORTED<string Kind, string ScaleVecSize> {
@@ -2404,12 +2407,12 @@ class NVVM_LDMATRIX<WMMA_REGS Frag, int Transposed>
: Intrinsic<Frag.regs, [llvm_anyptr_ty],
[IntrReadMem, IntrArgMemOnly, IntrNoCallback, ReadOnly<ArgIndex<0>>,
NoCapture<ArgIndex<0>>],
- LDMATRIX_NAME<Frag, Transposed>.intr>;
+ LDMATRIX_NAME<Frag, Transposed>.intr_name>;
foreach transposed = [0, 1] in {
foreach frag = NVVM_MMA_OPS.all_ldmatrix_ops in {
if NVVM_LDMATRIX_SUPPORTED<frag, transposed>.ret then {
- def LDMATRIX_NAME<frag, transposed>.record
+ def LDMATRIX_NAME<frag, transposed>.record_name
: NVVM_LDMATRIX<frag, transposed>;
}
}
@@ -2421,12 +2424,12 @@ class NVVM_STMATRIX<WMMA_REGS Frag, int Transposed>
!listconcat([llvm_anyptr_ty], Frag.regs),
[IntrWriteMem, IntrArgMemOnly, IntrNoCallback,
WriteOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>],
- STMATRIX_NAME<Frag, Transposed>.intr>;
+ STMATRIX_NAME<Frag, Transposed>.intr_name>;
foreach transposed = [0, 1] in {
foreach frag = NVVM_MMA_OPS.all_stmatrix_ops in {
if NVVM_STMATRIX_SUPPORTED<frag, transposed>.ret then {
- def STMATRIX_NAME<frag, transposed>.record
+ def STMATRIX_NAME<frag, transposed>.record_name
: NVVM_STMATRIX<frag, transposed>;
}
}
@@ -2782,12 +2785,12 @@ foreach cta_group = ["cg1", "cg2"] in {
defvar name = "llvm.nvvm.tcgen05.cp." #
StrJoin<".", [shape, src_fmt, cta_group]>.ret;
- defvar intname = IntrinsicName<name>;
- def intname.record : Intrinsic<[],
+ defvar intrinsic_name = IntrinsicName<name>;
+ def intrinsic_name.record_name : Intrinsic<[],
[llvm_tmem_ptr_ty, // tmem_addr
llvm_i64_ty], // smem descriptor
[IntrConvergent, IntrInaccessibleMemOrArgMemOnly, NoCapture<ArgIndex<0>>],
- intname.intr>;
+ intrinsic_name.intr_name>;
}
}
}
@@ -2894,9 +2897,9 @@ foreach sp = [0, 1] in {
]
);
- def mma.record:
+ def mma.record_name:
DefaultAttrsIntrinsicFlags<[], args, flags, intrinsic_properties,
- mma.intr>;
+ mma.intr_name>;
}
}
}
@@ -2931,8 +2934,8 @@ foreach sp = [0, 1] in {
Range<ArgIndex<!add(nargs, 1)>, 0, !if(!eq(ashift, 1), 2, 4)>]
);
- def mma.record: DefaultAttrsIntrinsicFlags<[], args, flags, intrinsic_properties,
- mma.intr>;
+ def mma.record_name : DefaultAttrsIntrinsicFlags<[], args, flags,
+ intrinsic_properties, mma.intr_name>;
} // ashift
} // scale_d
} // cta_group
@@ -2957,11 +2960,11 @@ foreach sp = [0, 1] in {
defvar collector_usage = ArgIndex<!add(nargs, 1)>;
if NVVM_TCGEN05_MMA_BLOCKSCALE_SUPPORTED<kind, scale_vec_size>.ret then {
- def mma.record: DefaultAttrsIntrinsicFlags<[], args, flags,
+ def mma.record_name : DefaultAttrsIntrinsicFlags<[], args, flags,
!listconcat(mma.common_intr_props,
[Range<cta_group, 1, 3>,
Range<collector_usage, 0, 4>]),
- mma.intr>;
+ mma.intr_name>;
}
}
}
@@ -2990,9 +2993,9 @@ foreach sp = [0, 1] in {
Range<ArgIndex<!add(nargs, 2)>, 0, 4>]
);
- def mma.record:
+ def mma.record_name:
DefaultAttrsIntrinsicFlags<[], args, flags, intrinsic_properties,
- mma.intr>;
+ mma.intr_name>;
}
}
}
diff --git a/llvm/test/TableGen/intrinsic-manual-name.td b/llvm/test/TableGen/intrinsic-manual-name.td
new file mode 100644
index 0000000000000..5751fc2874b97
--- /dev/null
+++ b/llvm/test/TableGen/intrinsic-manual-name.td
@@ -0,0 +1,6 @@
+// RUN: llvm-tblgen -gen-intrinsic-impl -I %p/../../include %s -DTEST_INTRINSICS_SUPPRESS_DEFS 2>&1 | FileCheck %s -DFILE=%s
+
+include "llvm/IR/Intrinsics.td"
+
+// CHECK: [[FILE]]:[[@LINE+1]]:5: note: Explicitly specified name matches default name, consider dropping it
+def int_foo0 : Intrinsic<[llvm_anyint_ty], [], [], "llvm.foo0">;
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
index d0993e8d7cd91..f36bf05661782 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
@@ -292,7 +292,7 @@ CodeGenIntrinsic::CodeGenIntrinsic(const Record *R,
if (Name == DerivedName)
PrintNote(DefLoc, "Explicitly specified name matches default name, "
- "consider dropping it.");
+ "consider dropping it");
}
// If TargetPrefix is specified, make sure that Name starts with
>From 132075fbf3dc3f58a5b2958166974ae317f69a4d Mon Sep 17 00:00:00 2001
From: Rahul Joshi <rjoshi at nvidia.com>
Date: Thu, 23 Oct 2025 09:37:09 -0700
Subject: [PATCH 3/5] Rename DerivedName to DefaultName
---
llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp | 8 ++++----
1 file changed, 4 insertions(+), 4 deletions(-)
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
index f36bf05661782..cd866469792a2 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
@@ -278,19 +278,19 @@ CodeGenIntrinsic::CodeGenIntrinsic(const Record *R,
TargetPrefix = R->getValueAsString("TargetPrefix");
Name = R->getValueAsString("LLVMName").str();
- std::string DerivedName = "llvm." + EnumName.str();
- llvm::replace(DerivedName, '_', '.');
+ std::string DefaultName = "llvm." + EnumName.str();
+ llvm::replace(DefaultName, '_', '.');
if (Name == "") {
// If an explicit name isn't specified, derive one from the DefName.
- Name = std::move(DerivedName);
+ Name = std::move(DefaultName);
} else {
// Verify it starts with "llvm.".
if (!StringRef(Name).starts_with("llvm."))
PrintFatalError(DefLoc, "Intrinsic '" + DefName +
"'s name does not start with 'llvm.'!");
- if (Name == DerivedName)
+ if (Name == DefaultName)
PrintNote(DefLoc, "Explicitly specified name matches default name, "
"consider dropping it");
}
>From c70902a6bc6114b95b765b46feab868d8f9027df Mon Sep 17 00:00:00 2001
From: Rahul Joshi <rjoshi at nvidia.com>
Date: Thu, 23 Oct 2025 10:02:04 -0700
Subject: [PATCH 4/5] Fix NVPTXIntrinsics.td references
---
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 10 +++++-----
1 file changed, 5 insertions(+), 5 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 22cf3a7eef2c1..774ab140aef4b 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -4946,7 +4946,7 @@ defset list<WMMA_INSTR> MMA_SPs = {
// ldmatrix.sync.aligned.m8n8[|.trans][|.shared].b16
//
class LDMATRIX<WMMA_REGINFO Frag, bit Transposed, string Space>
- : WMMA_INSTR<LDMATRIX_NAME<Frag, Transposed>.record, [(ins ADDR:$src)]>,
+ : WMMA_INSTR<LDMATRIX_NAME<Frag, Transposed>.record_name, [(ins ADDR:$src)]>,
Requires<Frag.Predicates> {
// Build PatFrag that only matches particular address space.
PatFrag IntrFrag = PatFrag<(ops node:$src), (Intr node:$src),
@@ -4981,7 +4981,7 @@ defset list<WMMA_INSTR> LDMATRIXs = {
// stmatrix.sync.aligned.m8n8[|.trans][|.shared].b16
//
class STMATRIX<WMMA_REGINFO Frag, bit Transposed, string Space>
- : WMMA_INSTR<STMATRIX_NAME<Frag, Transposed>.record, [!con((ins ADDR:$dst), Frag.Ins)]>,
+ : WMMA_INSTR<STMATRIX_NAME<Frag, Transposed>.record_name, [!con((ins ADDR:$dst), Frag.Ins)]>,
Requires<Frag.Predicates> {
// Build PatFrag that only matches particular address space.
dag PFOperands = !con((ops node:$dst),
@@ -5376,7 +5376,7 @@ class Tcgen05MMAInst<bit Sp, string KindStr, string ASpace,
Requires<PTXPredicates> {
Intrinsic Intrin = !cast<Intrinsic>(
- NVVM_TCGEN05_MMA<Sp, ASpace, AShift, ScaleInputD>.record
+ NVVM_TCGEN05_MMA<Sp, ASpace, AShift, ScaleInputD>.record_name
);
dag ScaleInpIns = !if(!eq(ScaleInputD, 1), (ins i64imm:$scale_input_d), (ins));
@@ -5618,7 +5618,7 @@ class Tcgen05MMABlockScaleInst<bit Sp, string ASpace, string KindStr,
Requires<[hasTcgen05Instructions, PTXPredicate]> {
Intrinsic Intrin = !cast<Intrinsic>(
- NVVM_TCGEN05_MMA_BLOCKSCALE<Sp, ASpace, KindStr, ScaleVecSize>.record);
+ NVVM_TCGEN05_MMA_BLOCKSCALE<Sp, ASpace, KindStr, ScaleVecSize>.record_name);
dag SparseMetadataIns = !if(!eq(Sp, 1), (ins B32:$spmetadata), (ins));
dag SparseMetadataIntr = !if(!eq(Sp, 1), (Intrin i32:$spmetadata), (Intrin));
@@ -5702,7 +5702,7 @@ class Tcgen05MMAWSInst<bit Sp, string ASpace, string KindStr,
Requires<[hasTcgen05Instructions]> {
Intrinsic Intrin = !cast<Intrinsic>(
- NVVM_TCGEN05_MMA_WS<Sp, ASpace, HasZeroColMask>.record);
+ NVVM_TCGEN05_MMA_WS<Sp, ASpace, HasZeroColMask>.record_name);
dag ZeroColMaskIns = !if(!eq(HasZeroColMask, 1),
(ins B64:$zero_col_mask), (ins));
>From 6dcaade392fcd9db8db2791279c7853ec7ed41e9 Mon Sep 17 00:00:00 2001
From: Rahul Joshi <rjoshi at nvidia.com>
Date: Thu, 23 Oct 2025 10:07:18 -0700
Subject: [PATCH 5/5] Rename more .intr/.record in NVVM
---
llvm/include/llvm/IR/IntrinsicsNVVM.td | 24 ++++++++++++------------
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 10 +++++-----
2 files changed, 17 insertions(+), 17 deletions(-)
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index a6511a567cfbf..c9df6c43fd396 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -456,7 +456,7 @@ class WMMA_REGS<string Geom, string Frag, string PtxEltType, bit IsSparse = fals
}
class WMMA_NAME_LDST<string Op, WMMA_REGS Frag, string Layout, int WithStride> {
- string intr = "llvm.nvvm.wmma."
+ string intr_name = "llvm.nvvm.wmma."
# Frag.geom
# "." # Op
# "." # Frag.frag
@@ -467,7 +467,7 @@ class WMMA_NAME_LDST<string Op, WMMA_REGS Frag, string Layout, int WithStride> {
// TODO(tra): record name should ideally use the same field order as the intrinsic.
// E.g. string record = !subst("llvm", "int",
// !subst(".", "_", llvm));
- string record = "int_nvvm_wmma_"
+ string record_name = "int_nvvm_wmma_"
# Frag.geom
# "_" # Op
# "_" # Frag.frag
@@ -496,7 +496,7 @@ class MMA_SIGNATURE<WMMA_REGS A, WMMA_REGS B, WMMA_REGS C, WMMA_REGS D> {
class WMMA_NAME<string ALayout, string BLayout, int Satfinite, string Rnd, string b1op,
WMMA_REGS A, WMMA_REGS B, WMMA_REGS C, WMMA_REGS D> {
string signature = MMA_SIGNATURE<A, B, C, D>.ret;
- string record = "int_nvvm_wmma_"
+ string record_name = "int_nvvm_wmma_"
# A.geom
# "_mma"
# !subst(".", "_", b1op)
@@ -510,7 +510,7 @@ class WMMA_NAME<string ALayout, string BLayout, int Satfinite, string Rnd, strin
class MMA_NAME<string ALayout, string BLayout, int Satfinite, string b1op, string Kind,
WMMA_REGS A, WMMA_REGS B, WMMA_REGS C, WMMA_REGS D> {
string signature = MMA_SIGNATURE<A, B, C, D>.ret;
- string record = "int_nvvm_mma"
+ string record_name = "int_nvvm_mma"
# !subst(".", "_", b1op)
# "_" # A.geom
# "_" # ALayout
@@ -524,7 +524,7 @@ class MMA_SP_NAME<string Metadata, string Kind, int Satfinite,
WMMA_REGS A, WMMA_REGS B,
WMMA_REGS C, WMMA_REGS D> {
string signature = MMA_SIGNATURE<A, B, C, D>.ret;
- string record = "int_nvvm_mma"
+ string record_name = "int_nvvm_mma"
# "_" # !subst("::", "_", Metadata)
# "_" # A.geom
# "_row_col"
@@ -2288,7 +2288,7 @@ class NVVM_WMMA_LD<WMMA_REGS Frag, string Layout, int WithStride>
: Intrinsic<Frag.regs,
!if(WithStride, [llvm_anyptr_ty, llvm_i32_ty], [llvm_anyptr_ty]),
[IntrWillReturn, IntrReadMem, IntrArgMemOnly, IntrNoCallback, ReadOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>],
- WMMA_NAME_LDST<"load", Frag, Layout, WithStride>.intr>;
+ WMMA_NAME_LDST<"load", Frag, Layout, WithStride>.intr_name>;
// WMMA.STORE.D
class NVVM_WMMA_ST<WMMA_REGS Frag, string Layout, int WithStride>
@@ -2298,18 +2298,18 @@ class NVVM_WMMA_ST<WMMA_REGS Frag, string Layout, int WithStride>
Frag.regs,
!if(WithStride, [llvm_i32_ty], [])),
[IntrWriteMem, IntrArgMemOnly, IntrNoCallback, WriteOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>],
- WMMA_NAME_LDST<"store", Frag, Layout, WithStride>.intr>;
+ WMMA_NAME_LDST<"store", Frag, Layout, WithStride>.intr_name>;
// Create all load/store variants
foreach layout = ["row", "col"] in {
foreach stride = [0, 1] in {
foreach frag = NVVM_MMA_OPS.all_ld_ops in
if NVVM_WMMA_LDST_SUPPORTED<frag, layout>.ret then
- def WMMA_NAME_LDST<"load", frag, layout, stride>.record
+ def WMMA_NAME_LDST<"load", frag, layout, stride>.record_name
: NVVM_WMMA_LD<frag, layout, stride>;
foreach frag = NVVM_MMA_OPS.all_st_ops in
if NVVM_WMMA_LDST_SUPPORTED<frag, layout>.ret then
- def WMMA_NAME_LDST<"store", frag, layout, stride>.record
+ def WMMA_NAME_LDST<"store", frag, layout, stride>.record_name
: NVVM_WMMA_ST<frag, layout, stride>;
}
}
@@ -2328,7 +2328,7 @@ foreach layout_a = ["row", "col"] in {
foreach b1op = NVVM_MMA_B1OPS<op>.ret in {
if NVVM_WMMA_SUPPORTED<op, layout_a, layout_b, satf, rnd>.ret then {
def WMMA_NAME<layout_a, layout_b, satf, rnd, b1op,
- op[0], op[1], op[2], op[3]>.record
+ op[0], op[1], op[2], op[3]>.record_name
: NVVM_MMA<op[0], op[1], op[2], op[3]>;
}
} // b1op
@@ -2345,7 +2345,7 @@ foreach layout_a = ["row", "col"] in {
foreach b1op = NVVM_MMA_B1OPS<op>.ret in {
foreach kind = ["", "kind::f8f6f4"] in {
if NVVM_MMA_SUPPORTED<op, layout_a, layout_b, kind, satf>.ret then {
- def MMA_NAME<layout_a, layout_b, satf, b1op, kind, op[0], op[1], op[2], op[3]>.record
+ def MMA_NAME<layout_a, layout_b, satf, b1op, kind, op[0], op[1], op[2], op[3]>.record_name
: NVVM_MMA<op[0], op[1], op[2], op[3]>;
}
} // kind
@@ -2394,7 +2394,7 @@ foreach metadata = ["sp", "sp::ordered_metadata"] in {
foreach op = NVVM_MMA_OPS.all_mma_sp_ops in {
if NVVM_MMA_SP_SUPPORTED<op, metadata, kind, satf>.ret then {
def MMA_SP_NAME<metadata, kind, satf,
- op[0], op[1], op[2], op[3]>.record
+ op[0], op[1], op[2], op[3]>.record_name
: NVVM_MMA_SP<op[0], op[1], op[2], op[3]>;
}
} // op
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 774ab140aef4b..598735f5972bc 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -4675,7 +4675,7 @@ class WMMA_INSTR<string _Intr, list<dag> _Args>
//
class WMMA_LOAD<WMMA_REGINFO Frag, string Layout, string Space, bit WithStride>
- : WMMA_INSTR<WMMA_NAME_LDST<"load", Frag, Layout, WithStride>.record,
+ : WMMA_INSTR<WMMA_NAME_LDST<"load", Frag, Layout, WithStride>.record_name,
[!con((ins ADDR:$src),
!if(WithStride, (ins B32:$ldm), (ins)))]>,
Requires<Frag.Predicates> {
@@ -4714,7 +4714,7 @@ class WMMA_LOAD<WMMA_REGINFO Frag, string Layout, string Space, bit WithStride>
//
class WMMA_STORE_D<WMMA_REGINFO Frag, string Layout, string Space,
bit WithStride>
- : WMMA_INSTR<WMMA_NAME_LDST<"store", Frag, Layout, WithStride>.record,
+ : WMMA_INSTR<WMMA_NAME_LDST<"store", Frag, Layout, WithStride>.record_name,
[!con((ins ADDR:$dst),
Frag.Ins,
!if(WithStride, (ins B32:$ldm), (ins)))]>,
@@ -4778,7 +4778,7 @@ class MMA_OP_PREDICATES<WMMA_REGINFO FragA, string b1op> {
class WMMA_MMA<WMMA_REGINFO FragA, WMMA_REGINFO FragB,
WMMA_REGINFO FragC, WMMA_REGINFO FragD,
string ALayout, string BLayout, int Satfinite, string rnd, string b1op>
- : WMMA_INSTR<WMMA_NAME<ALayout, BLayout, Satfinite, rnd, b1op, FragA, FragB, FragC, FragD>.record,
+ : WMMA_INSTR<WMMA_NAME<ALayout, BLayout, Satfinite, rnd, b1op, FragA, FragB, FragC, FragD>.record_name,
[FragA.Ins, FragB.Ins, FragC.Ins]>,
// Requires does not seem to have effect on Instruction w/o Patterns.
// We set it here anyways and propagate to the Pat<> we construct below.
@@ -4837,7 +4837,7 @@ defset list<WMMA_INSTR> WMMAs = {
class MMA<WMMA_REGINFO FragA, WMMA_REGINFO FragB,
WMMA_REGINFO FragC, WMMA_REGINFO FragD,
string ALayout, string BLayout, int Satfinite, string b1op, string Kind>
- : WMMA_INSTR<MMA_NAME<ALayout, BLayout, Satfinite, b1op, Kind, FragA, FragB, FragC, FragD>.record,
+ : WMMA_INSTR<MMA_NAME<ALayout, BLayout, Satfinite, b1op, Kind, FragA, FragB, FragC, FragD>.record_name,
[FragA.Ins, FragB.Ins, FragC.Ins]>,
// Requires does not seem to have effect on Instruction w/o Patterns.
// We set it here anyways and propagate to the Pat<> we construct below.
@@ -4891,7 +4891,7 @@ class MMA_SP<WMMA_REGINFO FragA, WMMA_REGINFO FragB,
WMMA_REGINFO FragC, WMMA_REGINFO FragD,
string Metadata, string Kind, int Satfinite>
: WMMA_INSTR<MMA_SP_NAME<Metadata, Kind, Satfinite,
- FragA, FragB, FragC, FragD>.record,
+ FragA, FragB, FragC, FragD>.record_name,
[FragA.Ins, FragB.Ins, FragC.Ins,
(ins B32:$metadata, i32imm:$selector)]>,
// Requires does not seem to have effect on Instruction w/o Patterns.
More information about the llvm-commits
mailing list