[NVPTX, CUDA] barrier intrinsics and builtins for sm_90

[tra]

Re-enabled .relaxed test, now that ptx80 is available.

[tra]

whitespace fix.

clang/include/clang/Basic/BuiltinsNVPTX.def

 BUILTIN(__nvvm_bar0_popc, "ii", "")
 BUILTIN(__nvvm_bar0_and, "ii", "")
 BUILTIN(__nvvm_bar0_or, "ii", "")
 BUILTIN(__nvvm_bar_sync, "vi", "n")
 TARGET_BUILTIN(__nvvm_bar_warp_sync, "vUi", "n", PTX60)
 TARGET_BUILTIN(__nvvm_barrier_sync, "vUi", "n", PTX60)
 TARGET_BUILTIN(__nvvm_barrier_sync_cnt, "vUiUi", "n", PTX60)
 
+TARGET_BUILTIN(__nvvm_barrier_cluster_arrive, "v", "n", AND(SM_90,PTX78))
+TARGET_BUILTIN(__nvvm_barrier_cluster_arrive_relaxed, "v", "n", AND(SM_90,PTX80))
+TARGET_BUILTIN(__nvvm_barrier_cluster_wait, "v", "n", AND(SM_90,PTX78))
+TARGET_BUILTIN(__nvvm_fence_sc_cluster, "v", "n", AND(SM_90,PTX78))
+
 // Shuffle
 
 BUILTIN(__nvvm_shfl_down_i32, "iiii", "")
 BUILTIN(__nvvm_shfl_down_f32, "ffii", "")
 BUILTIN(__nvvm_shfl_up_i32, "iiii", "")
 BUILTIN(__nvvm_shfl_up_f32, "ffii", "")
 BUILTIN(__nvvm_shfl_bfly_i32, "iiii", "")
 BUILTIN(__nvvm_shfl_bfly_f32, "ffii", "")

clang/lib/CodeGen/CGBuiltin.cpp

Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID,
   case NVPTX::BI__nvvm_getctarank:
     return Builder.CreateCall(
         CGM.getIntrinsic(Intrinsic::nvvm_getctarank),
         EmitScalarExpr(E->getArg(0)));
   case NVPTX::BI__nvvm_getctarank_shared_cluster:
     return Builder.CreateCall(
         CGM.getIntrinsic(Intrinsic::nvvm_getctarank_shared_cluster),
         EmitScalarExpr(E->getArg(0)));
