[clang] [CUDA] Change '__activemask' to use '__nvvm_activemask()' (PR #79892)

via cfe-commits cfe-commits at lists.llvm.org
Mon Jan 29 12:30:02 PST 2024


llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-clang

Author: Joseph Huber (jhuber6)

<details>
<summary>Changes</summary>

Summary:
We recently added builitin support for this function.


---
Full diff: https://github.com/llvm/llvm-project/pull/79892.diff


1 Files Affected:

- (modified) clang/lib/Headers/__clang_cuda_intrinsics.h (+1-3) 


``````````diff
diff --git a/clang/lib/Headers/__clang_cuda_intrinsics.h b/clang/lib/Headers/__clang_cuda_intrinsics.h
index 3c3948863c1d453..a04e8b6de44d053 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
 }
 

``````````

</details>


https://github.com/llvm/llvm-project/pull/79892


More information about the cfe-commits mailing list