diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp index f09b3b92c4ea0..4b695f970e248 100644 --- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp @@ -274,7 +274,7 @@ llvm::Value *CodeGenFunction::EmitScalarOrConstFoldImmArg(unsigned ICEArguments, void CodeGenFunction::AddAMDGPUFenceAddressSpaceMMRA(llvm::Instruction *Inst, const CallExpr *E) { - constexpr const char *Tag = "amdgpu-as"; + constexpr const char *Tag = "amdgpu-synchronize-as"; LLVMContext &Ctx = Inst->getContext(); SmallVector MMRAs; diff --git a/clang/test/CodeGenCXX/builtin-amdgcn-fence.cpp b/clang/test/CodeGenCXX/builtin-amdgcn-fence.cpp index 3af5a21ba0cd5..1e977dd6420f4 100644 --- a/clang/test/CodeGenCXX/builtin-amdgcn-fence.cpp +++ b/clang/test/CodeGenCXX/builtin-amdgcn-fence.cpp @@ -105,7 +105,7 @@ void test_mixed() { __builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "local", "local", "global", "local", "local"); } //. -// CHECK: [[META3]] = !{!"amdgpu-as", !"local"} -// CHECK: [[META4]] = !{!"amdgpu-as", !"global"} +// CHECK: [[META3]] = !{!"amdgpu-synchronize-as", !"local"} +// CHECK: [[META4]] = !{!"amdgpu-synchronize-as", !"global"} // CHECK: [[META5]] = !{[[META4]], [[META3]]} //. diff --git a/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst index c5b9bd9de66e1..89e17d58116ea 100644 --- a/llvm/docs/AMDGPUUsage.rst +++ b/llvm/docs/AMDGPUUsage.rst @@ -6344,10 +6344,13 @@ also have to wait on all global memory operations, which is unnecessary. :doc:`Memory Model Relaxation Annotations ` can be used as an optimization hint for fences to solve this problem. -The AMDGPU backend recognizes the following tags on fences: +The AMDGPU backend recognizes the following tags on fences to control which address +space a fence can synchronize: -- ``amdgpu-as:local`` - fence only the local address space -- ``amdgpu-as:global``- fence only the global address space +- ``amdgpu-synchronize-as:local`` - for the local address space +- ``amdgpu-synchronize-as:global``- for the global address space + +Multiple tags can be used at the same time to synchronize with more than one address space. .. note:: diff --git a/llvm/lib/Target/AMDGPU/SIMemoryLegalizer.cpp b/llvm/lib/Target/AMDGPU/SIMemoryLegalizer.cpp index 3212060f303a5..123fbfe0b2a48 100644 --- a/llvm/lib/Target/AMDGPU/SIMemoryLegalizer.cpp +++ b/llvm/lib/Target/AMDGPU/SIMemoryLegalizer.cpp @@ -704,12 +704,12 @@ void diagnoseUnknownMMRAASName(const MachineInstr &MI, StringRef AS) { DiagnosticInfoUnsupported(Fn, Str.str(), MI.getDebugLoc(), DS_Warning)); } -/// Reads \p MI's MMRAs to parse the "amdgpu-as" MMRA. +/// Reads \p MI's MMRAs to parse the "amdgpu-synchronize-as" MMRA. /// If this tag isn't present, or if it has no meaningful values, returns \p /// Default. Otherwise returns all the address spaces concerned by the MMRA. static SIAtomicAddrSpace getFenceAddrSpaceMMRA(const MachineInstr &MI, SIAtomicAddrSpace Default) { - static constexpr StringLiteral FenceASPrefix = "amdgpu-as"; + static constexpr StringLiteral FenceASPrefix = "amdgpu-synchronize-as"; auto MMRA = MMRAMetadata(MI.getMMRAMetadata()); if (!MMRA) diff --git a/llvm/test/CodeGen/AMDGPU/memory-legalizer-fence-mmra-global.ll b/llvm/test/CodeGen/AMDGPU/memory-legalizer-fence-mmra-global.ll index 1379eb61e0853..80445f793934b 100644 --- a/llvm/test/CodeGen/AMDGPU/memory-legalizer-fence-mmra-global.ll +++ b/llvm/test/CodeGen/AMDGPU/memory-legalizer-fence-mmra-global.ll @@ -79,7 +79,7 @@ define amdgpu_kernel void @workgroup_acquire_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("workgroup") acquire, !mmra !{!"amdgpu-as", !"global"} + fence syncscope("workgroup") acquire, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -146,7 +146,7 @@ define amdgpu_kernel void @workgroup_release_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("workgroup") release, !mmra !{!"amdgpu-as", !"global"} + fence syncscope("workgroup") release, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -218,7 +218,7 @@ define amdgpu_kernel void @workgroup_acq_rel_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("workgroup") acq_rel, !mmra !{!"amdgpu-as", !"global"} + fence syncscope("workgroup") acq_rel, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -290,7 +290,7 @@ define amdgpu_kernel void @workgroup_seq_cst_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("workgroup") seq_cst, !mmra !{!"amdgpu-as", !"global"} + fence syncscope("workgroup") seq_cst, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -360,7 +360,7 @@ define amdgpu_kernel void @workgroup_one_as_acquire_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("workgroup-one-as") acquire, !mmra !{!"amdgpu-as", !"global"} + fence syncscope("workgroup-one-as") acquire, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -427,7 +427,7 @@ define amdgpu_kernel void @workgroup_one_as_release_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("workgroup-one-as") release, !mmra !{!"amdgpu-as", !"global"} + fence syncscope("workgroup-one-as") release, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -499,7 +499,7 @@ define amdgpu_kernel void @workgroup_one_as_acq_rel_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("workgroup-one-as") acq_rel, !mmra !{!"amdgpu-as", !"global"} + fence syncscope("workgroup-one-as") acq_rel, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -571,7 +571,7 @@ define amdgpu_kernel void @workgroup_one_as_seq_cst_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("workgroup-one-as") seq_cst, !mmra !{!"amdgpu-as", !"global"} + fence syncscope("workgroup-one-as") seq_cst, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -663,7 +663,7 @@ define amdgpu_kernel void @agent_acquire_fence() { ; GFX12-CU-NEXT: global_inv scope:SCOPE_DEV ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("agent") acquire, !mmra !{!"amdgpu-as", !"global"} + fence syncscope("agent") acquire, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -745,7 +745,7 @@ define amdgpu_kernel void @agent_release_fence() { ; GFX12-CU-NEXT: s_wait_storecnt 0x0 ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("agent") release, !mmra !{!"amdgpu-as", !"global"} + fence syncscope("agent") release, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -843,7 +843,7 @@ define amdgpu_kernel void @agent_acq_rel_fence() { ; GFX12-CU-NEXT: global_inv scope:SCOPE_DEV ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("agent") acq_rel, !mmra !{!"amdgpu-as", !"global"} + fence syncscope("agent") acq_rel, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -941,7 +941,7 @@ define amdgpu_kernel void @agent_seq_cst_fence() { ; GFX12-CU-NEXT: global_inv scope:SCOPE_DEV ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("agent") seq_cst, !mmra !{!"amdgpu-as", !"global"} + fence syncscope("agent") seq_cst, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -1033,7 +1033,7 @@ define amdgpu_kernel void @agent_one_as_acquire_fence() { ; GFX12-CU-NEXT: global_inv scope:SCOPE_DEV ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("agent-one-as") acquire, !mmra !{!"amdgpu-as", !"global"} + fence syncscope("agent-one-as") acquire, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -1115,7 +1115,7 @@ define amdgpu_kernel void @agent_one_as_release_fence() { ; GFX12-CU-NEXT: s_wait_storecnt 0x0 ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("agent-one-as") release, !mmra !{!"amdgpu-as", !"global"} + fence syncscope("agent-one-as") release, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -1213,7 +1213,7 @@ define amdgpu_kernel void @agent_one_as_acq_rel_fence() { ; GFX12-CU-NEXT: global_inv scope:SCOPE_DEV ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("agent-one-as") acq_rel, !mmra !{!"amdgpu-as", !"global"} + fence syncscope("agent-one-as") acq_rel, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -1311,7 +1311,7 @@ define amdgpu_kernel void @agent_one_as_seq_cst_fence() { ; GFX12-CU-NEXT: global_inv scope:SCOPE_DEV ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("agent-one-as") seq_cst, !mmra !{!"amdgpu-as", !"global"} + fence syncscope("agent-one-as") seq_cst, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -1405,7 +1405,7 @@ define amdgpu_kernel void @system_acquire_fence() { ; GFX12-CU-NEXT: global_inv scope:SCOPE_SYS ; GFX12-CU-NEXT: s_endpgm entry: - fence acquire, !mmra !{!"amdgpu-as", !"global"} + fence acquire, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -1491,7 +1491,7 @@ define amdgpu_kernel void @system_release_fence() { ; GFX12-CU-NEXT: s_wait_storecnt 0x0 ; GFX12-CU-NEXT: s_endpgm entry: - fence release, !mmra !{!"amdgpu-as", !"global"} + fence release, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -1595,7 +1595,7 @@ define amdgpu_kernel void @system_acq_rel_fence() { ; GFX12-CU-NEXT: global_inv scope:SCOPE_SYS ; GFX12-CU-NEXT: s_endpgm entry: - fence acq_rel, !mmra !{!"amdgpu-as", !"global"} + fence acq_rel, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -1699,7 +1699,7 @@ define amdgpu_kernel void @system_seq_cst_fence() { ; GFX12-CU-NEXT: global_inv scope:SCOPE_SYS ; GFX12-CU-NEXT: s_endpgm entry: - fence seq_cst, !mmra !{!"amdgpu-as", !"global"} + fence seq_cst, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -1793,7 +1793,7 @@ define amdgpu_kernel void @system_one_as_acquire_fence() { ; GFX12-CU-NEXT: global_inv scope:SCOPE_SYS ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("one-as") acquire, !mmra !{!"amdgpu-as", !"global"} + fence syncscope("one-as") acquire, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -1879,7 +1879,7 @@ define amdgpu_kernel void @system_one_as_release_fence() { ; GFX12-CU-NEXT: s_wait_storecnt 0x0 ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("one-as") release, !mmra !{!"amdgpu-as", !"global"} + fence syncscope("one-as") release, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -1983,7 +1983,7 @@ define amdgpu_kernel void @system_one_as_acq_rel_fence() { ; GFX12-CU-NEXT: global_inv scope:SCOPE_SYS ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("one-as") acq_rel, !mmra !{!"amdgpu-as", !"global"} + fence syncscope("one-as") acq_rel, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } @@ -2087,6 +2087,6 @@ define amdgpu_kernel void @system_one_as_seq_cst_fence() { ; GFX12-CU-NEXT: global_inv scope:SCOPE_SYS ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("one-as") seq_cst, !mmra !{!"amdgpu-as", !"global"} + fence syncscope("one-as") seq_cst, !mmra !{!"amdgpu-synchronize-as", !"global"} ret void } diff --git a/llvm/test/CodeGen/AMDGPU/memory-legalizer-fence-mmra-local.ll b/llvm/test/CodeGen/AMDGPU/memory-legalizer-fence-mmra-local.ll index 971015b391ca8..7a419a5031ba9 100644 --- a/llvm/test/CodeGen/AMDGPU/memory-legalizer-fence-mmra-local.ll +++ b/llvm/test/CodeGen/AMDGPU/memory-legalizer-fence-mmra-local.ll @@ -77,7 +77,7 @@ define amdgpu_kernel void @workgroup_acquire_fence() { ; GFX12-CU-NEXT: s_wait_dscnt 0x0 ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("workgroup") acquire, !mmra !{!"amdgpu-as", !"local"} + fence syncscope("workgroup") acquire, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -143,7 +143,7 @@ define amdgpu_kernel void @workgroup_release_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("workgroup") release, !mmra !{!"amdgpu-as", !"local"} + fence syncscope("workgroup") release, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -209,7 +209,7 @@ define amdgpu_kernel void @workgroup_acq_rel_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("workgroup") acq_rel, !mmra !{!"amdgpu-as", !"local"} + fence syncscope("workgroup") acq_rel, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -275,7 +275,7 @@ define amdgpu_kernel void @workgroup_seq_cst_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("workgroup") seq_cst, !mmra !{!"amdgpu-as", !"local"} + fence syncscope("workgroup") seq_cst, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -332,7 +332,7 @@ define amdgpu_kernel void @workgroup_one_as_acquire_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("workgroup-one-as") acquire, !mmra !{!"amdgpu-as", !"local"} + fence syncscope("workgroup-one-as") acquire, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -389,7 +389,7 @@ define amdgpu_kernel void @workgroup_one_as_release_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("workgroup-one-as") release, !mmra !{!"amdgpu-as", !"local"} + fence syncscope("workgroup-one-as") release, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -446,7 +446,7 @@ define amdgpu_kernel void @workgroup_one_as_acq_rel_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("workgroup-one-as") acq_rel, !mmra !{!"amdgpu-as", !"local"} + fence syncscope("workgroup-one-as") acq_rel, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -503,7 +503,7 @@ define amdgpu_kernel void @workgroup_one_as_seq_cst_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("workgroup-one-as") seq_cst, !mmra !{!"amdgpu-as", !"local"} + fence syncscope("workgroup-one-as") seq_cst, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -571,7 +571,7 @@ define amdgpu_kernel void @agent_acquire_fence() { ; GFX12-CU-NEXT: s_wait_dscnt 0x0 ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("agent") acquire, !mmra !{!"amdgpu-as", !"local"} + fence syncscope("agent") acquire, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -637,7 +637,7 @@ define amdgpu_kernel void @agent_release_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("agent") release, !mmra !{!"amdgpu-as", !"local"} + fence syncscope("agent") release, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -703,7 +703,7 @@ define amdgpu_kernel void @agent_acq_rel_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("agent") acq_rel, !mmra !{!"amdgpu-as", !"local"} + fence syncscope("agent") acq_rel, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -769,7 +769,7 @@ define amdgpu_kernel void @agent_seq_cst_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("agent") seq_cst, !mmra !{!"amdgpu-as", !"local"} + fence syncscope("agent") seq_cst, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -826,7 +826,7 @@ define amdgpu_kernel void @agent_one_as_acquire_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("agent-one-as") acquire, !mmra !{!"amdgpu-as", !"local"} + fence syncscope("agent-one-as") acquire, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -883,7 +883,7 @@ define amdgpu_kernel void @agent_one_as_release_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("agent-one-as") release, !mmra !{!"amdgpu-as", !"local"} + fence syncscope("agent-one-as") release, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -940,7 +940,7 @@ define amdgpu_kernel void @agent_one_as_acq_rel_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("agent-one-as") acq_rel, !mmra !{!"amdgpu-as", !"local"} + fence syncscope("agent-one-as") acq_rel, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -997,7 +997,7 @@ define amdgpu_kernel void @agent_one_as_seq_cst_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("agent-one-as") seq_cst, !mmra !{!"amdgpu-as", !"local"} + fence syncscope("agent-one-as") seq_cst, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -1065,7 +1065,7 @@ define amdgpu_kernel void @system_acquire_fence() { ; GFX12-CU-NEXT: s_wait_dscnt 0x0 ; GFX12-CU-NEXT: s_endpgm entry: - fence acquire, !mmra !{!"amdgpu-as", !"local"} + fence acquire, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -1131,7 +1131,7 @@ define amdgpu_kernel void @system_release_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence release, !mmra !{!"amdgpu-as", !"local"} + fence release, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -1197,7 +1197,7 @@ define amdgpu_kernel void @system_acq_rel_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence acq_rel, !mmra !{!"amdgpu-as", !"local"} + fence acq_rel, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -1263,7 +1263,7 @@ define amdgpu_kernel void @system_seq_cst_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence seq_cst, !mmra !{!"amdgpu-as", !"local"} + fence seq_cst, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -1320,7 +1320,7 @@ define amdgpu_kernel void @system_one_as_acquire_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("one-as") acquire, !mmra !{!"amdgpu-as", !"local"} + fence syncscope("one-as") acquire, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -1377,7 +1377,7 @@ define amdgpu_kernel void @system_one_as_release_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("one-as") release, !mmra !{!"amdgpu-as", !"local"} + fence syncscope("one-as") release, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -1434,7 +1434,7 @@ define amdgpu_kernel void @system_one_as_acq_rel_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("one-as") acq_rel, !mmra !{!"amdgpu-as", !"local"} + fence syncscope("one-as") acq_rel, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void } @@ -1491,6 +1491,6 @@ define amdgpu_kernel void @system_one_as_seq_cst_fence() { ; GFX12-CU: ; %bb.0: ; %entry ; GFX12-CU-NEXT: s_endpgm entry: - fence syncscope("one-as") seq_cst, !mmra !{!"amdgpu-as", !"local"} + fence syncscope("one-as") seq_cst, !mmra !{!"amdgpu-synchronize-as", !"local"} ret void }