[LLVM][Intrinsics] Print note if manual name matches default name #164716
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
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.
Member
|
@llvm/pr-subscribers-tablegen @llvm/pr-subscribers-backend-nvptx Author: Rahul Joshi (jurahul) ChangesPrint 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. Full diff: https://github.com/llvm/llvm-project/pull/164716.diff 3 Files Affected:
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
|
Member
|
@llvm/pr-subscribers-llvm-ir Author: Rahul Joshi (jurahul) ChangesPrint 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. Full diff: https://github.com/llvm/llvm-project/pull/164716.diff 3 Files Affected:
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
|
arsenm
approved these changes
Oct 23, 2025
durga4github
approved these changes
Oct 23, 2025
Contributor
durga4github
left a comment
durga4github
left a comment
There was a problem hiding this comment.
I have a few minor comments, LGTM otherwise.
Contributor
Author
|
I also renamed the variable in CodeGenIntrinsics.cpp from |
Contributor
Author
|
@durga4github I renamed all instances of |
Contributor
Sure, that's a welcome change! |
dvbuka
pushed a commit
to dvbuka/llvm-project
that referenced
this pull request
Oct 27, 2025
…vm#164716) 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.
Lukacma
pushed a commit
to Lukacma/llvm-project
that referenced
this pull request
Oct 29, 2025
…vm#164716) 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.
aokblast
pushed a commit
to aokblast/llvm-project
that referenced
this pull request
Oct 30, 2025
…vm#164716) 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.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
4 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
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.