[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