[clang] [CUDA] Change '__activemask' to use '__nvvm_activemask()' (PR #79892)
Joseph Huber via cfe-commits
cfe-commits at lists.llvm.org
Mon Jan 29 12:29:34 PST 2024
https://github.com/jhuber6 created https://github.com/llvm/llvm-project/pull/79892
Summary:
We recently added builitin support for this function.
>From 5f316d30a179dd21cfadd50d232de622d394ccea Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Mon, 29 Jan 2024 14:28:35 -0600
Subject: [PATCH] [CUDA] Change '__activemask' to use '__nvvm_activemask()'
Summary:
We recently added builitin support for this function.
---
clang/lib/Headers/__clang_cuda_intrinsics.h | 4 +---
1 file changed, 1 insertion(+), 3 deletions(-)
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
}
More information about the cfe-commits
mailing list