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

[Artem Belevich via Phabricator via llvm-commits]

tra added inline comments.


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




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




Repository:
  rL LLVM

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

https://reviews.llvm.org/D67130