https://github.com/RiverDave created
https://github.com/llvm/llvm-project/pull/221255
This is from some NYI's I found while compiling MiniFE
__syncthreads_{and,or,count} path; match classic codegen (icmp ne,
cta.red.*.aligned.all, zext i1->i32 for and/or).
>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] [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;
+}
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
