[llvm] [NVPTX] Further refactor intrinsic definitions to remove redundancy (NFC) (PR #139924)

via llvm-commits llvm-commits at lists.llvm.org
Wed May 14 08:37:36 PDT 2025


llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-backend-nvptx

Author: Alex MacLean (AlexMaclean)

<details>
<summary>Changes</summary>

Note: the diff indicates this change has no impact on the intrinsic code generated by table-gen. 

---

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


1 Files Affected:

- (modified) llvm/include/llvm/IR/IntrinsicsNVVM.td (+643-2384) 


``````````diff
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 0b26bb9829005..3e3a55c05a9e0 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -357,38 +357,33 @@ class MMA_SIGNATURE<WMMA_REGS A, WMMA_REGS B, WMMA_REGS C, WMMA_REGS D> {
      !ne(A.ptx_elt_type, B.ptx_elt_type): [A, B],
      true: [A]
      );
-   string ret = !foldl("", id_frags, a, b, !strconcat(a, ".", b.ptx_elt_type));
+   string ret = !foldl("", id_frags, a, b, !strconcat(a, "_", b.ptx_elt_type));
 }
 
 class WMMA_NAME<string ALayout, string BLayout, int Satfinite, string Rnd, string b1op,
                 WMMA_REGS A, WMMA_REGS B, WMMA_REGS C, WMMA_REGS D> {
   string signature = MMA_SIGNATURE<A, B, C, D>.ret;
-  string llvm = "llvm.nvvm.wmma."
-                # A.geom
-                # ".mma"
-                # b1op
-                # "." # ALayout
-                # "." # BLayout
-                # !if(!ne(Rnd, ""), !strconcat(".", Rnd), "")
-                # signature
-                # !if(Satfinite, ".satfinite", "");
-
-  string record = !subst(".", "_",
-                  !subst("llvm.", "int_", llvm));
+  string record = "int_nvvm_wmma_"
+                  # A.geom
+                  # "_mma"
+                  # b1op
+                  # "_" # ALayout
+                  # "_" # BLayout
+                  # !if(!ne(Rnd, ""), !strconcat("_", Rnd), "")
+                  # signature
+                  # !if(Satfinite, "_satfinite", "");
 }
 
 class MMA_NAME<string ALayout, string BLayout, int Satfinite, string b1op,
                WMMA_REGS A, WMMA_REGS B, WMMA_REGS C, WMMA_REGS D> {
   string signature = MMA_SIGNATURE<A, B, C, D>.ret;
-  string llvm = "llvm.nvvm.mma"
-                # b1op
-                # "." # A.geom
-                # "." # ALayout
-                # "." # BLayout
-                # !if(Satfinite, ".satfinite", "")
-                # signature;
-  string record = !subst(".", "_",
-                  !subst("llvm.", "int_", llvm));
+  string record = "int_nvvm_mma"
+                  # b1op
+                  # "_" # A.geom
+                  # "_" # ALayout
+                  # "_" # BLayout
+                  # !if(Satfinite, "_satfinite", "")
+                  # signature;
 }
 
 class LDMATRIX_NAME<WMMA_REGS Frag, int Trans> {
@@ -602,7 +597,7 @@ class NVVM_WMMA_SUPPORTED<list<WMMA_REGS> frags, string layout_a, string layout_
 
 class NVVM_MMA_B1OPS<list<WMMA_REGS> frags> {
   list<string> ret = !cond(
-    !eq(frags[0].ptx_elt_type, "b1") : [".xor.popc", ".and.popc"],
+    !eq(frags[0].ptx_elt_type, "b1") : ["_xor_popc", "_and_popc"],
     true: [""]
   );
 }
@@ -696,101 +691,6 @@ class SHFL_INFO<bit sync, string mode, string type, bit return_pred> {
     [OpType, llvm_i32_ty, llvm_i32_ty]);
 }
 
