Skip to content

Commit b12df55

Browse files
vmaksimojsji
vmaksimo
authored and
jsji
committed
Fix mangling for atomic builtins used with SPV_KHR_untyped_pointers (#2771)
This change allows to preserve the correct builtin mangling in reverse translation. All the existing tests for atomics (except atomic flag instructions which are not covered by the extension) were updated to verify we get the same mangling with and without extension enabled. Original commit: KhronosGroup/SPIRV-LLVM-Translator@566023769b3ab6a
1 parent 1f8de71 commit b12df55

18 files changed

+134
-12
lines changed

‎llvm-spirv/lib/SPIRV/SPIRVInternal.h

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -556,6 +556,19 @@ inline unsigned findFirstPtr(const Container &Args) {
556556
return PtArg - Args.begin();
557557
}
558558

559+
// Utility function to check if a type is a TypedPointerType
560+
inline bool isTypedPointerType(llvm::Type *Ty) {
561+
return llvm::isa<llvm::TypedPointerType>(Ty);
562+
}
563+
564+
template <typename Container>
565+
inline unsigned findFirstPtrType(const Container &Args) {
566+
auto PtArg = std::find_if(Args.begin(), Args.end(), [](Type *T) {
567+
return T->isPointerTy() || isTypedPointerType(T);
568+
});
569+
return PtArg - Args.begin();
570+
}
571+
559572
bool isSupportedTriple(Triple T);
560573
void removeFnAttr(CallInst *Call, Attribute::AttrKind Attr);
561574
void addFnAttr(CallInst *Call, Attribute::AttrKind Attr);

‎llvm-spirv/lib/SPIRV/SPIRVReader.cpp

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3402,6 +3402,21 @@ Instruction *SPIRVToLLVM::transBuiltinFromInst(const std::string &FuncName,
34023402
transOCLBuiltinFromInstPreproc(BI, RetTy, Ops);
34033403
std::vector<Type *> ArgTys =
34043404
transTypeVector(SPIRVInstruction::getOperandTypes(Ops), true);
3405+
3406+
// Special handling for "truly" untyped pointers to preserve correct
3407+
// builtin mangling of atomic operations.
3408+
auto Ptr = findFirstPtrType(ArgTys);
3409+
if (Ptr < ArgTys.size() &&
3410+
BI->getValueType(Ops[Ptr]->getId())->isTypeUntypedPointerKHR()) {
3411+
if (isAtomicOpCodeUntypedPtrSupported(BI->getOpCode())) {
3412+
auto *AI = static_cast<SPIRVAtomicInstBase *>(BI);
3413+
ArgTys[Ptr] = TypedPointerType::get(
3414+
transType(AI->getSemanticType()),
3415+
SPIRSPIRVAddrSpaceMap::rmap(
3416+
BI->getValueType(Ops[Ptr]->getId())->getPointerStorageClass()));
3417+
}
3418+
}
3419+
34053420
for (auto &I : ArgTys) {
34063421
if (isa<FunctionType>(I)) {
34073422
I = TypedPointerType::get(I, SPIRAS_Private);

‎llvm-spirv/lib/SPIRV/libSPIRV/SPIRVInstruction.h

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2928,6 +2928,24 @@ class SPIRVAtomicInstBase : public SPIRVInstTemplateBase {
29282928
assert(this->getModule()->getSPIRVVersion() < VersionNumber::SPIRV_1_4 &&
29292929
"OpAtomicCompareExchangeWeak is removed starting from SPIR-V 1.4");
29302930
}
2931+
2932+
// This method is needed for correct translation of atomic instructions when
2933+
// SPV_KHR_untyped_pointers is enabled.
2934+
// The interpreted data type for untyped pointers is specified by the Result
2935+
// Type if it exists, or from the type of the object being stored in other
2936+
// case.
2937+
SPIRVType *getSemanticType() {
2938+
switch (OpCode) {
2939+
case OpAtomicStore:
2940+
// Get type of Value operand
2941+
return getOperand(3)->getType();
2942+
default: {
2943+
if (hasType())
2944+
return getType();
2945+
return nullptr;
2946+
}
2947+
}
2948+
}
29312949
};
29322950

29332951
class SPIRVAtomicStoreInst : public SPIRVAtomicInstBase {

‎llvm-spirv/lib/SPIRV/libSPIRV/SPIRVOpCode.h

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -69,6 +69,13 @@ inline bool isAtomicOpCode(Op OpCode) {
6969
OpCode == OpAtomicFlagTestAndSet || OpCode == OpAtomicFlagClear ||
7070
isFPAtomicOpCode(OpCode);
7171
}
72+
inline bool isAtomicOpCodeUntypedPtrSupported(Op OpCode) {
73+
static_assert(OpAtomicLoad < OpAtomicXor, "");
74+
return ((unsigned)OpCode >= OpAtomicLoad &&
75+
(unsigned)OpCode <= OpAtomicXor) ||
76+
isFPAtomicOpCode(OpCode);
77+
}
78+
7279
inline bool isBinaryOpCode(Op OpCode) {
7380
return ((unsigned)OpCode >= OpIAdd && (unsigned)OpCode <= OpFMod) ||
7481
OpCode == OpDot || OpCode == OpIAddCarry || OpCode == OpISubBorrow ||

‎llvm-spirv/test/AtomicBuiltinsFloat.ll

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4,6 +4,10 @@
44
; RUN: llvm-spirv %t.bc -o %t.spv
55
; RUN: spirv-val %t.spv
66

7+
; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_KHR_untyped_pointers -o - | FileCheck %s
8+
; RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
9+
; RUN: spirv-val %t.spv
10+
711
; CHECK-LABEL: Label
812
; CHECK: Store
913
; CHECK-COUNT-3: AtomicStore

‎llvm-spirv/test/AtomicCompareExchange.ll

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3,6 +3,10 @@
33
; RUN: llvm-spirv -to-text %t.spv -o - | FileCheck %s --check-prefix=CHECK-SPIRV
44
; RUN: spirv-val %t.spv
55

6+
; RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
7+
; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_KHR_untyped_pointers -o - | FileCheck %s --check-prefix=CHECK-SPIRV
8+
; RUN: spirv-val %t.spv
9+
610
; CHECK-SPIRV: TypeInt [[Int:[0-9]+]] 32 0
711
; CHECK-SPIRV: Constant [[Int]] [[MemScope_CrossDevice:[0-9]+]] 0
812
; CHECK-SPIRV: Constant [[Int]] [[MemSemEqual_SeqCst:[0-9]+]] 16

‎llvm-spirv/test/AtomicCompareExchange_cl20.ll

Lines changed: 7 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,12 @@
11
; RUN: llvm-as %s -o %t.bc
2-
; RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s
2+
; RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s --check-prefixes=CHECK,CHECK-TYPED-PTR
33
; RUN: llvm-spirv %t.bc -o %t.spv
44
; RUN: spirv-val %t.spv
55

6+
; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_KHR_untyped_pointers -o - | FileCheck %s --check-prefixes=CHECK,CHECK-UNTYPED-PTR
7+
; RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
8+
; RUN: spirv-val %t.spv
9+
610
target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
711
target triple = "spir-unknown-unknown"
812

@@ -14,7 +18,8 @@ target triple = "spir-unknown-unknown"
1418
; CHECK: 4 TypeInt [[int:[0-9]+]] 32 0
1519
; CHECK: Constant [[int]] [[DeviceScope:[0-9]+]] 1
1620
; CHECK: Constant [[int]] [[SequentiallyConsistent_MS:[0-9]+]] 16
17-
; CHECK: 4 TypePointer [[int_ptr:[0-9]+]] 8 [[int]]
21+
; CHECK-TYPED-PTR: 4 TypePointer [[int_ptr:[0-9]+]] 8 [[int]]
22+
; CHECK-UNTYPED-PTR: 3 TypeUntypedPointerKHR [[int_ptr:[0-9]+]] 8
1823
; CHECK: 2 TypeBool [[bool:[0-9]+]]
1924

2025
; Function Attrs: nounwind

‎llvm-spirv/test/atomic-load-store.ll

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3,6 +3,10 @@
33
; RUN: spirv-val %t.spv
44
; RUN: llvm-spirv -to-text %t.spv -o - | FileCheck %s
55

6+
; RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
7+
; RUN: spirv-val %t.spv
8+
; RUN: llvm-spirv -to-text %t.spv -o - | FileCheck %s
9+
610
; CHECK-DAG: Constant [[#]] [[#CrossDeviceScope:]] 0
711
; CHECK-DAG: Constant [[#]] [[#Release:]] 4
812
; CHECK-DAG: Constant [[#]] [[#SequentiallyConsistent:]] 16
@@ -14,7 +18,7 @@ target triple = "spir64"
1418
; Function Attrs: nounwind
1519
define dso_local spir_func void @test() {
1620
entry:
17-
; CHECK: Variable [[#]] [[#PTR:]]
21+
; CHECK: {{(Variable|UntypedVariableKHR)}} [[#]] [[#PTR:]]
1822
%0 = alloca i32
1923

2024
; CHECK: AtomicStore [[#PTR]] [[#CrossDeviceScope]] {{.+}} [[#]]

‎llvm-spirv/test/atomicrmw.ll

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3,6 +3,10 @@
33
; RUN: spirv-val %t.spv
44
; RUN: llvm-spirv -to-text %t.spv -o - | FileCheck %s
55

6+
; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_KHR_untyped_pointers -o %t.spv
7+
; RUN: spirv-val %t.spv
8+
; RUN: llvm-spirv -to-text %t.spv -o - | FileCheck %s
9+
610
; CHECK: TypeInt [[Int:[0-9]+]] 32 0
711
; CHECK-DAG: Constant [[Int]] [[MemSem_Relaxed:[0-9]+]] 0
812
; CHECK-DAG: Constant [[Int]] [[MemSem_Acquire:[0-9]+]] 2
@@ -11,8 +15,8 @@
1115
; CHECK-DAG: Constant [[Int]] [[MemSem_SequentiallyConsistent:[0-9]+]] 16
1216
; CHECK-DAG: Constant [[Int]] [[Value:[0-9]+]] 42
1317
; CHECK: TypeFloat [[Float:[0-9]+]] 32
14-
; CHECK: Variable {{[0-9]+}} [[Pointer:[0-9]+]]
15-
; CHECK: Variable {{[0-9]+}} [[FPPointer:[0-9]+]]
18+
; CHECK: {{(Variable|UntypedVariableKHR)}} {{[0-9]+}} [[Pointer:[0-9]+]]
19+
; CHECK: {{(Variable|UntypedVariableKHR)}} {{[0-9]+}} [[FPPointer:[0-9]+]]
1620
; CHECK: Constant [[Float]] [[FPValue:[0-9]+]] 1109917696
1721

1822
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"

‎llvm-spirv/test/transcoding/AtomicCompareExchangeExplicit_cl20.cl

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,13 @@
66
// RUN: llvm-spirv -r --spirv-target-env=CL2.0 %t.spv -o %t.rev.bc
77
// RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
88

9+
// RUN: llvm-spirv %t.bc -spirv-text -o %t.txt --spirv-ext=+SPV_KHR_untyped_pointers
10+
// RUN: FileCheck < %t.txt %s --check-prefix=CHECK-SPIRV
11+
// RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
12+
// RUN: spirv-val %t.spv
13+
// RUN: llvm-spirv -r --spirv-target-env=CL2.0 %t.spv -o %t.rev.bc
14+
// RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
15+
916
#define DEFINE_KERNEL(TYPE) \
1017
__kernel void testAtomicCompareExchangeExplicit_cl20_##TYPE( \
1118
volatile global atomic_##TYPE* object, \

‎llvm-spirv/test/transcoding/AtomicCompareExchange_cl20.ll

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4,6 +4,12 @@ target triple = "spir-unknown-unknown"
44

55
; RUN: llvm-as %s -o %t.bc
66
; RUN: llvm-spirv %t.bc -o %t.spv
7+
; RUN: spirv-val %t.spv
8+
; RUN: llvm-spirv -r --spirv-target-env=CL2.0 %t.spv -o %t.bc
9+
; RUN: llvm-dis < %t.bc | FileCheck %s
10+
11+
; RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
12+
; RUN: spirv-val %t.spv
713
; RUN: llvm-spirv -r --spirv-target-env=CL2.0 %t.spv -o %t.bc
814
; RUN: llvm-dis < %t.bc | FileCheck %s
915

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

Lines changed: 8 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,12 @@
11
// RUN: %clang_cc1 %s -triple spir -cl-std=CL1.2 -emit-llvm-bc -fdeclare-opencl-builtins -o %t.bc
22
// RUN: llvm-spirv %t.bc -o %t.spv
33
// RUN: spirv-val %t.spv
4-
// RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
4+
// RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV-TYPED-PTRS
5+
// RUN: llvm-spirv %t.spv -r --spirv-target-env=CL1.2 -o - | llvm-dis -o - | FileCheck %s --check-prefix=CHECK-LLVM
6+
7+
// RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
8+
// RUN: spirv-val %t.spv
9+
// RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV-UNTYPED-PTRS
510
// RUN: llvm-spirv %t.spv -r --spirv-target-env=CL1.2 -o - | llvm-dis -o - | FileCheck %s --check-prefix=CHECK-LLVM
611

712
// This test checks that the translator is capable to correctly translate
@@ -19,7 +24,8 @@ __kernel void test_atomic_cmpxchg(__global int *p, int cmp, int val) {
1924

2025
// CHECK-SPIRV: EntryPoint [[#]] [[TEST:[0-9]+]] "test_atomic_cmpxchg"
2126
// CHECK-SPIRV-DAG: TypeInt [[UINT:[0-9]+]] 32 0
22-
// CHECK-SPIRV-DAG: TypePointer [[UINT_PTR:[0-9]+]] 5 [[UINT]]
27+
// CHECK-SPIRV-TYPED-PTRS-DAG: TypePointer [[UINT_PTR:[0-9]+]] 5 [[UINT]]
28+
// CHECK-SPIRV-UNTYPED-PTRS-DAG: TypeUntypedPointerKHR [[UINT_PTR:[0-9]+]] 5
2329
//
2430
// In SPIR-V, atomic_cmpxchg is represented as OpAtomicCompareExchange [2],
2531
// which also includes memory scope and two memory semantic arguments. The

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

Lines changed: 8 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,12 @@
11
// RUN: %clang_cc1 %s -triple spir -cl-std=CL1.2 -emit-llvm-bc -fdeclare-opencl-builtins -o %t.bc
22
// RUN: llvm-spirv %t.bc -o %t.spv
33
// RUN: spirv-val %t.spv
4-
// RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
4+
// RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV-TYPED-PTRS
5+
// RUN: llvm-spirv %t.spv -r --spirv-target-env=CL1.2 -o - | llvm-dis -o - | FileCheck %s --check-prefix=CHECK-LLVM
6+
7+
// RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
8+
// RUN: spirv-val %t.spv
9+
// RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV-UNTYPED-PTRS
510
// RUN: llvm-spirv %t.spv -r --spirv-target-env=CL1.2 -o - | llvm-dis -o - | FileCheck %s --check-prefix=CHECK-LLVM
611

712
// This test checks that the translator is capable to correctly translate
@@ -15,7 +20,8 @@ __kernel void test_legacy_atomics(__global int *p, int val) {
1520

1621
// CHECK-SPIRV: EntryPoint [[#]] [[TEST:[0-9]+]] "test_legacy_atomics"
1722
// CHECK-SPIRV-DAG: TypeInt [[UINT:[0-9]+]] 32 0
18-
// CHECK-SPIRV-DAG: TypePointer [[UINT_PTR:[0-9]+]] 5 [[UINT]]
23+
// CHECK-SPIRV-TYPED-PTRS-DAG: TypePointer [[UINT_PTR:[0-9]+]] 5 [[UINT]]
24+
// CHECK-SPIRV-UNTYPED-PTRS-DAG: TypeUntypedPointerKHR [[UINT_PTR:[0-9]+]] 5
1925
//
2026
// In SPIR-V, atomic_add is represented as OpAtomicIAdd [2], which also includes
2127
// memory scope and memory semantic arguments. The translator applies a default

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

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,13 @@
66
; RUN: llvm-spirv %t.spv -r --spirv-target-env=CL2.0 -o - | llvm-dis -o %t.rev.ll
77
; RUN: FileCheck < %t.rev.ll %s -check-prefix=CHECK-LLVM
88

9+
; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_add,+SPV_KHR_untyped_pointers -o %t.spv
10+
; RUN: spirv-val %t.spv
11+
; RUN: llvm-spirv %t.spv -to-text -o %t.spt
12+
; RUN: FileCheck < %t.spt %s -check-prefix=CHECK-SPIRV
13+
; RUN: llvm-spirv %t.spv -r --spirv-target-env=CL2.0 -o - | llvm-dis -o %t.rev.ll
14+
; RUN: FileCheck < %t.rev.ll %s -check-prefix=CHECK-LLVM
15+
916
target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
1017
target triple = "spir64"
1118

‎llvm-spirv/test/transcoding/atomic_explicit_arguments.cl

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -5,6 +5,10 @@
55
// RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
66
// RUN: llvm-spirv %t.spv -r --spirv-target-env=CL2.0 -o - | llvm-dis -o - | FileCheck %s --check-prefix=CHECK-LLVM
77

8+
// RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
9+
// RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
10+
// RUN: llvm-spirv %t.spv -r --spirv-target-env=CL2.0 -o - | llvm-dis -o - | FileCheck %s --check-prefix=CHECK-LLVM
11+
812
int load (volatile atomic_int* obj, memory_order order, memory_scope scope) {
913
return atomic_load_explicit(obj, order, scope);
1014
}
@@ -33,7 +37,7 @@ int load (volatile atomic_int* obj, memory_order order, memory_scope scope) {
3337

3438
// CHECK-SPIRV: Function [[int]] [[TRANS_MEM_SCOPE]]
3539
// CHECK-SPIRV: FunctionParameter [[int]] [[KEY:[0-9]+]]
36-
// CHECK-SPIRV: Variable {{[0-9]+}} [[RES:[0-9]+]]
40+
// CHECK-SPIRV: {{(Variable|UntypedVariableKHR)}} {{[0-9]+}} [[RES:[0-9]+]]
3741
// CHECK-SPIRV: Switch [[KEY]] [[CASE_2:[0-9]+]] 0 [[CASE_0:[0-9]+]] 1 [[CASE_1:[0-9]+]] 2 [[CASE_2]] 3 [[CASE_3:[0-9]+]] 4 [[CASE_4:[0-9]+]]
3842
// CHECK-SPIRV: Label [[CASE_0]]
3943
// CHECK-SPIRV: Store [[RES]] [[FOUR]]
@@ -57,7 +61,7 @@ int load (volatile atomic_int* obj, memory_order order, memory_scope scope) {
5761

5862
// CHECK-SPIRV: Function [[int]] [[TRANS_MEM_ORDER]]
5963
// CHECK-SPIRV: FunctionParameter [[int]] [[KEY:[0-9]+]]
60-
// CHECK-SPIRV: Variable {{[0-9]+}} [[RES:[0-9]+]]
64+
// CHECK-SPIRV: {{(Variable|UntypedVariableKHR)}} {{[0-9]+}} [[RES:[0-9]+]]
6165
// CHECK-SPIRV: Switch [[KEY]] [[CASE_5:[0-9]+]] 0 [[CASE_0:[0-9]+]] 2 [[CASE_2:[0-9]+]] 3 [[CASE_3:[0-9]+]] 4 [[CASE_4:[0-9]+]] 5 [[CASE_5]]
6266
// CHECK-SPIRV: Label [[CASE_0]]
6367
// CHECK-SPIRV: Store [[RES]] [[ZERO]]

‎llvm-spirv/test/transcoding/atomic_flag.cl

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4,7 +4,6 @@
44
// RUN: spirv-val %t.spv
55
// RUN: llvm-spirv -r --spirv-target-env=CL2.0 %t.spv -o %t.rev.bc
66
// RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
7-
87
kernel void testAtomicFlag(global int *res) {
98
atomic_flag f;
109

‎llvm-spirv/test/transcoding/atomic_load_store.ll

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,13 @@ target triple = "spir-unknown-unknown"
1010
; RUN: llvm-spirv -r --spirv-target-env=CL2.0 %t.spv -o %t.bc
1111
; RUN: llvm-dis < %t.bc | FileCheck %s --check-prefix=CHECK-LLVM
1212

13+
; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_KHR_untyped_pointers -spirv-text -o %t.spt
14+
; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
15+
; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_KHR_untyped_pointers -o %t.spv
16+
; RUN: spirv-val %t.spv
17+
; RUN: llvm-spirv -r --spirv-target-env=CL2.0 %t.spv -o %t.bc
18+
; RUN: llvm-dis < %t.bc | FileCheck %s --check-prefix=CHECK-LLVM
19+
1320
; Check 'LLVM ==> SPIR-V ==> LLVM' conversion of atomic_load and atomic_store.
1421

1522

‎llvm-spirv/test/transcoding/atomics_1.2.ll

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -7,6 +7,12 @@ target triple = "spir64-unknown-unknown"
77
; RUN: llvm-spirv -r %t.spv -o %t.bc
88
; RUN: llvm-dis < %t.bc | FileCheck %s
99

10+
; RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
11+
; TODO: investigate why function parameters are decorated with Volatile multiple times.
12+
; R/UN: spirv-val %t.spv
13+
; RUN: llvm-spirv -r %t.spv -o %t.bc
14+
; RUN: llvm-dis < %t.bc | FileCheck %s
15+
1016
; Most of atomics lost information about the sign of the integer operand
1117
; but since this concerns only built-ins with two-complement's arithmetics
1218
; it shouldn't cause any problems.

0 commit comments

Comments
 (0)