Skip to content

Commit d8a5a97

Browse files
Pierre-vhPierre-vh
committed
Refactor
1 parent a608563 commit d8a5a97

File tree

11 files changed

+240
-238
lines changed

11 files changed

+240
-238
lines changed

‎clang/docs/LanguageExtensions.rst

Lines changed: 14 additions & 26 deletions
Original file line numberDiff line numberDiff line change
@@ -4419,47 +4419,35 @@ AMDGPU Language Extensions
44194419
__builtin_amdgcn_fence
44204420
^^^^^^^^^^^^^^^^^^^^^^
44214421
4422-
``__builtin_amdgcn_fence`` emits a fence for all address spaces
4423-
and takes the following arguments:
4422+
``__builtin_amdgcn_fence`` emits a fence.
44244423
44254424
* ``unsigned`` atomic ordering, e.g. ``__ATOMIC_ACQUIRE``
44264425
* ``const char *`` synchronization scope, e.g. ``workgroup``
4427-
4428-
.. code-block:: c++
4429-
4430-
__builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup");
4431-
__builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent");
4432-
4433-
__builtin_amdgcn_masked_fence
4434-
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
4435-
4436-
``__builtin_amdgcn_masked_fence`` emits a fence for one or more address
4437-
spaces and takes the following arguments:
4438-
4439-
* ``unsigned`` atomic ordering, e.g. ``__ATOMIC_ACQUIRE``
4440-
* ``const char *`` synchronization scope, e.g. ``workgroup``
4441-
* Zero or more ``const char *`` address spaces.
4426+
* Zero or more ``const char *`` address spaces names.
44424427
44434428
The address spaces arguments must be string literals with known values, such as:
44444429
44454430
* ``"local"``
44464431
* ``"global"``
44474432
* ``"image"``
44484433
4449-
If there are no address spaces specified, this fence behaves like
4450-
``__builtin_amdgcn_fence``.
4434+
If one or more address space name are provided, the code generator will attempt
4435+
to emit potentially faster instructions that only fence those address spaces.
4436+
Emitting such instructions may not always be possible and the compiler is free
4437+
to fence more aggressively.
44514438
4452-
Examples:
4439+
If no address spaces names are provided, all address spaces are fenced.
44534440
44544441
.. code-block:: c++
44554442
4456-
__builtin_amdgcn_masked_fence(__ATOMIC_SEQ_CST, "workgroup", "local")
4457-
__builtin_amdgcn_masked_fence(__ATOMIC_SEQ_CST, "workgroup", "local", "global")
4443+
// Fence all address spaces.
4444+
__builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup");
4445+
__builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent");
4446+
4447+
// Fence only requested address spaces.
4448+
__builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup", "local")
4449+
__builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup", "local", "global")
44584450
4459-
Note that this fence may affect more than just the address spaces
4460-
specified; in some cases, the address space mask may
4461-
be lost during optimization and a normal fence for all address
4462-
spaces (``__builtin_amdgcn_fence``) will be emitted instead.
44634451
44644452
ARM/AArch64 Language Extensions
44654453
-------------------------------

‎clang/include/clang/Basic/BuiltinsAMDGPU.def

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -68,8 +68,7 @@ BUILTIN(__builtin_amdgcn_sched_group_barrier, "vIiIiIi", "n")
6868
BUILTIN(__builtin_amdgcn_iglp_opt, "vIi", "n")
6969
BUILTIN(__builtin_amdgcn_s_dcache_inv, "v", "n")
7070
BUILTIN(__builtin_amdgcn_buffer_wbinvl1, "v", "n")
71-
BUILTIN(__builtin_amdgcn_fence, "vUicC*", "n")
72-
BUILTIN(__builtin_amdgcn_masked_fence, "vUicC*.", "n")
71+
BUILTIN(__builtin_amdgcn_fence, "vUicC*.", "n")
7372
BUILTIN(__builtin_amdgcn_groupstaticsize, "Ui", "n")
7473
BUILTIN(__builtin_amdgcn_wavefrontsize, "Ui", "nc")
7574

‎clang/lib/CodeGen/CGBuiltin.cpp

Lines changed: 7 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -18414,14 +18414,13 @@ Value *CodeGenFunction::EmitHLSLBuiltinExpr(unsigned BuiltinID,
1841418414
return nullptr;
1841518415
}
1841618416