-class CP_ASYNC_BULK_TENSOR_G2S_INTR<int dim, string mode> {
-  string Name = "int_nvvm_cp_async_bulk_tensor_g2s_" # mode # "_" # dim # "d";
-
-  bit IsIm2Col = !if(!eq(mode, "im2col"), 1, 0);
-  int NumIm2ColOffsets = !if(IsIm2Col, !add(dim, -2), 0);
-  list<LLVMType> Im2ColOffsetsTy = !listsplat(llvm_i16_ty, NumIm2ColOffsets);
-  list<LLVMType> TensorDimsTy = !listsplat(llvm_i32_ty, dim);
-  list<LLVMType> ArgsTy = !listconcat(
-                          [llvm_shared_cluster_ptr_ty,  // dst_shared_cluster_ptr
-                           llvm_shared_ptr_ty,          // mbarrier_smem_ptr
-                           llvm_ptr_ty],                // tensormap_ptr
-                           TensorDimsTy,                // actual tensor dims
-                           Im2ColOffsetsTy,             // im2col offsets
-                          [llvm_i16_ty,                 // cta_mask
-                           llvm_i64_ty,                 // cache_hint
-                           llvm_i1_ty,                  // Flag for cta_mask
-                           llvm_i1_ty]                  // Flag for cache_hint
-                          );
-
-  int TempFlagsStartIdx = !add(dim, 5);
-  int FlagsStartIdx = !add(TempFlagsStartIdx, NumIm2ColOffsets);
-  list<IntrinsicProperty> IntrProp = [IntrConvergent,
-        WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<2>>,
-        NoCapture<ArgIndex<0>>, NoCapture<ArgIndex<1>>, NoCapture<ArgIndex<2>>,
-        ImmArg<ArgIndex<FlagsStartIdx>>,
-        ImmArg<ArgIndex<!add(FlagsStartIdx, 1)>>];
-}
-
-class CP_ASYNC_BULK_TENSOR_S2G_INTR<int dim, string mode> {
-  string Name = "int_nvvm_cp_async_bulk_tensor_s2g_" # mode # "_" # dim # "d";
-
-  list<LLVMType> TensorDimsTy = !listsplat(llvm_i32_ty, dim);
-  list<LLVMType> ArgsTy = !listconcat(
-                          [llvm_shared_ptr_ty,  // src_smem_ptr
-                           llvm_ptr_ty],        // tensormap_ptr
-                           TensorDimsTy,        // actual tensor dims
-                          [llvm_i64_ty,         // cache_hint
-                           llvm_i1_ty]          // Flag for cache_hint
-                          );
-  int FlagsStartIdx = !add(dim, 3);
-  list<IntrinsicProperty> IntrProp = [IntrConvergent,
-        ReadOnly<ArgIndex<0>>, ReadOnly<ArgIndex<1>>,
-        NoCapture<ArgIndex<0>>, NoCapture<ArgIndex<1>>,
-        ImmArg<ArgIndex<FlagsStartIdx>>];
-}
-
-class CP_ASYNC_BULK_TENSOR_PREFETCH_INTR<int dim, string mode> {
-  string Name = "int_nvvm_cp_async_bulk_tensor_prefetch_" # mode # "_" # dim # "d";
-
-  bit IsIm2Col = !if(!eq(mode, "im2col"), 1, 0);
-  int NumIm2ColOffsets = !if(IsIm2Col, !add(dim, -2), 0);
-  list<LLVMType> Im2ColOffsetsTy = !listsplat(llvm_i16_ty, NumIm2ColOffsets);
-  list<LLVMType> TensorDimsTy = !listsplat(llvm_i32_ty, dim);
-  list<LLVMType> ArgsTy = !listconcat(
-                          [llvm_ptr_ty],     // tensormap_ptr
-                           TensorDimsTy,     // actual tensor dims
-                           Im2ColOffsetsTy,  // im2col offsets
-                          [llvm_i64_ty,      // cache_hint
-                           llvm_i1_ty]       // Flag for cache_hint
-                          );
-
-  int TempFlagsStartIdx = !add(dim, 2);
-  int FlagsStartIdx = !add(TempFlagsStartIdx, NumIm2ColOffsets);
-  list<IntrinsicProperty> IntrProp = [IntrConvergent,
-        ReadOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>,
-        ImmArg<ArgIndex<FlagsStartIdx>>];
-}
-
-class CP_ASYNC_BULK_TENSOR_REDUCE_INTR<int dim, string mode, string op> {
-  string Suffix = op # "_" # mode # "_" # dim # "d";
-  string Name = "int_nvvm_cp_async_bulk_tensor_reduce_" # Suffix;
-
-  list<LLVMType> TensorDimsTy = !listsplat(llvm_i32_ty, dim);
-  list<LLVMType> ArgsTy = !listconcat(
-                          [llvm_shared_ptr_ty,  // src_smem_ptr
-                           llvm_ptr_ty],        // tensormap_ptr
-                           TensorDimsTy,        // actual tensor dims
-                          [llvm_i64_ty,         // cache_hint
-                           llvm_i1_ty]          // Flag for cache_hint
-                          );
-  int FlagsStartIdx = !add(dim, 3);
-  list<IntrinsicProperty> IntrProp = [IntrConvergent,
-        ReadOnly<ArgIndex<0>>, ReadOnly<ArgIndex<1>>,
-        NoCapture<ArgIndex<0>>, NoCapture<ArgIndex<1>>,
-        ImmArg<ArgIndex<FlagsStartIdx>>];
-}
-
-class NVVM_TCGEN05_LDST_NAME<string Op, string Shape, int Num> {
-  string intr =   "llvm.nvvm.tcgen05." # Op
-                  # "." # Shape
-                  # "." # "x" # !shl(1, 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,
@@ -810,6 +710,28 @@ class NVVM_TCGEN05_LDST_ACCESS_SIZE<string Shape, int Num> {
                         true : llvm_void_ty);
 }
 
