Skip to content

[NFC][AMDGPU] Rename "amdgpu-as" to "amdgpu-synchronize-as" #148627

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 1 commit into
base: main
Choose a base branch
from
+59 −56
Open

[NFC][AMDGPU] Rename "amdgpu-as" to "amdgpu-synchronize-as" #148627

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<MMRAMetadata::TagT, 3> MMRAs;
Expand Down
4 changes: 2 additions & 2 deletions clang/test/CodeGenCXX/builtin-amdgcn-fence.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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]]}
//.
9 changes: 6 additions & 3 deletions llvm/docs/AMDGPUUsage.rst
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6344,10 +6344,13 @@ also have to wait on all global memory operations, which is unnecessary.

:doc:`Memory Model Relaxation Annotations <MemoryModelRelaxationAnnotations>` 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::

Expand Down
4 changes: 2 additions & 2 deletions llvm/lib/Target/AMDGPU/SIMemoryLegalizer.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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";
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

If we want to phase out amdgpu-as, I think it'd be better to still support/auto upgrade the old one.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

it's droppable metadata, we don't need to worry about auto upgrade


auto MMRA = MMRAMetadata(MI.getMMRAMetadata());
if (!MMRA)
Expand Down
48 changes: 24 additions & 24 deletions llvm/test/CodeGen/AMDGPU/memory-legalizer-fence-mmra-global.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}

Expand Down Expand Up @@ -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
}
Loading
Loading