llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-amdgpu @llvm/pr-subscribers-clang-codegen Author: Pierre van Houtryve (Pierre-vh) <details> <summary>Changes</summary> Load monitor operations make more sense as atomic operations, as non-atomic operations cannot be used for inter-thread communication w/o additional synchronization. The previous built-in made it work because one could just override the CPol bits, but that bypasses the memory model and forces the user to learn about ISA bits encoding. Making load monitor an atomic operation has a couple of advantages. First, the memory model foundation for it is stronger. We just lean on the existing rules for atomic operations. Second, the CPol bits are abstracted away from the user, which avoids leaking ISA details into the API. This patch also adds supporting memory model and intrinsics documentation to AMDGPUUsage. Solves SWDEV-516398. --- Patch is 64.90 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/177343.diff 16 Files Affected: - (modified) clang/include/clang/Basic/BuiltinsAMDGPU.td (+6-6) - (modified) clang/include/clang/Sema/SemaAMDGPU.h (+12) - (modified) clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp (+31-24) - (modified) clang/lib/Sema/SemaAMDGPU.cpp (+48-21) - (modified) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-load-monitor.cl (+24-24) - (modified) clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl (+30-8) - (modified) llvm/docs/AMDGPUUsage.rst (+91-17) - (modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+10-11) - (modified) llvm/lib/Target/AMDGPU/AMDGPUGISel.td (+4) - (modified) llvm/lib/Target/AMDGPU/AMDGPUInstrInfo.td (+9) - (modified) llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp (+20) - (modified) llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp (+16-6) - (modified) llvm/lib/Target/AMDGPU/FLATInstructions.td (+6-6) - (modified) llvm/lib/Target/AMDGPU/SIISelLowering.cpp (+58-28) - (modified) llvm/lib/Target/AMDGPU/SIInstructions.td (+10) - (modified) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.load.monitor.gfx1250.ll (+74-54) ``````````diff diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td index 12ffad305e7c0..3032b387d91ab 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.td +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td @@ -707,12 +707,12 @@ def __builtin_amdgcn_s_cluster_barrier : AMDGPUBuiltin<"void()", [], "gfx1250-in def __builtin_amdgcn_flat_prefetch : AMDGPUBuiltin<"void(void const address_space<0> *, _Constant int)", [Const], "vmem-pref-insts">; def __builtin_amdgcn_global_prefetch : AMDGPUBuiltin<"void(void const address_space<1> *, _Constant int)", [Const], "vmem-pref-insts">; -def __builtin_amdgcn_global_load_monitor_b32 : AMDGPUBuiltin<"int(int address_space<1> *, _Constant int)", [Const], "gfx1250-insts">; -def __builtin_amdgcn_global_load_monitor_b64 : AMDGPUBuiltin<"_ExtVector<2, int>(_ExtVector<2, int address_space<1> *>, _Constant int)", [Const], "gfx1250-insts">; -def __builtin_amdgcn_global_load_monitor_b128 : AMDGPUBuiltin<"_ExtVector<4, int>(_ExtVector<4, int address_space<1> *>, _Constant int)", [Const], "gfx1250-insts">; -def __builtin_amdgcn_flat_load_monitor_b32 : AMDGPUBuiltin<"int(int address_space<0> *, _Constant int)", [Const], "gfx1250-insts">; -def __builtin_amdgcn_flat_load_monitor_b64 : AMDGPUBuiltin<"_ExtVector<2, int>(_ExtVector<2, int address_space<0> *>, _Constant int)", [Const], "gfx1250-insts">; -def __builtin_amdgcn_flat_load_monitor_b128 : AMDGPUBuiltin<"_ExtVector<4, int>(_ExtVector<4, int address_space<0> *>, _Constant int)", [Const], "gfx1250-insts">; +def __builtin_amdgcn_global_atomic_load_monitor_b32 : AMDGPUBuiltin<"int(int address_space<1> *, _Constant int, char const *)", [Const], "gfx1250-insts">; +def __builtin_amdgcn_global_atomic_load_monitor_b64 : AMDGPUBuiltin<"_ExtVector<2, int>(_ExtVector<2, int address_space<1> *>, _Constant int, char const *)", [Const], "gfx1250-insts">; +def __builtin_amdgcn_global_atomic_load_monitor_b128 : AMDGPUBuiltin<"_ExtVector<4, int>(_ExtVector<4, int address_space<1> *>, _Constant int, char const *)", [Const], "gfx1250-insts">; +def __builtin_amdgcn_flat_atomic_load_monitor_b32 : AMDGPUBuiltin<"int(int address_space<0> *, _Constant int, char const *)", [Const], "gfx1250-insts">; +def __builtin_amdgcn_flat_atomic_load_monitor_b64 : AMDGPUBuiltin<"_ExtVector<2, int>(_ExtVector<2, int address_space<0> *>, _Constant int, char const *)", [Const], "gfx1250-insts">; +def __builtin_amdgcn_flat_atomic_load_monitor_b128 : AMDGPUBuiltin<"_ExtVector<4, int>(_ExtVector<4, int address_space<0> *>, _Constant int, char const *)", [Const], "gfx1250-insts">; def __builtin_amdgcn_cluster_load_b32 : AMDGPUBuiltin<"int(int address_space<1> *, _Constant int, int)", [Const], "mcast-load-insts,wavefrontsize32">; def __builtin_amdgcn_cluster_load_b64 : AMDGPUBuiltin<"_ExtVector<2, int>(_ExtVector<2, int address_space<1> *>, _Constant int, int)", [Const], "mcast-load-insts,wavefrontsize32">; def __builtin_amdgcn_cluster_load_b128 : AMDGPUBuiltin<"_ExtVector<4, int>(_ExtVector<4, int address_space<1> *>, _Constant int, int)", [Const], "mcast-load-insts,wavefrontsize32">; diff --git a/clang/include/clang/Sema/SemaAMDGPU.h b/clang/include/clang/Sema/SemaAMDGPU.h index bac812a9d4fcf..eb6e73dd7322f 100644 --- a/clang/include/clang/Sema/SemaAMDGPU.h +++ b/clang/include/clang/Sema/SemaAMDGPU.h @@ -26,7 +26,19 @@ class SemaAMDGPU : public SemaBase { bool CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID, CallExpr *TheCall); + /// Emits a diagnostic if the arg at \p ArgIdx of \p TheCall is not a string + /// literal. \returns true if a diagnostic was emitted. + bool checkStringLiteralArg(CallExpr *TheCall, unsigned ArgIdx); + + /// Emits a diagnostic if the arg at \p ArgIdx of \p TheCall is not atomic + /// ordering encoded in the C ABI format, or if the atomic ordering is not + /// valid for the operation type as defined by \p MayLoad and \p MayStore. + /// \returns true if a diagnostic was emitted. + bool checkAtomicOrderingCABIArg(CallExpr *TheCall, unsigned ArgIdx, + bool MayLoad, bool MayStore); + bool checkCoopAtomicFunctionCall(CallExpr *TheCall, bool IsStore); + bool checkAtomicMonitorLoad(CallExpr *TheCall); bool checkMovDPPFunctionCall(CallExpr *TheCall, unsigned NumArgs, unsigned NumDataArgs); diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp index a8a5bc348f00c..5c720619ec7db 100644 --- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp @@ -366,6 +366,14 @@ void CodeGenFunction::AddAMDGPUFenceAddressSpaceMMRA(llvm::Instruction *Inst, Inst->setMetadata(LLVMContext::MD_mmra, MMRAMetadata::getMD(Ctx, MMRAs)); } +static llvm::MetadataAsValue *getStringAsMDValue(llvm::LLVMContext &Ctx, + const clang::Expr *E) { + StringRef Arg = + cast<clang::StringLiteral>(E->IgnoreParenCasts())->getString(); + llvm::MDNode *MD = llvm::MDNode::get(Ctx, {llvm::MDString::get(Ctx, Arg)}); + return llvm::MetadataAsValue::get(Ctx, MD); +} + static Intrinsic::ID getIntrinsicIDforWaveReduction(unsigned BuiltinID) { switch (BuiltinID) { default: @@ -781,40 +789,42 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, llvm::Function *F = CGM.getIntrinsic(IID, {LoadTy}); return Builder.CreateCall(F, {Addr}); } - case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b32: - case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b64: - case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b128: - case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b32: - case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b64: - case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b128: { + case AMDGPU::BI__builtin_amdgcn_global_atomic_load_monitor_b32: + case AMDGPU::BI__builtin_amdgcn_global_atomic_load_monitor_b64: + case AMDGPU::BI__builtin_amdgcn_global_atomic_load_monitor_b128: + case AMDGPU::BI__builtin_amdgcn_flat_atomic_load_monitor_b32: + case AMDGPU::BI__builtin_amdgcn_flat_atomic_load_monitor_b64: + case AMDGPU::BI__builtin_amdgcn_flat_atomic_load_monitor_b128: { Intrinsic::ID IID; switch (BuiltinID) { - case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b32: - IID = Intrinsic::amdgcn_global_load_monitor_b32; + case AMDGPU::BI__builtin_amdgcn_global_atomic_load_monitor_b32: + IID = Intrinsic::amdgcn_global_atomic_load_monitor_b32; break; - case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b64: - IID = Intrinsic::amdgcn_global_load_monitor_b64; + case AMDGPU::BI__builtin_amdgcn_global_atomic_load_monitor_b64: + IID = Intrinsic::amdgcn_global_atomic_load_monitor_b64; break; - case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b128: - IID = Intrinsic::amdgcn_global_load_monitor_b128; + case AMDGPU::BI__builtin_amdgcn_global_atomic_load_monitor_b128: + IID = Intrinsic::amdgcn_global_atomic_load_monitor_b128; break; - case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b32: - IID = Intrinsic::amdgcn_flat_load_monitor_b32; + case AMDGPU::BI__builtin_amdgcn_flat_atomic_load_monitor_b32: + IID = Intrinsic::amdgcn_flat_atomic_load_monitor_b32; break; - case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b64: - IID = Intrinsic::amdgcn_flat_load_monitor_b64; + case AMDGPU::BI__builtin_amdgcn_flat_atomic_load_monitor_b64: + IID = Intrinsic::amdgcn_flat_atomic_load_monitor_b64; break; - case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b128: - IID = Intrinsic::amdgcn_flat_load_monitor_b128; + case AMDGPU::BI__builtin_amdgcn_flat_atomic_load_monitor_b128: + IID = Intrinsic::amdgcn_flat_atomic_load_monitor_b128; break; } + LLVMContext &Ctx = CGM.getLLVMContext(); llvm::Type *LoadTy = ConvertType(E->getType()); llvm::Value *Addr = EmitScalarExpr(E->getArg(0)); - llvm::Value *Val = EmitScalarExpr(E->getArg(1)); + llvm::Value *AO = EmitScalarExpr(E->getArg(1)); + llvm::Value *Scope = getStringAsMDValue(Ctx, E->getArg(2)); llvm::Function *F = CGM.getIntrinsic(IID, {LoadTy}); - return Builder.CreateCall(F, {Addr, Val}); + return Builder.CreateCall(F, {Addr, AO, Scope}); } case AMDGPU::BI__builtin_amdgcn_cluster_load_b32: case AMDGPU::BI__builtin_amdgcn_cluster_load_b64: @@ -876,10 +886,7 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, const unsigned ScopeArg = E->getNumArgs() - 1; for (unsigned i = 0; i != ScopeArg; ++i) Args.push_back(EmitScalarExpr(E->getArg(i))); - StringRef Arg = cast<StringLiteral>(E->getArg(ScopeArg)->IgnoreParenCasts()) - ->getString(); - llvm::MDNode *MD = llvm::MDNode::get(Ctx, {llvm::MDString::get(Ctx, Arg)}); - Args.push_back(llvm::MetadataAsValue::get(Ctx, MD)); + Args.push_back(getStringAsMDValue(Ctx, E->getArg(ScopeArg))); // Intrinsic is typed based on the pointer AS. Pointer is always the first // argument. llvm::Function *F = CGM.getIntrinsic(IID, {Args[0]->getType()}); diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index 4261e1849133f..a53cadd27a184 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -119,6 +119,13 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID, case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_16x8B: case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_8x16B: return checkCoopAtomicFunctionCall(TheCall, /*IsStore=*/true); + case AMDGPU::BI__builtin_amdgcn_flat_atomic_load_monitor_b32: + case AMDGPU::BI__builtin_amdgcn_flat_atomic_load_monitor_b64: + case AMDGPU::BI__builtin_amdgcn_flat_atomic_load_monitor_b128: + case AMDGPU::BI__builtin_amdgcn_global_atomic_load_monitor_b32: + case AMDGPU::BI__builtin_amdgcn_global_atomic_load_monitor_b64: + case AMDGPU::BI__builtin_amdgcn_global_atomic_load_monitor_b128: + return checkAtomicMonitorLoad(TheCall); case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32: case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32: case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32: @@ -341,22 +348,9 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID, return false; } -bool SemaAMDGPU::checkCoopAtomicFunctionCall(CallExpr *TheCall, bool IsStore) { - bool Fail = false; - - // First argument is a global or generic pointer. - Expr *PtrArg = TheCall->getArg(0); - QualType PtrTy = PtrArg->getType()->getPointeeType(); - unsigned AS = getASTContext().getTargetAddressSpace(PtrTy.getAddressSpace()); - if (AS != llvm::AMDGPUAS::FLAT_ADDRESS && - AS != llvm::AMDGPUAS::GLOBAL_ADDRESS) { - Fail = true; - Diag(TheCall->getBeginLoc(), diag::err_amdgcn_coop_atomic_invalid_as) - << PtrArg->getSourceRange(); - } - - // Check atomic ordering - Expr *AtomicOrdArg = TheCall->getArg(IsStore ? 2 : 1); +bool SemaAMDGPU::checkAtomicOrderingCABIArg(CallExpr *TheCall, unsigned ArgIdx, + bool MayLoad, bool MayStore) { + Expr *AtomicOrdArg = TheCall->getArg(ArgIdx); Expr::EvalResult AtomicOrdArgRes; if (!AtomicOrdArg->EvaluateAsInt(AtomicOrdArgRes, getASTContext())) llvm_unreachable("Intrinsic requires imm for atomic ordering argument!"); @@ -366,22 +360,55 @@ bool SemaAMDGPU::checkCoopAtomicFunctionCall(CallExpr *TheCall, bool IsStore) { // Atomic ordering cannot be acq_rel in any case, acquire for stores or // release for loads. if (!llvm::isValidAtomicOrderingCABI((unsigned)Ord) || - (Ord == llvm::AtomicOrderingCABI::acq_rel) || - Ord == (IsStore ? llvm::AtomicOrderingCABI::acquire - : llvm::AtomicOrderingCABI::release)) { + (!(MayLoad && MayStore) && (Ord == llvm::AtomicOrderingCABI::acq_rel)) || + (!MayLoad && Ord == llvm::AtomicOrderingCABI::acquire) || + (!MayStore && Ord == llvm::AtomicOrderingCABI::release)) { return Diag(AtomicOrdArg->getBeginLoc(), diag::warn_atomic_op_has_invalid_memory_order) << 0 << AtomicOrdArg->getSourceRange(); } - // Last argument is a string literal + return false; +} + +bool SemaAMDGPU::checkStringLiteralArg(CallExpr *TheCall, unsigned ArgIdx) { Expr *Arg = TheCall->getArg(TheCall->getNumArgs() - 1); if (!isa<StringLiteral>(Arg->IgnoreParenImpCasts())) { - Fail = true; Diag(TheCall->getBeginLoc(), diag::err_expr_not_string_literal) << Arg->getSourceRange(); + return true; + } + return false; +} + +bool SemaAMDGPU::checkCoopAtomicFunctionCall(CallExpr *TheCall, bool IsStore) { + bool Fail = false; + + // First argument is a global or generic pointer. + Expr *PtrArg = TheCall->getArg(0); + QualType PtrTy = PtrArg->getType()->getPointeeType(); + unsigned AS = getASTContext().getTargetAddressSpace(PtrTy.getAddressSpace()); + if (AS != llvm::AMDGPUAS::FLAT_ADDRESS && + AS != llvm::AMDGPUAS::GLOBAL_ADDRESS) { + Fail = true; + Diag(TheCall->getBeginLoc(), diag::err_amdgcn_coop_atomic_invalid_as) + << PtrArg->getSourceRange(); } + // Check atomic ordering + Fail |= checkAtomicOrderingCABIArg( + TheCall, IsStore ? 2 : 1, /*MayLoad=*/!IsStore, /*MayStore=*/IsStore); + // Last argument is the syncscope as a string literal. + Fail |= checkStringLiteralArg(TheCall, TheCall->getNumArgs() - 1); + + return Fail; +} + +bool SemaAMDGPU::checkAtomicMonitorLoad(CallExpr *TheCall) { + bool Fail = false; + Fail |= checkAtomicOrderingCABIArg(TheCall, 1, /*MayLoad=*/true, + /*MayStore=*/false); + Fail |= checkStringLiteralArg(TheCall, 2); return Fail; } diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-load-monitor.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-load-monitor.cl index f2552d40fa273..efdbfc25714fb 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-load-monitor.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-load-monitor.cl @@ -5,62 +5,62 @@ typedef int v2i __attribute__((ext_vector_type(2))); typedef int v4i __attribute__((ext_vector_type(4))); -// CHECK-GFX1250-LABEL: @test_amdgcn_global_load_monitor_b32( +// CHECK-GFX1250-LABEL: @test_amdgcn_global_atomic_load_monitor_b32( // CHECK-GFX1250-NEXT: entry: -// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.global.load.monitor.b32.i32(ptr addrspace(1) [[INPTR:%.*]], i32 1) +// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.global.atomic.load.monitor.b32.i32(ptr addrspace(1) [[INPTR:%.*]], i32 0, metadata [[META8:![0-9]+]]) // CHECK-GFX1250-NEXT: ret i32 [[TMP0]] // -int test_amdgcn_global_load_monitor_b32(global int* inptr) +int test_amdgcn_global_atomic_load_monitor_b32(global int* inptr) { - return __builtin_amdgcn_global_load_monitor_b32(inptr, 1); + return __builtin_amdgcn_global_atomic_load_monitor_b32(inptr, __ATOMIC_RELAXED, ""); } -// CHECK-GFX1250-LABEL: @test_amdgcn_global_load_monitor_b64( +// CHECK-GFX1250-LABEL: @test_amdgcn_global_atomic_load_monitor_b64( // CHECK-GFX1250-NEXT: entry: -// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.global.load.monitor.b64.v2i32(ptr addrspace(1) [[INPTR:%.*]], i32 10) +// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.global.atomic.load.monitor.b64.v2i32(ptr addrspace(1) [[INPTR:%.*]], i32 2, metadata [[META9:![0-9]+]]) // CHECK-GFX1250-NEXT: ret <2 x i32> [[TMP0]] // -v2i test_amdgcn_global_load_monitor_b64(global v2i* inptr) +v2i test_amdgcn_global_atomic_load_monitor_b64(global v2i* inptr) { - return __builtin_amdgcn_global_load_monitor_b64(inptr, 10); + return __builtin_amdgcn_global_atomic_load_monitor_b64(inptr, __ATOMIC_ACQUIRE, "agent"); } -// CHECK-GFX1250-LABEL: @test_amdgcn_global_load_monitor_b128( +// CHECK-GFX1250-LABEL: @test_amdgcn_global_atomic_load_monitor_b128( // CHECK-GFX1250-NEXT: entry: -// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.global.load.monitor.b128.v4i32(ptr addrspace(1) [[INPTR:%.*]], i32 22) +// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.global.atomic.load.monitor.b128.v4i32(ptr addrspace(1) [[INPTR:%.*]], i32 2, metadata [[META10:![0-9]+]]) // CHECK-GFX1250-NEXT: ret <4 x i32> [[TMP0]] // -v4i test_amdgcn_global_load_monitor_b128(global v4i* inptr) +v4i test_amdgcn_global_atomic_load_monitor_b128(global v4i* inptr) { - return __builtin_amdgcn_global_load_monitor_b128(inptr, 22); + return __builtin_amdgcn_global_atomic_load_monitor_b128(inptr, __ATOMIC_ACQUIRE, "workgroup"); } -// CHECK-GFX1250-LABEL: @test_amdgcn_flat_load_monitor_b32( +// CHECK-GFX1250-LABEL: @test_amdgcn_flat_atomic_load_monitor_b32( // CHECK-GFX1250-NEXT: entry: -// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.flat.load.monitor.b32.i32(ptr [[INPTR:%.*]], i32 27) +// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.flat.atomic.load.monitor.b32.i32(ptr [[INPTR:%.*]], i32 0, metadata [[META8]]) // CHECK-GFX1250-NEXT: ret i32 [[TMP0]] // -int test_amdgcn_flat_load_monitor_b32(int* inptr) +int test_amdgcn_flat_atomic_load_monitor_b32(int* inptr) { - return __builtin_amdgcn_flat_load_monitor_b32(inptr, 27); + return __builtin_amdgcn_flat_atomic_load_monitor_b32(inptr, __ATOMIC_RELAXED, ""); } -// CHECK-GFX1250-LABEL: @test_amdgcn_flat_load_monitor_b64( +// CHECK-GFX1250-LABEL: @test_amdgcn_flat_atomic_load_monitor_b64( // CHECK-GFX1250-NEXT: entry: -// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.flat.load.monitor.b64.v2i32(ptr [[INPTR:%.*]], i32 1) +// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.flat.atomic.load.monitor.b64.v2i32(ptr [[INPTR:%.*]], i32 5, metadata [[META11:![0-9]+]]) // CHECK-GFX1250-NEXT: ret <2 x i32> [[TMP0]] // -v2i test_amdgcn_flat_load_monitor_b64(v2i* inptr) +v2i test_amdgcn_flat_atomic_load_monitor_b64(v2i* inptr) { - return __builtin_amdgcn_flat_load_monitor_b64(inptr, 1); + return __builtin_amdgcn_flat_atomic_load_monitor_b64(inptr, __ATOMIC_SEQ_CST, "cluster"); } -// CHECK-GFX1250-LABEL: @test_amdgcn_flat_load_monitor_b128( +// CHECK-GFX1250-LABEL: @test_amdgcn_flat_atomic_load_monitor_b128( // CHECK-GFX1250-NEXT: entry: -// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.flat.load.monitor.b128.v4i32(ptr [[INPTR:%.*]], i32 0) +// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.flat.atomic.load.monitor.b128.v4i32(ptr [[INPTR:%.*]], i32 0, metadata [[META8]]) // CHECK-GFX1250-NEXT: ret <4 x i32> [[TMP0]] // -v4i test_amdgcn_flat_load_monitor_b128(v4i* inptr) +v4i test_amdgcn_flat_atomic_load_monitor_b128(v4i* inptr) { - return __builtin_amdgcn_flat_load_monitor_b128(inptr, 0); + return __builtin_amdgcn_flat_atomic_load_monitor_b128(inptr, __ATOMIC_RELAXED, ""); } diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl index da6a03bc93eeb..dfbe5e3b30396 100644 --- a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl +++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl @@ -94,15 +94,37 @@ void test_cvt_scale_pk(global half8 *outh8, global bfloat8 *outy8, uint2 src2, *outf16 = __builtin_amdgcn_cvt_scale_pk16_f32_bf6(src3, scale, 16); // expected-error {{argument value 16 is outside the valid range [0, 15]}} } -void test_amdgcn_load_monitor(global int* b32gaddr, global v2i* b64gaddr, global v4i* b128gaddr, int *b32faddr, v2i* b64faddr, v4i *b128faddr, - global int* b32out, global v2i* b64out, global v4i* b128out, int cpol) +void test_amdgcn_atomic_load_monitor_ao_constant(global int* b32gaddr, global v2i* b64gaddr, global v4i* b128gaddr, int *b32faddr, v2i* b64faddr, v4i *b128faddr, + global int* b32out, global v2i* b64out, global v4i* b128out, int ao) { - *b32out = __builtin_amdgcn_global_load_monitor_b32(b32gaddr, cpol); // expected-error {{'__builtin_amdgcn_global_load_monitor_b32' must be a constant integer}} - *b64out = __builtin_amdgcn_global_load_monitor_b64(b64gaddr, cpol); // expected-error {{'__builtin_amdgcn_global_load_monitor_b64' must be a constant integer}} - *b128out = __built... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/177343 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