+class TexVector<string name, list<LLVMType> types> {
+  string Name = name;
+  list<LLVMType> Types = types;
+}
+
+def TV_I8 : TexVector<"i8", [llvm_i16_ty]>;
+def TV_I16 : TexVector<"i16", [llvm_i16_ty]>;
+def TV_I32 : TexVector<"i32", [llvm_i32_ty]>;
+def TV_I64 : TexVector<"i64", [llvm_i64_ty]>;
+def TV_V2I8 : TexVector<"v2i8", !listsplat(llvm_i16_ty, 2)>;
+def TV_V2I16 : TexVector<"v2i16", !listsplat(llvm_i16_ty, 2)>;
+def TV_V2I32 : TexVector<"v2i32", !listsplat(llvm_i32_ty, 2)>;
+def TV_V2I64 : TexVector<"v2i64", !listsplat(llvm_i64_ty, 2)>;
+def TV_V4I8 : TexVector<"v4i8", !listsplat(llvm_i16_ty, 4)>;
+def TV_V4I16 : TexVector<"v4i16", !listsplat(llvm_i16_ty, 4)>;
+def TV_V4I32 : TexVector<"v4i32", !listsplat(llvm_i32_ty, 4)>;
+
+
+def V4F32 : TexVector<"v4f32", !listsplat(llvm_float_ty, 4)>;
+def V4S32 : TexVector<"v4s32", !listsplat(llvm_i32_ty, 4)>;
+def V4U32 : TexVector<"v4u32", !listsplat(llvm_i32_ty, 4)>;
+
 class NVVMBuiltin :
   ClangBuiltin<!strconcat("__", !substr(NAME, !size("int_")))> {
     assert !eq(!substr(NAME, 0, !size("int_nvvm_")), "int_nvvm_"),
@@ -828,131 +750,116 @@ let TargetPrefix = "nvvm" in {
 //
 // Min Max
 //
+  let IntrProperties = [IntrNoMem, IntrSpeculatable, Commutative] in {
+    foreach operation = ["min", "max"] in {
+      def int_nvvm_f # operation # _d : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty]>;
 
-  foreach operation = ["min", "max"] in {
-    def int_nvvm_f # operation # _d : NVVMBuiltin,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_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]>;
 
-    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]>;
+            def int_nvvm_f # operation # ftz # nan # variant # _f16 :
+              DefaultAttrsIntrinsic<[llvm_half_ty], [llvm_half_ty, llvm_half_ty]>;
 
-          def int_nvvm_f # operation # ftz # nan # variant # _f16 :
-            DefaultAttrsIntrinsic<[llvm_half_ty], [llvm_half_ty, llvm_half_ty],
-              [IntrNoMem, IntrSpeculatable, Commutative]>;
+            def int_nvvm_f # operation # ftz # nan # variant # _f16x2 :
+              DefaultAttrsIntrinsic<[llvm_v2f16_ty], [llvm_v2f16_ty, llvm_v2f16_ty]>;
 
-          def int_nvvm_f # operation # ftz # nan # variant # _f16x2 :
-            DefaultAttrsIntrinsic<[llvm_v2f16_ty], [llvm_v2f16_ty, llvm_v2f16_ty],
-              [IntrNoMem, IntrSpeculatable, Commutative]>;
+            def int_nvvm_f # operation # ftz # nan # variant # _bf16 : NVVMBuiltin,
+              DefaultAttrsIntrinsic<[llvm_bfloat_ty], [llvm_bfloat_ty, llvm_bfloat_ty]>;
 
-          def int_nvvm_f # operation # ftz # nan # variant # _bf16 : NVVMBuiltin,
-            DefaultAttrsIntrinsic<[llvm_bfloat_ty], [llvm_bfloat_ty, llvm_bfloat_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]>;
-        } // ftz
-      } // nan
-    } // variant
-  } // operation
+            def int_nvvm_f # operation # ftz # nan # variant # _bf16x2 : NVVMBuiltin,
+              DefaultAttrsIntrinsic<[llvm_v2bf16_ty], [llvm_v2bf16_ty, llvm_v2bf16_ty]>;
+          } // ftz
+        } // nan
+      } // variant
+    } // operation
+  }
 
 //
 // Multiplication
 //
