[clang] fb77712 - [clang][AMDGPU] Widen ballot for read_exec_lo/hi to wavefront size (#212813)

[via cfe-commits]

Author: Arseniy Obolenskiy
Date: 2026-08-03T11:50:09+02:00
New Revision: fb77712fd1dd911330bd432d1fd9313081b40c59

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

LOG: [clang][AMDGPU] Widen ballot for read_exec_lo/hi to wavefront size (#212813)

GlobalISel cannot select a ballot narrower than the wavefront width,
since it can't represent one bit per lane

Widen the ballot to the wave size and narrow the result afterwards

This is a prerequisite for relanding
https://github.com/llvm/llvm-project/pull/211493 (reverted in
https://github.com/llvm/llvm-project/pull/212628 to unblock buildbot) to
prevent device libs side failures

---------

Co-authored-by: Matt Arsenault <arsenm2 at gmail.com>

Added: 
    

Modified: 
    clang/lib/Basic/Targets/SPIR.h
    clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
    clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl
    clang/test/CodeGenOpenCL/builtins-amdgcn.cl

Removed: 
    


################################################################################
diff  --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h
index a5d58dcf2c835..9240bd940fa1b 100644
--- a/clang/lib/Basic/Targets/SPIR.h
+++ b/clang/lib/Basic/Targets/SPIR.h
@@ -454,6 +454,10 @@ class LLVM_LIBRARY_VISIBILITY SPIRV64AMDGCNTargetInfo final
   void getTargetDefines(const LangOptions &Opts,
                         MacroBuilder &Builder) const override;
 
+  const llvm::omp::GV &getGridValue() const override {
+    return llvm::omp::SPIRVGridValues;
+  }
+
   void setAuxTarget(const TargetInfo *Aux) override;
 
   void adjust(DiagnosticsEngine &Diags, LangOptions &Opts,

diff  --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
index 29199f1726c1a..720f956130cc1 100644
--- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
@@ -274,13 +274,21 @@ Value *EmitAMDGPUGridSize(CodeGenFunction &CGF, unsigned Index) {
 
 // Generates the IR for __builtin_read_exec_*.
 // Lowers the builtin to amdgcn_ballot intrinsic.
+//
+// The ballot must be taken at the wavefront width: a ballot narrower than the
+// wave size cannot represent one bit per lane and fails to select. Request the
+// mask at the wave width and narrow it afterwards for the _lo and _hi halves.
 static Value *EmitAMDGCNBallotForExec(CodeGenFunction &CGF, const CallExpr *E,
                                       llvm::Type *RegisterType,
                                       llvm::Type *ValueType, bool isExecHi) {
   CodeGen::CGBuilderTy &Builder = CGF.Builder;
   CodeGen::CodeGenModule &CGM = CGF.CGM;
 
-  Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_ballot, {RegisterType});
+  unsigned WaveSize = CGF.getTarget().getGridValue().GV_Warp_Size;
+  unsigned BallotSize = std::max(WaveSize, RegisterType->getIntegerBitWidth());
+  llvm::Type *BallotType = Builder.getIntNTy(BallotSize);
+
+  Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_ballot, {BallotType});
   llvm::Value *Call = Builder.CreateCall(F, {Builder.getInt1(true)});
 
   if (isExecHi) {
@@ -289,7 +297,7 @@ static Value *EmitAMDGCNBallotForExec(CodeGenFunction &CGF, const CallExpr *E,
     return Rt2;
   }
 
-  return Call;
+  return Builder.CreateTrunc(Call, ValueType);
 }
 
 static llvm::Value *loadTextureDescPtorAsVec8I32(CodeGenFunction &CGF,

diff  --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl
index ef39558f94f2c..b049031156335 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl
@@ -37,13 +37,12 @@ void test_read_exec(global ulong* out) {
 }
 
 // CHECK-LABEL: @test_read_exec_lo(
-// CHECK: call i32 @llvm.amdgcn.ballot.i32(i1 true)
+// CHECK: call i64 @llvm.amdgcn.ballot.i64(i1 true)
+// CHECK: and i64 [[A:%.*]], 4294967295
 void test_read_exec_lo(global ulong* out) {
   *out = __builtin_amdgcn_read_exec_lo();
 }
 
-// CHECK: declare i32 @llvm.amdgcn.ballot.i32(i1) #[[$NOUNWIND_READONLY:[0-9]+]]
-
 // CHECK-LABEL: @test_read_exec_hi(
 // CHECK: call i64 @llvm.amdgcn.ballot.i64(i1 true)
 // CHECK: lshr i64 [[A:%.*]], 32

diff  --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
index 28c420a5760f4..78a4d7af386c2 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -1048,13 +1048,12 @@ void test_read_exec(global ulong* out) {
 }
 
 // CHECK-LABEL: @test_read_exec_lo(
-// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.ballot.i32(i1 true)
+// CHECK: [[A:%.*]] = {{.*}}call{{.*}} i64 @llvm.amdgcn.ballot.i64(i1 true)
+// CHECK: trunc i64 [[A]] to i32
 void test_read_exec_lo(global uint* out) {
   *out = __builtin_amdgcn_read_exec_lo();
 }
 
-// CHECK: declare i32 @llvm.amdgcn.ballot.i32(i1){{.*}} #[[$NOUNWIND_READONLY_NOPOISON:[0-9]+]]
-
 // CHECK-LABEL: @test_read_exec_hi(
 // CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.ballot.i64(i1 true)
 // CHECK: lshr i64 [[A:%.*]], 32