Skip to content

Commit cd1b84c

Browse files
Pierre-vhPierre-vh
authored
[NFC][AMDGPU] Rename "amdgpu-as" to "amdgpu-synchronize-as" (llvm#148627)
"amdgpu-as" is way too vague and doesn't give enough context. We may want to support it on normal atomics too, to control the synchronized (ordered) AS. If we do that, the name has to be less vague.
1 parent 1a32bcb commit cd1b84c

File tree

6 files changed

+59
-56
lines changed

6 files changed

+59
-56
lines changed

‎clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -274,7 +274,7 @@ llvm::Value *CodeGenFunction::EmitScalarOrConstFoldImmArg(unsigned ICEArguments,
274274

275275
void CodeGenFunction::AddAMDGPUFenceAddressSpaceMMRA(llvm::Instruction *Inst,
276276
const CallExpr *E) {
277-
constexpr const char *Tag = "amdgpu-as";
277+
constexpr const char *Tag = "amdgpu-synchronize-as";
278278

279279
LLVMContext &Ctx = Inst->getContext();
280280
SmallVector<MMRAMetadata::TagT, 3> MMRAs;

‎clang/test/CodeGenCXX/builtin-amdgcn-fence.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -105,7 +105,7 @@ void test_mixed() {
105105
__builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "local", "local", "global", "local", "local");
106106
}
107107
//.
108-
// CHECK: [[META3]] = !{!"amdgpu-as", !"local"}
109-
// CHECK: [[META4]] = !{!"amdgpu-as", !"global"}
108+
// CHECK: [[META3]] = !{!"amdgpu-synchronize-as", !"local"}
109+
// CHECK: [[META4]] = !{!"amdgpu-synchronize-as", !"global"}
110110
// CHECK: [[META5]] = !{[[META4]], [[META3]]}
111111
//.

‎llvm/docs/AMDGPUUsage.rst‎

Lines changed: 6 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -6358,10 +6358,13 @@ also have to wait on all global memory operations, which is unnecessary.
63586358

63596359
:doc:`Memory Model Relaxation Annotations <MemoryModelRelaxationAnnotations>` can
63606360
be used as an optimization hint for fences to solve this problem.
6361-
The AMDGPU backend recognizes the following tags on fences:
6361+
The AMDGPU backend recognizes the following tags on fences to control which address
6362+
space a fence can synchronize:
63626363

6363-
- ``amdgpu-as:local`` - fence only the local address space
6364-
- ``amdgpu-as:global``- fence only the global address space
6364+
- ``amdgpu-synchronize-as:local`` - for the local address space
6365+
- ``amdgpu-synchronize-as:global``- for the global address space
6366+
6367+
Multiple tags can be used at the same time to synchronize with more than one address space.
63656368

63666369
.. note::
63676370

‎llvm/lib/Target/AMDGPU/SIMemoryLegalizer.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -704,12 +704,12 @@ void diagnoseUnknownMMRAASName(const MachineInstr &MI, StringRef AS) {
704704
DiagnosticInfoUnsupported(Fn, Str.str(), MI.getDebugLoc(), DS_Warning));
705705
}
706706

707-
/// Reads \p MI's MMRAs to parse the "amdgpu-as" MMRA.
707+
/// Reads \p MI's MMRAs to parse the "amdgpu-synchronize-as" MMRA.
708708
/// If this tag isn't present, or if it has no meaningful values, returns \p
709709
/// Default. Otherwise returns all the address spaces concerned by the MMRA.
710710
static SIAtomicAddrSpace getFenceAddrSpaceMMRA(const MachineInstr &MI,
711711
SIAtomicAddrSpace Default) {
712-
static constexpr StringLiteral FenceASPrefix = "amdgpu-as";
712+
static constexpr StringLiteral FenceASPrefix = "amdgpu-synchronize-as";
713713

714714
auto MMRA = MMRAMetadata(MI.getMMRAMetadata());
715715
if (!MMRA)

‎llvm/test/CodeGen/AMDGPU/memory-legalizer-fence-mmra-global.ll‎

Lines changed: 24 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -79,7 +79,7 @@ define amdgpu_kernel void @workgroup_acquire_fence() {
7979
; GFX12-CU: ; %bb.0: ; %entry
8080
; GFX12-CU-NEXT: s_endpgm
8181
entry:
82-
fence syncscope("workgroup") acquire, !mmra !{!"amdgpu-as", !"global"}
82+
fence syncscope("workgroup") acquire, !mmra !{!"amdgpu-synchronize-as", !"global"}
8383
ret void
8484
}
8585

@@ -146,7 +146,7 @@ define amdgpu_kernel void @workgroup_release_fence() {
146146
; GFX12-CU: ; %bb.0: ; %entry
147147
; GFX12-CU-NEXT: s_endpgm
148148
entry:
149-
fence syncscope("workgroup") release, !mmra !{!"amdgpu-as", !"global"}
149+
fence syncscope("workgroup") release, !mmra !{!"amdgpu-synchronize-as", !"global"}
150150
ret void
151151
}
152152

@@ -218,7 +218,7 @@ define amdgpu_kernel void @workgroup_acq_rel_fence() {
218218
; GFX12-CU: ; %bb.0: ; %entry
219219
; GFX12-CU-NEXT: s_endpgm
220220
entry:
221-
fence syncscope("workgroup") acq_rel, !mmra !{!"amdgpu-as", !"global"}
221+
fence syncscope("workgroup") acq_rel, !mmra !{!"amdgpu-synchronize-as", !"global"}
222222
ret void
223223
}
224224

@@ -290,7 +290,7 @@ define amdgpu_kernel void @workgroup_seq_cst_fence() {
290290
; GFX12-CU: ; %bb.0: ; %entry
291291
; GFX12-CU-NEXT: s_endpgm
292292
entry:
293-
fence syncscope("workgroup") seq_cst, !mmra !{!"amdgpu-as", !"global"}
293+
fence syncscope("workgroup") seq_cst, !mmra !{!"amdgpu-synchronize-as", !"global"}
294294
ret void
295295
}
296296

@@ -360,7 +360,7 @@ define amdgpu_kernel void @workgroup_one_as_acquire_fence() {
360360
; GFX12-CU: ; %bb.0: ; %entry
361361
; GFX12-CU-NEXT: s_endpgm
362362
entry:
363-
fence syncscope("workgroup-one-as") acquire, !mmra !{!"amdgpu-as", !"global"}
363+
fence syncscope("workgroup-one-as") acquire, !mmra !{!"amdgpu-synchronize-as", !"global"}
364364
ret void
365365
}
366366

@@ -427,7 +427,7 @@ define amdgpu_kernel void @workgroup_one_as_release_fence() {
427427
; GFX12-CU: ; %bb.0: ; %entry
428428
; GFX12-CU-NEXT: s_endpgm
429429
entry:
430-
fence syncscope("workgroup-one-as") release, !mmra !{!"amdgpu-as", !"global"}
430+
fence syncscope("workgroup-one-as") release, !mmra !{!"amdgpu-synchronize-as", !"global"}
431431
ret void
432432
}
433433

@@ -499,7 +499,7 @@ define amdgpu_kernel void @workgroup_one_as_acq_rel_fence() {
499499
; GFX12-CU: ; %bb.0: ; %entry
500500
; GFX12-CU-NEXT: s_endpgm
501501
entry:
502-
fence syncscope("workgroup-one-as") acq_rel, !mmra !{!"amdgpu-as", !"global"}
502+
fence syncscope("workgroup-one-as") acq_rel, !mmra !{!"amdgpu-synchronize-as", !"global"}
503503
ret void
504504
}
505505

@@ -571,7 +571,7 @@ define amdgpu_kernel void @workgroup_one_as_seq_cst_fence() {
571571
; GFX12-CU: ; %bb.0: ; %entry
572572
; GFX12-CU-NEXT: s_endpgm
573573
entry:
574-
fence syncscope("workgroup-one-as") seq_cst, !mmra !{!"amdgpu-as", !"global"}
574+
fence syncscope("workgroup-one-as") seq_cst, !mmra !{!"amdgpu-synchronize-as", !"global"}
575575
ret void
576576
}
577577

@@ -663,7 +663,7 @@ define amdgpu_kernel void @agent_acquire_fence() {
663663
; GFX12-CU-NEXT: global_inv scope:SCOPE_DEV
664664
; GFX12-CU-NEXT: s_endpgm
665665
entry:
666-
fence syncscope("agent") acquire, !mmra !{!"amdgpu-as", !"global"}
666+
fence syncscope("agent") acquire, !mmra !{!"amdgpu-synchronize-as", !"global"}
667667
ret void
668668
}
669669

@@ -745,7 +745,7 @@ define amdgpu_kernel void @agent_release_fence() {
745745
; GFX12-CU-NEXT: s_wait_storecnt 0x0
746746
; GFX12-CU-NEXT: s_endpgm
747747
entry:
748-
fence syncscope("agent") release, !mmra !{!"amdgpu-as", !"global"}
748+
fence syncscope("agent") release, !mmra !{!"amdgpu-synchronize-as", !"global"}
749749
ret void
750750
}
751751

@@ -843,7 +843,7 @@ define amdgpu_kernel void @agent_acq_rel_fence() {
843843
; GFX12-CU-NEXT: global_inv scope:SCOPE_DEV
844844
; GFX12-CU-NEXT: s_endpgm
845845
entry:
846-
fence syncscope("agent") acq_rel, !mmra !{!"amdgpu-as", !"global"}
846+
fence syncscope("agent") acq_rel, !mmra !{!"amdgpu-synchronize-as", !"global"}
847847
ret void
848848
}
849849

@@ -941,7 +941,7 @@ define amdgpu_kernel void @agent_seq_cst_fence() {
941941
; GFX12-CU-NEXT: global_inv scope:SCOPE_DEV
942942
; GFX12-CU-NEXT: s_endpgm
943943
entry:
944-
fence syncscope("agent") seq_cst, !mmra !{!"amdgpu-as", !"global"}
944+
fence syncscope("agent") seq_cst, !mmra !{!"amdgpu-synchronize-as", !"global"}
945945
ret void
946946
}
947947

@@ -1033,7 +1033,7 @@ define amdgpu_kernel void @agent_one_as_acquire_fence() {
10331033
; GFX12-CU-NEXT: global_inv scope:SCOPE_DEV
10341034
; GFX12-CU-NEXT: s_endpgm
10351035
entry:
1036-
fence syncscope("agent-one-as") acquire, !mmra !{!"amdgpu-as", !"global"}
1036+
fence syncscope("agent-one-as") acquire, !mmra !{!"amdgpu-synchronize-as", !"global"}
10371037
ret void
10381038
}
10391039

@@ -1115,7 +1115,7 @@ define amdgpu_kernel void @agent_one_as_release_fence() {
11151115
; GFX12-CU-NEXT: s_wait_storecnt 0x0
11161116
; GFX12-CU-NEXT: s_endpgm
11171117
entry:
1118-
fence syncscope("agent-one-as") release, !mmra !{!"amdgpu-as", !"global"}
1118+
fence syncscope("agent-one-as") release, !mmra !{!"amdgpu-synchronize-as", !"global"}
11191119
ret void
11201120
}
11211121

@@ -1213,7 +1213,7 @@ define amdgpu_kernel void @agent_one_as_acq_rel_fence() {
12131213
; GFX12-CU-NEXT: global_inv scope:SCOPE_DEV
12141214
; GFX12-CU-NEXT: s_endpgm
12151215
entry:
1216-
fence syncscope("agent-one-as") acq_rel, !mmra !{!"amdgpu-as", !"global"}
1216+
fence syncscope("agent-one-as") acq_rel, !mmra !{!"amdgpu-synchronize-as", !"global"}
12171217
ret void
12181218
}
12191219

@@ -1311,7 +1311,7 @@ define amdgpu_kernel void @agent_one_as_seq_cst_fence() {
13111311
; GFX12-CU-NEXT: global_inv scope:SCOPE_DEV
13121312
; GFX12-CU-NEXT: s_endpgm
13131313
entry:
1314-
fence syncscope("agent-one-as") seq_cst, !mmra !{!"amdgpu-as", !"global"}
1314+
fence syncscope("agent-one-as") seq_cst, !mmra !{!"amdgpu-synchronize-as", !"global"}
13151315
ret void
13161316
}
13171317

@@ -1405,7 +1405,7 @@ define amdgpu_kernel void @system_acquire_fence() {
14051405
; GFX12-CU-NEXT: global_inv scope:SCOPE_SYS
14061406
; GFX12-CU-NEXT: s_endpgm
14071407
entry:
1408-
fence acquire, !mmra !{!"amdgpu-as", !"global"}
1408+
fence acquire, !mmra !{!"amdgpu-synchronize-as", !"global"}
14091409
ret void
14101410
}
14111411

@@ -1491,7 +1491,7 @@ define amdgpu_kernel void @system_release_fence() {
14911491
; GFX12-CU-NEXT: s_wait_storecnt 0x0
14921492
; GFX12-CU-NEXT: s_endpgm
14931493
entry:
1494-
fence release, !mmra !{!"amdgpu-as", !"global"}
1494+
fence release, !mmra !{!"amdgpu-synchronize-as", !"global"}
14951495
ret void
14961496
}
14971497

@@ -1595,7 +1595,7 @@ define amdgpu_kernel void @system_acq_rel_fence() {
15951595
; GFX12-CU-NEXT: global_inv scope:SCOPE_SYS
15961596
; GFX12-CU-NEXT: s_endpgm
15971597
entry:
1598-
fence acq_rel, !mmra !{!"amdgpu-as", !"global"}
1598+
fence acq_rel, !mmra !{!"amdgpu-synchronize-as", !"global"}
15991599
ret void
16001600
}
16011601

@@ -1699,7 +1699,7 @@ define amdgpu_kernel void @system_seq_cst_fence() {
16991699
; GFX12-CU-NEXT: global_inv scope:SCOPE_SYS
17001700
; GFX12-CU-NEXT: s_endpgm
17011701
entry:
1702-
fence seq_cst, !mmra !{!"amdgpu-as", !"global"}
1702+
fence seq_cst, !mmra !{!"amdgpu-synchronize-as", !"global"}
17031703
ret void
17041704
}
17051705

@@ -1793,7 +1793,7 @@ define amdgpu_kernel void @system_one_as_acquire_fence() {
17931793
; GFX12-CU-NEXT: global_inv scope:SCOPE_SYS
17941794
; GFX12-CU-NEXT: s_endpgm
17951795
entry:
1796-
fence syncscope("one-as") acquire, !mmra !{!"amdgpu-as", !"global"}
1796+
fence syncscope("one-as") acquire, !mmra !{!"amdgpu-synchronize-as", !"global"}
17971797
ret void
17981798
}
17991799

@@ -1879,7 +1879,7 @@ define amdgpu_kernel void @system_one_as_release_fence() {
18791879
; GFX12-CU-NEXT: s_wait_storecnt 0x0
18801880
; GFX12-CU-NEXT: s_endpgm
18811881
entry:
1882-
fence syncscope("one-as") release, !mmra !{!"amdgpu-as", !"global"}
1882+
fence syncscope("one-as") release, !mmra !{!"amdgpu-synchronize-as", !"global"}
18831883
ret void
18841884
}
18851885

@@ -1983,7 +1983,7 @@ define amdgpu_kernel void @system_one_as_acq_rel_fence() {
19831983
; GFX12-CU-NEXT: global_inv scope:SCOPE_SYS
19841984
; GFX12-CU-NEXT: s_endpgm
19851985
entry:
1986-
fence syncscope("one-as") acq_rel, !mmra !{!"amdgpu-as", !"global"}
1986+
fence syncscope("one-as") acq_rel, !mmra !{!"amdgpu-synchronize-as", !"global"}
19871987
ret void
19881988
}
19891989

@@ -2087,6 +2087,6 @@ define amdgpu_kernel void @system_one_as_seq_cst_fence() {
20872087
; GFX12-CU-NEXT: global_inv scope:SCOPE_SYS
20882088
; GFX12-CU-NEXT: s_endpgm
20892089
entry:
2090-
fence syncscope("one-as") seq_cst, !mmra !{!"amdgpu-as", !"global"}
2090+
fence syncscope("one-as") seq_cst, !mmra !{!"amdgpu-synchronize-as", !"global"}
20912091
ret void
20922092
}

0 commit comments

Comments
 (0)