[clang] 51379a9 - [CUDA] Change '__activemask' to use '__nvvm_activemask()' (#79892)

[via cfe-commits]

Author: Joseph Huber
Date: 2024-01-29T17:34:45-06:00
New Revision: 51379a982efc64eb319ba6f4c1af4dce31af499a

URL: https://github.com/llvm/llvm-project/commit/51379a982efc64eb319ba6f4c1af4dce31af499a
DIFF: https://github.com/llvm/llvm-project/commit/51379a982efc64eb319ba6f4c1af4dce31af499a.diff

LOG: [CUDA] Change '__activemask' to use '__nvvm_activemask()' (#79892)

Summary:
We recently added builitin support for this function.

Added: 
    

Modified: 
    clang/lib/Headers/__clang_cuda_intrinsics.h

Removed: 
    


################################################################################
diff  --git a/clang/lib/Headers/__clang_cuda_intrinsics.h b/clang/lib/Headers/__clang_cuda_intrinsics.h
index 3c3948863c1d45..a04e8b6de44d05 100644
--- a/clang/lib/Headers/__clang_cuda_intrinsics.h
+++ b/clang/lib/Headers/__clang_cuda_intrinsics.h
@@ -215,9 +215,7 @@ 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;
+  return __nvvm_activemask();
 #endif
 }