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

[Alex MacLean via llvm-commits]

================
@@ -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>;
----------------
AlexMaclean wrote:

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. 

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