Skip to content

[NVPTX] Refactor intrinsic definitions with loops and classes to remove redundancy (NFC) #139611

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom
Open
wants to merge 3 commits into
base: main
Choose a base branch
from
Open

[NVPTX] Refactor intrinsic definitions with loops and classes to remove redundancy (NFC) #139611

wants to merge 3 commits into from
+1,352 −2,844

Conversation

AlexMaclean
Copy link
Member

No description provided.

@AlexMaclean AlexMaclean self-assigned this May 12, 2025
@llvmbot
Copy link
Member

llvmbot commented May 12, 2025

@llvm/pr-subscribers-backend-nvptx

Author: Alex MacLean (AlexMaclean)

Changes

Patch is 286.33 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/139611.diff

1 Files Affected:

  • (modified) llvm/include/llvm/IR/IntrinsicsNVVM.td (+1352-2844)
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 640fdf3f86326..93c7db40b1c6e 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -686,10 +686,7 @@ class SHFL_INFO<bit sync, string mode, string type, bit return_pred> {
                   # !if(return_pred, "p", "");
 
   string Name = "int_nvvm_shfl_" # Suffix;
-  string Builtin = "__nvvm_shfl_" # Suffix;
-  string IntrName = "llvm.nvvm.shfl." # !subst("_",".", Suffix);
   bit withGccBuiltin = !not(return_pred);
-  bit withoutGccBuiltin = return_pred;
   LLVMType OpType = !cond(
     !eq(type,"i32"): llvm_i32_ty,
     !eq(type,"f32"): llvm_float_ty);
@@ -794,8 +791,6 @@ class NVVM_TCGEN05_LDST_NAME<string Op, string Shape, int Num> {
   string record = !subst(".", "_",
                   !subst("llvm.", "int_", intr));
 }
-
-
 class NVVM_TCGEN05_LDST_ACCESS_SIZE<string Shape, int Num> {
   int shift = !cond(!eq(Shape, "16x128b"): 1,
                     !eq(Shape, "16x256b"): 2,
@@ -815,12 +810,18 @@ class NVVM_TCGEN05_LDST_ACCESS_SIZE<string Shape, int Num> {
                         true : llvm_void_ty);
 }
 
+class NVVMBuiltin :
+  ClangBuiltin<!strconcat("__", !substr(NAME, !size("int_")))> {
+    assert !eq(!substr(NAME, 0, !size("int_nvvm_")), "int_nvvm_"),
+           "NVVMBuiltin must be a NVVM intrinsic starting with 'int_nvvm_'";
+}
+
 let TargetPrefix = "nvvm" in {
-  def int_nvvm_prmt : ClangBuiltin<"__nvvm_prmt">,
+  def int_nvvm_prmt : NVVMBuiltin,
       DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
         [IntrNoMem, IntrSpeculatable]>;
 
-  def int_nvvm_nanosleep : ClangBuiltin<"__nvvm_nanosleep">,
+  def int_nvvm_nanosleep : NVVMBuiltin,
       DefaultAttrsIntrinsic<[], [llvm_i32_ty],
                             [IntrConvergent, IntrNoMem, IntrHasSideEffects]>;
 
@@ -829,52 +830,34 @@ let TargetPrefix = "nvvm" in {
 //
 
   foreach operation = ["min", "max"] in {
-    def int_nvvm_f # operation # _d :
-      ClangBuiltin<!strconcat("__nvvm_f", operation, "_d")>,
+    def int_nvvm_f # operation # _d : NVVMBuiltin,
       DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
         [IntrNoMem, IntrSpeculatable, Commutative]>;
 
-    foreach variant = ["_f", "_ftz_f", "_nan_f", "_ftz_nan_f",
-      "_xorsign_abs_f", "_ftz_xorsign_abs_f", "_nan_xorsign_abs_f",
-      "_ftz_nan_xorsign_abs_f"] in {
-      def int_nvvm_f # operation # variant :
-        ClangBuiltin<!strconcat("__nvvm_f", operation, variant)>,
-        DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-          [IntrNoMem, IntrSpeculatable, Commutative]>;
-    }
+    foreach variant = ["", "_xorsign_abs"] in {
+      foreach nan = ["", "_nan"] in {
+        foreach ftz = ["", "_ftz"] in {
+          def int_nvvm_f # operation # ftz # nan # variant # _f : NVVMBuiltin,
+            DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
+              [IntrNoMem, IntrSpeculatable, Commutative]>;
 
-    foreach variant = ["_f16", "_ftz_f16", "_nan_f16", "_ftz_nan_f16",
-      "_xorsign_abs_f16", "_ftz_xorsign_abs_f16", "_nan_xorsign_abs_f16",
-      "_ftz_nan_xorsign_abs_f16"] in {
-      def int_nvvm_f # operation # variant :
-        DefaultAttrsIntrinsic<[llvm_half_ty], [llvm_half_ty, llvm_half_ty],
-          [IntrNoMem, IntrSpeculatable, Commutative]>;
-    }
+          def int_nvvm_f # operation # ftz # nan # variant # _f16 :
+            DefaultAttrsIntrinsic<[llvm_half_ty], [llvm_half_ty, llvm_half_ty],
+              [IntrNoMem, IntrSpeculatable, Commutative]>;
 
-    foreach variant = ["_f16x2", "_ftz_f16x2", "_nan_f16x2",
-      "_ftz_nan_f16x2", "_xorsign_abs_f16x2", "_ftz_xorsign_abs_f16x2",
-      "_nan_xorsign_abs_f16x2", "_ftz_nan_xorsign_abs_f16x2"] in {
-      def int_nvvm_f # operation # variant :
-        DefaultAttrsIntrinsic<[llvm_v2f16_ty], [llvm_v2f16_ty, llvm_v2f16_ty],
-          [IntrNoMem, IntrSpeculatable, Commutative]>;
-    }
+          def int_nvvm_f # operation # ftz # nan # variant # _f16x2 :
+            DefaultAttrsIntrinsic<[llvm_v2f16_ty], [llvm_v2f16_ty, llvm_v2f16_ty],
+              [IntrNoMem, IntrSpeculatable, Commutative]>;
 
-    foreach variant = ["_bf16", "_ftz_bf16", "_nan_bf16", "_ftz_nan_bf16",
-      "_xorsign_abs_bf16", "_ftz_xorsign_abs_bf16", "_nan_xorsign_abs_bf16",
-      "_ftz_nan_xorsign_abs_bf16"] in {
-      def int_nvvm_f # operation # variant :
-        ClangBuiltin<!strconcat("__nvvm_f", operation, variant)>,
-        DefaultAttrsIntrinsic<[llvm_bfloat_ty], [llvm_bfloat_ty, llvm_bfloat_ty],
-          [IntrNoMem, IntrSpeculatable, Commutative]>;
-    }
+          def int_nvvm_f # operation # ftz # nan # variant # _bf16 : NVVMBuiltin,
+            DefaultAttrsIntrinsic<[llvm_bfloat_ty], [llvm_bfloat_ty, llvm_bfloat_ty],
+              [IntrNoMem, IntrSpeculatable, Commutative]>;
 
-    foreach variant = ["_bf16x2", "_ftz_bf16x2", "_nan_bf16x2",
-      "_ftz_nan_bf16x2", "_xorsign_abs_bf16x2", "_ftz_xorsign_abs_bf16x2",
-      "_nan_xorsign_abs_bf16x2", "_ftz_nan_xorsign_abs_bf16x2"]  in {
-      def int_nvvm_f # operation # variant :
-        ClangBuiltin<!strconcat("__nvvm_f", operation, variant)>,
-        DefaultAttrsIntrinsic<[llvm_v2bf16_ty], [llvm_v2bf16_ty, llvm_v2bf16_ty],
-          [IntrNoMem, IntrSpeculatable, Commutative]>;
+          def int_nvvm_f # operation # ftz # nan # variant # _bf16x2 : NVVMBuiltin,
+            DefaultAttrsIntrinsic<[llvm_v2bf16_ty], [llvm_v2bf16_ty, llvm_v2bf16_ty],
+              [IntrNoMem, IntrSpeculatable, Commutative]>;
+        }
+      }
     }
   }
 
@@ -882,315 +865,208 @@ let TargetPrefix = "nvvm" in {
 // Multiplication
 //
 
-  def int_nvvm_mulhi_s : ClangBuiltin<"__nvvm_mulhi_s">,
-      DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mulhi_us : ClangBuiltin<"__nvvm_mulhi_us">,
-      DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
+  foreach sign = ["", "u"] in {
+    def int_nvvm_mulhi_ # sign # s : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty],
+          [IntrNoMem, IntrSpeculatable, Commutative]>;
 
-  def int_nvvm_mulhi_i : ClangBuiltin<"__nvvm_mulhi_i">,
-      DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mulhi_ui : ClangBuiltin<"__nvvm_mulhi_ui">,
-      DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
+    def int_nvvm_mulhi_ # sign # i : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
+          [IntrNoMem, IntrSpeculatable, Commutative]>;
 
-  def int_nvvm_mulhi_ll : ClangBuiltin<"__nvvm_mulhi_ll">,
-      DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mulhi_ull : ClangBuiltin<"__nvvm_mulhi_ull">,
-      DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
+    def int_nvvm_mulhi_ # sign # ll : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty],
+          [IntrNoMem, IntrSpeculatable, Commutative]>;
 
-  def int_nvvm_mul_rn_ftz_f : ClangBuiltin<"__nvvm_mul_rn_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mul_rn_f : ClangBuiltin<"__nvvm_mul_rn_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mul_rz_ftz_f : ClangBuiltin<"__nvvm_mul_rz_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mul_rz_f : ClangBuiltin<"__nvvm_mul_rz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mul_rm_ftz_f : ClangBuiltin<"__nvvm_mul_rm_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mul_rm_f : ClangBuiltin<"__nvvm_mul_rm_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mul_rp_ftz_f : ClangBuiltin<"__nvvm_mul_rp_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mul_rp_f : ClangBuiltin<"__nvvm_mul_rp_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
+    def int_nvvm_mul24_ # sign # i : NVVMBuiltin,
+      DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
         [IntrNoMem, IntrSpeculatable, Commutative]>;
+  }
 
-  def int_nvvm_mul_rn_d : ClangBuiltin<"__nvvm_mul_rn_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mul_rz_d : ClangBuiltin<"__nvvm_mul_rz_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mul_rm_d : ClangBuiltin<"__nvvm_mul_rm_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mul_rp_d : ClangBuiltin<"__nvvm_mul_rp_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
+  foreach rnd = ["rn", "rz", "rm", "rp"] in {
+    foreach ftz = ["", "_ftz"] in
+      def int_nvvm_mul_ # rnd # ftz # _f : NVVMBuiltin,
+          DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
+            [IntrNoMem, IntrSpeculatable, Commutative]>;
 
-  def int_nvvm_mul24_i : ClangBuiltin<"__nvvm_mul24_i">,
-      DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mul24_ui : ClangBuiltin<"__nvvm_mul24_ui">,
-      DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
+    def int_nvvm_mul_ # rnd # _d : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
+          [IntrNoMem, IntrSpeculatable, Commutative]>;
+  }
 
 //
 // Div
 //
 
-  def int_nvvm_div_approx_ftz_f : ClangBuiltin<"__nvvm_div_approx_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
-  def int_nvvm_div_approx_f : ClangBuiltin<"__nvvm_div_approx_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
-
-  def int_nvvm_div_rn_ftz_f : ClangBuiltin<"__nvvm_div_rn_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
-  def int_nvvm_div_rn_f : ClangBuiltin<"__nvvm_div_rn_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
-
-  def int_nvvm_div_rz_ftz_f : ClangBuiltin<"__nvvm_div_rz_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
-  def int_nvvm_div_rz_f : ClangBuiltin<"__nvvm_div_rz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
-
-  def int_nvvm_div_rm_ftz_f : ClangBuiltin<"__nvvm_div_rm_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
-  def int_nvvm_div_rm_f : ClangBuiltin<"__nvvm_div_rm_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
-
-  def int_nvvm_div_rp_ftz_f : ClangBuiltin<"__nvvm_div_rp_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
-  def int_nvvm_div_rp_f : ClangBuiltin<"__nvvm_div_rp_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
-
-  def int_nvvm_div_rn_d : ClangBuiltin<"__nvvm_div_rn_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
-        [IntrNoMem]>;
-  def int_nvvm_div_rz_d : ClangBuiltin<"__nvvm_div_rz_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
-        [IntrNoMem]>;
-  def int_nvvm_div_rm_d : ClangBuiltin<"__nvvm_div_rm_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
-        [IntrNoMem]>;
-  def int_nvvm_div_rp_d : ClangBuiltin<"__nvvm_div_rp_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
-        [IntrNoMem]>;
+  foreach ftz = ["", "_ftz"] in {
+    def int_nvvm_div_approx # ftz # _f : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
+          [IntrNoMem]>;
 
-  def int_nvvm_div_full : ClangBuiltin<"__nvvm_div_full">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
-  def int_nvvm_div_full_ftz : ClangBuiltin<"__nvvm_div_full_ftz">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
+    def int_nvvm_div_full # ftz : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
+          [IntrNoMem]>;
+  }
+
+  foreach rnd = ["rn", "rz", "rm", "rp"] in {
+    foreach ftz = ["", "_ftz"] in
+      def int_nvvm_div_ # rnd # ftz # _f : NVVMBuiltin,
+          DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
+            [IntrNoMem]>;
+
+    def int_nvvm_div_ # rnd # _d : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
+          [IntrNoMem]>;
+  }
 
 //
 // Sad
 //
 
-  def int_nvvm_sad_s : ClangBuiltin<"__nvvm_sad_s">,
-      DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
-        [IntrNoMem, Commutative, IntrSpeculatable]>;
-  def int_nvvm_sad_us : ClangBuiltin<"__nvvm_sad_us">,
-      DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
-        [IntrNoMem, Commutative, IntrSpeculatable]>;
-
-  def int_nvvm_sad_i : ClangBuiltin<"__nvvm_sad_i">,
-      DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
-        [IntrNoMem, Commutative, IntrSpeculatable]>;
-  def int_nvvm_sad_ui : ClangBuiltin<"__nvvm_sad_ui">,
-      DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
-        [IntrNoMem, Commutative, IntrSpeculatable]>;
+  foreach sign = ["", "u"] in {
+    def int_nvvm_sad_ # sign # s : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
+          [IntrNoMem, Commutative, IntrSpeculatable]>;
 
-  def int_nvvm_sad_ll : ClangBuiltin<"__nvvm_sad_ll">,
-      DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i64_ty],
-        [IntrNoMem, Commutative, IntrSpeculatable]>;
-  def int_nvvm_sad_ull : ClangBuiltin<"__nvvm_sad_ull">,
-      DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i64_ty],
-        [IntrNoMem, Commutative, IntrSpeculatable]>;
+    def int_nvvm_sad_ # sign # i : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
+          [IntrNoMem, Commutative, IntrSpeculatable]>;
 
+    def int_nvvm_sad_ # sign # ll : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i64_ty],
+          [IntrNoMem, Commutative, IntrSpeculatable]>;
+  }
 
 //
 // Floor  Ceil
 //
 
-  def int_nvvm_floor_ftz_f : ClangBuiltin<"__nvvm_floor_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_floor_f : ClangBuiltin<"__nvvm_floor_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_floor_d : ClangBuiltin<"__nvvm_floor_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem, IntrSpeculatable]>;
-
-  def int_nvvm_ceil_ftz_f : ClangBuiltin<"__nvvm_ceil_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_ceil_f : ClangBuiltin<"__nvvm_ceil_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_ceil_d : ClangBuiltin<"__nvvm_ceil_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem, IntrSpeculatable]>;
+  foreach op = ["floor", "ceil"] in {
+    foreach ftz = ["", "_ftz"] in
+      def int_nvvm_ # op # ftz # _f : NVVMBuiltin,
+          DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
+    def int_nvvm_ # op # _d : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem, IntrSpeculatable]>;
+  }
 
 //
 // Abs
 //
 
-  def int_nvvm_fabs_ftz :
+  foreach ftz = ["", "_ftz"] in
+    def int_nvvm_fabs # ftz :
       DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>],
                             [IntrNoMem, IntrSpeculatable]>;
 
-  def int_nvvm_fabs :
-      DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>],
-                            [IntrNoMem, IntrSpeculatable]>;
 //
 // Abs, Neg bf16, bf16x2
 //
 
-  foreach unary = ["neg"] in {
-    def int_nvvm_ # unary # _bf16 :
-      ClangBuiltin<!strconcat("__nvvm_", unary, "_bf16")>,
-      DefaultAttrsIntrinsic<[llvm_bfloat_ty], [llvm_bfloat_ty], [IntrNoMem]>;
-    def int_nvvm_ # unary # _bf16x2 :
-      ClangBuiltin<!strconcat("__nvvm_", unary, "_bf16x2")>,
-      DefaultAttrsIntrinsic<[llvm_v2bf16_ty], [llvm_v2bf16_ty], [IntrNoMem]>;
-  }
+  def int_nvvm_neg_bf16 : NVVMBuiltin,
+    DefaultAttrsIntrinsic<[llvm_bfloat_ty], [llvm_bfloat_ty], [IntrNoMem]>;
+  def int_nvvm_neg_bf16x2 : NVVMBuiltin,
+    DefaultAttrsIntrinsic<[llvm_v2bf16_ty], [llvm_v2bf16_ty], [IntrNoMem]>;
 
 //
 // Round
 //
 
-  def int_nvvm_round_ftz_f : ClangBuiltin<"__nvvm_round_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_round_f : ClangBuiltin<"__nvvm_round_f">,
+  foreach ftz = ["", "_ftz"] in
+    def int_nvvm_round # ftz # _f : NVVMBuiltin,
       DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
 
-  def int_nvvm_round_d : ClangBuiltin<"__nvvm_round_d">,
+  def int_nvvm_round_d : NVVMBuiltin,
       DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem, IntrSpeculatable]>;
 
 //
 // Trunc
 //
 
-  def int_nvvm_trunc_ftz_f : ClangBuiltin<"__nvvm_trunc_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_trunc_f : ClangBuiltin<"__nvvm_trunc_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
+  foreach ftz = ["", "_ftz"] in
+    def int_nvvm_trunc # ftz # _f : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
 
-  def int_nvvm_trunc_d : ClangBuiltin<"__nvvm_trunc_d">,
+  def int_nvvm_trunc_d : NVVMBuiltin,
       DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem, IntrSpeculatable]>;
 
 //
 // Saturate
 //
 
