Skip to content

Commit 4019d91

Browse files
MrSidimsjsji
authored andcommitted
Add memory semantics to atomic builtins (#3375)
It's deduced from a pointer's address space. Resolves: KhronosGroup/SPIRV-LLVM-Translator#3371 Signed-off-by: Sidorov, Dmitry <dmitry.sidorov@intel.com> Original commit: KhronosGroup/SPIRV-LLVM-Translator@630f8ef0759944e
1 parent f30e016 commit 4019d91

File tree

9 files changed

+156
-30
lines changed

9 files changed

+156
-30
lines changed

llvm-spirv/lib/SPIRV/OCLToSPIRV.cpp

Lines changed: 107 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -42,15 +42,21 @@
4242
#include "SPIRVInternal.h"
4343
#include "libSPIRV/SPIRVDebug.h"
4444

45+
#include "llvm/ADT/SmallPtrSet.h"
46+
#include "llvm/ADT/SmallVector.h"
4547
#include "llvm/ADT/StringSwitch.h"
4648
#include "llvm/Analysis/ValueTracking.h"
49+
#include "llvm/IR/Constants.h"
4750
#include "llvm/IR/IRBuilder.h"
4851
#include "llvm/IR/Instruction.h"
4952
#include "llvm/IR/Instructions.h"
53+
#include "llvm/IR/Operator.h"
5054
#include "llvm/IR/PatternMatch.h"
55+
#include "llvm/IR/TypedPointerType.h"
5156
#include "llvm/Support/Debug.h"
5257

5358
#include <algorithm>
59+
#include <optional>
5460
#include <regex>
5561
#include <set>
5662

@@ -62,6 +68,88 @@ using namespace SPIRV;
6268
using namespace OCLUtil;
6369

6470
namespace SPIRV {
71+
72+
static unsigned getAddressSpaceFromType(const Type *Ty) {
73+
assert(Ty && "Can't deduce pointer AS");
74+
if (auto *TypedPtr = dyn_cast<TypedPointerType>(Ty))
75+
return TypedPtr->getAddressSpace();
76+
if (auto *Ptr = dyn_cast<PointerType>(Ty))
77+
return Ptr->getAddressSpace();
78+
llvm_unreachable("Can't deduce pointer AS");
79+
}
80+
81+
// Performs an address space inference analysis.
82+
static unsigned getAddressSpaceFromValue(const Value *Ptr) {
83+
assert(Ptr && "Can't deduce pointer AS");
84+
85+
SmallPtrSet<const Value *, 8> Visited;
86+
SmallVector<const Value *, 8> Worklist;
87+
Worklist.push_back(Ptr);
88+
unsigned AS = SPIRAS_Generic;
89+
90+
while (!Worklist.empty()) {
91+
const Value *Current = Worklist.pop_back_val();
92+
if (!Visited.insert(Current).second)
93+
continue;
94+
95+
unsigned DeducedAS = getAddressSpaceFromType(Current->getType());
96+
if (DeducedAS != SPIRAS_Generic)
97+
return DeducedAS;
98+
AS = DeducedAS;
99+
100+
// Find origins of the pointer and add to the worklist.
101+
if (auto *Op = dyn_cast<Operator>(Current)) {
102+
switch (Op->getOpcode()) {
103+
case Instruction::AddrSpaceCast:
104+
case Instruction::BitCast:
105+
case Instruction::GetElementPtr:
106+
Worklist.push_back(Op->getOperand(0));
107+
break;
108+
case Instruction::Select:
109+
Worklist.push_back(Op->getOperand(1));
110+
Worklist.push_back(Op->getOperand(2));
111+
break;
112+
case Instruction::PHI: {
113+
auto *Phi = cast<PHINode>(Op);
114+
for (Value *Incoming : Phi->incoming_values())
115+
Worklist.push_back(Incoming);
116+
break;
117+
}
118+
default:
119+
break;
120+
}
121+
}
122+
}
123+
124+
return AS;
125+
}
126+
127+
// Sets memory semantic mask of an atomic depending on a pointer argument
128+
// address space.
129+
static unsigned
130+
getAtomicPointerMemorySemanticsMemoryMask(const Value *Ptr,
131+
const Type *RecordedType) {
132+
assert((Ptr && RecordedType) &&
133+
"Can't evaluate atomic builtin's memory semantic");
134+
unsigned AddrSpace = getAddressSpaceFromType(RecordedType);
135+
if (AddrSpace == SPIRAS_Generic)
136+
AddrSpace = getAddressSpaceFromValue(Ptr);
137+
138+
switch (AddrSpace) {
139+
case SPIRAS_Global:
140+
case SPIRAS_GlobalDevice:
141+
case SPIRAS_GlobalHost:
142+
return MemorySemanticsCrossWorkgroupMemoryMask;
143+
case SPIRAS_Local:
144+
return MemorySemanticsWorkgroupMemoryMask;
145+
case SPIRAS_Generic:
146+
return MemorySemanticsCrossWorkgroupMemoryMask |
147+
MemorySemanticsWorkgroupMemoryMask;
148+
default:
149+
return MemorySemanticsMaskNone;
150+
}
151+
}
152+
65153
static size_t getOCLCpp11AtomicMaxNumOps(StringRef Name) {
66154
return StringSwitch<size_t>(Name)
67155
.Cases({"load", "flag_test_and_set", "flag_clear"}, 3)
@@ -700,6 +788,11 @@ void OCLToSPIRVBase::transAtomicBuiltin(CallInst *CI,
700788
const size_t ScopeIdx = ArgsCount - 1;
701789
const size_t OrderIdx = ScopeIdx - NumOrder;
702790

791+
unsigned PtrMemSemantics = MemorySemanticsMaskNone;
792+
if (Mutator.arg_size() > 0)
793+
PtrMemSemantics = getAtomicPointerMemorySemanticsMemoryMask(
794+
Mutator.getArg(0), Mutator.getType(0));
795+
703796
if (NeedsNegate) {
704797
Mutator.mapArg(1, [=](Value *V) {
705798
IRBuilder<> IRB(CI);
@@ -710,9 +803,20 @@ void OCLToSPIRVBase::transAtomicBuiltin(CallInst *CI,
710803
return transOCLMemScopeIntoSPIRVScope(V, OCLMS_device, CI);
711804
});
712805
for (size_t I = 0; I < NumOrder; ++I) {
713-
Mutator.mapArg(OrderIdx + I, [=](Value *V) {
714-
return transOCLMemOrderIntoSPIRVMemorySemantics(V, OCLMO_seq_cst, CI);
715-
});
806+
Mutator.mapArg(
807+
OrderIdx + I, [=](IRBuilder<> &Builder, Value *V) -> Value * {
808+
Value *MemSem =
809+
transOCLMemOrderIntoSPIRVMemorySemantics(V, OCLMO_seq_cst, CI);
810+
if (PtrMemSemantics == MemorySemanticsMaskNone)
811+
return MemSem;
812+
813+
auto *MemSemTy = cast<IntegerType>(MemSem->getType());
814+
auto *Mask = ConstantInt::get(MemSemTy, PtrMemSemantics);
815+
if (auto *Const = dyn_cast<ConstantInt>(MemSem))
816+
return static_cast<Value *>(ConstantInt::get(
817+
MemSemTy, Const->getZExtValue() | PtrMemSemantics));
818+
return Builder.CreateOr(MemSem, Mask);
819+
});
716820
}
717821

718822
// Order of args in SPIR-V:

llvm-spirv/test/AtomicCompareExchangeExplicit.ll

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,9 @@ target triple = "spir64"
88

99
; CHECK-DAG: 4 TypeInt [[#int:]] 32 0
1010
; CHECK-DAG: Constant [[#int]] [[#DeviceScope:]] 4
11-
; CHECK-DAG: Constant [[#int]] [[#SequentiallyConsistent_MS:]] 0
11+
; Memory semantics: 256 = WorkgroupMemory (256) | SequentiallyConsistent (0)
12+
; Local address space (3) maps to WorkgroupMemory storage class
13+
; CHECK-DAG: Constant [[#int]] [[#SequentiallyConsistent_MS:]] 256
1214
; CHECK-DAG: 4 TypePointer [[#int_ptr:]] 4 [[#int]]
1315
; CHECK-DAG: 2 TypeBool [[#bool:]]
1416

llvm-spirv/test/AtomicCompareExchange_cl20.ll

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,9 @@ target triple = "spir-unknown-unknown"
1717

1818
; CHECK: 4 TypeInt [[int:[0-9]+]] 32 0
1919
; CHECK: Constant [[int]] [[DeviceScope:[0-9]+]] 1
20-
; CHECK: Constant [[int]] [[SequentiallyConsistent_MS:[0-9]+]] 16
20+
; For generic AS with SequentiallyConsistent: 784 = 768 (storage class) + 16 (SeqCst)
21+
; Where 768 = CrossWorkgroupMemory (512) | WorkgroupMemory (256)
22+
; CHECK: Constant [[int]] [[SequentiallyConsistent_MS:[0-9]+]] 784
2123
; CHECK-TYPED-PTR: 4 TypePointer [[int_ptr:[0-9]+]] 8 [[int]]
2224
; CHECK-UNTYPED-PTR: 3 TypeUntypedPointerKHR [[int_ptr:[0-9]+]] 8
2325
; CHECK: 2 TypeBool [[bool:[0-9]+]]

llvm-spirv/test/extensions/EXT/SPV_EXT_shader_atomic_float_/AtomicFSubEXTForOCL.ll

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -29,7 +29,8 @@ define dso_local spir_func void @test_atomic_float(ptr addrspace(1) %a) local_un
2929
entry:
3030
; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_32]] 13 7 10 11 [[NEGATIVE_229]]
3131
; CHECK-LLVM-CL20: call spir_func float @_Z25atomic_fetch_add_explicitPU3AS4VU7_Atomicff12memory_order12memory_scope(ptr addrspace(4) %a.as, float -2.290000e+02, i32 0, i32 1) #0
32-
; CHECK-LLVM-SPV: call spir_func float @_Z21__spirv_AtomicFAddEXTPU3AS1fiif(ptr addrspace(1) %a, i32 2, i32 0, float -2.290000e+02) #0
32+
; Memory semantics: 512 = CrossWorkgroupMemory (512) | Relaxed (0) for global AS
33+
; CHECK-LLVM-SPV: call spir_func float @_Z21__spirv_AtomicFAddEXTPU3AS1fiif(ptr addrspace(1) %a, i32 2, i32 512, float -2.290000e+02) #0
3334
%call2 = tail call spir_func float @_Z25atomic_fetch_sub_explicitPU3AS1VU7_Atomicff12memory_order12memory_scope(ptr addrspace(1) noundef %a, float noundef 2.290000e+02, i32 noundef 0, i32 noundef 1) #2
3435
ret void
3536
}
@@ -43,7 +44,8 @@ define dso_local spir_func void @test_atomic_double(ptr addrspace(1) %a) local_u
4344
entry:
4445
; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_64]] 21 18 10 11 [[NEGATIVE_334]]
4546
; CHECK-LLVM-CL20: call spir_func double @_Z25atomic_fetch_add_explicitPU3AS4VU7_Atomicdd12memory_order12memory_scope(ptr addrspace(4) %a.as, double -3.340000e+02, i32 0, i32 1) #0
46-
; CHECK-LLVM-SPV: call spir_func double @_Z21__spirv_AtomicFAddEXTPU3AS1diid(ptr addrspace(1) %a, i32 2, i32 0, double -3.340000e+02) #0
47+
; Memory semantics: 512 = CrossWorkgroupMemory (512) | Relaxed (0) for global AS
48+
; CHECK-LLVM-SPV: call spir_func double @_Z21__spirv_AtomicFAddEXTPU3AS1diid(ptr addrspace(1) %a, i32 2, i32 512, double -3.340000e+02) #0
4749
%call = tail call spir_func double @_Z25atomic_fetch_sub_explicitPU3AS1VU7_Atomicdd12memory_order12memory_scope(ptr addrspace(1) noundef %a, double noundef 3.340000e+02, i32 noundef 0, i32 noundef 1) #2
4850
ret void
4951
}

llvm-spirv/test/transcoding/AtomicCompareExchangeExplicit_cl20.cl

Lines changed: 7 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -49,9 +49,13 @@ DEFINE_KERNEL(double)
4949
//; Constants below correspond to the SPIR-V spec
5050
//CHECK-SPIRV-DAG: Constant [[int32]] [[DeviceScope:[0-9]+]] 1
5151
//CHECK-SPIRV-DAG: Constant [[int32]] [[WorkgroupScope:[0-9]+]] 2
52-
//CHECK-SPIRV-DAG: Constant [[int32]] [[ReleaseMemSem:[0-9]+]] 4
53-
//CHECK-SPIRV-DAG: Constant [[int32]] [[RelaxedMemSem:[0-9]+]] 0
54-
//CHECK-SPIRV-DAG: Constant [[int32]] [[AcqRelMemSem:[0-9]+]] 8
52+
//; Memory semantics include both memory order and storage class bits
53+
//; 516 = CrossWorkgroupMemory (512) | Release (4)
54+
//; 512 = CrossWorkgroupMemory (512) | Relaxed (0)
55+
//; 520 = CrossWorkgroupMemory (512) | AcqRel (8)
56+
//CHECK-SPIRV-DAG: Constant [[int32]] [[ReleaseMemSem:[0-9]+]] 516
57+
//CHECK-SPIRV-DAG: Constant [[int32]] [[RelaxedMemSem:[0-9]+]] 512
58+
//CHECK-SPIRV-DAG: Constant [[int32]] [[AcqRelMemSem:[0-9]+]] 520
5559

5660
//CHECK-SPIRV: AtomicCompareExchange [[int32]] {{[0-9]+}} {{[0-9]+}} [[DeviceScope]] [[ReleaseMemSem]] [[RelaxedMemSem]]
5761
//CHECK-SPIRV: AtomicCompareExchange [[int32]] {{[0-9]+}} {{[0-9]+}} [[WorkgroupScope]] [[AcqRelMemSem]] [[RelaxedMemSem]]

llvm-spirv/test/transcoding/OpenCL/atomic_cmpxchg.cl

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -35,9 +35,8 @@ __kernel void test_atomic_cmpxchg(__global int *p, int cmp, int val) {
3535
// 0x2 Workgroup
3636
// CHECK-SPIRV-DAG: Constant [[UINT]] [[WORKGROUP_SCOPE:[0-9]+]] 2
3737
//
38-
// 0x0 Relaxed
39-
// TODO: do we need CrossWorkgroupMemory here as well?
40-
// CHECK-SPIRV-DAG: Constant [[UINT]] [[RELAXED:[0-9]+]] 0
38+
// 0x0 Relaxed | 0x200 CrossWorkgroupMemory
39+
// CHECK-SPIRV-DAG: Constant [[UINT]] [[RELAXED:[0-9]+]] 512
4140
//
4241
// CHECK-SPIRV: Function {{[0-9]+}} [[TEST]]
4342
// CHECK-SPIRV: FunctionParameter [[UINT_PTR]] [[PTR:[0-9]+]]

llvm-spirv/test/transcoding/OpenCL/atomic_legacy.cl

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -31,8 +31,9 @@ __kernel void test_legacy_atomics(__global int *p, int val) {
3131
// 0x2 Workgroup
3232
// CHECK-SPIRV-DAG: Constant [[UINT]] [[WORKGROUP_SCOPE:[0-9]+]] 2
3333
//
34-
// 0x0 Relaxed
35-
// CHECK-SPIRV-DAG: Constant [[UINT]] [[RELAXED:[0-9]+]] 0
34+
// 0x200 CrossWorkgroupMemory | 0x0 Relaxed = 512
35+
// Global address space (AS 1) maps to CrossWorkgroupMemory storage class
36+
// CHECK-SPIRV-DAG: Constant [[UINT]] [[RELAXED:[0-9]+]] 512
3637
//
3738
// CHECK-SPIRV: Function {{[0-9]+}} [[TEST]]
3839
// CHECK-SPIRV: FunctionParameter [[UINT_PTR]] [[PTR:[0-9]+]]

llvm-spirv/test/transcoding/OpenCL/atomic_syncscope_test.ll

Lines changed: 19 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -33,19 +33,23 @@ target triple = "spir64"
3333
; 4 - sub_group
3434

3535
; CHECK-SPIRV-DAG: Constant [[#]] [[#ConstInt0:]] 0
36-
; CHECK-SPIRV-DAG: Constant [[#]] [[#SequentiallyConsistent:]] 16
36+
; CHECK-SPIRV-DAG: Constant [[#]] [[#SCPrivate:]] 16
3737
; CHECK-SPIRV-DAG: Constant [[#]] [[#ConstInt1:]] 1
3838
; CHECK-SPIRV-DAG: Constant [[#]] [[#ConstInt2:]] 2
3939
; CHECK-SPIRV-DAG: Constant [[#]] [[#ConstInt3:]] 3
4040
; CHECK-SPIRV-DAG: Constant [[#]] [[#ConstInt4:]] 4
4141
; CHECK-SPIRV-DAG: Constant [[#]] [[#Const2Power30:]] 1073741824
4242
; CHECK-SPIRV-DAG: Constant [[#]] [[#ConstInt42:]] 42
43+
; Note: Storage class bits (SCGlobal, SCLocal, etc.) are not added for plain LLVM IR atomics
44+
; Only OpenCL builtin atomics get the storage class memory semantics bits from the patch
4345

4446
; AtomicLoad ResTypeId ResId PtrId MemScopeId MemSemanticsId
45-
; CHECK-SPIRV: AtomicLoad [[#]] [[#]] [[#]] [[#ConstInt2]] [[#SequentiallyConsistent]]
46-
; CHECK-SPIRV: AtomicLoad [[#]] [[#]] [[#]] [[#ConstInt1]] [[#SequentiallyConsistent]]
47-
; CHECK-SPIRV: AtomicLoad [[#]] [[#]] [[#]] [[#ConstInt0]] [[#SequentiallyConsistent]]
48-
; CHECK-SPIRV: AtomicLoad [[#]] [[#]] [[#]] [[#ConstInt3]] [[#SequentiallyConsistent]]
47+
; Note: Plain LLVM atomic loads don't get storage class bits added (only OpenCL builtins do)
48+
; These use SCPrivate (16) which is SequentiallyConsistent without storage class bits
49+
; CHECK-SPIRV: AtomicLoad [[#]] [[#]] [[#]] [[#ConstInt2]] [[#SCPrivate]]
50+
; CHECK-SPIRV: AtomicLoad [[#]] [[#]] [[#]] [[#ConstInt1]] [[#SCPrivate]]
51+
; CHECK-SPIRV: AtomicLoad [[#]] [[#]] [[#]] [[#ConstInt0]] [[#SCPrivate]]
52+
; CHECK-SPIRV: AtomicLoad [[#]] [[#]] [[#]] [[#ConstInt3]] [[#SCPrivate]]
4953

5054
; CHECK-LLVM: call spir_func i32 @_Z20atomic_load_explicitPU3AS4VU7_Atomici12memory_order12memory_scope(ptr{{.*}}, i32 5, i32 1)
5155
; CHECK-LLVM: call spir_func i32 @_Z20atomic_load_explicitPU3AS4VU7_Atomici12memory_order12memory_scope(ptr{{.*}}, i32 5, i32 2)
@@ -62,8 +66,9 @@ entry:
6266
}
6367

6468
; AtomicStore PtrId MemScopeId MemSemanticsId ValueId
65-
; CHECK-SPIRV: AtomicStore [[#]] [[#ConstInt3]] [[#SequentiallyConsistent]] [[#ConstInt1]]
66-
; CHECK-SPIRV: AtomicStore [[#]] [[#ConstInt2]] [[#SequentiallyConsistent]] [[#ConstInt1]]
69+
; Plain LLVM IR store atomic instructions don't get storage class bits
70+
; CHECK-SPIRV: AtomicStore [[#]] [[#ConstInt3]] [[#SCPrivate]] [[#ConstInt1]]
71+
; CHECK-SPIRV: AtomicStore [[#]] [[#ConstInt2]] [[#SCPrivate]] [[#ConstInt1]]
6772
; CHECK-LLVM: call spir_func void @_Z21atomic_store_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(ptr{{.*}}, i32 5, i32 4)
6873
; CHECK-LLVM: call spir_func void @_Z21atomic_store_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(ptr{{.*}}, i32 5, i32 1)
6974

@@ -75,11 +80,11 @@ entry:
7580
}
7681

7782
; Atomic* ResTypeId ResId PtrId MemScopeId MemSemanticsId ValueId
78-
; CHECK-SPIRV: AtomicAnd [[#]] [[#]] [[#]] [[#ConstInt4]] [[#SequentiallyConsistent]] [[#ConstInt1]]
79-
; CHECK-SPIRV: AtomicSMin [[#]] [[#]] [[#]] [[#ConstInt0]] [[#SequentiallyConsistent]] [[#ConstInt1]]
80-
; CHECK-SPIRV: AtomicSMax [[#]] [[#]] [[#]] [[#ConstInt0]] [[#SequentiallyConsistent]] [[#ConstInt1]]
81-
; CHECK-SPIRV: AtomicUMin [[#]] [[#]] [[#]] [[#ConstInt2]] [[#SequentiallyConsistent]] [[#ConstInt1]]
82-
; CHECK-SPIRV: AtomicUMax [[#]] [[#]] [[#]] [[#ConstInt2]] [[#SequentiallyConsistent]] [[#ConstInt1]]
83+
; CHECK-SPIRV: AtomicAnd [[#]] [[#]] [[#]] [[#ConstInt4]] [[#SCPrivate]] [[#ConstInt1]]
84+
; CHECK-SPIRV: AtomicSMin [[#]] [[#]] [[#]] [[#ConstInt0]] [[#SCPrivate]] [[#ConstInt1]]
85+
; CHECK-SPIRV: AtomicSMax [[#]] [[#]] [[#]] [[#ConstInt0]] [[#SCPrivate]] [[#ConstInt1]]
86+
; CHECK-SPIRV: AtomicUMin [[#]] [[#]] [[#]] [[#ConstInt2]] [[#SCPrivate]] [[#ConstInt1]]
87+
; CHECK-SPIRV: AtomicUMax [[#]] [[#]] [[#]] [[#ConstInt2]] [[#SCPrivate]] [[#ConstInt1]]
8388

8489
; CHECK-LLVM: call spir_func i32 @_Z25atomic_fetch_and_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(ptr{{.*}}, i32 1, i32 5, i32 0)
8590
; CHECK-LLVM: call spir_func i32 @_Z25atomic_fetch_min_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(ptr{{.*}}, i32 1, i32 5, i32 3)
@@ -109,7 +114,7 @@ entry:
109114
}
110115

111116
; AtomicExchange ResTypeId ResId PtrId MemScopeId MemSemanticsId ValueId
112-
; CHECK-SPIRV: AtomicExchange [[#]] [[#]] [[#]] [[#ConstInt2]] [[#SequentiallyConsistent]] [[#Const2Power30]]
117+
; CHECK-SPIRV: AtomicExchange [[#]] [[#]] [[#]] [[#ConstInt2]] [[#SCPrivate]] [[#Const2Power30]]
113118
; CHECK-LLVM: call spir_func i32 @_Z24atomic_exchange_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(ptr{{.*}}, i32 1073741824, i32 5, i32 1)
114119

115120
define dso_local float @ff3(ptr captures(none) noundef %d) local_unnamed_addr #0 {
@@ -120,6 +125,7 @@ entry:
120125
}
121126

122127
; AtomicFAddEXT ResTypeId ResId PtrId MemScopeId MemSemanticsId ValueId
128+
; Plain LLVM atomicrmw fadd doesn't get storage class bits
123129
; CHECK-SPIRV: AtomicFAddEXT [[#]] [[#]] [[#]] [[#ConstInt2]] [[#ConstInt0]] [[#]]
124130
; CHECK-LLVM: call spir_func float @_Z25atomic_fetch_add_explicitPU3AS4VU7_Atomicff12memory_order12memory_scope(ptr{{.*}}, i32 0, i32 1)
125131

llvm-spirv/test/transcoding/atomic_explicit_arguments.cl

Lines changed: 8 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -18,6 +18,8 @@ int load (volatile atomic_int* obj, memory_order order, memory_scope scope) {
1818
// CHECK-SPIRV: Name [[TRANS_MEM_ORDER:[0-9]+]] "__translate_ocl_memory_order"
1919

2020
// CHECK-SPIRV: TypeInt [[int:[0-9]+]] 32 0
21+
// Memory semantics for generic AS: 768 = CrossWorkgroupMemory | WorkgroupMemory
22+
// CHECK-SPIRV-DAG: Constant [[int]] [[GENERIC_STORAGE_MASK:[0-9]+]] 768
2123
// CHECK-SPIRV-DAG: Constant [[int]] [[ZERO:[0-9]+]] 0
2224
// CHECK-SPIRV-DAG: Constant [[int]] [[ONE:[0-9]+]] 1
2325
// CHECK-SPIRV-DAG: Constant [[int]] [[TWO:[0-9]+]] 2
@@ -31,8 +33,11 @@ int load (volatile atomic_int* obj, memory_order order, memory_scope scope) {
3133
// CHECK-SPIRV: FunctionParameter {{[0-9]+}} [[OCL_ORDER:[0-9]+]]
3234
// CHECK-SPIRV: FunctionParameter {{[0-9]+}} [[OCL_SCOPE:[0-9]+]]
3335

36+
//
3437
// CHECK-SPIRV: FunctionCall [[int]] [[SPIRV_SCOPE:[0-9]+]] [[TRANS_MEM_SCOPE]] [[OCL_SCOPE]]
35-
// CHECK-SPIRV: FunctionCall [[int]] [[SPIRV_ORDER:[0-9]+]] [[TRANS_MEM_ORDER]] [[OCL_ORDER]]
38+
// CHECK-SPIRV: FunctionCall [[int]] [[SPIRV_ORDER_BASE:[0-9]+]] [[TRANS_MEM_ORDER]] [[OCL_ORDER]]
39+
// The translated memory order is combined with storage class semantics for generic AS
40+
// CHECK-SPIRV: BitwiseOr [[int]] [[SPIRV_ORDER:[0-9]+]] [[SPIRV_ORDER_BASE]] [[GENERIC_STORAGE_MASK]]
3641
// CHECK-SPIRV: AtomicLoad [[int]] {{[0-9]+}} [[OBJECT]] [[SPIRV_SCOPE]] [[SPIRV_ORDER]]
3742

3843
// CHECK-SPIRV: Function [[int]] [[TRANS_MEM_SCOPE]]
@@ -86,5 +91,6 @@ int load (volatile atomic_int* obj, memory_order order, memory_scope scope) {
8691

8792
// CHECK-LLVM: define spir_func i32 @load(ptr addrspace(4) %[[obj:[0-9a-zA-Z._]+]], i32 %[[order:[0-9a-zA-Z._]+]], i32 %[[scope:[0-9a-zA-Z._]+]]) #0 {
8893
// CHECK-LLVM: entry:
89-
// CHECK-LLVM: call spir_func i32 @_Z20atomic_load_explicitPU3AS4VU7_Atomici12memory_order12memory_scope(ptr addrspace(4) %[[obj]], i32 %[[order]], i32 %[[scope]])
94+
// CHECK-LLVM: %[[#]] = or i32 %{{[0-9a-zA-Z._]+}}, 768
95+
// CHECK-LLVM: call spir_func i32 @_Z20atomic_load_explicitPU3AS4VU7_Atomici12memory_order12memory_scope(ptr addrspace(4) %[[obj]], i32 %{{[0-9a-zA-Z._]+}}, i32 %[[scope]])
9096
// CHECK-LLVM: }

0 commit comments

Comments
 (0)