+  let IntrProperties = [IntrNoMem, IntrSpeculatable, Commutative] in {
+    foreach sign = ["", "u"] in {
+      def int_nvvm_mulhi_ # sign # s : NVVMBuiltin,
+          DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty]>;
 
-  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_ # sign # i : NVVMBuiltin,
-        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]>;
 
-    def int_nvvm_mulhi_ # sign # ll : NVVMBuiltin,
-        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]>;
 
-    def int_nvvm_mul24_ # sign # i : NVVMBuiltin,
-      DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
-        [IntrNoMem, IntrSpeculatable, Commutative]>;
-  }
+      def int_nvvm_mul24_ # sign # i : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty]>;
+    }
 
-  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]>;
+    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]>;
 
-    def int_nvvm_mul_ # rnd # _d : NVVMBuiltin,
-        DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
-          [IntrNoMem, IntrSpeculatable, Commutative]>;
+      def int_nvvm_mul_ # rnd # _d : NVVMBuiltin,
+          DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty]>;
+    }
   }
 
 //
 // Div
 //
+  let IntrProperties = [IntrNoMem] in {
+    foreach ftz = ["", "_ftz"] in {
+      def int_nvvm_div_approx # ftz # _f : NVVMBuiltin,
+          DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty]>;
 
-  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 # ftz : NVVMBuiltin,
-        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]>;
+    }
 
-  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]>;
+    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]>;
 
-    def int_nvvm_div_ # rnd # _d : NVVMBuiltin,
-        DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
-          [IntrNoMem]>;
+      def int_nvvm_div_ # rnd # _d : NVVMBuiltin,
+          DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty]>;
+    }
   }
 
 //
 // Sad
 //