-  def int_nvvm_saturate_ftz_f : ClangBuiltin<"__nvvm_saturate_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_saturate_f : ClangBuiltin<"__nvvm_saturate_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
+  foreach ftz = ["", "_ftz"] in
+    def int_nvvm_saturate # ftz # _f : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
 
-  def int_nvvm_saturate_d : ClangBuiltin<"__nvvm_saturate_d">,
+  def int_nvvm_saturate_d : NVVMBuiltin,
       DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem, I...
[truncated]

@llvmbot
Copy link
Member

llvmbot commented May 12, 2025

@llvm/pr-subscribers-llvm-ir

Author: Alex MacLean (AlexMaclean)

Changes

Patch is 286.33 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/139611.diff

1 Files Affected:

  • (modified) llvm/include/llvm/IR/IntrinsicsNVVM.td (+1352-2844)
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 640fdf3f86326..93c7db40b1c6e 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -686,10 +686,7 @@ class SHFL_INFO<bit sync, string mode, string type, bit return_pred> {
                   # !if(return_pred, "p", "");
 
   string Name = "int_nvvm_shfl_" # Suffix;
-  string Builtin = "__nvvm_shfl_" # Suffix;
-  string IntrName = "llvm.nvvm.shfl." # !subst("_",".", Suffix);
   bit withGccBuiltin = !not(return_pred);
-  bit withoutGccBuiltin = return_pred;
   LLVMType OpType = !cond(
     !eq(type,"i32"): llvm_i32_ty,
     !eq(type,"f32"): llvm_float_ty);
@@ -794,8 +791,6 @@ class NVVM_TCGEN05_LDST_NAME<string Op, string Shape, int Num> {
   string record = !subst(".", "_",
                   !subst("llvm.", "int_", intr));
 }
-
-
 class NVVM_TCGEN05_LDST_ACCESS_SIZE<string Shape, int Num> {
   int shift = !cond(!eq(Shape, "16x128b"): 1,
                     !eq(Shape, "16x256b"): 2,
@@ -815,12 +810,18 @@ class NVVM_TCGEN05_LDST_ACCESS_SIZE<string Shape, int Num> {
                         true : llvm_void_ty);
 }
 
+class NVVMBuiltin :
+  ClangBuiltin<!strconcat("__", !substr(NAME, !size("int_")))> {
+    assert !eq(!substr(NAME, 0, !size("int_nvvm_")), "int_nvvm_"),
+           "NVVMBuiltin must be a NVVM intrinsic starting with 'int_nvvm_'";
+}
+
 let TargetPrefix = "nvvm" in {
-  def int_nvvm_prmt : ClangBuiltin<"__nvvm_prmt">,
+  def int_nvvm_prmt : NVVMBuiltin,
       DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
         [IntrNoMem, IntrSpeculatable]>;
 
-  def int_nvvm_nanosleep : ClangBuiltin<"__nvvm_nanosleep">,
+  def int_nvvm_nanosleep : NVVMBuiltin,
       DefaultAttrsIntrinsic<[], [llvm_i32_ty],
                             [IntrConvergent, IntrNoMem, IntrHasSideEffects]>;
 
@@ -829,52 +830,34 @@ let TargetPrefix = "nvvm" in {
 //
 
   foreach operation = ["min", "max"] in {
-    def int_nvvm_f # operation # _d :
-      ClangBuiltin<!strconcat("__nvvm_f", operation, "_d")>,
+    def int_nvvm_f # operation # _d : NVVMBuiltin,
       DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
         [IntrNoMem, IntrSpeculatable, Commutative]>;
 
-    foreach variant = ["_f", "_ftz_f", "_nan_f", "_ftz_nan_f",
-      "_xorsign_abs_f", "_ftz_xorsign_abs_f", "_nan_xorsign_abs_f",
-      "_ftz_nan_xorsign_abs_f"] in {
-      def int_nvvm_f # operation # variant :
-        ClangBuiltin<!strconcat("__nvvm_f", operation, variant)>,
-        DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-          [IntrNoMem, IntrSpeculatable, Commutative]>;
-    }
+    foreach variant = ["", "_xorsign_abs"] in {
+      foreach nan = ["", "_nan"] in {
+        foreach ftz = ["", "_ftz"] in {
+          def int_nvvm_f # operation # ftz # nan # variant # _f : NVVMBuiltin,
+            DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
+              [IntrNoMem, IntrSpeculatable, Commutative]>;
 
-    foreach variant = ["_f16", "_ftz_f16", "_nan_f16", "_ftz_nan_f16",
-      "_xorsign_abs_f16", "_ftz_xorsign_abs_f16", "_nan_xorsign_abs_f16",
-      "_ftz_nan_xorsign_abs_f16"] in {
-      def int_nvvm_f # operation # variant :
-        DefaultAttrsIntrinsic<[llvm_half_ty], [llvm_half_ty, llvm_half_ty],
-          [IntrNoMem, IntrSpeculatable, Commutative]>;
-    }
+          def int_nvvm_f # operation # ftz # nan # variant # _f16 :
+            DefaultAttrsIntrinsic<[llvm_half_ty], [llvm_half_ty, llvm_half_ty],
+              [IntrNoMem, IntrSpeculatable, Commutative]>;
 
-    foreach variant = ["_f16x2", "_ftz_f16x2", "_nan_f16x2",
-      "_ftz_nan_f16x2", "_xorsign_abs_f16x2", "_ftz_xorsign_abs_f16x2",
-      "_nan_xorsign_abs_f16x2", "_ftz_nan_xorsign_abs_f16x2"] in {
-      def int_nvvm_f # operation # variant :
-        DefaultAttrsIntrinsic<[llvm_v2f16_ty], [llvm_v2f16_ty, llvm_v2f16_ty],
-          [IntrNoMem, IntrSpeculatable, Commutative]>;
-    }
+          def int_nvvm_f # operation # ftz # nan # variant # _f16x2 :
+            DefaultAttrsIntrinsic<[llvm_v2f16_ty], [llvm_v2f16_ty, llvm_v2f16_ty],
+              [IntrNoMem, IntrSpeculatable, Commutative]>;
 
-    foreach variant = ["_bf16", "_ftz_bf16", "_nan_bf16", "_ftz_nan_bf16",
-      "_xorsign_abs_bf16", "_ftz_xorsign_abs_bf16", "_nan_xorsign_abs_bf16",
-      "_ftz_nan_xorsign_abs_bf16"] in {
-      def int_nvvm_f # operation # variant :
-        ClangBuiltin<!strconcat("__nvvm_f", operation, variant)>,
-        DefaultAttrsIntrinsic<[llvm_bfloat_ty], [llvm_bfloat_ty, llvm_bfloat_ty],
-          [IntrNoMem, IntrSpeculatable, Commutative]>;
-    }
+          def int_nvvm_f # operation # ftz # nan # variant # _bf16 : NVVMBuiltin,
+            DefaultAttrsIntrinsic<[llvm_bfloat_ty], [llvm_bfloat_ty, llvm_bfloat_ty],
+              [IntrNoMem, IntrSpeculatable, Commutative]>;
 
-    foreach variant = ["_bf16x2", "_ftz_bf16x2", "_nan_bf16x2",
-      "_ftz_nan_bf16x2", "_xorsign_abs_bf16x2", "_ftz_xorsign_abs_bf16x2",
-      "_nan_xorsign_abs_bf16x2", "_ftz_nan_xorsign_abs_bf16x2"]  in {
-      def int_nvvm_f # operation # variant :
-        ClangBuiltin<!strconcat("__nvvm_f", operation, variant)>,
-        DefaultAttrsIntrinsic<[llvm_v2bf16_ty], [llvm_v2bf16_ty, llvm_v2bf16_ty],
-          [IntrNoMem, IntrSpeculatable, Commutative]>;
+          def int_nvvm_f # operation # ftz # nan # variant # _bf16x2 : NVVMBuiltin,
+            DefaultAttrsIntrinsic<[llvm_v2bf16_ty], [llvm_v2bf16_ty, llvm_v2bf16_ty],
+              [IntrNoMem, IntrSpeculatable, Commutative]>;
+        }
+      }
     }
   }
 
@@ -882,315 +865,208 @@ let TargetPrefix = "nvvm" in {
 // Multiplication
 //
 
-  def int_nvvm_mulhi_s : ClangBuiltin<"__nvvm_mulhi_s">,
-      DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mulhi_us : ClangBuiltin<"__nvvm_mulhi_us">,
-      DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
+  foreach sign = ["", "u"] in {
+    def int_nvvm_mulhi_ # sign # s : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty],
+          [IntrNoMem, IntrSpeculatable, Commutative]>;
 
-  def int_nvvm_mulhi_i : ClangBuiltin<"__nvvm_mulhi_i">,
-      DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mulhi_ui : ClangBuiltin<"__nvvm_mulhi_ui">,
-      DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
+    def int_nvvm_mulhi_ # sign # i : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
+          [IntrNoMem, IntrSpeculatable, Commutative]>;
 
-  def int_nvvm_mulhi_ll : ClangBuiltin<"__nvvm_mulhi_ll">,
-      DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mulhi_ull : ClangBuiltin<"__nvvm_mulhi_ull">,
-      DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
+    def int_nvvm_mulhi_ # sign # ll : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty],
+          [IntrNoMem, IntrSpeculatable, Commutative]>;
 
-  def int_nvvm_mul_rn_ftz_f : ClangBuiltin<"__nvvm_mul_rn_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mul_rn_f : ClangBuiltin<"__nvvm_mul_rn_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mul_rz_ftz_f : ClangBuiltin<"__nvvm_mul_rz_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mul_rz_f : ClangBuiltin<"__nvvm_mul_rz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mul_rm_ftz_f : ClangBuiltin<"__nvvm_mul_rm_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mul_rm_f : ClangBuiltin<"__nvvm_mul_rm_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mul_rp_ftz_f : ClangBuiltin<"__nvvm_mul_rp_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mul_rp_f : ClangBuiltin<"__nvvm_mul_rp_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
+    def int_nvvm_mul24_ # sign # i : NVVMBuiltin,
+      DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
         [IntrNoMem, IntrSpeculatable, Commutative]>;
+  }
 
-  def int_nvvm_mul_rn_d : ClangBuiltin<"__nvvm_mul_rn_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mul_rz_d : ClangBuiltin<"__nvvm_mul_rz_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mul_rm_d : ClangBuiltin<"__nvvm_mul_rm_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mul_rp_d : ClangBuiltin<"__nvvm_mul_rp_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
+  foreach rnd = ["rn", "rz", "rm", "rp"] in {
+    foreach ftz = ["", "_ftz"] in
+      def int_nvvm_mul_ # rnd # ftz # _f : NVVMBuiltin,
+          DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
+            [IntrNoMem, IntrSpeculatable, Commutative]>;
 
-  def int_nvvm_mul24_i : ClangBuiltin<"__nvvm_mul24_i">,
-      DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_mul24_ui : ClangBuiltin<"__nvvm_mul24_ui">,
-      DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
+    def int_nvvm_mul_ # rnd # _d : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
+          [IntrNoMem, IntrSpeculatable, Commutative]>;
+  }
 
 //
 // Div
 //
 
-  def int_nvvm_div_approx_ftz_f : ClangBuiltin<"__nvvm_div_approx_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
-  def int_nvvm_div_approx_f : ClangBuiltin<"__nvvm_div_approx_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
-
-  def int_nvvm_div_rn_ftz_f : ClangBuiltin<"__nvvm_div_rn_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
-  def int_nvvm_div_rn_f : ClangBuiltin<"__nvvm_div_rn_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
-
-  def int_nvvm_div_rz_ftz_f : ClangBuiltin<"__nvvm_div_rz_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
-  def int_nvvm_div_rz_f : ClangBuiltin<"__nvvm_div_rz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
-
-  def int_nvvm_div_rm_ftz_f : ClangBuiltin<"__nvvm_div_rm_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
-  def int_nvvm_div_rm_f : ClangBuiltin<"__nvvm_div_rm_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
-
-  def int_nvvm_div_rp_ftz_f : ClangBuiltin<"__nvvm_div_rp_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
-  def int_nvvm_div_rp_f : ClangBuiltin<"__nvvm_div_rp_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
-
-  def int_nvvm_div_rn_d : ClangBuiltin<"__nvvm_div_rn_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
-        [IntrNoMem]>;
-  def int_nvvm_div_rz_d : ClangBuiltin<"__nvvm_div_rz_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
-        [IntrNoMem]>;
-  def int_nvvm_div_rm_d : ClangBuiltin<"__nvvm_div_rm_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
-        [IntrNoMem]>;
-  def int_nvvm_div_rp_d : ClangBuiltin<"__nvvm_div_rp_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
-        [IntrNoMem]>;
+  foreach ftz = ["", "_ftz"] in {
+    def int_nvvm_div_approx # ftz # _f : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
+          [IntrNoMem]>;
 
-  def int_nvvm_div_full : ClangBuiltin<"__nvvm_div_full">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
-  def int_nvvm_div_full_ftz : ClangBuiltin<"__nvvm_div_full_ftz">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-        [IntrNoMem]>;
+    def int_nvvm_div_full # ftz : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
+          [IntrNoMem]>;
+  }
+
+  foreach rnd = ["rn", "rz", "rm", "rp"] in {
+    foreach ftz = ["", "_ftz"] in
+      def int_nvvm_div_ # rnd # ftz # _f : NVVMBuiltin,
+          DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
+            [IntrNoMem]>;
+
+    def int_nvvm_div_ # rnd # _d : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
+          [IntrNoMem]>;
+  }
 
 //
 // Sad
 //
 
-  def int_nvvm_sad_s : ClangBuiltin<"__nvvm_sad_s">,
-      DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
-        [IntrNoMem, Commutative, IntrSpeculatable]>;
-  def int_nvvm_sad_us : ClangBuiltin<"__nvvm_sad_us">,
-      DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
-        [IntrNoMem, Commutative, IntrSpeculatable]>;
-
-  def int_nvvm_sad_i : ClangBuiltin<"__nvvm_sad_i">,
-      DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
-        [IntrNoMem, Commutative, IntrSpeculatable]>;
-  def int_nvvm_sad_ui : ClangBuiltin<"__nvvm_sad_ui">,
-      DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
-        [IntrNoMem, Commutative, IntrSpeculatable]>;
+  foreach sign = ["", "u"] in {
+    def int_nvvm_sad_ # sign # s : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
+          [IntrNoMem, Commutative, IntrSpeculatable]>;
 
-  def int_nvvm_sad_ll : ClangBuiltin<"__nvvm_sad_ll">,
-      DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i64_ty],
-        [IntrNoMem, Commutative, IntrSpeculatable]>;
-  def int_nvvm_sad_ull : ClangBuiltin<"__nvvm_sad_ull">,
-      DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i64_ty],
-        [IntrNoMem, Commutative, IntrSpeculatable]>;
+    def int_nvvm_sad_ # sign # i : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
+          [IntrNoMem, Commutative, IntrSpeculatable]>;
 
+    def int_nvvm_sad_ # sign # ll : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i64_ty],
+          [IntrNoMem, Commutative, IntrSpeculatable]>;
+  }
 
 //
 // Floor  Ceil
 //
 
-  def int_nvvm_floor_ftz_f : ClangBuiltin<"__nvvm_floor_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_floor_f : ClangBuiltin<"__nvvm_floor_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_floor_d : ClangBuiltin<"__nvvm_floor_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem, IntrSpeculatable]>;
-
-  def int_nvvm_ceil_ftz_f : ClangBuiltin<"__nvvm_ceil_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_ceil_f : ClangBuiltin<"__nvvm_ceil_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_ceil_d : ClangBuiltin<"__nvvm_ceil_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem, IntrSpeculatable]>;
+  foreach op = ["floor", "ceil"] in {
+    foreach ftz = ["", "_ftz"] in
+      def int_nvvm_ # op # ftz # _f : NVVMBuiltin,
+          DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
+    def int_nvvm_ # op # _d : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem, IntrSpeculatable]>;
+  }
 
 //
 // Abs
 //
 
-  def int_nvvm_fabs_ftz :
+  foreach ftz = ["", "_ftz"] in
+    def int_nvvm_fabs # ftz :
       DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>],
                             [IntrNoMem, IntrSpeculatable]>;
 
-  def int_nvvm_fabs :
-      DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>],
-                            [IntrNoMem, IntrSpeculatable]>;
 //
 // Abs, Neg bf16, bf16x2
 //
 
-  foreach unary = ["neg"] in {
-    def int_nvvm_ # unary # _bf16 :
-      ClangBuiltin<!strconcat("__nvvm_", unary, "_bf16")>,
-      DefaultAttrsIntrinsic<[llvm_bfloat_ty], [llvm_bfloat_ty], [IntrNoMem]>;
-    def int_nvvm_ # unary # _bf16x2 :
-      ClangBuiltin<!strconcat("__nvvm_", unary, "_bf16x2")>,
-      DefaultAttrsIntrinsic<[llvm_v2bf16_ty], [llvm_v2bf16_ty], [IntrNoMem]>;
-  }
+  def int_nvvm_neg_bf16 : NVVMBuiltin,
+    DefaultAttrsIntrinsic<[llvm_bfloat_ty], [llvm_bfloat_ty], [IntrNoMem]>;
+  def int_nvvm_neg_bf16x2 : NVVMBuiltin,
+    DefaultAttrsIntrinsic<[llvm_v2bf16_ty], [llvm_v2bf16_ty], [IntrNoMem]>;
 
 //
 // Round
 //
 
-  def int_nvvm_round_ftz_f : ClangBuiltin<"__nvvm_round_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_round_f : ClangBuiltin<"__nvvm_round_f">,
+  foreach ftz = ["", "_ftz"] in
+    def int_nvvm_round # ftz # _f : NVVMBuiltin,
       DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
 
-  def int_nvvm_round_d : ClangBuiltin<"__nvvm_round_d">,
+  def int_nvvm_round_d : NVVMBuiltin,
       DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem, IntrSpeculatable]>;
 
 //
 // Trunc
 //
 
-  def int_nvvm_trunc_ftz_f : ClangBuiltin<"__nvvm_trunc_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_trunc_f : ClangBuiltin<"__nvvm_trunc_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
+  foreach ftz = ["", "_ftz"] in
+    def int_nvvm_trunc # ftz # _f : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
 
-  def int_nvvm_trunc_d : ClangBuiltin<"__nvvm_trunc_d">,
+  def int_nvvm_trunc_d : NVVMBuiltin,
       DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem, IntrSpeculatable]>;
 
 //
 // Saturate
 //
 
