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, \

0 commit comments

Comments
 (0)