From 2bb4226c7c6da0edf40b4f1b87e9a625ff2a0e31 Mon Sep 17 00:00:00 2001 From: Rahul Joshi Date: Thu, 23 Oct 2025 12:00:03 -0700 Subject: [PATCH] [LLVM][Intrinsics] Print note if manual name matches default name (#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. --- llvm/include/llvm/IR/Intrinsics.td | 12 +-- llvm/include/llvm/IR/IntrinsicsNVVM.td | 100 ++++++++++-------- llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 20 ++-- llvm/test/TableGen/intrinsic-manual-name.td | 6 ++ .../TableGen/Basic/CodeGenIntrinsics.cpp | 10 +- 5 files changed, 87 insertions(+), 61 deletions(-) create mode 100644 llvm/test/TableGen/intrinsic-manual-name.td diff --git a/llvm/include/llvm/IR/Intrinsics.td b/llvm/include/llvm/IR/Intrinsics.td index e6cce9a4eea1..4d59ee8676b9 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 3af1750ffcf3..c9df6c43fd39 100644 --- a/llvm/include/llvm/IR/IntrinsicsNVVM.td +++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td @@ -456,7 +456,7 @@ class WMMA_REGS { - string intr = "llvm.nvvm.wmma." + string intr_name = "llvm.nvvm.wmma." # Frag.geom # "." # Op # "." # Frag.frag @@ -467,7 +467,7 @@ class WMMA_NAME_LDST { // 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 { class WMMA_NAME { string signature = MMA_SIGNATURE.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 signature = MMA_SIGNATURE.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 signature = MMA_SIGNATURE.ret; - string record = "int_nvvm_mma" + string record_name = "int_nvvm_mma" # "_" # !subst("::", "_", Metadata) # "_" # A.geom # "_row_col" @@ -533,26 +533,37 @@ class MMA_SP_NAME { + 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_name = !if(!ne(!find(name, "_"), -1), name, ""); +} + class LDMATRIX_NAME { - 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_name = IntrinsicName.intr_name; + string record_name = IntrinsicName.record_name; } class STMATRIX_NAME { - 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_name = IntrinsicName.intr_name; + string record_name = IntrinsicName.record_name; } // Generates list of 4-tuples of WMMA_REGS representing a valid MMA op. @@ -1042,45 +1053,49 @@ class NVVM_TCGEN05_MMA_BASE { class NVVM_TCGEN05_MMA: NVVM_TCGEN05_MMA_BASE { - 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_name = IntrinsicName.intr_name; + string record_name = IntrinsicName.record_name; } class NVVM_TCGEN05_MMA_BLOCKSCALE: NVVM_TCGEN05_MMA_BASE { - 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_name = IntrinsicName.intr_name; + string record_name = IntrinsicName.record_name; } class NVVM_TCGEN05_MMA_WS: NVVM_TCGEN05_MMA_BASE { - 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_name = IntrinsicName.intr_name; + string record_name = IntrinsicName.record_name; } class NVVM_TCGEN05_MMA_DISABLE_OUTPUT_LANE: NVVM_TCGEN05_MMA_BASE { - 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_name = IntrinsicName.intr_name; + string record_name = IntrinsicName.record_name; } class NVVM_TCGEN05_MMA_BLOCKSCALE_SUPPORTED { @@ -2273,7 +2288,7 @@ class NVVM_WMMA_LD : Intrinsic>, NoCapture>], - WMMA_NAME_LDST<"load", Frag, Layout, WithStride>.intr>; + WMMA_NAME_LDST<"load", Frag, Layout, WithStride>.intr_name>; // WMMA.STORE.D class NVVM_WMMA_ST @@ -2283,18 +2298,18 @@ class NVVM_WMMA_ST Frag.regs, !if(WithStride, [llvm_i32_ty], [])), [IntrWriteMem, IntrArgMemOnly, IntrNoCallback, WriteOnly>, NoCapture>], - 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.ret then - def WMMA_NAME_LDST<"load", frag, layout, stride>.record + def WMMA_NAME_LDST<"load", frag, layout, stride>.record_name : NVVM_WMMA_LD; foreach frag = NVVM_MMA_OPS.all_st_ops in if NVVM_WMMA_LDST_SUPPORTED.ret then - def WMMA_NAME_LDST<"store", frag, layout, stride>.record + def WMMA_NAME_LDST<"store", frag, layout, stride>.record_name : NVVM_WMMA_ST; } } @@ -2313,7 +2328,7 @@ foreach layout_a = ["row", "col"] in { foreach b1op = NVVM_MMA_B1OPS.ret in { if NVVM_WMMA_SUPPORTED.ret then { def WMMA_NAME.record + op[0], op[1], op[2], op[3]>.record_name : NVVM_MMA; } } // b1op @@ -2330,7 +2345,7 @@ foreach layout_a = ["row", "col"] in { foreach b1op = NVVM_MMA_B1OPS.ret in { foreach kind = ["", "kind::f8f6f4"] in { if NVVM_MMA_SUPPORTED.ret then { - def MMA_NAME.record + def MMA_NAME.record_name : NVVM_MMA; } } // kind @@ -2379,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.ret then { def MMA_SP_NAME.record + op[0], op[1], op[2], op[3]>.record_name : NVVM_MMA_SP; } } // op @@ -2392,12 +2407,12 @@ class NVVM_LDMATRIX : Intrinsic>, NoCapture>], - LDMATRIX_NAME.intr>; + LDMATRIX_NAME.intr_name>; foreach transposed = [0, 1] in { foreach frag = NVVM_MMA_OPS.all_ldmatrix_ops in { if NVVM_LDMATRIX_SUPPORTED.ret then { - def LDMATRIX_NAME.record + def LDMATRIX_NAME.record_name : NVVM_LDMATRIX; } } @@ -2409,12 +2424,12 @@ class NVVM_STMATRIX !listconcat([llvm_anyptr_ty], Frag.regs), [IntrWriteMem, IntrArgMemOnly, IntrNoCallback, WriteOnly>, NoCapture>], - STMATRIX_NAME.intr>; + STMATRIX_NAME.intr_name>; foreach transposed = [0, 1] in { foreach frag = NVVM_MMA_OPS.all_stmatrix_ops in { if NVVM_STMATRIX_SUPPORTED.ret then { - def STMATRIX_NAME.record + def STMATRIX_NAME.record_name : NVVM_STMATRIX; } } @@ -2767,14 +2782,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 intrinsic_name = IntrinsicName; + def intrinsic_name.record_name : Intrinsic<[], [llvm_tmem_ptr_ty, // tmem_addr llvm_i64_ty], // smem descriptor [IntrConvergent, IntrInaccessibleMemOrArgMemOnly, NoCapture>], - "llvm.nvvm.tcgen05.cp." # name_suffix>; + intrinsic_name.intr_name>; } } } @@ -2881,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>; } } } @@ -2918,8 +2934,8 @@ foreach sp = [0, 1] in { Range, 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 @@ -2944,11 +2960,11 @@ foreach sp = [0, 1] in { defvar collector_usage = ArgIndex; if NVVM_TCGEN05_MMA_BLOCKSCALE_SUPPORTED.ret then { - def mma.record: DefaultAttrsIntrinsicFlags<[], args, flags, + def mma.record_name : DefaultAttrsIntrinsicFlags<[], args, flags, !listconcat(mma.common_intr_props, [Range, Range]), - mma.intr>; + mma.intr_name>; } } } @@ -2977,9 +2993,9 @@ foreach sp = [0, 1] in { Range, 0, 4>] ); - def mma.record: + def mma.record_name: DefaultAttrsIntrinsicFlags<[], args, flags, intrinsic_properties, - mma.intr>; + mma.intr_name>; } } } diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td index 22cf3a7eef2c..598735f5972b 100644 --- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td +++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td @@ -4675,7 +4675,7 @@ class WMMA_INSTR _Args> // class WMMA_LOAD - : WMMA_INSTR.record, + : WMMA_INSTR.record_name, [!con((ins ADDR:$src), !if(WithStride, (ins B32:$ldm), (ins)))]>, Requires { @@ -4714,7 +4714,7 @@ class WMMA_LOAD // class WMMA_STORE_D - : WMMA_INSTR.record, + : WMMA_INSTR.record_name, [!con((ins ADDR:$dst), Frag.Ins, !if(WithStride, (ins B32:$ldm), (ins)))]>, @@ -4778,7 +4778,7 @@ class MMA_OP_PREDICATES { class WMMA_MMA - : WMMA_INSTR.record, + : WMMA_INSTR.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 WMMAs = { class MMA - : WMMA_INSTR.record, + : WMMA_INSTR.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_INSTR.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. @@ -4946,7 +4946,7 @@ defset list MMA_SPs = { // ldmatrix.sync.aligned.m8n8[|.trans][|.shared].b16 // class LDMATRIX - : WMMA_INSTR.record, [(ins ADDR:$src)]>, + : WMMA_INSTR.record_name, [(ins ADDR:$src)]>, Requires { // Build PatFrag that only matches particular address space. PatFrag IntrFrag = PatFrag<(ops node:$src), (Intr node:$src), @@ -4981,7 +4981,7 @@ defset list LDMATRIXs = { // stmatrix.sync.aligned.m8n8[|.trans][|.shared].b16 // class STMATRIX - : WMMA_INSTR.record, [!con((ins ADDR:$dst), Frag.Ins)]>, + : WMMA_INSTR.record_name, [!con((ins ADDR:$dst), Frag.Ins)]>, Requires { // Build PatFrag that only matches particular address space. dag PFOperands = !con((ops node:$dst), @@ -5376,7 +5376,7 @@ class Tcgen05MMAInst { Intrinsic Intrin = !cast( - NVVM_TCGEN05_MMA.record + NVVM_TCGEN05_MMA.record_name ); dag ScaleInpIns = !if(!eq(ScaleInputD, 1), (ins i64imm:$scale_input_d), (ins)); @@ -5618,7 +5618,7 @@ class Tcgen05MMABlockScaleInst { Intrinsic Intrin = !cast( - NVVM_TCGEN05_MMA_BLOCKSCALE.record); + NVVM_TCGEN05_MMA_BLOCKSCALE.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 { Intrinsic Intrin = !cast( - NVVM_TCGEN05_MMA_WS.record); + NVVM_TCGEN05_MMA_WS.record_name); dag ZeroColMaskIns = !if(!eq(HasZeroColMask, 1), (ins B64:$zero_col_mask), (ins)); diff --git a/llvm/test/TableGen/intrinsic-manual-name.td b/llvm/test/TableGen/intrinsic-manual-name.td new file mode 100644 index 000000000000..5751fc2874b9 --- /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 be7537c83da3..cd866469792a 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 DefaultName = "llvm." + EnumName.str(); + llvm::replace(DefaultName, '_', '.'); + if (Name == "") { // If an explicit name isn't specified, derive one from the DefName. - Name = "llvm." + EnumName.str(); - llvm::replace(Name, '_', '.'); + 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 == DefaultName) + PrintNote(DefLoc, "Explicitly specified name matches default name, " + "consider dropping it"); } // If TargetPrefix is specified, make sure that Name starts with