[PATCH] D67130: [NVPTX] Add activemask intrinsic.

[Alexey Bataev via Phabricator via llvm-commits]

ABataev marked 2 inline comments as done.
ABataev added inline comments.


================
Comment at: include/llvm/IR/IntrinsicsNVVM.td:4096
+  Intrinsic<[llvm_i32_ty], [],
+            [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.activemask">,
+  GCCBuiltin<"__nvvm_vote_activemask">;
----------------
tra wrote:
> 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.
> 
> 
Will add a test


================
Comment at: include/llvm/IR/IntrinsicsNVVM.td:4097
+            [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.activemask">,
+  GCCBuiltin<"__nvvm_vote_activemask">;
+
----------------
tra wrote:
> AFAICT NVCC does not provide __nvvm_vote_activemask builtin.
> 
> 
Will remove it, copy-paste.


Repository:
  rL LLVM

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D67130/new/

https://reviews.llvm.org/D67130