[NVPTX] Add activemask intrinsic.

[ABataev]

Summary

Since PTX62 vote.ballot cannot be used to get the mask of active
threads, instead activemask.b32 instruction must be used. Required for
Cuda10.

LLVM part of the patches to fix PR43156.

[tra]

include/llvm/IR/IntrinsicsNVVM.td

4096

Are these attribute sufficient to prevent CSE'ing out of divergent branches.

E.g. we must not allow transforming this:

if (cond(threadIdx.x))
  foo(activemask());
else
  bar(activemask());

into that:

int x = activemask();
if(cond(threadIdx.x)
  foo(x);
else
  bar(x);

It would be great to add a test for that.

4097

AFAICT NVCC does not provide __nvvm_vote_activemask builtin.

[ABataev]

include/llvm/IR/IntrinsicsNVVM.td

4096

Will add a test

4097

Will remove it, copy-paste.

include/llvm/IR/IntrinsicsNVVM.td

Intrinsic<[llvm_i1_ty], [llvm_i32_ty, llvm_i1_ty],
             [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.vote.uni.sync">,
   GCCBuiltin<"__nvvm_vote_uni_sync">;
 // vote.sync.ballot mask, pred
 def int_nvvm_vote_ballot_sync :
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i1_ty],
             [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.vote.ballot.sync">,
   GCCBuiltin<"__nvvm_vote_ballot_sync">;
 
+// activemask
+def int_nvvm_activemask :
+  Intrinsic<[llvm_i32_ty], [],
+            [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.activemask">,
+  GCCBuiltin<"__nvvm_vote_activemask">;
+
 //
 // MATCH.SYNC
 //
 // match.any.sync.b32 mask, value
 def int_nvvm_match_any_sync_i32 :
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
             [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.match.any.sync.i32">,
   GCCBuiltin<"__nvvm_match_any_sync_i32">;

lib/Target/NVPTX/NVPTX.td

 def PTX43 : SubtargetFeature<"ptx43", "PTXVersion", "43",
                              "Use PTX version 4.3">;
 def PTX50 : SubtargetFeature<"ptx50", "PTXVersion", "50",
                              "Use PTX version 5.0">;
 def PTX60 : SubtargetFeature<"ptx60", "PTXVersion", "60",
                              "Use PTX version 6.0">;
 def PTX61 : SubtargetFeature<"ptx61", "PTXVersion", "61",
                              "Use PTX version 6.1">;
+def PTX62 : SubtargetFeature<"ptx62", "PTXVersion", "62",
+                             "Use PTX version 6.2">;
 def PTX63 : SubtargetFeature<"ptx63", "PTXVersion", "63",
                              "Use PTX version 6.3">;
 def PTX64 : SubtargetFeature<"ptx64", "PTXVersion", "64",
                              "Use PTX version 6.4">;
 
 //===----------------------------------------------------------------------===//
 // NVPTX supported processors.
 //===----------------------------------------------------------------------===//

lib/Target/NVPTX/NVPTXInstrInfo.td

 def hasHWROT32 : Predicate<"Subtarget->hasHWROT32()">;
 def noHWROT32 : Predicate<"!Subtarget->hasHWROT32()">;
 
 def true : Predicate<"true">;
 
 def hasPTX31 : Predicate<"Subtarget->getPTXVersion() >= 31">;
 def hasPTX60 : Predicate<"Subtarget->getPTXVersion() >= 60">;
 def hasPTX61 : Predicate<"Subtarget->getPTXVersion() >= 61">;
+def hasPTX62 : Predicate<"Subtarget->getPTXVersion() >= 62">;
 def hasPTX63 : Predicate<"Subtarget->getPTXVersion() >= 63">;
 
 def hasSM30 : Predicate<"Subtarget->getSmVersion() >= 30">;
 def hasSM70 : Predicate<"Subtarget->getSmVersion() >= 70">;
 def hasSM72 : Predicate<"Subtarget->getSmVersion() >= 72">;
 def hasSM75 : Predicate<"Subtarget->getSmVersion() >= 75">;
 
 def useShortPtr : Predicate<"useShortPointers()">;

lib/Target/NVPTX/NVPTXIntrinsics.td

def r : NVPTXInst<(outs regclass:$dest), (ins Int32Regs:$mask, Int1Regs:$pred),
           Requires<[hasPTX60, hasSM30]>;
 }
 
 defm VOTE_SYNC_ALL : VOTE_SYNC<Int1Regs, "all.pred", int_nvvm_vote_all_sync>;
 defm VOTE_SYNC_ANY : VOTE_SYNC<Int1Regs, "any.pred", int_nvvm_vote_any_sync>;
 defm VOTE_SYNC_UNI : VOTE_SYNC<Int1Regs, "uni.pred", int_nvvm_vote_uni_sync>;
 defm VOTE_SYNC_BALLOT : VOTE_SYNC<Int32Regs, "ballot.b32", int_nvvm_vote_ballot_sync>;
 
+// activemask.b32
+def ACTIVEMASK : NVPTXInst<(outs Int32Regs:$dest), (ins), "activemask.b32 \t$dest;",
+                             [(set Int32Regs:$dest, (int_nvvm_activemask))]>,
+        Requires<[hasPTX62, hasSM30]>;
+
 multiclass MATCH_ANY_SYNC<NVPTXRegClass regclass, string ptxtype, Intrinsic IntOp,
                           Operand ImmOp> {
   def ii : NVPTXInst<(outs regclass:$dest), (ins i32imm:$mask, ImmOp:$value),
               "match.any.sync." # ptxtype # " \t$dest, $value, $mask;",
               [(set regclass:$dest, (IntOp imm:$mask, imm:$value))]>,
            Requires<[hasPTX60, hasSM70]>;
   def ir : NVPTXInst<(outs regclass:$dest), (ins Int32Regs:$mask, ImmOp:$value),
               "match.any.sync." # ptxtype # " \t$dest, $value, $mask;",

test/CodeGen/NVPTX/activemask.ll

+; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 -mattr=+ptx62 | FileCheck %s
+
+declare i32 @llvm.nvvm.activemask()
+; CHECK-LABEL: .func{{.*}}activemask
+define i32 @activemask() {
+  ; CHECK: activemask.b32
+  %val = call i32 @llvm.nvvm.activemask()
+  ret i32 %val
+}