[libc-commits] [PATCH] D148810: [libc] Add more utility functions for the GPU

[Joseph Huber via Phabricator via libc-commits]

jhuber6 added a comment.

In D148810#4290519 <https://reviews.llvm.org/D148810#4290519>, @JonChesterfield wrote:

> Using volatile to model convergence/divergence sounds wrong. Perhaps we're missing an intrinsic for activemask.

There isn't as far as I know, this is how CUDA in clang defines `__activemaks()`

  inline __device__ unsigned int __activemask() {                                     
  #if CUDA_VERSION < 9020                                                             
    return __nvvm_vote_ballot(1);                                                     
  #else                                                                               
    unsigned int mask;                                                                
    asm volatile("activemask.b32 %0;" : "=r"(mask));                                  
    return mask;                                                                      
  #endif                                                                              
  } 

The volatile is necessary because the activemask is sensitive to control flow so we can't change its location.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D148810