+  case NVPTX::BI__nvvm_barrier_cluster_arrive:
+    return Builder.CreateCall(
+        CGM.getIntrinsic(Intrinsic::nvvm_barrier_cluster_arrive));
+  case NVPTX::BI__nvvm_barrier_cluster_arrive_relaxed:
+    return Builder.CreateCall(
+        CGM.getIntrinsic(Intrinsic::nvvm_barrier_cluster_arrive_relaxed));
+  case NVPTX::BI__nvvm_barrier_cluster_wait:
+    return Builder.CreateCall(
+        CGM.getIntrinsic(Intrinsic::nvvm_barrier_cluster_wait));
+  case NVPTX::BI__nvvm_fence_sc_cluster:
+    return Builder.CreateCall(
+        CGM.getIntrinsic(Intrinsic::nvvm_fence_sc_cluster));
   default:
     return nullptr;
   }
 }
 
 namespace {
 struct BuiltinAlignArgs {
   llvm::Value *Src = nullptr;

clang/test/CodeGenCUDA/builtins-sm90.cu

-// RUN: %clang_cc1 "-triple" "nvptx64-nvidia-cuda" "-target-feature" "+ptx78" "-target-cpu" "sm_90" -emit-llvm -fcuda-is-device -o - %s | FileCheck %s
+// RUN: %clang_cc1 "-triple" "nvptx64-nvidia-cuda" "-target-feature" "+ptx80" "-target-cpu" "sm_90" -emit-llvm -fcuda-is-device -o - %s | FileCheck %s
 
 // CHECK: define{{.*}} void @_Z6kernelPlPvj(
 __attribute__((global)) void kernel(long *out, void *ptr, unsigned u) {
   int i = 0;
   // CHECK: call i1 @llvm.nvvm.isspacep.shared.cluster
   out[i++] = __nvvm_isspacep_shared_cluster(ptr);
 
   // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.clusterid.x()
__attribute__((global)) void kernel(long *out, void *ptr, unsigned u) {
   out[i++] = (long) __nvvm_mapa(ptr, u);
   // CHECK: call ptr addrspace(3) @llvm.nvvm.mapa.shared.cluster(ptr addrspace(3) %{{.*}}, i32 %{{.*}})
   out[i++] = (long) __nvvm_mapa_shared_cluster(sptr, u);
   // CHECK: call i32 @llvm.nvvm.getctarank(ptr {{.*}})
   out[i++] = __nvvm_getctarank(ptr);
   // CHECK: call i32 @llvm.nvvm.getctarank.shared.cluster(ptr addrspace(3) {{.*}})
   out[i++] = __nvvm_getctarank_shared_cluster(sptr);
 
+  // CHECK: call void @llvm.nvvm.barrier.cluster.arrive()
+  __nvvm_barrier_cluster_arrive();
+  // CHECK: call void @llvm.nvvm.barrier.cluster.arrive.relaxed()
+  __nvvm_barrier_cluster_arrive_relaxed();
+  // CHECK: call void @llvm.nvvm.barrier.cluster.wait()
+  __nvvm_barrier_cluster_wait();
+  // CHECK: call void @llvm.nvvm.fence.sc.cluster()
+  __nvvm_fence_sc_cluster();
+
   // CHECK: ret void
 }

llvm/include/llvm/IR/IntrinsicsNVVM.td

// Bar.Sync
   // barrier.sync id[, cnt]
   def int_nvvm_barrier_sync :
       Intrinsic<[], [llvm_i32_ty], [IntrConvergent, IntrNoCallback]>,
       ClangBuiltin<"__nvvm_barrier_sync">;
   def int_nvvm_barrier_sync_cnt :
       Intrinsic<[], [llvm_i32_ty, llvm_i32_ty], [IntrConvergent, IntrNoCallback]>,
       ClangBuiltin<"__nvvm_barrier_sync_cnt">;
 
+  // barrier.cluster.[wait, arrive, arrive.relaxed]
+  def int_nvvm_barrier_cluster_arrive :
+      Intrinsic<[], [], [IntrConvergent, IntrNoCallback]>;
+  def int_nvvm_barrier_cluster_arrive_relaxed :
+      Intrinsic<[], [], [IntrConvergent, IntrNoCallback]>;
+  def int_nvvm_barrier_cluster_wait :
+      Intrinsic<[], [], [IntrConvergent, IntrNoCallback]>;
+
   // Membar
   def int_nvvm_membar_cta : ClangBuiltin<"__nvvm_membar_cta">,
       Intrinsic<[], [], [IntrNoCallback]>;
   def int_nvvm_membar_gl : ClangBuiltin<"__nvvm_membar_gl">,
       Intrinsic<[], [], [IntrNoCallback]>;
   def int_nvvm_membar_sys : ClangBuiltin<"__nvvm_membar_sys">,
       Intrinsic<[], [], [IntrNoCallback]>;
+  def int_nvvm_fence_sc_cluster:
+      Intrinsic<[], [], [IntrNoCallback]>;
 
 // Async Copy
 def int_nvvm_cp_async_mbarrier_arrive :
     ClangBuiltin<"__nvvm_cp_async_mbarrier_arrive">,
     Intrinsic<[],[llvm_i64ptr_ty],[IntrConvergent, IntrNoCallback]>;
 def int_nvvm_cp_async_mbarrier_arrive_shared :
     ClangBuiltin<"__nvvm_cp_async_mbarrier_arrive_shared">,
     Intrinsic<[],[llvm_shared_i64ptr_ty],[IntrConvergent, IntrNoCallback]>;

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

 def INT_BARRIER_SYNC_CNT_IR : NVPTXInst<(outs), (ins i32imm:$id, Int32Regs:$cnt),
                  "barrier.sync \t$id, $cnt;",
                  [(int_nvvm_barrier_sync_cnt imm:$id, Int32Regs:$cnt)]>,
         Requires<[hasPTX<60>, hasSM<30>]>;
 def INT_BARRIER_SYNC_CNT_II : NVPTXInst<(outs), (ins i32imm:$id, i32imm:$cnt),
                  "barrier.sync \t$id, $cnt;",
                  [(int_nvvm_barrier_sync_cnt imm:$id, imm:$cnt)]>,
         Requires<[hasPTX<60>, hasSM<30>]>;
+class INT_BARRIER_CLUSTER<string variant, Intrinsic Intr,
+                          list<Predicate> Preds = [hasPTX<78>, hasSM<90>]>:
+        NVPTXInst<(outs), (ins), "barrier.cluster."# variant #";", [(Intr)]>,
+        Requires<Preds>;
+
+def barrier_cluster_arrive:
+        INT_BARRIER_CLUSTER<"arrive", int_nvvm_barrier_cluster_arrive>;
+def barrier_cluster_arrive_relaxed:
+        INT_BARRIER_CLUSTER<"arrive.relaxed",
+        int_nvvm_barrier_cluster_arrive_relaxed, [hasPTX<80>, hasSM<90>]>;
+def barrier_cluster_wait:
+        INT_BARRIER_CLUSTER<"wait", int_nvvm_barrier_cluster_wait>;
 
 class SHFL_INSTR<bit sync, string mode, string reg, bit return_pred,
                  bit offset_imm, bit mask_imm, bit threadmask_imm>
       : NVPTXInst<(outs), (ins), "?", []> {
   NVPTXRegClass rc = !cond(
     !eq(reg, "i32"): Int32Regs,
     !eq(reg, "f32"): Float32Regs);
   string IntrName = "int_nvvm_shfl_"

 class MEMBAR<string StrOp, Intrinsic IntOP> :
               NVPTXInst<(outs), (ins),
             StrOp, [(IntOP)]>;
 
 def INT_MEMBAR_CTA : MEMBAR<"membar.cta;", int_nvvm_membar_cta>;
 def INT_MEMBAR_GL  : MEMBAR<"membar.gl;",  int_nvvm_membar_gl>;
 def INT_MEMBAR_SYS : MEMBAR<"membar.sys;", int_nvvm_membar_sys>;
 
+def INT_FENCE_SC_CLUSTER:
+       MEMBAR<"fence.sc.cluster;", int_nvvm_fence_sc_cluster>,
+       Requires<[hasPTX<78>, hasSM<90>]>;
 
 //-----------------------------------
 // Async Copy Functions
 //-----------------------------------
 
 multiclass CP_ASYNC_MBARRIER_ARRIVE<string NoInc, string AddrSpace, Intrinsic Intrin> {
   def _32 : NVPTXInst<(outs), (ins Int32Regs:$addr),
             !strconcat("cp.async.mbarrier.arrive", NoInc, AddrSpace, ".b64 [$addr];"),

llvm/test/CodeGen/NVPTX/intrinsics-sm90.ll

-; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx78| FileCheck --check-prefixes=CHECK %s
-; RUN: %if ptxas-11.8 %{ llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx78| %ptxas-verify -arch=sm_90 %}
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx80| FileCheck --check-prefixes=CHECK %s
+; RUN: %if ptxas-11.8 %{ llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx80| %ptxas-verify -arch=sm_90 %}
 
 ; CHECK-LABEL: test_isspacep
 define i1 @test_isspacep_shared_cluster(ptr %p) {
 ; CHECK: isspacep.shared::cluster
   %a = tail call i1 @llvm.nvvm.isspacep.shared.cluster(ptr %p)
 ; CHECK: ret
   ret i1 %a
 }

 ; CHECK-LABEL: test_is_explicit_cluster(
 define i1 @test_is_explicit_cluster() {
 ; CHECK: mov.pred %p{{[0-9]+}}, %is_explicit_cluster;
 ; CHECK: ret;
         %x = call i1 @llvm.nvvm.is_explicit_cluster()
         ret i1 %x
 }
 
+; CHECK-LABEL: test_barrier_cluster(
+define void @test_barrier_cluster() {
+; CHECK: barrier.cluster.arrive;
+       call void @llvm.nvvm.barrier.cluster.arrive()
+; CHECK: barrier.cluster.arrive.relaxed;
+       call void @llvm.nvvm.barrier.cluster.arrive.relaxed()
+; CHECK: barrier.cluster.wait;
+       call void @llvm.nvvm.barrier.cluster.wait()
+; CHECK: fence.sc.cluster
+       call void @llvm.nvvm.fence.sc.cluster()
+       ret void
+}
+
 
 declare i1 @llvm.nvvm.isspacep.shared.cluster(ptr %p);
 declare ptr @llvm.nvvm.mapa(ptr %p, i32 %r);
 declare ptr addrspace(3) @llvm.nvvm.mapa.shared.cluster(ptr addrspace(3) %p, i32 %r);
 declare i32 @llvm.nvvm.getctarank(ptr %p);
 declare i32 @llvm.nvvm.getctarank.shared.cluster(ptr addrspace(3) %p);
 declare i32 @llvm.nvvm.read.ptx.sreg.clusterid.x()
 declare i32 @llvm.nvvm.read.ptx.sreg.clusterid.y()
 declare i32 @llvm.nvvm.read.ptx.sreg.clusterid.z()
 declare i32 @llvm.nvvm.read.ptx.sreg.clusterid.w()
 declare i32 @llvm.nvvm.read.ptx.sreg.nclusterid.x()
 declare i32 @llvm.nvvm.read.ptx.sreg.nclusterid.y()
 declare i32 @llvm.nvvm.read.ptx.sreg.nclusterid.z()
 declare i32 @llvm.nvvm.read.ptx.sreg.nclusterid.w()
 declare i32 @llvm.nvvm.read.ptx.sreg.cluster.ctarank()
 declare i32 @llvm.nvvm.read.ptx.sreg.cluster.nctarank()
 declare i1 @llvm.nvvm.is_explicit_cluster()
+declare void @llvm.nvvm.barrier.cluster.arrive()
+declare void @llvm.nvvm.barrier.cluster.arrive.relaxed()
+declare void @llvm.nvvm.barrier.cluster.wait()
+declare void @llvm.nvvm.fence.sc.cluster()