Skip to content

Commit 85815e7

Browse files
airliedairlied
authored
Add an entry point wrapper around functions (llvm pass) (#1149)
SPIR-V spec states: "It is invalid for any function to be targeted by both an OpEntryPoint instruction and an OpFunctionCall instruction." In order to satisfy SPIR-V that entrypoints and functions must be different, this introduces an entrypoint wrapper around functions at the LLVM IR level, then fixes up a few things like naming at the SPIRV translation.
1 parent 2db19de commit 85815e7

25 files changed

+155
-36
lines changed

‎lib/SPIRV/SPIRVInternal.h

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -377,6 +377,7 @@ const static char TranslateOCLMemScope[] = "__translate_ocl_memory_scope";
377377
const static char TranslateSPIRVMemOrder[] = "__translate_spirv_memory_order";
378378
const static char TranslateSPIRVMemScope[] = "__translate_spirv_memory_scope";
379379
const static char TranslateSPIRVMemFence[] = "__translate_spirv_memory_fence";
380+
const static char EntrypointPrefix[] = "__spirv_entry_";
380381
} // namespace kSPIRVName
381382

382383
namespace kSPIRVPostfix {

‎lib/SPIRV/SPIRVReader.cpp

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2768,6 +2768,24 @@ Function *SPIRVToLLVM::transFunction(SPIRVFunction *BF) {
27682768
return Loc->second;
27692769

27702770
auto IsKernel = isKernel(BF);
2771+
2772+
if (IsKernel) {
2773+
// search for a previous function with the same name
2774+
// upgrade it to a kernel and drop this if it's found
2775+
for (auto &I : FuncMap) {
2776+
auto BFName = I.getFirst()->getName();
2777+
if (BF->getName() == BFName) {
2778+
auto *F = I.getSecond();
2779+
F->setCallingConv(CallingConv::SPIR_KERNEL);
2780+
F->setLinkage(GlobalValue::ExternalLinkage);
2781+
F->setDSOLocal(false);
2782+
F = cast<Function>(mapValue(BF, F));
2783+
mapFunction(BF, F);
2784+
return F;
2785+
}
2786+
}
2787+
}
2788+
27712789
auto Linkage = IsKernel ? GlobalValue::ExternalLinkage : transLinkageType(BF);
27722790
FunctionType *FT = dyn_cast<FunctionType>(transType(BF->getFunctionType()));
27732791
std::string FuncName = BF->getName();

‎lib/SPIRV/SPIRVRegularizeLLVM.cpp

Lines changed: 70 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -39,6 +39,7 @@
3939

4040
#include "OCLUtil.h"
4141
#include "SPIRVInternal.h"
42+
#include "SPIRVMDWalker.h"
4243
#include "libSPIRV/SPIRVDebug.h"
4344

4445
#include "llvm/ADT/StringExtras.h" // llvm::isDigit
@@ -72,6 +73,11 @@ class SPIRVRegularizeLLVMBase {
7273
// Lower functions
7374
bool regularize();
7475

76+
// SPIR-V disallows functions being entrypoints and called
77+
// LLVM doesn't. This adds a wrapper around the entry point
78+
// that later SPIR-V writer renames.
79+
void addKernelEntryPoint(Module *M);
80+
7581
/// Erase cast inst of function and replace with the function.
7682
/// Assuming F is a SPIR-V builtin function with op code \param OC.
7783
void lowerFuncPtr(Function *F, Op OC);
@@ -437,6 +443,7 @@ bool SPIRVRegularizeLLVMBase::runRegularizeLLVM(Module &Module) {
437443
bool SPIRVRegularizeLLVMBase::regularize() {
438444
eraseUselessFunctions(M);
439445
lowerFuncPtr(M);
446+
addKernelEntryPoint(M);
440447

441448
for (auto I = M->begin(), E = M->end(); I != E;) {
442449
Function *F = &(*I++);
@@ -605,6 +612,69 @@ void SPIRVRegularizeLLVMBase::lowerFuncPtr(Module *M) {
605612
lowerFuncPtr(I.first, I.second);
606613
}
607614

615+
void SPIRVRegularizeLLVMBase::addKernelEntryPoint(Module *M) {
616+
std::vector<Function *> Work;
617+
618+
// Get a list of all functions that have SPIR kernel calling conv
619+
for (auto &F : *M) {
620+
if (F.getCallingConv() == CallingConv::SPIR_KERNEL)
621+
Work.push_back(&F);
622+
}
623+
for (auto &F : Work) {
624+
// for declarations just make them into SPIR functions.
625+
F->setCallingConv(CallingConv::SPIR_FUNC);
626+
if (F->isDeclaration())
627+
continue;
628+
629+
// Otherwise add a wrapper around the function to act as an entry point.
630+
FunctionType *FType = F->getFunctionType();
631+
std::string WrapName =
632+
kSPIRVName::EntrypointPrefix + static_cast<std::string>(F->getName());
633+
Function *WrapFn =
634+
getOrCreateFunction(M, F->getReturnType(), FType->params(), WrapName);
635+
636+
auto *CallBB = BasicBlock::Create(M->getContext(), "", WrapFn);
637+
IRBuilder<> Builder(CallBB);
638+
639+
Function::arg_iterator DestI = WrapFn->arg_begin();
640+
for (const Argument &I : F->args()) {
641+
DestI->setName(I.getName());
642+
DestI++;
643+
}
644+
SmallVector<Value *, 1> Args;
645+
for (Argument &I : WrapFn->args()) {
646+
Args.emplace_back(&I);
647+
}
648+
auto *CI = CallInst::Create(F, ArrayRef<Value *>(Args), "", CallBB);
649+
CI->setCallingConv(F->getCallingConv());
650+
CI->setAttributes(F->getAttributes());
651+
652+
// copy over all the metadata (should it be removed from F?)
653+
SmallVector<std::pair<unsigned, MDNode *>> MDs;
654+
F->getAllMetadata(MDs);
655+
WrapFn->setAttributes(F->getAttributes());
656+
for (auto MD = MDs.begin(), End = MDs.end(); MD != End; ++MD) {
657+
WrapFn->addMetadata(MD->first, *MD->second);
658+
}
659+
WrapFn->setCallingConv(CallingConv::SPIR_KERNEL);
660+
WrapFn->setLinkage(llvm::GlobalValue::InternalLinkage);
661+
662+
Builder.CreateRet(F->getReturnType()->isVoidTy() ? nullptr : CI);
663+
664+
// Have to find the spir-v metadata for execution mode and transfer it to
665+
// the wrapper.
666+
if (auto NMD = SPIRVMDWalker(*M).getNamedMD(kSPIRVMD::ExecutionMode)) {
667+
while (!NMD.atEnd()) {
668+
Function *MDF = nullptr;
669+
auto N = NMD.nextOp(); /* execution mode MDNode */
670+
N.get(MDF);
671+
if (MDF == F)
672+
N.M->replaceOperandWith(0, ValueAsMetadata::get(WrapFn));
673+
}
674+
}
675+
}
676+
}
677+
608678
} // namespace SPIRV
609679

610680
INITIALIZE_PASS(SPIRVRegularizeLLVMLegacy, "spvregular",

‎lib/SPIRV/SPIRVWriter.cpp

Lines changed: 13 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -638,8 +638,15 @@ SPIRVFunction *LLVMToSPIRVBase::transFunctionDecl(Function *F) {
638638
SPIRVFunction *BF =
639639
static_cast<SPIRVFunction *>(mapValue(F, BM->addFunction(BFT)));
640640
BF->setFunctionControlMask(transFunctionControlMask(F));
641-
if (F->hasName())
642-
BM->setName(BF, F->getName().str());
641+
if (F->hasName()) {
642+
if (isKernel(F)) {
643+
/* strip the prefix as the runtime will be looking for this name */
644+
std::string Prefix = kSPIRVName::EntrypointPrefix;
645+
std::string Name = F->getName().str();
646+
BM->setName(BF, Name.substr(Prefix.size()));
647+
} else
648+
BM->setName(BF, F->getName().str());
649+
}
643650
if (!isKernel(F) && F->getLinkage() != GlobalValue::InternalLinkage)
644651
BF->setLinkageType(transLinkageType(F));
645652

@@ -3735,7 +3742,7 @@ void LLVMToSPIRVBase::transFunction(Function *I) {
37353742

37363743
if (isKernel(I)) {
37373744
auto Interface = collectEntryPointInterfaces(BF, I);
3738-
BM->addEntryPoint(ExecutionModelKernel, BF->getId(), I->getName().str(),
3745+
BM->addEntryPoint(ExecutionModelKernel, BF->getId(), BF->getName(),
37393746
Interface);
37403747
}
37413748
}
@@ -4064,8 +4071,9 @@ bool LLVMToSPIRVBase::transMetadata() {
40644071
// Work around to translate kernel_arg_type and kernel_arg_type_qual metadata
40654072
static void transKernelArgTypeMD(SPIRVModule *BM, Function *F, MDNode *MD,
40664073
std::string MDName) {
4067-
std::string KernelArgTypesMDStr =
4068-
std::string(MDName) + "." + F->getName().str() + ".";
4074+
std::string Prefix = kSPIRVName::EntrypointPrefix;
4075+
std::string Name = F->getName().str().substr(Prefix.size());
4076+
std::string KernelArgTypesMDStr = std::string(MDName) + "." + Name + ".";
40694077
for (const auto &TyOp : MD->operands())
40704078
KernelArgTypesMDStr += cast<MDString>(TyOp)->getString().str() + ",";
40714079
BM->getString(KernelArgTypesMDStr);

‎test/entry_point_func.ll

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,21 @@
1+
;; Test to check that an LLVM spir_kernel gets translated into an
2+
;; Entrypoint wrapper and Function with LinkageAttributes
3+
; RUN: llvm-as %s -o %t.bc
4+
; RUN: llvm-spirv %t.bc -o - -spirv-text | FileCheck %s --check-prefix=CHECK-SPIRV
5+
; RUN: llvm-spirv %t.bc -o %t.spv
6+
; RUN: spirv-val %t.spv
7+
8+
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
9+
target triple = "spir64-unknown-unknown"
10+
11+
define spir_kernel void @testfunction() {
12+
ret void
13+
}
14+
15+
; Check there is an entrypoint and a function produced.
16+
; CHECK-SPIRV: EntryPoint 6 [[EP:[0-9]+]] "testfunction"
17+
; CHECK-SPIRV: Name [[FUNC:[0-9]+]] "testfunction"
18+
; CHECK-SPIRV: Decorate [[FUNC]] LinkageAttributes "testfunction" Export
19+
; CHECK-SPIRV: Function 2 [[FUNC]] 0 3
20+
; CHECK-SPIRV: Function 2 [[EP]] 0 3
21+
; CHECK-SPIRV: FunctionCall 2 8 [[FUNC]]

‎test/mem2reg.cl

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,10 +1,11 @@
11
// RUN: %clang_cc1 -O0 -S -triple spir-unknown-unknown -cl-std=CL2.0 -x cl -disable-O0-optnone %s -emit-llvm-bc -o %t.bc
22
// RUN: llvm-spirv -s %t.bc
3-
// RUN: llvm-dis < %t.bc | FileCheck %s --check-prefixes=CHECK,CHECK-WO
3+
// RUN: llvm-dis < %t.bc | FileCheck %s --check-prefixes=CHECK-WO
44
// RUN: llvm-spirv -s -spirv-mem2reg %t.bc -o %t.opt.bc
5-
// RUN: llvm-dis < %t.opt.bc | FileCheck %s --check-prefixes=CHECK,CHECK-W
6-
// CHECK-LABEL: spir_kernel void @foo
5+
// RUN: llvm-dis < %t.opt.bc | FileCheck %s --check-prefixes=CHECK-W
6+
// CHECK-W-LABEL: spir_func void @foo
77
// CHECK-W-NOT: alloca i32
8+
// CHECK-WO-LABEL: spir_kernel void @foo
89
// CHECK-WO: alloca i32
910
__kernel void foo(__global int *a) {
1011
*a = *a + 1;

‎test/transcoding/FPGAUnstructuredLoopAttr.ll

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -9,10 +9,10 @@
99
; CHECK-SPIRV: 2 Capability FPGALoopControlsINTEL
1010
; CHECK-SPIRV: 9 Extension "SPV_INTEL_fpga_loop_controls"
1111
; CHECK-SPIRV: 11 Extension "SPV_INTEL_unstructured_loop_controls"
12-
; CHECK-SPIRV: 4 EntryPoint 6 [[FOO:[0-9]+]] "foo"
13-
; CHECK-SPIRV: 4 EntryPoint 6 [[BOO:[0-9]+]] "boo"
12+
; CHECK-SPIRV: 3 Name [[FOO:[0-9]+]] "foo"
1413
; CHECK-SPIRV: 4 Name [[ENTRY_1:[0-9]+]] "entry"
1514
; CHECK-SPIRV: 5 Name [[FOR:[0-9]+]] "for.cond"
15+
; CHECK-SPIRV: 3 Name [[BOO:[0-9]+]] "boo"
1616
; CHECK-SPIRV: 4 Name [[ENTRY_2:[0-9]+]] "entry"
1717
; CHECK-SPIRV: 5 Name [[WHILE:[0-9]+]] "while.body"
1818

‎test/transcoding/KernelArgTypeInOpString.ll

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -39,8 +39,8 @@
3939
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"
4040
target triple = "spir-unknown-unknown"
4141

42-
; CHECK-SPIRV-WORKAROUND: String 14 "kernel_arg_type.foo.image_kernel_data*,myInt,struct struct_name*,"
43-
; CHECK-SPIRV-WORKAROUND-NEGATIVE-NOT: String 14 "kernel_arg_type.foo.image_kernel_data*,myInt,struct struct_name*,"
42+
; CHECK-SPIRV-WORKAROUND: String 20 "kernel_arg_type.foo.image_kernel_data*,myInt,struct struct_name*,"
43+
; CHECK-SPIRV-WORKAROUND-NEGATIVE-NOT: String 20 "kernel_arg_type.foo.image_kernel_data*,myInt,struct struct_name*,"
4444

4545
; CHECK-LLVM-WORKAROUND: !kernel_arg_type [[TYPE:![0-9]+]]
4646
; CHECK-LLVM-WORKAROUND: [[TYPE]] = !{!"image_kernel_data*", !"myInt", !"struct struct_name*"}

‎test/transcoding/KernelArgTypeInOpString2.ll

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -41,8 +41,8 @@
4141
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"
4242
target triple = "spir"
4343

44-
; CHECK-SPIRV-WORKAROUND: String 17 "kernel_arg_type.foo.cl::tt::vec<float, 4>*,"
45-
; CHECK-SPIRV-WORKAROUND-NEGATIVE-NOT: String 17 "kernel_arg_type.foo.cl::tt::vec<float, 4>*,"
44+
; CHECK-SPIRV-WORKAROUND: String 21 "kernel_arg_type.foo.cl::tt::vec<float, 4>*,"
45+
; CHECK-SPIRV-WORKAROUND-NEGATIVE-NOT: String 21 "kernel_arg_type.foo.cl::tt::vec<float, 4>*,"
4646

4747
; CHECK-LLVM-WORKAROUND: !kernel_arg_type [[TYPE:![0-9]+]]
4848
; CHECK-LLVM-WORKAROUND: [[TYPE]] = !{!"cl::tt::vec<float, 4>*"}

‎test/transcoding/OpenCL/atomic_cmpxchg.cl

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,7 @@ __kernel void test_atomic_cmpxchg(__global int *p, int cmp, int val) {
1717
atomic_cmpxchg(up, ucmp, uval);
1818
}
1919

20-
// CHECK-SPIRV: EntryPoint {{[0-9]+}} [[TEST:[0-9]+]] "test_atomic_cmpxchg"
20+
// CHECK-SPIRV: Name [[TEST:[0-9]+]] "test_atomic_cmpxchg"
2121
// CHECK-SPIRV-DAG: TypeInt [[UINT:[0-9]+]] 32 0
2222
// CHECK-SPIRV-DAG: TypePointer [[UINT_PTR:[0-9]+]] 5 [[UINT]]
2323
//

‎test/transcoding/OpenCL/atomic_legacy.cl

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -13,7 +13,7 @@ __kernel void test_legacy_atomics(__global int *p, int val) {
1313
atomic_add(p, val); // from OpenCL C 1.1
1414
}
1515

16-
// CHECK-SPIRV: EntryPoint {{[0-9]+}} [[TEST:[0-9]+]] "test_legacy_atomics"
16+
// CHECK-SPIRV: Name [[TEST:[0-9]+]] "test_legacy_atomics"
1717
// CHECK-SPIRV-DAG: TypeInt [[UINT:[0-9]+]] 32 0
1818
// CHECK-SPIRV-DAG: TypePointer [[UINT_PTR:[0-9]+]] 5 [[UINT]]
1919
//

‎test/transcoding/OpenCL/atomic_work_item_fence.cl

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -23,7 +23,7 @@ __kernel void test_mem_fence_non_const_flags(cl_mem_fence_flags flags, memory_or
2323
// atomic_work_item_fence(flags, order, scope);
2424
}
2525

26-
// CHECK-SPIRV: EntryPoint {{[0-9]+}} [[TEST_CONST_FLAGS:[0-9]+]] "test_mem_fence_const_flags"
26+
// CHECK-SPIRV: Name [[TEST_CONST_FLAGS:[0-9]+]] "test_mem_fence_const_flags"
2727
// CHECK-SPIRV: TypeInt [[UINT:[0-9]+]] 32 0
2828
//
2929
// 0x0 Relaxed + 0x100 WorkgroupMemory

‎test/transcoding/OpenCL/barrier.cl

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -28,7 +28,7 @@ __kernel void test_barrier_non_const_flags(cl_mem_fence_flags flags) {
2828
// barrier(flags);
2929
}
3030

31-
// CHECK-SPIRV: EntryPoint {{[0-9]+}} [[TEST_CONST_FLAGS:[0-9]+]] "test_barrier_const_flags"
31+
// CHECK-SPIRV: Name [[TEST_CONST_FLAGS:[0-9]+]] "test_barrier_const_flags"
3232
// CHECK-SPIRV: TypeInt [[UINT:[0-9]+]] 32 0
3333
//
3434
// In SPIR-V, barrier is represented as OpControlBarrier [3] and OpenCL

‎test/transcoding/OpenCL/mem_fence.cl

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -34,7 +34,7 @@ __kernel void test_mem_fence_non_const_flags(cl_mem_fence_flags flags) {
3434
// mem_fence(flags);
3535
}
3636

37-
// CHECK-SPIRV: EntryPoint {{[0-9]+}} [[TEST_CONST_FLAGS:[0-9]+]] "test_mem_fence_const_flags"
37+
// CHECK-SPIRV: Name [[TEST_CONST_FLAGS:[0-9]+]] "test_mem_fence_const_flags"
3838
// CHECK-SPIRV: TypeInt [[UINT:[0-9]+]] 32 0
3939
//
4040
// In SPIR-V, mem_fence is represented as OpMemoryBarrier [2] and OpenCL

‎test/transcoding/OpenCL/sub_group_barrier.cl

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -31,7 +31,7 @@ __kernel void test_barrier_non_const_flags(cl_mem_fence_flags flags, memory_scop
3131
// sub_group_barrier(flags, scope);
3232
}
3333

34-
// CHECK-SPIRV: EntryPoint {{[0-9]+}} [[TEST_CONST_FLAGS:[0-9]+]] "test_barrier_const_flags"
34+
// CHECK-SPIRV: Name [[TEST_CONST_FLAGS:[0-9]+]] "test_barrier_const_flags"
3535
// CHECK-SPIRV: TypeInt [[UINT:[0-9]+]] 32 0
3636
//
3737
// In SPIR-V, barrier is represented as OpControlBarrier [2] and OpenCL

‎test/transcoding/OpenCL/work_group_barrier.cl

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -33,7 +33,7 @@ __kernel void test_barrier_non_const_flags(cl_mem_fence_flags flags, memory_scop
3333
// work_group_barrier(flags, scope);
3434
}
3535

36-
// CHECK-SPIRV: EntryPoint {{[0-9]+}} [[TEST_CONST_FLAGS:[0-9]+]] "test_barrier_const_flags"
36+
// CHECK-SPIRV: Name [[TEST_CONST_FLAGS:[0-9]+]] "test_barrier_const_flags"
3737
// CHECK-SPIRV: TypeInt [[UINT:[0-9]+]] 32 0
3838
//
3939
// In SPIR-V, barrier is represented as OpControlBarrier [2] and OpenCL

‎test/transcoding/SPV_INTEL_function_pointers/fp-from-host.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,7 @@
1717
; CHECK-SPIRV: Capability FunctionPointersINTEL
1818
; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers"
1919
;
20-
; CHECK-SPIRV: EntryPoint {{[0-9]+}} [[KERNEL_ID:[0-9]+]] "test"
20+
; CHECK-SPIRV: Name [[KERNEL_ID:[0-9]+]] "test"
2121
; CHECK-SPIRV: TypeInt [[INT32_TYPE_ID:[0-9]+]] 32
2222
; CHECK-SPIRV: TypePointer [[INT_PTR:[0-9]+]] 5 [[INT32_TYPE_ID]]
2323
; CHECK-SPIRV: TypeFunction [[FOO_TYPE_ID:[0-9]+]] [[INT32_TYPE_ID]] [[INT32_TYPE_ID]]

‎test/transcoding/SPV_INTEL_function_pointers/function-pointer-as-function-arg.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -33,7 +33,7 @@
3333
; CHECK-SPIRV: Capability FunctionPointersINTEL
3434
; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers"
3535
;
36-
; CHECK-SPIRV: EntryPoint 6 [[KERNEL_ID:[0-9]+]] "test"
36+
; CHECK-SPIRV: Name [[KERNEL_ID:[0-9]+]] "test"
3737
; CHECK-SPIRV: TypeInt [[TYPE_INT32_ID:[0-9]+]] 32
3838
; CHECK-SPIRV: TypeFunction [[FOO_TYPE_ID:[0-9]+]] [[TYPE_INT32_ID]] [[TYPE_INT32_ID]]
3939
; CHECK-SPIRV: TypePointer [[FOO_PTR_TYPE_ID:[0-9]+]] {{[0-9]+}} [[FOO_TYPE_ID]]

‎test/transcoding/SPV_INTEL_function_pointers/function-pointer.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,7 @@
1919
;
2020
; CHECK-SPIRV: Capability FunctionPointersINTEL
2121
; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers"
22-
; CHECK-SPIRV: EntryPoint 6 [[KERNEL_ID:[0-9]+]] "test"
22+
; CHECK-SPIRV: Name [[KERNEL_ID:[0-9]+]] "test"
2323
; CHECK-SPIRV: TypeInt [[TYPE_INT_ID:[0-9]+]]
2424
; CHECK-SPIRV: TypeFunction [[FOO_TYPE_ID:[0-9]+]] [[TYPE_INT_ID]] [[TYPE_INT_ID]]
2525
; CHECK-SPIRV: TypePointer [[FOO_PTR_ID:[0-9]+]] {{[0-9]+}} [[FOO_TYPE_ID]]

‎test/transcoding/SPV_INTEL_function_pointers/non-uniform-function-pointer.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -29,7 +29,7 @@
2929
; CHECK-SPIRV: Capability FunctionPointersINTEL
3030
; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers"
3131
;
32-
; CHECK-SPIRV: EntryPoint 6 [[KERNEL_ID:[0-9]+]] "test"
32+
; CHECK-SPIRV: Name [[KERNEL_ID:[0-9]+]] "test"
3333
; CHECK-SPIRV: TypeInt [[TYPE_INT32_ID:[0-9+]]] 32
3434
; CHECK-SPIRV: TypeFunction [[FOO_TYPE_ID:[0-9]+]] [[TYPE_INT32_ID]] [[TYPE_INT32_ID]]
3535
; CHECK-SPIRV: TypePointer [[FOO_PTR_TYPE_ID:[0-9]+]] {{[0-9]+}} [[FOO_TYPE_ID]]

‎test/transcoding/SPV_INTEL_function_pointers/select.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -6,7 +6,7 @@
66
; RUN: llvm-dis %t.r.bc -o %t.r.ll
77
; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM
88

9-
; CHECK-SPIRV: EntryPoint 6 [[#KERNEL_ID:]] "_ZTS6kernel"
9+
; CHECK-SPIRV: Name [[#KERNEL_ID:]] "_ZTS6kernel"
1010
; CHECK-SPIRV-DAG: Name [[#BAR:]] "_Z3barii"
1111
; CHECK-SPIRV-DAG: Name [[#BAZ:]] "_Z3bazii"
1212
; CHECK-SPIRV: TypeInt [[#INT32:]] 32

‎test/transcoding/SPV_INTEL_joint_matrix/joint_matrix.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -13,7 +13,7 @@
1313

1414
; CHECK-SPIRV: Capability JointMatrixINTEL
1515
; CHECK-SPIRV: Extension "SPV_INTEL_joint_matrix"
16-
; CHECK-SPIRV: EntryPoint 6 [[#Kernel:]]
16+
; CHECK-SPIRV: Name [[#Kernel:]] "_ZTSZ4mainE11matrix_test"
1717

1818
; CHECK-SPIRV-DAG: TypeInt [[#ShortTy:]] 16 0
1919
; CHECK-SPIRV-DAG: TypeInt [[#CharTy:]] 8 0

‎test/transcoding/SampledImage.cl

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -27,8 +27,8 @@ void sample_kernel_int(image2d_t input, float2 coords, global int4 *results, sam
2727
}
2828

2929
// CHECK-SPIRV: Capability LiteralSampler
30-
// CHECK-SPIRV: EntryPoint 6 [[sample_kernel_float:[0-9]+]] "sample_kernel_float"
31-
// CHECK-SPIRV: EntryPoint 6 [[sample_kernel_int:[0-9]+]] "sample_kernel_int"
30+
// CHECK-SPIRV: Name [[sample_kernel_float:[0-9]+]] "sample_kernel_float"
31+
// CHECK-SPIRV: Name [[sample_kernel_int:[0-9]+]] "sample_kernel_int"
3232

3333
// CHECK-SPIRV: TypeSampler [[TypeSampler:[0-9]+]]
3434
// CHECK-SPIRV: TypeSampledImage [[SampledImageTy:[0-9]+]]
@@ -81,4 +81,4 @@ void sample_kernel_int(image2d_t input, float2 coords, global int4 *results, sam
8181
// CHECK-SPIRV: ImageSampleExplicitLod {{.*}} [[SampledImage6]]
8282
// CHECK-LLVM: call spir_func <4 x i32> @_Z11read_imagei14ocl_image2d_ro11ocl_samplerDv2_f(%opencl.image2d_ro_t addrspace(1)* %input, %opencl.sampler_t addrspace(2)* %1, <2 x float> %coords)
8383
// CHECK-SPV-IR: call spir_func %spirv.SampledImage._void_1_0_0_0_0_0_0 addrspace(1)* @_Z20__spirv_SampledImagePU3AS133__spirv_Image__void_1_0_0_0_0_0_0PU3AS215__spirv_Sampler(%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %input, %spirv.Sampler addrspace(2)* %1)
84-
// CHECK-SPV-IR: call spir_func <4 x i32> @_Z36__spirv_ImageSampleExplicitLod_Rint4PU3AS140__spirv_SampledImage__void_1_0_0_0_0_0_0Dv2_fif(%spirv.SampledImage._void_1_0_0_0_0_0_0 addrspace(1)* %TempSampledImage6, <2 x float> %coords, i32 2, float 0.000000e+00)
84+
// CHECK-SPV-IR: call spir_func <4 x i32> @_Z36__spirv_ImageSampleExplicitLod_Rint4PU3AS140__spirv_SampledImage__void_1_0_0_0_0_0_0Dv2_fif(%spirv.SampledImage._void_1_0_0_0_0_0_0 addrspace(1)* %TempSampledImage6, <2 x float> %coords, i32 2, float 0.000000e+00)

‎test/transcoding/kernel_arg_type_qual.ll

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -13,8 +13,8 @@ source_filename = "test.cl"
1313
target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
1414
target triple = "spir64-unknown-unknown."
1515

16-
; CHECK-SPIRV: String 12 "kernel_arg_type_qual.test.volatile,const,,"
17-
; CHECK-SPIRV: Name [[ARG:[0-9]+]] "g"
16+
; CHECK-SPIRV: String 18 "kernel_arg_type_qual.test.volatile,const,,"
17+
; CHECK-SPIRV: Name [[ARG:1[0-9]+]] "g"
1818
; CHECK-SPIRV: Decorate [[ARG]] Volatile
1919
; CHECK-SPIRV-NEGATIVE-NOT: String 12 "kernel_arg_type_qual.test.volatile,const,,"
2020

0 commit comments

Comments
 (0)