llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-amdgpu Author: Arseniy Obolenskiy (aobolensk) <details> <summary>Changes</summary> Allow vectors of floats in cir.atomic.fetch Bitcast v2bf16 operands from `<2 x i16>` to `<2 x bfloat>` and back, matching classic codegen --- Full diff: https://github.com/llvm/llvm-project/pull/231555.diff 6 Files Affected: - (modified) clang/include/clang/CIR/Dialect/IR/CIROps.td (+6-5) - (modified) clang/include/clang/CIR/Dialect/IR/CIRTypeConstraints.td (+7) - (modified) clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp (+16-12) - (modified) clang/lib/CIR/Dialect/IR/CIRDialect.cpp (+1-1) - (modified) clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp (+1-1) - (modified) clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-atomic.hip (+65) ``````````diff diff --git a/clang/include/clang/CIR/Dialect/IR/CIROps.td b/clang/include/clang/CIR/Dialect/IR/CIROps.td index 6d62fb3d93357..7651d15998faa 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIROps.td +++ b/clang/include/clang/CIR/Dialect/IR/CIROps.td @@ -9893,8 +9893,9 @@ def CIR_AtomicFetchOp : CIR_Op<"atomic.fetch", [ This operation takes 2 arguments: a pointer `ptr` and a value `val`. The type of `val` must match the pointee type of `ptr`. If the binary operation - is `add`, `sub`, `max`, or `min`, the type of `val` may either be an integer - type or a floating-point type. Otherwise, `val` must be an integer. + is `add`, `sub`, `max`, or `min`, the type of `val` may be an integer type, + a floating-point type, or a vector of floating-point type. Otherwise, `val` + must be an integer. This operation atomically loads the value from `ptr`, performs the binary operation as indicated by `binop` on the loaded value and `val`, and stores @@ -9921,10 +9922,10 @@ def CIR_AtomicFetchOp : CIR_Op<"atomic.fetch", [ %res = cir.atomic.fetch add seq_cst %ptr, %val : (!cir.ptr<!s32i>, !s32i) -> !s32i }]; - let results = (outs CIR_AnyIntOrFloatType:$result); + let results = (outs CIR_AnyIntOrFloatOrVecOfFloatType:$result); let arguments = (ins - Arg<CIR_PtrToIntOrFloatType, "", [MemRead, MemWrite]>:$ptr, - CIR_AnyIntOrFloatType:$val, + Arg<CIR_PtrToIntOrFloatOrVecOfFloatType, "", [MemRead, MemWrite]>:$ptr, + CIR_AnyIntOrFloatOrVecOfFloatType:$val, CIR_AtomicFetchKindAttr:$binop, Arg<CIR_MemOrderAttr, "memory order">:$mem_order, Arg<CIR_SyncScopeKindAttr, "synchronization scope">:$sync_scope, diff --git a/clang/include/clang/CIR/Dialect/IR/CIRTypeConstraints.td b/clang/include/clang/CIR/Dialect/IR/CIRTypeConstraints.td index 97d7b30236b82..8dd1d1ebb4b69 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRTypeConstraints.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRTypeConstraints.td @@ -379,6 +379,13 @@ def CIR_AnyFloatOrVecOfFloatType let cppFunctionName = "isFPOrVectorOfFPType"; } +def CIR_AnyIntOrFloatOrVecOfFloatType + : AnyTypeOf<[CIR_AnyIntOrFloatType, CIR_VectorOfFloatType], + "integer, floating point or vector of floating point type">; + +def CIR_PtrToIntOrFloatOrVecOfFloatType + : CIR_PtrToType<CIR_AnyIntOrFloatOrVecOfFloatType>; + // Types valid for bitwise ops (and/or/xor): // integer, boolean, or vector of bool or integer (no floating-point). def CIR_AnyBitwiseType diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index f28bfc67d9cc5..ccb78816e171e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -103,12 +103,21 @@ static mlir::Value emitAMDGPUAtomicRMW(CIRGenFunction &cgf, scope = decodeAMDGPUSyncScope(expr->getArg(3)); } + mlir::Type origTy = val.getType(); + if (auto vecTy = mlir::dyn_cast<cir::VectorType>(origTy); + vecTy && mlir::isa<cir::IntType>(vecTy.getElementType())) { + auto bf16VecTy = + cir::VectorType::get(builder.getBfloat16Ty(), vecTy.getSize()); + val = builder.createBitcast(loc, val, bf16VecTy); + ptr = ptr.withElementType(builder, bf16VecTy); + } + auto rmw = cir::AtomicFetchOp::create(builder, loc, ptr.emitRawPointer(), val, binOp, order, scope, isVolatile, /*fetch_first=*/true); rmw->setAttr(cir::CIRDialect::getAMDGPURawAtomicAttrName(), builder.getUnitAttr()); - return rmw->getResult(0); + return builder.createBitcast(loc, rmw->getResult(0), origTy); } // Emit the `amdgcn.dispatch.ptr` intrinsic, address-space-casting the @@ -1267,6 +1276,12 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f64: case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f32: case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f64: + case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2f16: + case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2bf16: + case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2f16: + case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2f16: + case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16: + case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16: return emitAMDGPUAtomicRMW(*this, expr, cir::AtomicFetchKind::Add, /*hasVolatileArg=*/false); case AMDGPU::BI__builtin_amdgcn_global_atomic_fmin_f64: @@ -1288,17 +1303,6 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, case AMDGPU::BI__builtin_amdgcn_ds_fmaxf: return emitAMDGPUAtomicRMW(*this, expr, cir::AtomicFetchKind::Max, /*hasVolatileArg=*/true); - case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2f16: - case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2bf16: - case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2f16: - case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2f16: - case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16: - case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; - } case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtn: case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtnl: return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.s.sendmsg.rtn", diff --git a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp index b40c0b965ff73..bff810c606896 100644 --- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp @@ -5034,7 +5034,7 @@ LogicalResult cir::AtomicFetchOp::verify() { getBinop() == cir::AtomicFetchKind::Minimum || getBinop() == cir::AtomicFetchKind::MaximumNum || getBinop() == cir::AtomicFetchKind::MinimumNum) && - !mlir::isa<cir::FPTypeInterface>(getVal().getType())) + !cir::isFPOrVectorOfFPType(getVal().getType())) return emitError("atomic maximum, minimum, maximum_num, and minimum_num " "operation could only operate on floating-point values"); diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index 868bcfc206578..3ec633aac98fc 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -1568,7 +1568,7 @@ mlir::LogicalResult CIRToLLVMAtomicFetchOpLowering::matchAndRewrite( if (auto intTy = mlir::dyn_cast<cir::IntType>(op.getVal().getType())) { isInt = true; isSignedInt = intTy.isSigned(); - } else if (mlir::isa<cir::SingleType, cir::DoubleType>( + } else if (mlir::isa<cir::SingleType, cir::DoubleType, cir::VectorType>( op.getVal().getType())) { isInt = false; } else { diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-atomic.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-atomic.hip index 9525994fc6800..091acaac5d1ec 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-atomic.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-atomic.hip @@ -12,6 +12,12 @@ typedef __attribute__((address_space(3))) float *float_lds_ptr; typedef __attribute__((address_space(3))) double *double_lds_ptr; typedef __attribute__((address_space(1))) float *float_global_ptr; typedef __attribute__((address_space(1))) double *double_global_ptr; +typedef _Float16 half2 __attribute__((ext_vector_type(2))); +typedef short short2 __attribute__((ext_vector_type(2))); +typedef __attribute__((address_space(3))) half2 *half2_lds_ptr; +typedef __attribute__((address_space(1))) half2 *half2_global_ptr; +typedef __attribute__((address_space(3))) short2 *short2_lds_ptr; +typedef __attribute__((address_space(1))) short2 *short2_global_ptr; // CIR-LABEL: @_Z5inc32PVjj // CIR: cir.atomic.fetch uinc_wrap seq_cst syncscope(device) fetch_first %{{.+}}, %{{.+}} volatile : (!cir.ptr<!u32i>, !u32i) -> !u32i {cir.amdgpu_raw_atomic} @@ -171,6 +177,65 @@ __device__ double flat_atomic_fmax_f64(double *p, double v) { return __builtin_amdgcn_flat_atomic_fmax_f64(p, v); } +// CIR-LABEL: @_Z20ds_atomic_fadd_v2f16 +// CIR: cir.atomic.fetch add relaxed syncscope(device) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.vector<2 x !cir.f16>, target_address_space(3)>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x !cir.f16> {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z20ds_atomic_fadd_v2f16 +// LLVM: atomicrmw fadd ptr addrspace(3) %{{.+}}, <2 x half> %{{.+}} syncscope("agent") monotonic, align 4 +// LLVM-NOT: amdgpu.no.fine.grained.memory +__device__ half2 ds_atomic_fadd_v2f16(half2_lds_ptr p, half2 v) { + return __builtin_amdgcn_ds_atomic_fadd_v2f16(p, v); +} + +// CIR-LABEL: @_Z24global_atomic_fadd_v2f16 +// CIR: cir.atomic.fetch add relaxed syncscope(device) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.vector<2 x !cir.f16>, target_address_space(1)>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x !cir.f16> {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z24global_atomic_fadd_v2f16 +// LLVM: atomicrmw fadd ptr addrspace(1) %{{.+}}, <2 x half> %{{.+}} syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}} +__device__ half2 global_atomic_fadd_v2f16(half2_global_ptr p, half2 v) { + return __builtin_amdgcn_global_atomic_fadd_v2f16(p, v); +} + +// CIR-LABEL: @_Z22flat_atomic_fadd_v2f16 +// CIR: cir.atomic.fetch add relaxed syncscope(device) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.vector<2 x !cir.f16>>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x !cir.f16> {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z22flat_atomic_fadd_v2f16 +// LLVM: atomicrmw fadd ptr %{{.+}}, <2 x half> %{{.+}} syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}} +__device__ half2 flat_atomic_fadd_v2f16(half2 *p, half2 v) { + return __builtin_amdgcn_flat_atomic_fadd_v2f16(p, v); +} + +// CIR-LABEL: @_Z21ds_atomic_fadd_v2bf16 +// CIR: %[[VAL:.+]] = cir.cast bitcast %{{.+}} : !cir.vector<2 x !s16i> -> !cir.vector<2 x !cir.bf16> +// CIR: %[[PTR:.+]] = cir.cast bitcast %{{.+}} : !cir.ptr<!cir.vector<2 x !s16i>, target_address_space(3)> -> !cir.ptr<!cir.vector<2 x !cir.bf16>, target_address_space(3)> +// CIR: %[[RMW:.+]] = cir.atomic.fetch add relaxed syncscope(device) fetch_first %[[PTR]], %[[VAL]] : (!cir.ptr<!cir.vector<2 x !cir.bf16>, target_address_space(3)>, !cir.vector<2 x !cir.bf16>) -> !cir.vector<2 x !cir.bf16> {cir.amdgpu_raw_atomic} +// CIR: cir.cast bitcast %[[RMW]] : !cir.vector<2 x !cir.bf16> -> !cir.vector<2 x !s16i> +// LLVM-LABEL: @_Z21ds_atomic_fadd_v2bf16 +// LLVM: %[[VAL:.+]] = bitcast <2 x i16> %{{.+}} to <2 x bfloat> +// LLVM: %[[RMW:.+]] = atomicrmw fadd ptr addrspace(3) %{{.+}}, <2 x bfloat> %[[VAL]] syncscope("agent") monotonic, align 4 +// LLVM-NOT: amdgpu.no.fine.grained.memory +// LLVM: bitcast <2 x bfloat> %[[RMW]] to <2 x i16> +__device__ short2 ds_atomic_fadd_v2bf16(short2_lds_ptr p, short2 v) { + return __builtin_amdgcn_ds_atomic_fadd_v2bf16(p, v); +} + +// CIR-LABEL: @_Z25global_atomic_fadd_v2bf16 +// CIR: cir.atomic.fetch add relaxed syncscope(device) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.vector<2 x !cir.bf16>, target_address_space(1)>, !cir.vector<2 x !cir.bf16>) -> !cir.vector<2 x !cir.bf16> {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z25global_atomic_fadd_v2bf16 +// LLVM: %[[VAL:.+]] = bitcast <2 x i16> %{{.+}} to <2 x bfloat> +// LLVM: %[[RMW:.+]] = atomicrmw fadd ptr addrspace(1) %{{.+}}, <2 x bfloat> %[[VAL]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}} +// LLVM: bitcast <2 x bfloat> %[[RMW]] to <2 x i16> +__device__ short2 global_atomic_fadd_v2bf16(short2_global_ptr p, short2 v) { + return __builtin_amdgcn_global_atomic_fadd_v2bf16(p, v); +} + +// CIR-LABEL: @_Z23flat_atomic_fadd_v2bf16 +// CIR: cir.atomic.fetch add relaxed syncscope(device) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.vector<2 x !cir.bf16>>, !cir.vector<2 x !cir.bf16>) -> !cir.vector<2 x !cir.bf16> {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z23flat_atomic_fadd_v2bf16 +// LLVM: %[[VAL:.+]] = bitcast <2 x i16> %{{.+}} to <2 x bfloat> +// LLVM: %[[RMW:.+]] = atomicrmw fadd ptr %{{.+}}, <2 x bfloat> %[[VAL]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}} +// LLVM: bitcast <2 x bfloat> %[[RMW]] to <2 x i16> +__device__ short2 flat_atomic_fadd_v2bf16(short2 *p, short2 v) { + return __builtin_amdgcn_flat_atomic_fadd_v2bf16(p, v); +} + // CIR-LABEL: @_Z5fencev // CIR: cir.atomic.fence syncscope(device) seq_cst // LLVM-LABEL: @_Z5fencev `````````` </details> https://github.com/llvm/llvm-project/pull/231555 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
