llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clangir Author: Ayokunle Amodu (ayokunle321) <details> <summary>Changes</summary> Adds codegen for the following NVVM atomic add builtins: - __nvvm_atom_add_gen_i (int) - __nvvm_atom_add_gen_l (long) - __nvvm_atom_add_gen_ll (long long) These are lowered to `atomicrmw` add instructions. --- Full diff: https://github.com/llvm/llvm-project/pull/210534.diff 4 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp (+7-8) - (modified) clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp (+4-4) - (modified) clang/lib/CIR/CodeGen/CIRGenFunction.h (+6) - (added) clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu (+37) ``````````diff diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp index 3c48b8b67d9ee..9673b38fc9f5c 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp @@ -163,11 +163,10 @@ static Address checkAtomicAlignment(CIRGenFunction &cgf, const CallExpr *e) { /// Utility to insert an atomic instruction based on Intrinsic::ID /// and the expression node. -static mlir::Value makeBinaryAtomicValue( - CIRGenFunction &cgf, cir::AtomicFetchKind kind, const CallExpr *expr, - mlir::Type *originalArgType = nullptr, - mlir::Value *emittedArgValue = nullptr, - cir::MemOrder ordering = cir::MemOrder::SequentiallyConsistent) { +mlir::Value CIRGenFunction::makeBinaryAtomicValue( + cir::AtomicFetchKind kind, const CallExpr *expr, mlir::Type *originalArgType, + mlir::Value *emittedArgValue, cir::MemOrder ordering) { + CIRGenFunction &cgf = *this; QualType type = expr->getType(); QualType ptrType = expr->getArg(0)->getType(); @@ -224,7 +223,7 @@ static mlir::Value makeBinaryAtomicValue( static RValue emitBinaryAtomic(CIRGenFunction &cgf, cir::AtomicFetchKind atomicOpkind, const CallExpr *e) { - return RValue::get(makeBinaryAtomicValue(cgf, atomicOpkind, e)); + return RValue::get(cgf.makeBinaryAtomicValue(atomicOpkind, e)); } template <typename BinOp> @@ -234,8 +233,8 @@ static RValue emitBinaryAtomicPost(CIRGenFunction &cgf, mlir::Value emittedArgValue; mlir::Type originalArgType; clang::QualType typ = e->getType(); - mlir::Value result = makeBinaryAtomicValue( - cgf, atomicOpkind, e, &originalArgType, &emittedArgValue); + mlir::Value result = cgf.makeBinaryAtomicValue( + atomicOpkind, e, &originalArgType, &emittedArgValue); clang::CIRGen::CIRGenBuilderTy &builder = cgf.getBuilder(); result = BinOp::create(builder, result.getLoc(), result, emittedArgValue); diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp index 4db2d7259c6ba..c8469d3a575b4 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp @@ -39,10 +39,10 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) { case NVPTX::BI__nvvm_atom_add_gen_i: case NVPTX::BI__nvvm_atom_add_gen_l: case NVPTX::BI__nvvm_atom_add_gen_ll: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeBinaryAtomicValue(cir::AtomicFetchKind::Add, expr, + /*originalArgType=*/nullptr, + /*emittedArgValue=*/nullptr, + cir::MemOrder::Relaxed); case NVPTX::BI__nvvm_atom_sub_gen_i: case NVPTX::BI__nvvm_atom_sub_gen_l: case NVPTX::BI__nvvm_atom_sub_gen_ll: diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h index 322355fde3957..fa926ef54f7e2 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunction.h +++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h @@ -1648,6 +1648,12 @@ class CIRGenFunction : public CIRGenTypeCache { const Expr *memOrder, bool isStore, bool isLoad, bool isFence, llvm::function_ref<void(cir::MemOrder)> emitAtomicOp); + mlir::Value makeBinaryAtomicValue( + cir::AtomicFetchKind kind, const clang::CallExpr *expr, + mlir::Type *originalArgType = nullptr, + mlir::Value *emittedArgValue = nullptr, + cir::MemOrder ordering = cir::MemOrder::SequentiallyConsistent); + mlir::LogicalResult emitAttributedStmt(const AttributedStmt &s); AutoVarEmission emitAutoVarAlloca(const clang::VarDecl &d, diff --git a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu new file mode 100644 index 0000000000000..f532b3fedf46c --- /dev/null +++ b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu @@ -0,0 +1,37 @@ +#include "Inputs/cuda.h" + +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_80 -x cuda \ +// RUN: -fcuda-is-device -fclangir -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_80 -x cuda \ +// RUN: -fcuda-is-device -fclangir -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s + +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_80 -x cuda \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// CIR-LABEL: @_Z19test_atom_add_gen_iPii +// CIR: cir.atomic.fetch add relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i +// LLVM-LABEL: @_Z19test_atom_add_gen_iPii +// LLVM: atomicrmw add ptr %{{.*}}, i32 %{{.*}} monotonic, align 4 +__device__ void test_atom_add_gen_i(int *p, int val) { + __nvvm_atom_add_gen_i(p, val); +} + +// CIR-LABEL: @_Z19test_atom_add_gen_lPll +// CIR: cir.atomic.fetch add relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z19test_atom_add_gen_lPll +// LLVM: atomicrmw add ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_add_gen_l(long *p, long val) { + __nvvm_atom_add_gen_l(p, val); +} + +// CIR-LABEL: @_Z20test_atom_add_gen_llPxx +// CIR: cir.atomic.fetch add relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z20test_atom_add_gen_llPxx +// LLVM: atomicrmw add ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_add_gen_ll(long long *p, long long val) { + __nvvm_atom_add_gen_ll(p, val); +} `````````` </details> https://github.com/llvm/llvm-project/pull/210534 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