18417-
void CodeGenFunction::AddAMDGCNAddressSpaceMMRA(llvm::Instruction *Inst,
18418-
const CallExpr *E,
18419-
unsigned FirstASNameIdx) {
18420-
constexpr const char *Tag = "opencl-fence-mem";
18417+
void CodeGenFunction::AddAMDGCNFenceAddressSpaceMMRA(llvm::Instruction *Inst,
18418+
const CallExpr *E) {
18419+
constexpr const char *Tag = "amdgpu-as";
1842118420

1842218421
LLVMContext &Ctx = Inst->getContext();
1842318422
SmallVector<MMRAMetadata::TagT, 3> MMRAs;
18424-
for (unsigned K = FirstASNameIdx; K < E->getNumArgs(); ++K) {
18423+
for (unsigned K = 2; K < E->getNumArgs(); ++K) {
1842518424
llvm::Value *V = EmitScalarExpr(E->getArg(K));
1842618425
StringRef AS;
1842718426
if (llvm::getConstantStringInfo(V, AS) &&
@@ -19105,14 +19104,12 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
1910519104
Function *F = CGM.getIntrinsic(Intrinsic::fshr, Src0->getType());
1910619105
return Builder.CreateCall(F, { Src0, Src1, Src2 });
1910719106
}
19108-
case AMDGPU::BI__builtin_amdgcn_fence:
19109-
case AMDGPU::BI__builtin_amdgcn_masked_fence: {
19107+
case AMDGPU::BI__builtin_amdgcn_fence: {
1911019108
ProcessOrderScopeAMDGCN(EmitScalarExpr(E->getArg(0)),
1911119109
EmitScalarExpr(E->getArg(1)), AO, SSID);
1911219110
FenceInst *Fence = Builder.CreateFence(AO, SSID);
19113-
if (BuiltinID == AMDGPU::BI__builtin_amdgcn_masked_fence &&
19114-
E->getNumArgs() > 2)
19115-
AddAMDGCNAddressSpaceMMRA(Fence, E, 2);
19111+
if (E->getNumArgs() > 2)
19112+
AddAMDGCNFenceAddressSpaceMMRA(Fence, E);
1911619113
return Fence;
1911719114
}
1911819115
case AMDGPU::BI__builtin_amdgcn_atomic_inc32:

‎clang/lib/CodeGen/CodeGenFunction.h

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -4632,8 +4632,8 @@ class CodeGenFunction : public CodeGenTypeCache {
46324632
llvm::Value *EmitRISCVBuiltinExpr(unsigned BuiltinID, const CallExpr *E,
46334633
ReturnValueSlot ReturnValue);
46344634

4635-
void AddAMDGCNAddressSpaceMMRA(llvm::Instruction *Inst, const CallExpr *E,
4636-
unsigned FirstASNameIdx);
4635+
void AddAMDGCNFenceAddressSpaceMMRA(llvm::Instruction *Inst,
4636+
const CallExpr *E);
46374637
void ProcessOrderScopeAMDGCN(llvm::Value *Order, llvm::Value *Scope,
46384638
llvm::AtomicOrdering &AO,
46394639
llvm::SyncScope::ID &SSID);

‎clang/lib/Sema/SemaChecking.cpp

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -5704,7 +5704,6 @@ bool Sema::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
57045704
ScopeIndex = 3;
57055705
break;
57065706
case AMDGPU::BI__builtin_amdgcn_fence:
5707-
case AMDGPU::BI__builtin_amdgcn_masked_fence:
57085707
OrderIndex = 0;
57095708
ScopeIndex = 1;
57105709
break;
@@ -5730,8 +5729,7 @@ bool Sema::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
57305729
switch (static_cast<llvm::AtomicOrderingCABI>(Ord)) {
57315730
case llvm::AtomicOrderingCABI::relaxed:
57325731
case llvm::AtomicOrderingCABI::consume:
5733-
if (BuiltinID == AMDGPU::BI__builtin_amdgcn_fence ||
5734-
BuiltinID == AMDGPU::BI__builtin_amdgcn_masked_fence)
5732+
if (BuiltinID == AMDGPU::BI__builtin_amdgcn_fence)
57355733
return Diag(ArgExpr->getBeginLoc(),
57365734
diag::warn_atomic_op_has_invalid_memory_order)
57375735
<< 0 << ArgExpr->getSourceRange();

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

Lines changed: 0 additions & 100 deletions
This file was deleted.

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

Lines changed: 98 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -1,22 +1,113 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 4
12
// REQUIRES: amdgpu-registered-target
23
// RUN: %clang_cc1 %s -emit-llvm -O0 -o - \
3-
// RUN: -triple=amdgcn-amd-amdhsa | opt -S | FileCheck %s
4+
// RUN: -triple=amdgcn-amd-amdhsa | FileCheck %s
45

6+
// CHECK-LABEL: define dso_local void @_Z25test_memory_fence_successv(
7+
// CHECK-SAME: ) #[[ATTR0:[0-9]+]] {
8+
// CHECK-NEXT: entry:
9+
// CHECK-NEXT: fence syncscope("workgroup") seq_cst
10+
// CHECK-NEXT: fence syncscope("agent") acquire
11+
// CHECK-NEXT: fence seq_cst
12+
// CHECK-NEXT: fence syncscope("agent") acq_rel
13+
// CHECK-NEXT: fence syncscope("workgroup") release
14+
// CHECK-NEXT: ret void
15+
//
516
void test_memory_fence_success() {
6-
// CHECK-LABEL: test_memory_fence_success
717

8-
// CHECK: fence syncscope("workgroup") seq_cst
918
__builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup");
1019

11-
// CHECK: fence syncscope("agent") acquire
1220
__builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent");
1321

14-
// CHECK: fence seq_cst
1522
__builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "");
1623

17-
// CHECK: fence syncscope("agent") acq_rel
1824
__builtin_amdgcn_fence(4, "agent");
1925

20-
// CHECK: fence syncscope("workgroup") release
2126
__builtin_amdgcn_fence(3, "workgroup");
2227
}
28+
29+
// CHECK-LABEL: define dso_local void @_Z10test_localv(
30+
// CHECK-SAME: ) #[[ATTR0]] {
31+
// CHECK-NEXT: entry:
32+
// CHECK-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META3:![0-9]+]]
33+
// CHECK-NEXT: fence syncscope("agent") acquire, !mmra [[META3]]
34+
// CHECK-NEXT: fence seq_cst, !mmra [[META3]]
35+
// CHECK-NEXT: fence syncscope("agent") acq_rel, !mmra [[META3]]
36+
// CHECK-NEXT: fence syncscope("workgroup") release, !mmra [[META3]]
37+
// CHECK-NEXT: ret void
38+
//
39+
void test_local() {
40+
__builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "local");
41+
42+
__builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent", "local");
43+
44+
__builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "", "local");
45+
46+
__builtin_amdgcn_fence(4, "agent", "local");
47+
48+
__builtin_amdgcn_fence(3, "workgroup", "local");
49+
}
50+
51+
52+
// CHECK-LABEL: define dso_local void @_Z11test_globalv(
53+
// CHECK-SAME: ) #[[ATTR0]] {
54+
// CHECK-NEXT: entry:
55+
// CHECK-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META4:![0-9]+]]
56+
// CHECK-NEXT: fence syncscope("agent") acquire, !mmra [[META4]]
57+
// CHECK-NEXT: fence seq_cst, !mmra [[META4]]
58+
// CHECK-NEXT: fence syncscope("agent") acq_rel, !mmra [[META4]]
59+
// CHECK-NEXT: fence syncscope("workgroup") release, !mmra [[META4]]
60+
// CHECK-NEXT: ret void
61+
//
62+
void test_global() {
63+
__builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "global");
64+
65+
__builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent", "global");
66+
67+
__builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "", "global");
68+
69+
__builtin_amdgcn_fence(4, "agent", "global");
70+
71+
__builtin_amdgcn_fence(3, "workgroup", "global");
72+
}
73+
74+
// CHECK-LABEL: define dso_local void @_Z10test_imagev(
75+
// CHECK-SAME: ) #[[ATTR0]] {
76+
// CHECK-NEXT: entry:
77+
// CHECK-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META5:![0-9]+]]
78+
// CHECK-NEXT: fence syncscope("agent") acquire, !mmra [[META5]]
79+
// CHECK-NEXT: fence seq_cst, !mmra [[META5]]
80+
// CHECK-NEXT: fence syncscope("agent") acq_rel, !mmra [[META5]]
81+
// CHECK-NEXT: fence syncscope("workgroup") release, !mmra [[META5]]
82+
// CHECK-NEXT: ret void
83+
//
84+
void test_image() {
85+
__builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "image");
86+
87+
__builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent", "image");
88+
89+
__builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "", "image");
90+
91+
__builtin_amdgcn_fence(4, "agent", "image");
92+
93+
__builtin_amdgcn_fence(3, "workgroup", "image");
94+
}
95+
96+
// CHECK-LABEL: define dso_local void @_Z10test_mixedv(
97+
// CHECK-SAME: ) #[[ATTR0]] {
98+
// CHECK-NEXT: entry:
99+
// CHECK-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META6:![0-9]+]]
100+
// CHECK-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META7:![0-9]+]]
101+
// CHECK-NEXT: ret void
102+
//
103+
void test_mixed() {
104+
__builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "image", "global");
105+
__builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "image", "local", "global");
106+
}
107+
//.
108+
// CHECK: [[META3]] = !{!"amdgpu-as", !"local"}
109+
// CHECK: [[META4]] = !{!"amdgpu-as", !"global"}
110+
// CHECK: [[META5]] = !{!"amdgpu-as", !"image"}
111+
// CHECK: [[META6]] = !{[[META5]], [[META4]]}
112+
// CHECK: [[META7]] = !{[[META5]], [[META3]], [[META4]]}
113+
//.

0 commit comments

Comments
 (0)