diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
index ae994005c588a..c5042f2621517 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
@@ -69,6 +69,23 @@ static mlir::Value emitUnaryNVVMIntrinsic(CIRGenFunction &cgf,
       .getResult();
 }
 
+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::Type si32Ty = builder.getSInt32Ty();
+  mlir::Value zero = builder.getNullValue(si32Ty, loc);
+  mlir::Value pred = builder.createCompare(
+      loc, cir::CmpOpKind::ne, cgf.emitScalarExpr(expr->getArg(0)), zero);
+  mlir::Type resultTy = returnsPred ? mlir::Type(builder.getBoolTy()) : si32Ty;
+  mlir::Value result = builder.emitIntrinsicCallOp(
+      loc, intrinsicName, resultTy, mlir::ValueRange{zero, pred});
+  if (returnsPred)
+    result = builder.createBoolToInt(result, si32Ty);
+  return result;
+}
+
 static mlir::Value makeScopedAtomicRMW(CIRGenFunction &cgf,
                                        const CallExpr *expr,
                                        cir::AtomicFetchKind kind,
@@ -997,20 +1014,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;
+}