+  let IntrProperties = [IntrNoMem, Commutative, IntrSpeculatable] in {
+    foreach sign = ["", "u"] in {
+      def int_nvvm_sad_ # sign # s : NVVMBuiltin,
+          DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty, llvm_i16_ty]>;
 
-  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_ # sign # i : NVVMBuiltin,
-        DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
-          [IntrNoMem, Commutative, IntrSpeculatable]>;
+      def int_nvvm_sad_ # sign # i : NVVMBuiltin,
+          DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;
 
-    def int_nvvm_sad_ # sign # ll : NVVMBuiltin,
-        DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i64_ty],
-          [IntrNoMem, Commutative, IntrSpeculatable]>;
+      def int_nvvm_sad_ # sign # ll : NVVMBuiltin,
+          DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i64_ty]>;
+    }
   }
 
 //
 // Floor  Ceil
 //
-
-  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]>;
+  let IntrProperties = [IntrNoMem, IntrSpeculatable] in {
+    foreach op = ["floor", "ceil"] in {
+      foreach ftz = ["", "_ftz"] in
+        def int_nvvm_ # op # ftz # _f : NVVMBuiltin,
+            DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty]>;
+      def int_nvvm_ # op # _d : NVVMBuiltin,
+          DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty]>;
+    }
   }
 
 //
 // Abs
 //
-
   foreach ftz = ["", "_ftz"] in
     def int_nvvm_fabs # ftz :
       DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>],
@@ -961,7 +868,6 @@ let TargetPrefix = "nvvm" in {
 //
 // Abs, Neg bf16, bf16x2
 //
-
   def int_nvvm_neg_bf16 : NVVMBuiltin,
     DefaultAttrsIntrinsic<[llvm_bfloat_ty], [llvm_bfloat_ty], [IntrNoMem]>;
   def int_nvvm_neg_bf16x2 : NVVMBuiltin,
@@ -970,62 +876,65 @@ let TargetPrefix = "nvvm" in {
 //
 // Round
 //
+  let IntrProperties = [IntrNoMem, IntrSpeculatable] in {
+    foreach ftz = ["", "_ftz"] in
+      def int_nvvm_round # ftz # _f : NVVMBuiltin,
+          DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty]>;
 
-  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 : NVVMBuiltin,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem, IntrSpeculatable]>;
+    def int_nvvm_round_d : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty]>;
+  }
 
 //
 // Trunc
 //
+  let IntrProperties = [IntrNoMem, IntrSpeculatable] in {
+    foreach ftz = ["", "_ftz"] in
+      def int_nvvm_trunc # ftz # _f : NVVMBuiltin,
+          DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty]>;
 
-  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 : NVVMBuiltin,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem, IntrSpeculatable]>;
+    def int_nvvm_trunc_d : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty]>;
+  }
 
 //
 // Saturate
 //
+  let IntrProperties = [IntrNoMem, IntrSpeculatable] in {
+    foreach ftz = ["", "_ftz"] in
+      def int_nvvm_saturate # ftz # _f : NVVMBuiltin,
+          DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty]>;
 
-  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 : NVVMBuiltin,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem, IntrSpeculatable]>;
+    def int_nvvm_saturate_d : NVVMBuiltin,
+        DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty]>;
+  }
 
 //
 // Exp2  Log2
 //
+  let IntrProperties = [IntrNoMem] in {
+    foreach ftz = ["", "_ftz"] in
+      def int_nvvm_ex2_approx # ftz # _f : NVVMBuiltin,
+          DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty]>;
 
-  foreach ftz = ["", "_ftz"] in
-    def int_nvvm_ex2_approx # ftz # _f : NVVMBuiltin,
-        DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
-
-  def int_n...
[truncated]

``````````

</details>


https://github.com/llvm/llvm-project/pull/139924


More information about the llvm-commits mailing list