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 <[email protected]> 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
