https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/212813
>From db065b623203ef797a37d310029c2f5cc1f744b4 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Wed, 29 Jul 2026 18:25:18 +0200 Subject: [PATCH 1/4] [clang][AMDGPU] Widen ballot for read_exec_lo/hi to wavefront size 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 --- clang/lib/Basic/Targets/SPIR.h | 4 ++++ clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 13 +++++++++++-- clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl | 5 ++--- clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 5 ++--- 4 files changed, 19 insertions(+), 8 deletions(-) diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h index a16c8ec79d2f2..7cfe927ac6463 100644 --- a/clang/lib/Basic/Targets/SPIR.h +++ b/clang/lib/Basic/Targets/SPIR.h @@ -453,6 +453,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 72d6f536165c5..0dfb5e91f72f9 100644 --- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp @@ -275,13 +275,22 @@ 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; + llvm::Type *BallotType = RegisterType; + if (BallotType->getIntegerBitWidth() < WaveSize) + BallotType = Builder.getIntNTy(WaveSize); + + Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_ballot, {BallotType}); llvm::Value *Call = Builder.CreateCall(F, {Builder.getInt1(true)}); if (isExecHi) { @@ -290,7 +299,7 @@ static Value *EmitAMDGCNBallotForExec(CodeGenFunction &CGF, const CallExpr *E, return Rt2; } - return Call; + return Builder.CreateTrunc(Call, CGF.ConvertType(E->getType())); } 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 >From 4f2e870b86d9cfdaed6c59dccd03e332f3a7e9e1 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Thu, 30 Jul 2026 06:55:54 +0200 Subject: [PATCH 2/4] Address comment --- clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp index 0dfb5e91f72f9..8023703c93d4e 100644 --- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp @@ -299,7 +299,7 @@ static Value *EmitAMDGCNBallotForExec(CodeGenFunction &CGF, const CallExpr *E, return Rt2; } - return Builder.CreateTrunc(Call, CGF.ConvertType(E->getType())); + return Builder.CreateTrunc(Call, ValueType); } static llvm::Value *loadTextureDescPtorAsVec8I32(CodeGenFunction &CGF, >From 45ed07df335c590e9152c89078d1e1424dc7fb3f Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Mon, 3 Aug 2026 10:13:54 +0200 Subject: [PATCH 3/4] Update clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp Co-authored-by: Matt Arsenault <[email protected]> --- clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp index 8023703c93d4e..93620c63a6e0a 100644 --- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp @@ -286,9 +286,7 @@ static Value *EmitAMDGCNBallotForExec(CodeGenFunction &CGF, const CallExpr *E, CodeGen::CodeGenModule &CGM = CGF.CGM; unsigned WaveSize = CGF.getTarget().getGridValue().GV_Warp_Size; - llvm::Type *BallotType = RegisterType; - if (BallotType->getIntegerBitWidth() < WaveSize) - BallotType = Builder.getIntNTy(WaveSize); + llvm::Type *BallotType = Builder.getIntNTy(WaveSize); Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_ballot, {BallotType}); llvm::Value *Call = Builder.CreateCall(F, {Builder.getInt1(true)}); >From 8eab44458826c2e6c187c871db6e2e774004ca94 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Mon, 3 Aug 2026 10:51:53 +0200 Subject: [PATCH 4/4] fix wave size guard --- clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp index 93620c63a6e0a..547cd987af35d 100644 --- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp @@ -286,7 +286,8 @@ static Value *EmitAMDGCNBallotForExec(CodeGenFunction &CGF, const CallExpr *E, CodeGen::CodeGenModule &CGM = CGF.CGM; unsigned WaveSize = CGF.getTarget().getGridValue().GV_Warp_Size; - llvm::Type *BallotType = Builder.getIntNTy(WaveSize); + 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)}); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
