Skip to content

Navigation Menu

Sign in

Search code, repositories, users, issues, pull requests...

Provide feedback

We read every piece of feedback, and take your input very seriously.

Saved searches

Use saved searches to filter your results more quickly
Appearance settings

[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
Merged
merged 4 commits into from
May 13, 2025
Merged

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

merged 4 commits into from
May 13, 2025

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 Outdated
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.

Copy link
Member Author

Choose a reason for hiding this comment

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

Added. I've added such comments on each of the longer loops that I added as well.

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?

Copy link
Member Author

Choose a reason for hiding this comment

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

Yep. I think it only makes sense to explicitly specify the name when it differs from the name which will be derived by default. Otherwise we're just adding the complexity of composing essentially the same name twice.

llvm/include/llvm/IR/IntrinsicsNVVM.td Outdated

// 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.

Copy link
Member Author

Choose a reason for hiding this comment

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

Fixed! I've fixed a few more cases I could find, but I assume there are some I missed as well. It's a shame there isn't an automatic formatter for td.

Copy link
Member

Choose a reason for hiding this comment

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

Someone was supposedly working on one:
https://llvm.org/devmtg/2023-10/slides/quicktalks/Shah-TableGenFormatter.pdf

I'm not sure there's a consensus on what a "well formatted" tablegen file looks like, but it would be nice to have something. clang-format has some rudimentary table-gen support, but it never worked well.

@AlexMaclean AlexMaclean merged commit af54c70 into llvm:main May 13, 2025
11 checks passed
@llvm-ci
Copy link
Collaborator

llvm-ci commented May 13, 2025

LLVM Buildbot has detected a new failure on builder llvm-clang-x86_64-gcc-ubuntu running on sie-linux-worker3 while building llvm at step 6 "test-build-unified-tree-check-all".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/174/builds/17713

Here is the relevant piece of the build log for the reference 
Step 6 (test-build-unified-tree-check-all) failure: test (failure)
******************** TEST 'AddressSanitizer-x86_64-linux-dynamic :: TestCases/asan_lsan_deadlock.cpp' FAILED ********************
Exit Code: 1

Command Output (stderr):
--
/home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/build/./bin/clang  --driver-mode=g++ -fsanitize=address -mno-omit-leaf-frame-pointer -fno-omit-frame-pointer -fno-optimize-sibling-calls -gline-tables-only  -m64  -shared-libasan -O0 /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/compiler-rt/test/asan/TestCases/asan_lsan_deadlock.cpp -o /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/build/runtimes/runtimes-bins/compiler-rt/test/asan/X86_64LinuxDynamicConfig/TestCases/Output/asan_lsan_deadlock.cpp.tmp # RUN: at line 4
+ /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/build/./bin/clang --driver-mode=g++ -fsanitize=address -mno-omit-leaf-frame-pointer -fno-omit-frame-pointer -fno-optimize-sibling-calls -gline-tables-only -m64 -shared-libasan -O0 /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/compiler-rt/test/asan/TestCases/asan_lsan_deadlock.cpp -o /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/build/runtimes/runtimes-bins/compiler-rt/test/asan/X86_64LinuxDynamicConfig/TestCases/Output/asan_lsan_deadlock.cpp.tmp
env ASAN_OPTIONS=detect_leaks=1 not  /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/build/runtimes/runtimes-bins/compiler-rt/test/asan/X86_64LinuxDynamicConfig/TestCases/Output/asan_lsan_deadlock.cpp.tmp 2>&1 | FileCheck /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/compiler-rt/test/asan/TestCases/asan_lsan_deadlock.cpp # RUN: at line 5
+ env ASAN_OPTIONS=detect_leaks=1 not /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/build/runtimes/runtimes-bins/compiler-rt/test/asan/X86_64LinuxDynamicConfig/TestCases/Output/asan_lsan_deadlock.cpp.tmp
+ FileCheck /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/compiler-rt/test/asan/TestCases/asan_lsan_deadlock.cpp
�[1m/home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/compiler-rt/test/asan/TestCases/asan_lsan_deadlock.cpp:58:12: �[0m�[0;1;31merror: �[0m�[1mCHECK: expected string not found in input
�[0m // CHECK: SUMMARY: AddressSanitizer: stack-buffer-overflow
�[0;1;32m           ^
�[0m�[1m<stdin>:1:1: �[0m�[0;1;30mnote: �[0m�[1mscanning from here
�[0m=================================================================
�[0;1;32m^
�[0m�[1m<stdin>:2:10: �[0m�[0;1;30mnote: �[0m�[1mpossible intended match here
�[0m==2845085==ERROR: AddressSanitizer: stack-buffer-overflow on address 0x7bc3464de034 at pc 0x5564f989f220 bp 0x7bc3446fdce0 sp 0x7bc3446fdcd8
�[0;1;32m         ^
�[0m
Input file: <stdin>
Check file: /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/compiler-rt/test/asan/TestCases/asan_lsan_deadlock.cpp

-dump-input=help explains the following input dump.

Input was:
<<<<<<
�[1m�[0m�[0;1;30m            1: �[0m�[1m�[0;1;46m================================================================= �[0m
�[0;1;31mcheck:58'0     X~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ error: no match found
�[0m�[0;1;30m            2: �[0m�[1m�[0;1;46m==2845085==ERROR: AddressSanitizer: stack-buffer-overflow on address 0x7bc3464de034 at pc 0x5564f989f220 bp 0x7bc3446fdce0 sp 0x7bc3446fdcd8 �[0m
�[0;1;31mcheck:58'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
�[0m�[0;1;35mcheck:58'1              ?                                                                                                                                    possible intended match
�[0m�[0;1;30m            3: �[0m�[1m�[0;1;46mWRITE of size 4 at 0x7bc3464de034 thread T2 �[0m
�[0;1;31mcheck:58'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
�[0m>>>>>>

--

********************

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.
4 participants
@AlexMaclean @llvmbot @llvm-ci @Artem-B
Morty Proxy This is a proxified and sanitized view of the page, visit original site.