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

Artem Belevich via llvm-commits llvm-commits at lists.llvm.org
Mon May 12 17:47:04 PDT 2025


================
@@ -1812,45 +1420,36 @@ foreach scope = ["cta", "cluster", "gpu", "sys"] in {
 }
 
 // Async Copy
-def int_nvvm_cp_async_mbarrier_arrive :
-    ClangBuiltin<"__nvvm_cp_async_mbarrier_arrive">,
+def int_nvvm_cp_async_mbarrier_arrive : NVVMBuiltin,
     Intrinsic<[],[llvm_ptr_ty],[IntrConvergent, IntrNoCallback]>;
-def int_nvvm_cp_async_mbarrier_arrive_shared :
-    ClangBuiltin<"__nvvm_cp_async_mbarrier_arrive_shared">,
+def int_nvvm_cp_async_mbarrier_arrive_shared : NVVMBuiltin,
     Intrinsic<[],[llvm_shared_ptr_ty],[IntrConvergent, IntrNoCallback]>;
-def int_nvvm_cp_async_mbarrier_arrive_noinc :
-    ClangBuiltin<"__nvvm_cp_async_mbarrier_arrive_noinc">,
+def int_nvvm_cp_async_mbarrier_arrive_noinc : NVVMBuiltin,
     Intrinsic<[],[llvm_ptr_ty],[IntrConvergent, IntrNoCallback]>;
-def int_nvvm_cp_async_mbarrier_arrive_noinc_shared :
-    ClangBuiltin<"__nvvm_cp_async_mbarrier_arrive_noinc_shared">,
+def int_nvvm_cp_async_mbarrier_arrive_noinc_shared : NVVMBuiltin,
     Intrinsic<[],[llvm_shared_ptr_ty],[IntrConvergent, IntrNoCallback]>;
 
-multiclass CP_ASYNC_SHARED_GLOBAL<string n, string cc> {
+multiclass CP_ASYNC_SHARED_GLOBAL {
   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>;
----------------
Artem-B wrote:

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

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


More information about the llvm-commits mailing list