-  def int_nvvm_saturate_ftz_f : ClangBuiltin<"__nvvm_saturate_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_saturate_f : ClangBuiltin<"__nvvm_saturate_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
+  foreach ftz = ["", "_ftz"] in
+    def int_nvvm_saturate # ftz # _f : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
 
-  def int_nvvm_saturate_d : ClangBuiltin<"__nvvm_saturate_d">,
+  def int_nvvm_saturate_d : NVVMBuiltin,
       DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem, I...
[truncated]

Artem-B
Artem-B approved these changes May 13, 2025
View reviewed changes
Copy link
Member

@Artem-B Artem-B left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

1500 lines of tablegen gone while making the file more readable is a very nice cleanup.

llvm/include/llvm/IR/IntrinsicsNVVM.td
Comment on lines +1237 to +1238
}
}
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Nit: a comment mentioning which loop is closed by each '}' would be helpful.

llvm/include/llvm/IR/IntrinsicsNVVM.td
def NAME: Intrinsic<[],[llvm_shared_ptr_ty, llvm_global_ptr_ty],
[IntrArgMemOnly, IntrNoCallback, NoAlias<ArgIndex<0>>, NoAlias<ArgIndex<1>>,
WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<1>>],
"llvm.nvvm.cp.async." # cc # ".shared.global." # n>;
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Was the explicit name removed because it matches the default LLVM intrinsic name derived from the record name?

llvm/include/llvm/IR/IntrinsicsNVVM.td

// For getting the handle from a texture or surface variable
def int_nvvm_texsurf_handle
: Intrinsic<[llvm_i64_ty], [llvm_metadata_ty, llvm_anyptr_ty],
[IntrNoMem], "llvm.nvvm.texsurf.handle">;
[IntrNoMem]>;
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Nit: now that we have shorter list of record arguments, they can fit on one line. I think there are still a few more such short leftovers in the diff that could be joined with the line above them.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@Artem-B Artem-B Artem-B approved these changes

@justinfargnoli justinfargnoli Awaiting requested review from justinfargnoli

Assignees

@AlexMaclean AlexMaclean

Labels
backend:NVPTX llvm:ir
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
3 participants
@AlexMaclean @llvmbot @Artem-B