https://github.com/RiverDave updated https://github.com/llvm/llvm-project/pull/221255
>From 2b89a9eee6f11dccdb49ce9f24a268bf309a0012 Mon Sep 17 00:00:00 2001 From: David Rivera <[email protected]> Date: Fri, 4 Sep 2026 10:52:51 -0400 Subject: [PATCH 1/2] [CIR][NVPTX] Lower __nvvm_bar0_{and,or,popc} builtins These are the MiniFE __syncthreads_{and,or,count} path; match classic codegen (icmp ne, cta.red.*.aligned.all, zext i1->i32 for and/or). --- clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp | 40 ++++++++++++------ .../NVPTX/builtins-nvptx-sync.cu | 41 ++++++++++++++++++- 2 files changed, 68 insertions(+), 13 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp index ae994005c588a..1ae3d18760006 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp @@ -69,6 +69,26 @@ static mlir::Value emitUnaryNVVMIntrinsic(CIRGenFunction &cgf, .getResult(); } +/// Lower __nvvm_bar0_{and,or,popc} like classic codegen: +/// icmp ne %arg, 0; call nvvm.barrier.cta.red.*.aligned.all(0, pred); +/// zext i1 to i32 for and/or. +static mlir::Value emitBar0Reduction(CIRGenFunction &cgf, const CallExpr *expr, + llvm::StringRef intrinsicName, + bool returnsPred) { + CIRGenBuilderTy &builder = cgf.getBuilder(); + mlir::Location loc = cgf.getLoc(expr->getExprLoc()); + mlir::Value zero = builder.getConstInt(loc, builder.getSInt32Ty(), 0); + mlir::Value pred = builder.createCompare( + loc, cir::CmpOpKind::ne, cgf.emitScalarExpr(expr->getArg(0)), zero); + mlir::Type resultTy = + returnsPred ? mlir::Type(builder.getBoolTy()) : builder.getSInt32Ty(); + mlir::Value result = builder.emitIntrinsicCallOp( + loc, intrinsicName, resultTy, mlir::ValueRange{zero, pred}); + if (returnsPred) + result = builder.createBoolToInt(result, builder.getSInt32Ty()); + return result; +} + static mlir::Value makeScopedAtomicRMW(CIRGenFunction &cgf, const CallExpr *expr, cir::AtomicFetchKind kind, @@ -997,20 +1017,16 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) { mlir::ValueRange{emitScalarExpr(expr->getArg(0)), emitScalarExpr(expr->getArg(1))}); case NVPTX::BI__nvvm_bar0_and: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return emitBar0Reduction(*this, expr, + "nvvm.barrier.cta.red.and.aligned.all", + /*returnsPred=*/true); case NVPTX::BI__nvvm_bar0_or: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return emitBar0Reduction(*this, expr, "nvvm.barrier.cta.red.or.aligned.all", + /*returnsPred=*/true); case NVPTX::BI__nvvm_bar0_popc: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return emitBar0Reduction(*this, expr, + "nvvm.barrier.cta.red.popc.aligned.all", + /*returnsPred=*/false); default: return std::nullopt; diff --git a/clang/test/CIR/CodeGenBuiltins/NVPTX/builtins-nvptx-sync.cu b/clang/test/CIR/CodeGenBuiltins/NVPTX/builtins-nvptx-sync.cu index b1a57856dc392..b9277d3f587e3 100644 --- a/clang/test/CIR/CodeGenBuiltins/NVPTX/builtins-nvptx-sync.cu +++ b/clang/test/CIR/CodeGenBuiltins/NVPTX/builtins-nvptx-sync.cu @@ -15,7 +15,7 @@ #define __device__ __attribute__((device)) -// Tests CIR/LLVM lowering for NVPTX CTA-level sync barrier builtins. +// Tests CIR/LLVM lowering for NVPTX CTA-level sync and bar0 reduction builtins. // Mirrors the relevant slices of clang/test/CodeGen/builtins-nvptx.c and // clang/test/CodeGen/builtins-nvptx-ptx60.cu. @@ -43,3 +43,42 @@ __device__ void nvvm_sync(unsigned mask) { // OGCG: call void @llvm.nvvm.barrier.cta.sync.count(i32 %{{.*}}, i32 0) __nvvm_barrier_sync_cnt(mask, 0); } + +// CIR-LABEL: cir.func {{.*}} @_Z20nvvm_bar0_reductionsi +// LLVM-LABEL: define{{.*}} i32 @_Z20nvvm_bar0_reductionsi( +// OGCG-LABEL: define{{.*}} i32 @_Z20nvvm_bar0_reductionsi( +__device__ int nvvm_bar0_reductions(int i) { + int ret = 0; + + // CIR: %[[NE_AND:.*]] = cir.cmp ne {{.*}} : !s32i + // CIR: %[[AND:.*]] = cir.call_llvm_intrinsic "nvvm.barrier.cta.red.and.aligned.all" {{.*}} : (!s32i, !cir.bool) -> !cir.bool + // CIR: cir.cast bool_to_int %[[AND]] : !cir.bool -> !s32i + // LLVM: %[[NE_AND:.*]] = icmp ne i32 %{{.*}}, 0 + // LLVM: %[[AND:.*]] = call i1 @llvm.nvvm.barrier.cta.red.and.aligned.all(i32 0, i1 %[[NE_AND]]) + // LLVM: zext i1 %[[AND]] to i32 + // OGCG: %[[NE_AND:.*]] = icmp ne i32 %{{.*}}, 0 + // OGCG: %[[AND:.*]] = call i1 @llvm.nvvm.barrier.cta.red.and.aligned.all(i32 0, i1 %[[NE_AND]]) + // OGCG: zext i1 %[[AND]] to i32 + ret += __nvvm_bar0_and(i); + + // CIR: %[[NE_OR:.*]] = cir.cmp ne {{.*}} : !s32i + // CIR: %[[OR:.*]] = cir.call_llvm_intrinsic "nvvm.barrier.cta.red.or.aligned.all" {{.*}} : (!s32i, !cir.bool) -> !cir.bool + // CIR: cir.cast bool_to_int %[[OR]] : !cir.bool -> !s32i + // LLVM: %[[NE_OR:.*]] = icmp ne i32 %{{.*}}, 0 + // LLVM: %[[OR:.*]] = call i1 @llvm.nvvm.barrier.cta.red.or.aligned.all(i32 0, i1 %[[NE_OR]]) + // LLVM: zext i1 %[[OR]] to i32 + // OGCG: %[[NE_OR:.*]] = icmp ne i32 %{{.*}}, 0 + // OGCG: %[[OR:.*]] = call i1 @llvm.nvvm.barrier.cta.red.or.aligned.all(i32 0, i1 %[[NE_OR]]) + // OGCG: zext i1 %[[OR]] to i32 + ret += __nvvm_bar0_or(i); + + // CIR: %[[NE_POPC:.*]] = cir.cmp ne {{.*}} : !s32i + // CIR: cir.call_llvm_intrinsic "nvvm.barrier.cta.red.popc.aligned.all" {{.*}} : (!s32i, !cir.bool) -> !s32i + // LLVM: %[[NE_POPC:.*]] = icmp ne i32 %{{.*}}, 0 + // LLVM: call i32 @llvm.nvvm.barrier.cta.red.popc.aligned.all(i32 0, i1 %[[NE_POPC]]) + // OGCG: %[[NE_POPC:.*]] = icmp ne i32 %{{.*}}, 0 + // OGCG: call i32 @llvm.nvvm.barrier.cta.red.popc.aligned.all(i32 0, i1 %[[NE_POPC]]) + ret += __nvvm_bar0_popc(i); + + return ret; +} >From c8833275d0bb61c3aa7cfe7f0850d59657714e4f Mon Sep 17 00:00:00 2001 From: David Rivera <[email protected]> Date: Fri, 4 Sep 2026 11:37:10 -0400 Subject: [PATCH 2/2] [CIR][NVPTX] Drop emitBar0Reduction comment --- clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp | 3 --- 1 file changed, 3 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp index 1ae3d18760006..8f8c9f9cdd607 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp @@ -69,9 +69,6 @@ static mlir::Value emitUnaryNVVMIntrinsic(CIRGenFunction &cgf, .getResult(); } -/// Lower __nvvm_bar0_{and,or,popc} like classic codegen: -/// icmp ne %arg, 0; call nvvm.barrier.cta.red.*.aligned.all(0, pred); -/// zext i1 to i32 for and/or. static mlir::Value emitBar0Reduction(CIRGenFunction &cgf, const CallExpr *expr, llvm::StringRef intrinsicName, bool returnsPred) { _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
