From 5b4b5aaa04c7238948f46011d4d2b73b687d37e2 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Fri, 27 Sep 2024 10:49:56 +0200 Subject: [PATCH] [Backport to 15] Introduce CodeSectionINTEL storage class (#2728) This storage class is used for function pointers. It's added as based on cl_intel_function_pointers specification, it is not guaranteed that sizeof(void(*)(void) == sizeof(void *) - to allow consumers use this fact, we cannot say that function pointer belongs to the same storage class as data pointers. It wasn't added during initial implementation, now it's time to fill this gap. As it would be a breaking change its generation is added only under -spirv-emit-function-ptr-addr-space option. Also SPIR-V consumer may pass this option during reverse translation to get new address space even in a case, when OpConstantFunctionPointerINTEL doesn't reside in CodeSectionINTEL storage class. Expected behavior: No option is passed to the forward translation stage and function pointers are in addrspace(9): no CodeSectionINTEL storage class in SPIR-V The option is passed to the forward translation stage and function pointers are in addrepace(9): CodeSectionINTEL storage class is generated No option is passed to the reverse translation stage: function pointers are in private address space The option is passed to the reverse translation stage: function pointers are in addrspace(9) Spec: https://github.com/intel/llvm/blob/sycl/sycl/doc/design/spirv-extensions/SPV_INTEL_function_pointers.asciidoc The previous approach: #1392 --- include/LLVMSPIRVOpts.h | 12 ++ lib/SPIRV/SPIRVInternal.h | 4 + lib/SPIRV/SPIRVReader.cpp | 57 ++++-- lib/SPIRV/SPIRVReader.h | 2 +- lib/SPIRV/SPIRVWriter.cpp | 9 + lib/SPIRV/libSPIRV/SPIRVInstruction.cpp | 5 + lib/SPIRV/libSPIRV/SPIRVModule.cpp | 7 + lib/SPIRV/libSPIRV/SPIRVModule.h | 5 + .../CodeSectionINTEL/bitcast.ll | 51 +++++ .../const-function-pointer.ll | 67 +++++++ .../decor-func-ptr-arg-attr.ll | 67 +++++++ .../CodeSectionINTEL/fp-from-host.ll | 70 +++++++ .../function-pointer-as-function-arg.ll | 177 ++++++++++++++++++ .../function-pointer-dedicated-as.ll | 107 +++++++++++ .../CodeSectionINTEL/function-pointer.ll | 92 +++++++++ .../CodeSectionINTEL/gv-func-ptr.ll | 40 ++++ .../non-uniform-function-pointer.ll | 139 ++++++++++++++ .../CodeSectionINTEL/referenced-indirectly.ll | 82 ++++++++ .../CodeSectionINTEL/select.ll | 147 +++++++++++++++ tools/llvm-spirv/llvm-spirv.cpp | 7 + 20 files changed, 1135 insertions(+), 12 deletions(-) create mode 100644 test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/bitcast.ll create mode 100644 test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/const-function-pointer.ll create mode 100644 test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/decor-func-ptr-arg-attr.ll create mode 100644 test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/fp-from-host.ll create mode 100644 test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer-as-function-arg.ll create mode 100644 test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer-dedicated-as.ll create mode 100644 test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer.ll create mode 100644 test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/gv-func-ptr.ll create mode 100644 test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/non-uniform-function-pointer.ll create mode 100644 test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/referenced-indirectly.ll create mode 100644 test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/select.ll diff --git a/include/LLVMSPIRVOpts.h b/include/LLVMSPIRVOpts.h index c956eb7ab4..bd1c73f318 100644 --- a/include/LLVMSPIRVOpts.h +++ b/include/LLVMSPIRVOpts.h @@ -221,6 +221,14 @@ class TranslatorOpts { PreserveOCLKernelArgTypeMetadataThroughString = Value; } + bool shouldEmitFunctionPtrAddrSpace() const noexcept { + return EmitFunctionPtrAddrSpace; + } + + void setEmitFunctionPtrAddrSpace(bool Value) noexcept { + EmitFunctionPtrAddrSpace = Value; + } + private: // Common translation options VersionNumber MaxVersion = VersionNumber::MaximumVersion; @@ -262,6 +270,10 @@ class TranslatorOpts { // kernel_arg_type_qual metadata through OpString bool PreserveOCLKernelArgTypeMetadataThroughString = false; + // Controls if CodeSectionINTEL can be emitted and consumed with a dedicated + // address space + bool EmitFunctionPtrAddrSpace = false; + bool PreserveAuxData = false; }; diff --git a/lib/SPIRV/SPIRVInternal.h b/lib/SPIRV/SPIRVInternal.h index f0acfce59c..4c78d27b66 100644 --- a/lib/SPIRV/SPIRVInternal.h +++ b/lib/SPIRV/SPIRVInternal.h @@ -193,6 +193,7 @@ enum SPIRAddressSpace { SPIRAS_GlobalHost, SPIRAS_Input, SPIRAS_Output, + SPIRAS_CodeSectionINTEL, SPIRAS_Count, }; @@ -203,6 +204,8 @@ template <> inline void SPIRVMap::init() { add(SPIRAS_Local, "Local"); add(SPIRAS_Generic, "Generic"); add(SPIRAS_Input, "Input"); + add(SPIRAS_CodeSectionINTEL, "CodeSectionINTEL"); + add(SPIRAS_GlobalDevice, "GlobalDevice"); add(SPIRAS_GlobalHost, "GlobalHost"); } @@ -219,6 +222,7 @@ inline void SPIRVMap::init() { add(SPIRAS_Input, StorageClassInput); add(SPIRAS_GlobalDevice, StorageClassDeviceOnlyINTEL); add(SPIRAS_GlobalHost, StorageClassHostOnlyINTEL); + add(SPIRAS_CodeSectionINTEL, StorageClassCodeSectionINTEL); } typedef SPIRVMap SPIRSPIRVAddrSpaceMap; diff --git a/lib/SPIRV/SPIRVReader.cpp b/lib/SPIRV/SPIRVReader.cpp index 011ac07974..47315f77cb 100644 --- a/lib/SPIRV/SPIRVReader.cpp +++ b/lib/SPIRV/SPIRVReader.cpp @@ -381,11 +381,17 @@ Type *SPIRVToLLVM::transType(SPIRVType *T, bool IsClassMember) { } case internal::OpTypeTokenINTEL: return mapType(T, Type::getTokenTy(*Context)); - case OpTypePointer: + case OpTypePointer: { + unsigned AS = SPIRSPIRVAddrSpaceMap::rmap(T->getPointerStorageClass()); + if (AS == SPIRAS_CodeSectionINTEL && !BM->shouldEmitFunctionPtrAddrSpace()) + AS = SPIRAS_Private; + if (BM->shouldEmitFunctionPtrAddrSpace() && + T->getPointerElementType()->getOpCode() == OpTypeFunction) + AS = SPIRAS_CodeSectionINTEL; return mapType( T, PointerType::get( - transType(T->getPointerElementType(), IsClassMember), - SPIRSPIRVAddrSpaceMap::rmap(T->getPointerStorageClass()))); + transType(T->getPointerElementType(), IsClassMember), AS)); + } case OpTypeVector: return mapType(T, FixedVectorType::get(transType(T->getVectorComponentType()), @@ -1564,10 +1570,20 @@ Value *SPIRVToLLVM::transValueWithoutDecoration(SPIRVValue *BV, Function *F, case OpTypeVector: return mapValue(BV, ConstantVector::get(CV)); case OpTypeMatrix: - case OpTypeArray: - return mapValue( - BV, ConstantArray::get(dyn_cast(transType(BCC->getType())), - CV)); + case OpTypeArray: { + auto *AT = cast(transType(BCC->getType())); + for (size_t I = 0; I != AT->getNumElements(); ++I) { + auto *ElemTy = AT->getElementType(); + if (auto *ElemPtrTy = dyn_cast(ElemTy)) { + assert(isa(CV[I]->getType()) && + "Constant type doesn't match constexpr array element type"); + if (ElemPtrTy->getAddressSpace() != + cast(CV[I]->getType())->getAddressSpace()) + CV[I] = ConstantExpr::getAddrSpaceCast(CV[I], AT->getElementType()); + } + } + return mapValue(BV, ConstantArray::get(AT, CV)); + } case OpTypeStruct: { auto BCCTy = dyn_cast(transType(BCC->getType())); auto Members = BCCTy->getNumElements(); @@ -1582,7 +1598,12 @@ Value *SPIRVToLLVM::transValueWithoutDecoration(SPIRVValue *BV, Function *F, !BCCTy->getElementType(I)->isPointerTy()) continue; - CV[I] = ConstantExpr::getBitCast(CV[I], BCCTy->getElementType(I)); + if (cast(CV[I]->getType())->getAddressSpace() != + cast(BCCTy->getElementType(I))->getAddressSpace()) + CV[I] = + ConstantExpr::getAddrSpaceCast(CV[I], BCCTy->getElementType(I)); + else + CV[I] = ConstantExpr::getBitCast(CV[I], BCCTy->getElementType(I)); } } @@ -1620,7 +1641,10 @@ Value *SPIRVToLLVM::transValueWithoutDecoration(SPIRVValue *BV, Function *F, static_cast(BV); SPIRVFunction *F = BC->getFunction(); BV->setName(F->getName()); - return mapValue(BV, transFunction(F)); + const unsigned AS = BM->shouldEmitFunctionPtrAddrSpace() + ? SPIRAS_CodeSectionINTEL + : SPIRAS_Private; + return mapValue(BV, transFunction(F, AS)); } case OpUndef: @@ -3038,7 +3062,7 @@ void SPIRVToLLVM::transFunctionAttrs(SPIRVFunction *BF, Function *F) { }); } -Function *SPIRVToLLVM::transFunction(SPIRVFunction *BF) { +Function *SPIRVToLLVM::transFunction(SPIRVFunction *BF, unsigned AS) { auto Loc = FuncMap.find(BF); if (Loc != FuncMap.end()) return Loc->second; @@ -3087,7 +3111,7 @@ Function *SPIRVToLLVM::transFunction(SPIRVFunction *BF) { } Function *F = M->getFunction(FuncName); if (!F) - F = Function::Create(FT, Linkage, FuncName, M); + F = Function::Create(FT, Linkage, AS, FuncName, M); F = cast(mapValue(BF, F)); mapFunction(BF, F); @@ -3486,6 +3510,17 @@ bool SPIRVToLLVM::translate() { DbgTran->transDebugInst(EI); } + for (auto *FP : BM->getFunctionPointers()) { + SPIRVConstantFunctionPointerINTEL *BC = + static_cast(FP); + SPIRVFunction *F = BC->getFunction(); + FP->setName(F->getName()); + const unsigned AS = BM->shouldEmitFunctionPtrAddrSpace() + ? SPIRAS_CodeSectionINTEL + : SPIRAS_Private; + mapValue(FP, transFunction(F, AS)); + } + for (unsigned I = 0, E = BM->getNumFunctions(); I != E; ++I) { transFunction(BM->getFunction(I)); transUserSemantic(BM->getFunction(I)); diff --git a/lib/SPIRV/SPIRVReader.h b/lib/SPIRV/SPIRVReader.h index f1de7c6f52..ef4b6c3e1b 100644 --- a/lib/SPIRV/SPIRVReader.h +++ b/lib/SPIRV/SPIRVReader.h @@ -98,7 +98,7 @@ class SPIRVToLLVM { void transAuxDataInst(SPIRVExtInst *BC); std::vector transValue(const std::vector &, Function *F, BasicBlock *); - Function *transFunction(SPIRVFunction *F); + Function *transFunction(SPIRVFunction *F, unsigned AS = SPIRAS_Private); void transFunctionAttrs(SPIRVFunction *BF, Function *F); Value *transBlockInvoke(SPIRVValue *Invoke, BasicBlock *BB); Instruction *transWGSizeQueryBI(SPIRVInstruction *BI, BasicBlock *BB); diff --git a/lib/SPIRV/SPIRVWriter.cpp b/lib/SPIRV/SPIRVWriter.cpp index 54af00790d..c51ab56d9e 100644 --- a/lib/SPIRV/SPIRVWriter.cpp +++ b/lib/SPIRV/SPIRVWriter.cpp @@ -508,6 +508,11 @@ SPIRVType *LLVMToSPIRVBase::transPointerType(Type *ET, unsigned AddrSpc) { ((AddrSpc == SPIRAS_GlobalDevice) || (AddrSpc == SPIRAS_GlobalHost))) { return transPointerType(ET, SPIRAS_Global); } + // Lower function pointer address space to private if + // spirv-emit-function-ptr-addr-space is not passed + if (AddrSpc == SPIRAS_CodeSectionINTEL && + !BM->shouldEmitFunctionPtrAddrSpace()) + return transPointerType(ET, SPIRAS_Private); if (ST && !ST->isSized()) { Op OpCode; StringRef STName = ST->getName(); @@ -615,6 +620,10 @@ SPIRVType *LLVMToSPIRVBase::transPointerType(SPIRVType *ET, unsigned AddrSpc) { if (Loc != PointeeTypeMap.end()) return Loc->second; + if (AddrSpc == SPIRAS_CodeSectionINTEL && + !BM->shouldEmitFunctionPtrAddrSpace()) + return transPointerType(ET, SPIRAS_Private); + SPIRVType *TranslatedTy = BM->addPointerType( SPIRSPIRVAddrSpaceMap::map(static_cast(AddrSpc)), ET); PointeeTypeMap[TypeKey] = TranslatedTy; diff --git a/lib/SPIRV/libSPIRV/SPIRVInstruction.cpp b/lib/SPIRV/libSPIRV/SPIRVInstruction.cpp index 4612b4275c..61ad987baf 100644 --- a/lib/SPIRV/libSPIRV/SPIRVInstruction.cpp +++ b/lib/SPIRV/libSPIRV/SPIRVInstruction.cpp @@ -286,6 +286,11 @@ SPIRVInstruction *createInstFromSpecConstantOp(SPIRVSpecConstantOp *Inst) { auto OC = static_cast(Ops[0]); assert(isSpecConstantOpAllowedOp(OC) && "Op code not allowed for OpSpecConstantOp"); + auto *Const = Inst->getOperand(1); + // LLVM would eliminate a bitcast from a function pointer in a constexpr + // context. Cut this short here to avoid necessity to align address spaces + if (OC == OpBitcast && Const->getOpCode() == OpConstantFunctionPointerINTEL) + return static_cast(Const); Ops.erase(Ops.begin(), Ops.begin() + 1); auto *BM = Inst->getModule(); auto *RetInst = SPIRVInstTemplateBase::create( diff --git a/lib/SPIRV/libSPIRV/SPIRVModule.cpp b/lib/SPIRV/libSPIRV/SPIRVModule.cpp index 6829200d7c..1fc6ef7461 100644 --- a/lib/SPIRV/libSPIRV/SPIRVModule.cpp +++ b/lib/SPIRV/libSPIRV/SPIRVModule.cpp @@ -168,6 +168,13 @@ class SPIRVModuleImpl : public SPIRVModule { SPIRVConstant *getLiteralAsConstant(unsigned Literal) override; unsigned getNumFunctions() const override { return FuncVec.size(); } unsigned getNumVariables() const override { return VariableVec.size(); } + std::vector getFunctionPointers() const override { + std::vector Res; + for (auto *C : ConstVec) + if (C->getOpCode() == OpConstantFunctionPointerINTEL) + Res.emplace_back(C); + return Res; + } SourceLanguage getSourceLanguage(SPIRVWord *Ver = nullptr) const override { if (Ver) *Ver = SrcLangVer; diff --git a/lib/SPIRV/libSPIRV/SPIRVModule.h b/lib/SPIRV/libSPIRV/SPIRVModule.h index 6e69644473..41e18e5b31 100644 --- a/lib/SPIRV/libSPIRV/SPIRVModule.h +++ b/lib/SPIRV/libSPIRV/SPIRVModule.h @@ -140,6 +140,7 @@ class SPIRVModule { virtual SPIRVMemoryModelKind getMemoryModel() const = 0; virtual unsigned getNumFunctions() const = 0; virtual unsigned getNumVariables() const = 0; + virtual std::vector getFunctionPointers() const = 0; virtual SourceLanguage getSourceLanguage(SPIRVWord *) const = 0; virtual std::set &getSourceExtension() = 0; virtual SPIRVValue *getValue(SPIRVId TheId) const = 0; @@ -539,6 +540,10 @@ class SPIRVModule { .shouldPreserveOCLKernelArgTypeMetadataThroughString(); } + bool shouldEmitFunctionPtrAddrSpace() const noexcept { + return TranslationOpts.shouldEmitFunctionPtrAddrSpace(); + } + bool preserveAuxData() const noexcept { return TranslationOpts.preserveAuxData(); } diff --git a/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/bitcast.ll b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/bitcast.ll new file mode 100644 index 0000000000..7293f48c71 --- /dev/null +++ b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/bitcast.ll @@ -0,0 +1,51 @@ +; OpenCL C source: +; char foo(char a) { +; return a; +; } +; void bar() { +; int (*fun_ptr)(int) = &foo; +; } + +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv %t.bc -spirv-ext=+SPV_INTEL_function_pointers -o %t.spv +; RUN: llvm-spirv %t.spv -to-text -o %t.spt +; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV +; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.r.bc +; RUN: llvm-dis %t.r.bc -o %t.r.ll +; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM + +; CHECK-SPIRV: TypeFunction [[#FOO_TY:]] [[#]] [[#]] +; CHECK-SPIRV: TypeFunction [[#DEST_TY:]] [[#]] [[#]] +; CHECK-SPIRV: TypePointer [[#DEST_TY_PTR:]] [[#]] [[#DEST_TY]] +; CHECK-SPIRV: TypePointer [[#FOO_TY_PTR:]] [[#]] [[#FOO_TY]] +; CHECK-SPIRV: ConstantFunctionPointerINTEL [[#FOO_TY_PTR]] [[#FOO_PTR:]] [[#FOO:]] +; CHECK-SPIRV: Function [[#]] [[#FOO]] [[#]] [[#FOO_TY]] + +; CHECK-SPIRV: Bitcast [[#DEST_TY_PTR]] [[#]] [[#FOO_PTR]] + +; CHECK-LLVM: bitcast i8 (i8) addrspace(9)* @foo to i32 (i32) addrspace(9)* + +; ModuleID = './example.c' +source_filename = "./example.c" +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" +target triple = "spir" + +; Function Attrs: noinline nounwind optnone +define dso_local spir_func signext i8 @foo(i8 signext %0) #0 { + ret i8 %0 +} + +; Function Attrs: noinline nounwind optnone +define dso_local spir_func void @bar() #0 { + %1 = alloca i32 (i32)*, align 4 + store i32 (i32)* bitcast (i8 (i8)* @foo to i32 (i32)*), i32 (i32)** %1, align 4 + ret void +} + +attributes #0 = { noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } + +!llvm.module.flags = !{!0} +!llvm.ident = !{!1} + +!0 = !{i32 1, !"wchar_size", i32 4} +!1 = !{!"clang version 11.0.0 (https://github.com/llvm/llvm-project.git 0e1accd0f726eef2c47be9f37dd0a06cb50d207e)"} diff --git a/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/const-function-pointer.ll b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/const-function-pointer.ll new file mode 100644 index 0000000000..c937ae11b3 --- /dev/null +++ b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/const-function-pointer.ll @@ -0,0 +1,67 @@ +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -o %t.spv +; RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV +; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.r.bc +; RUN: llvm-dis %t.r.bc -o %t.r.ll +; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM + +; CHECK-SPIRV: Capability FunctionPointersINTEL +; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers" +; CHECK-SPIRV: Name [[F1Name:[0-9]+]] "f1" +; CHECK-SPIRV: Name [[F2Name:[0-9]+]] "f2" +; CHECK-SPIRV: TypeInt [[Int32:[0-9]+]] 32 +; CHECK-SPIRV: TypeInt [[Int64:[0-9]+]] 64 +; CHECK-SPIRV-DAG: Constant [[Int32]] [[XArg:[0-9]+]] 32 +; CHECK-SPIRV-DAG: Constant [[Int32]] [[YArg:[0-9]+]] 2 + +; CHECK-SPIRV: ConstantFunctionPointerINTEL {{[0-9]+}} [[F1:[0-9]+]] [[F1Name]] +; CHECK-SPIRV: ConstantFunctionPointerINTEL {{[0-9]+}} [[F2:[0-9]+]] [[F2Name]] +; CHECK-SPIRV: ConstantComposite {{[0-9]+}} [[ConstComp:[0-9]+]] [[F1]] [[F2]] +; CHECK-SPIRV: Variable {{[0-9]+}} [[Var:[0-9]+]] {{[0-9]+}} [[ConstComp]] + +; CHECK-SPIRV: InBoundsPtrAccessChain {{[0-9]+}} [[GEP:[0-9]+]] [[Var]] {{[0-9]+}} {{[0-9]+}} +; CHECK-SPIRV: Load {{[0-9]+}} [[FuncPtr:[0-9]+]] [[GEP]] +; CHECK-SPIRV: FunctionPointerCallINTEL [[Int32]] {{[0-9]+}} [[FuncPtr]] [[XArg]] [[YArg]] + +; CHECK-LLVM: @__const.main.funcs = internal constant [2 x i32 (i32, i32) addrspace(9)*] [i32 (i32, i32) addrspace(9)* @f1, i32 (i32, i32) addrspace(9)* @f2], align 16 +; CHECK-LLVM: %[[Idx:[a-z0-9]+]] = getelementptr inbounds [2 x i32 (i32, i32) addrspace(9)*], [2 x i32 (i32, i32) addrspace(9)*]* @__const.main.funcs, i64 0, i64 %{{[a-z0-9]+}} +; CHECK-LLVM: %[[FuncPtr:[a-z0-9]+]] = load i32 (i32, i32) addrspace(9)*, i32 (i32, i32) addrspace(9)** %[[Idx]], align 8 +; CHECK-LLVM: %{{[a-z0-9]+}} = call spir_func addrspace(9) i32 %[[FuncPtr]](i32 32, i32 2) + +target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128" +target triple = "spir-unknown-unknown" + +@__const.main.funcs = private unnamed_addr constant [2 x i32 (i32, i32)*] [i32 (i32, i32)* @f1, i32 (i32, i32)* @f2], align 16 + +; Function Attrs: norecurse nounwind readnone uwtable +define dso_local i32 @f1(i32 %a, i32 %b) #0 { +entry: + %add = add nsw i32 %b, %a + ret i32 %add +} + +; Function Attrs: norecurse nounwind readnone uwtable +define dso_local i32 @f2(i32 %a, i32 %b) #0 { +entry: + %sub = sub nsw i32 %a, %b + ret i32 %sub +} + +; Function Attrs: nounwind uwtable +define dso_local i32 @main() local_unnamed_addr #1 { +entry: + %call = tail call i32 @rand() #3 + %rem = srem i32 %call, 2 + %idxprom = sext i32 %rem to i64 + %arrayidx = getelementptr inbounds [2 x i32 (i32, i32)*], [2 x i32 (i32, i32)*]* @__const.main.funcs, i64 0, i64 %idxprom + %0 = load i32 (i32, i32)*, i32 (i32, i32)** %arrayidx, align 8 + %call1 = tail call i32 %0(i32 32, i32 2) #3 + ret i32 %call1 +} + +; Function Attrs: nounwind +declare dso_local i32 @rand() local_unnamed_addr #2 + +attributes #0 = { norecurse nounwind readnone uwtable "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #1 = { nounwind uwtable "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #2 = { nounwind "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "unsafe-fp-math"="false" "use-soft-float"="false" } diff --git a/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/decor-func-ptr-arg-attr.ll b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/decor-func-ptr-arg-attr.ll new file mode 100644 index 0000000000..edb8d910c2 --- /dev/null +++ b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/decor-func-ptr-arg-attr.ll @@ -0,0 +1,67 @@ +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv %t.bc -o %t.spt -spirv-text -spirv-ext=+SPV_INTEL_function_pointers +; RUN: FileCheck < %t.spt %s --check-prefix CHECK-SPIRV + +; RUN: llvm-spirv %t.spt -o %t.spv -to-binary +; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.rev.bc +; RUN: llvm-dis %t.rev.bc -o %t.rev.ll +; RUN: FileCheck < %t.rev.ll %s --check-prefix CHECK-LLVM + +; CHECK-SPIRV: Capability FunctionPointersINTEL +; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers" + +; CHECK-SPIRV: Decorate [[#TargetId:]] ArgumentAttributeINTEL 0 4 +; CHECK-SPIRV: Decorate [[#TargetId]] ArgumentAttributeINTEL 0 5 +; CHECK-SPIRV: Decorate [[#TargetId]] ArgumentAttributeINTEL 0 2 +; CHECK-SPIRV: FunctionPointerCallINTEL +; CHECK-SPIRV-SAME: [[#TargetId]] + +; CHECK-LLVM: call spir_func addrspace(9) void %cond.i.i(%multi_ptr* noalias nocapture byval(%multi_ptr) %agg.tmp.i.i) + +; ModuleID = 'sycl_test.cpp' +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" +target triple = "spir64-unknown-unknown" + +%"multi_ptr" = type { i32* } +%"range" = type { %"array" } +%"array" = type { [1 x i64] } +%wrapper_class = type { i32 addrspace(1)* } +%wrapper_class.0 = type { i32 addrspace(1)* } + +$RoundedRangeKernel = comdat any + +; Function Attrs: nounwind +define spir_func void @inc_function(%"multi_ptr"* byval(%"multi_ptr") noalias nocapture %ptr) #0 { +entry: + ret void +} + + +; Function Attrs: convergent norecurse +define weak_odr dso_local spir_kernel void @RoundedRangeKernel(%"range"* byval(%"range") align 8 %_arg_NumWorkItems, i1 zeroext %_arg_, %wrapper_class* byval(%wrapper_class) align 8 %_arg_1, %wrapper_class.0* byval(%wrapper_class.0) align 8 %_arg_2) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !6 { +entry: + %agg.tmp.i.i = alloca %"multi_ptr", align 8 + %cond.i.i = select i1 %_arg_, void (%"multi_ptr"*)* @inc_function, void (%"multi_ptr"*)* null + call spir_func void %cond.i.i(%"multi_ptr"* nonnull byval(%"multi_ptr") align 8 noalias nocapture %agg.tmp.i.i) #1, !callees !7 + ret void +} + +attributes #0 = { convergent norecurse "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee" "frame-pointer"="all" "min-legal-vector-width"="0" "no-infs-fp-math"="true" "no-nans-fp-math"="true" "no-signed-zeros-fp-math"="true" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="sycl_test.cpp" "uniform-work-group-size"="true" "unsafe-fp-math"="true" } +attributes #1 = { convergent } + +!llvm.module.flags = !{!0, !1} +!opencl.spir.version = !{!2} +!spirv.Source = !{!3} +!opencl.used.extensions = !{!4} +!opencl.used.optional.core.features = !{!4} +!opencl.compiler.options = !{!4} +!llvm.ident = !{!5} + +!0 = !{i32 1, !"wchar_size", i32 4} +!1 = !{i32 7, !"frame-pointer", i32 2} +!2 = !{i32 1, i32 2} +!3 = !{i32 4, i32 100000} +!4 = !{} +!5 = !{!"Compiler"} +!6 = !{i32 -1, i32 -1, i32 -1, i32 -1} +!7 = !{void (%"multi_ptr"*)* @inc_function} diff --git a/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/fp-from-host.ll b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/fp-from-host.ll new file mode 100644 index 0000000000..0c1394d48b --- /dev/null +++ b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/fp-from-host.ll @@ -0,0 +1,70 @@ +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_INTEL_function_pointers -o %t.spt +; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV +; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -o %t.spv +; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.r.bc +; RUN: llvm-dis %t.r.bc -o %t.r.ll +; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM +; +; Generated from: +; typedef int (*fp_t)(int); +; +; __kernel void test(__global int *fp, __global int *data) { +; +; data[0] = ((fp_t)(*fp))(data[1]); +; } +; +; CHECK-SPIRV: Capability FunctionPointersINTEL +; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers" +; +; CHECK-SPIRV: Name [[KERNEL_ID:[0-9]+]] "test" +; CHECK-SPIRV: TypeInt [[INT32_TYPE_ID:[0-9]+]] 32 +; CHECK-SPIRV: TypePointer [[INT_PTR:[0-9]+]] 5 [[INT32_TYPE_ID]] +; CHECK-SPIRV: TypeFunction [[FOO_TYPE_ID:[0-9]+]] [[INT32_TYPE_ID]] [[INT32_TYPE_ID]] +; CHECK-SPIRV: TypePointer [[FOO_TYPE_PTR_ID:[0-9]+]] {{[0-9]+}} [[FOO_TYPE_ID]] +; +; CHECK-SPIRV: Function {{[0-9]+}} [[KERNEL_ID]] +; CHECK-SPIRV: FunctionParameter [[INT_PTR]] [[FP:[0-9]+]] +; CHECK-SPIRV: Load [[INT32_TYPE_ID]] [[FUNC_ADDR:[0-9]+]] [[FP]] +; CHECK-SPIRV: ConvertUToPtr [[FOO_TYPE_PTR_ID]] [[FOO_PTR:[0-9]+]] [[FUNC_ADDR]] +; CHECK-SPIRV: FunctionPointerCallINTEL [[INT32_TYPE_ID]] {{[0-9]+}} [[FOO_PTR]] +; +; CHECK-LLVM: define spir_kernel void @test(i32 addrspace(1)* +; CHECK-LLVM: %{{.*}} = call spir_func addrspace(9) i32 %{{.*}}(i32 + +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" +target triple = "spir64-unknown-unknown" + +; Function Attrs: convergent nounwind +define spir_kernel void @test(i32 addrspace(1)* %fp, i32 addrspace(1)* %data) #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { +entry: + %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %data, i64 1 + %0 = load i32, i32 addrspace(1)* %arrayidx, align 4, !tbaa !8 + %1 = load i32, i32 addrspace(1)* %fp, align 4, !tbaa !8 + %2 = inttoptr i32 %1 to i32 (i32)* + %call = call spir_func i32 %2(i32 %0) #1 + %arrayidx1 = getelementptr inbounds i32, i32 addrspace(1)* %data, i64 0 + store i32 %call, i32 addrspace(1)* %arrayidx1, align 4, !tbaa !8 + ret void +} + +attributes #0 = { convergent nounwind "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #1 = { convergent } + +!llvm.module.flags = !{!0} +!opencl.ocl.version = !{!1} +!opencl.spir.version = !{!2} +!llvm.ident = !{!3} + +!0 = !{i32 1, !"wchar_size", i32 4} +!1 = !{i32 1, i32 0} +!2 = !{i32 1, i32 2} +!3 = !{!"clang version 7.1.0 "} +!4 = !{i32 1, i32 1} +!5 = !{!"none", !"none"} +!6 = !{!"int*", !"int*"} +!7 = !{!"", !""} +!8 = !{!9, !9, i64 0} +!9 = !{!"int", !10, i64 0} +!10 = !{!"omnipotent char", !11, i64 0} +!11 = !{!"Simple C/C++ TBAA"} diff --git a/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer-as-function-arg.ll b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer-as-function-arg.ll new file mode 100644 index 0000000000..7ee0e49daa --- /dev/null +++ b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer-as-function-arg.ll @@ -0,0 +1,177 @@ +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_INTEL_function_pointers -o %t.spt +; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV +; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -o %t.spv +; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.r.bc +; RUN: llvm-dis %t.r.bc -o %t.r.ll +; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM +; +; Generated from: +; int helper(int (*f)(int), int arg) { +; return f(arg); +; } +; +; int foo(int v) { +; return v + 1; +; } +; +; int bar(int v) { +; return v + 2; +; } +; +; __kernel void test(__global int *data, int control) { +; int (*fp)(int) = 0; +; +; if (get_global_id(0) % control == 0) +; fp = &foo; +; else +; fp = &bar; +; +; data[get_global_id(0)] = helper(fp, data[get_global_id(0)]); +; } +; +; CHECK-SPIRV: Capability FunctionPointersINTEL +; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers" +; +; CHECK-SPIRV: Name [[KERNEL_ID:[0-9]+]] "test" +; CHECK-SPIRV: TypeInt [[TYPE_INT32_ID:[0-9]+]] 32 +; CHECK-SPIRV: TypeFunction [[FOO_TYPE_ID:[0-9]+]] [[TYPE_INT32_ID]] [[TYPE_INT32_ID]] +; CHECK-SPIRV: TypePointer [[FOO_PTR_TYPE_ID:[0-9]+]] {{[0-9]+}} [[FOO_TYPE_ID]] +; CHECK-SPIRV: TypeFunction [[HELPER_TYPE_ID:[0-9]+]] [[TYPE_INT32_ID]] [[FOO_PTR_TYPE_ID]] [[TYPE_INT32_ID]] +; CHECK-SPIRV: TypePointer [[FOO_PTR_ALLOCA_TYPE_ID:[0-9]+]] {{[0-9]+}} [[FOO_PTR_TYPE_ID]] +; CHECK-SPIRV: TypePointer [[TYPE_INT32_ALLOCA_ID:[0-9]+]] {{[0-9]+}} [[TYPE_INT32_ID]] +; CHECK-SPIRV: FunctionPointerINTEL [[FOO_PTR_TYPE_ID]] [[FOO_PTR_ID:[0-9]+]] [[FOO_ID:[0-9]+]] +; CHECK-SPIRV: FunctionPointerINTEL [[FOO_PTR_TYPE_ID]] [[BAR_PTR_ID:[0-9]+]] [[BAR_ID:[0-9]+]] +; +; CHECK-SPIRV: Function {{[0-9]+}} [[HELPER_ID:[0-9]+]] {{[0-9]+}} [[HELPER_TYPE_ID]] +; CHECK-SPIRV: FunctionParameter [[FOO_PTR_TYPE_ID]] [[T_PTR_ARG_ID:[0-9]+]] +; CHECK-SPIRV: FunctionParameter [[TYPE_INT32_ID:[0-9]+]] [[INT_ARG_ID:[0-9]+]] +; CHECK-SPIRV: Variable [[FOO_PTR_ALLOCA_TYPE_ID]] [[T_PTR_ALLOCA_ID:[0-9]+]] +; CHECK-SPIRV: Variable [[TYPE_INT32_ALLOCA_ID]] [[INT_ALLOCA_ID:[0-9]+]] +; CHECK-SPIRV: Store [[T_PTR_ALLOCA_ID]] [[T_PTR_ARG_ID]] +; CHECK-SPIRV: Store [[INT_ALLOCA_ID]] [[INT_ARG_ID]] +; CHECK-SPIRV: Load [[FOO_PTR_TYPE_ID]] [[LOADED_T_PTR:[0-9]+]] [[T_PTR_ALLOCA_ID]] +; CHECK-SPIRV: Load [[TYPE_INT32_ID]] [[LOADED_INT:[0-9]+]] [[INT_ALLOCA_ID]] +; CHECK-SPIRV: FunctionPointerCallINTEL [[TYPE_INT32_ID]] [[RESULT:[0-9]+]] [[LOADED_T_PTR]] [[LOADED_INT]] +; CHECK-SPIRV: ReturnValue [[RESULT]] +; +; CHECK-SPIRV: Function {{[0-9]+}} [[FOO_ID]] {{[0-9]+}} [[FOO_TYPE_ID]] +; CHECK-SPIRV: Function {{[0-9]+}} [[BAR_ID]] {{[0-9]+}} [[FOO_TYPE_ID]] +; +; CHECK-SPIRV: Function {{[0-9]+}} [[KERNEL_ID]] +; CHECK-SPIRV: Variable [[FOO_PTR_ALLOCA_TYPE_ID]] [[F_PTR_ALLOCA_ID:[0-9]+]] +; CHECK-SPIRV: Store [[F_PTR_ALLOCA_ID]] [[FOO_PTR_ID]] +; CHECK-SPIRV: Store [[F_PTR_ALLOCA_ID]] [[BAR_PTR_ID]] +; CHECK-SPIRV: Load [[FOO_PTR_TYPE_ID]] [[LOADED_F_PTR:[0-9]+]] [[F_PTR_ALLOCA_ID]] +; CHECK-SPIRV: FunctionCall {{[0-9]+}} {{[0-9]+}} [[HELPER_ID]] [[LOADED_F_PTR]] +; +; CHECK-LLVM: define spir_func i32 @helper(i32 (i32) addrspace(9)* %[[F:.*]], +; CHECK-LLVM: %[[F_ADDR:.*]] = alloca i32 (i32) addrspace(9)* +; CHECK-LLVM: store i32 (i32) addrspace(9)* %[[F]], i32 (i32) addrspace(9)** %[[F_ADDR]] +; CHECK-LLVM: %[[F_LOADED:.*]] = load i32 (i32) addrspace(9)*, i32 (i32) addrspace(9)** %[[F_ADDR]] +; CHECK-LLVM: %[[CALL:.*]] = call spir_func addrspace(9) i32 %[[F_LOADED]] +; CHECK-LLVM: ret i32 %[[CALL]] +; +; CHECK-LLVM: define spir_kernel void @test +; CHECK-LLVM: %[[FP:.*]] = alloca i32 (i32) addrspace(9)* +; CHECK-LLVM: store i32 (i32) addrspace(9)* @foo, i32 (i32) addrspace(9)** %[[FP]] +; CHECK-LLVM: store i32 (i32) addrspace(9)* @bar, i32 (i32) addrspace(9)** %[[FP]] +; CHECK-LLVM: %[[FP_LOADED:.*]] = load i32 (i32) addrspace(9)*, i32 (i32) addrspace(9)** %[[FP]] +; CHECK-LLVM: call spir_func i32 @helper(i32 (i32) addrspace(9)* %[[FP_LOADED]] + + +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" +target triple = "spir64-unknown-unknown" + +; Function Attrs: convergent noinline nounwind optnone +define spir_func i32 @helper(i32 (i32)* %f, i32 %arg) #0 { +entry: + %f.addr = alloca i32 (i32)*, align 8 + %arg.addr = alloca i32, align 4 + store i32 (i32)* %f, i32 (i32)** %f.addr, align 8 + store i32 %arg, i32* %arg.addr, align 4 + %0 = load i32 (i32)*, i32 (i32)** %f.addr, align 8 + %1 = load i32, i32* %arg.addr, align 4 + %call = call spir_func i32 %0(i32 %1) #3 + ret i32 %call +} + +; Function Attrs: convergent noinline nounwind optnone +define spir_func i32 @foo(i32 %v) #0 { +entry: + %v.addr = alloca i32, align 4 + store i32 %v, i32* %v.addr, align 4 + %0 = load i32, i32* %v.addr, align 4 + %add = add nsw i32 %0, 1 + ret i32 %add +} + +; Function Attrs: convergent noinline nounwind optnone +define spir_func i32 @bar(i32 %v) #0 { +entry: + %v.addr = alloca i32, align 4 + store i32 %v, i32* %v.addr, align 4 + %0 = load i32, i32* %v.addr, align 4 + %add = add nsw i32 %0, 2 + ret i32 %add +} + +; Function Attrs: convergent noinline nounwind optnone +define spir_kernel void @test(i32 addrspace(1)* %data, i32 %control) #1 !kernel_arg_addr_space !1 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { +entry: + %data.addr = alloca i32 addrspace(1)*, align 8 + %control.addr = alloca i32, align 4 + %fp = alloca i32 (i32)*, align 8 + store i32 addrspace(1)* %data, i32 addrspace(1)** %data.addr, align 8 + store i32 %control, i32* %control.addr, align 4 + store i32 (i32)* null, i32 (i32)** %fp, align 8 + %call = call spir_func i64 @_Z13get_global_idj(i32 0) #4 + %0 = load i32, i32* %control.addr, align 4 + %conv = sext i32 %0 to i64 + %rem = urem i64 %call, %conv + %cmp = icmp eq i64 %rem, 0 + br i1 %cmp, label %if.then, label %if.else + +if.then: ; preds = %entry + store i32 (i32)* @foo, i32 (i32)** %fp, align 8 + br label %if.end + +if.else: ; preds = %entry + store i32 (i32)* @bar, i32 (i32)** %fp, align 8 + br label %if.end + +if.end: ; preds = %if.else, %if.then + %1 = load i32 (i32)*, i32 (i32)** %fp, align 8 + %2 = load i32 addrspace(1)*, i32 addrspace(1)** %data.addr, align 8 + %call2 = call spir_func i64 @_Z13get_global_idj(i32 0) #4 + %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %2, i64 %call2 + %3 = load i32, i32 addrspace(1)* %arrayidx, align 4 + %call3 = call spir_func i32 @helper(i32 (i32)* %1, i32 %3) #3 + %4 = load i32 addrspace(1)*, i32 addrspace(1)** %data.addr, align 8 + %call4 = call spir_func i64 @_Z13get_global_idj(i32 0) #4 + %arrayidx5 = getelementptr inbounds i32, i32 addrspace(1)* %4, i64 %call4 + store i32 %call3, i32 addrspace(1)* %arrayidx5, align 4 + ret void +} + +; Function Attrs: convergent nounwind readnone +declare spir_func i64 @_Z13get_global_idj(i32) #2 + +attributes #0 = { convergent noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #1 = { convergent noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #2 = { convergent nounwind readnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #3 = { convergent } +attributes #4 = { convergent nounwind readnone } + +!llvm.module.flags = !{!0} +!opencl.ocl.version = !{!1} +!opencl.spir.version = !{!2} +!llvm.ident = !{!3} + +!0 = !{i32 1, !"wchar_size", i32 4} +!1 = !{i32 1, i32 0} +!2 = !{i32 1, i32 2} +!3 = !{!"clang version 7.1.0 "} +!4 = !{!"none", !"none"} +!5 = !{!"int*", !"int"} +!6 = !{!"", !""} diff --git a/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer-dedicated-as.ll b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer-dedicated-as.ll new file mode 100644 index 0000000000..2ebbb429b3 --- /dev/null +++ b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer-dedicated-as.ll @@ -0,0 +1,107 @@ +; This test checks how a function pointer in a dedicated addr space would be +; translated with and without -spirv-emit-function-ptr-addr-space option. +; Expected behaviour: +; No option is passed to the forward translation stage - no CodeSectionINTEL storage class in SPIR-V +; The option is passed to the forward translation stage - CodeSectionINTEL storage class is generated +; No option is passed to the reverse translation stage - function pointers are in private address space +; The option is passed to the reverse translation stage - function pointers are in addrspace(9) +; +; Overall IR generation is tested elsewhere, here checks are very simple + +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_INTEL_function_pointers -spirv-emit-function-ptr-addr-space -o %t.spt +; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV-AS +; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -spirv-emit-function-ptr-addr-space -o %t.spv +; RUN: llvm-spirv -r %t.spv -o %t.r.bc +; RUN: llvm-dis %t.r.bc -o %t.r.ll +; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM-NO-AS + +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_INTEL_function_pointers -spirv-emit-function-ptr-addr-space -o %t.spt +; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV-AS +; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -spirv-emit-function-ptr-addr-space -o %t.spv +; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.r.bc +; RUN: llvm-dis %t.r.bc -o %t.r.ll +; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM-AS + +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_INTEL_function_pointers -o %t.spt +; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV-NO-AS +; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -o %t.spv +; RUN: llvm-spirv -r %t.spv -o %t.r.bc +; RUN: llvm-dis %t.r.bc -o %t.r.ll +; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM-NO-AS + +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_INTEL_function_pointers -o %t.spt +; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV-NO-AS +; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -o %t.spv +; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.r.bc +; RUN: llvm-dis %t.r.bc -o %t.r.ll +; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM-AS + +; CHECK-SPIRV-AS-DAG: TypePointer [[#PtrCodeTy:]] 5605 [[#]] +; CHECK-SPIRV-AS-DAG: TypePointer [[#PtrPrivTy:]] 7 [[#PtrCodeTy]] +; CHECK-SPIRV-AS-DAG: ConstantFunctionPointerINTEL [[#PtrCodeTy]] [[#FunPtr:]] +; CHECK-SPIRV-AS: Variable [[#PtrPrivTy]] [[#Var:]] 7 +; CHECK-SPIRV-AS: Store [[#Var]] [[#FunPtr]] +; CHECK-SPIRV-AS: Load [[#PtrCodeTy]] [[#Load:]] [[#Var]] +; CHECK-SPIRV-AS: FunctionPointerCallINTEL [[#]] [[#]] [[#Load]] [[#]] + +; CHECK-SPIRV-NO-AS-NOT: TypePointer [[#]] 5605 [[#]] + +; CHECK-LLVM-AS: define spir_func i32 @foo(i32 %{{.*}}) addrspace(9) + +; CHECK-LLVM-NO-AS-NOT: addrspace(9) + +; ModuleID = 'function-pointer-dedicated-as.bc' +target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024-G1" +target triple = "spir64-unknown-unknown" + +; Function Attrs: noinline nounwind +define spir_func i32 @foo(i32 %arg) addrspace(9) #0 { +entry: + %arg.addr = alloca i32, align 4 + store i32 %arg, i32* %arg.addr, align 4 + %0 = load i32, i32* %arg.addr, align 4 + %add = add nsw i32 %0, 10 + ret i32 %add +} + +; Function Attrs: noinline nounwind +define spir_kernel void @test(i32 addrspace(1)* %data, i32 %input) #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !6 !kernel_arg_type !7 !kernel_arg_type_qual !8 !kernel_arg_base_type !7 !spirv.ParameterDecorations !9 { +entry: + %data.addr = alloca i32 addrspace(1)*, align 8 + %input.addr = alloca i32, align 4 + %fp = alloca i32 (i32) addrspace(9)*, align 8 + store i32 addrspace(1)* %data, i32 addrspace(1)** %data.addr, align 8 + store i32 %input, i32* %input.addr, align 4 + store i32 (i32) addrspace(9)* @foo, i32 (i32) addrspace(9)** %fp, align 8 + %0 = load i32 (i32) addrspace(9)*, i32 (i32) addrspace(9)** %fp, align 8 + %1 = load i32, i32* %input.addr, align 4 + %call = call spir_func addrspace(9) i32 %0(i32 %1) + %2 = load i32 addrspace(1)*, i32 addrspace(1)** %data.addr, align 8 + store i32 %call, i32 addrspace(1)* %2, align 4 + ret void +} + +attributes #0 = { noinline nounwind } + +!spirv.MemoryModel = !{!0} +!spirv.Source = !{!1} +!opencl.spir.version = !{!2} +!opencl.ocl.version = !{!3} +!opencl.used.extensions = !{!4} +!opencl.used.optional.core.features = !{!4} +!spirv.Generator = !{!5} + +!0 = !{i32 2, i32 2} +!1 = !{i32 3, i32 100000} +!2 = !{i32 1, i32 2} +!3 = !{i32 1, i32 0} +!4 = !{} +!5 = !{i16 6, i16 14} +!6 = !{!"none", !"none"} +!7 = !{!"int*", !"int"} +!8 = !{!"", !""} +!9 = !{!4, !4} diff --git a/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer.ll b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer.ll new file mode 100644 index 0000000000..d43497b59f --- /dev/null +++ b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer.ll @@ -0,0 +1,92 @@ +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_INTEL_function_pointers -o %t.spt +; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV +; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -o %t.spv +; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.r.bc +; RUN: llvm-dis %t.r.bc -o %t.r.ll +; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM +; +; Generated from: +; int foo(int arg) { +; return arg + 10; +; } +; +; void __kernel test(__global int *data, int input) { +; int (__constant *fp)(int) = &foo; +; +; *data = fp(input); +; } +; +; CHECK-SPIRV: Capability FunctionPointersINTEL +; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers" +; CHECK-SPIRV: Name [[KERNEL_ID:[0-9]+]] "test" +; CHECK-SPIRV: TypeInt [[TYPE_INT_ID:[0-9]+]] +; CHECK-SPIRV: TypeFunction [[FOO_TYPE_ID:[0-9]+]] [[TYPE_INT_ID]] [[TYPE_INT_ID]] +; CHECK-SPIRV: TypePointer [[FOO_PTR_ID:[0-9]+]] {{[0-9]+}} [[FOO_TYPE_ID]] +; CHECK-SPIRV: TypePointer [[FOO_PTR_ALLOCA_ID:[0-9]+]] 7 [[FOO_PTR_ID]] +; CHECK-SPIRV: ConstantFunctionPointerINTEL [[FOO_PTR_ID]] [[FOO_PTR:[0-9]+]] [[FOO_ID:[0-9]+]] +; +; CHECK-SPIRV: Function {{[0-9]+}} [[FOO_ID]] {{[0-9]+}} [[FOO_TYPE_ID]] +; CHECK-SPIRV: Function {{[0-9]+}} [[KERNEL_ID]] +; CHECK-SPIRV: Variable [[FOO_PTR_ALLOCA_ID]] [[FOO_PTR_ALLOCA:[0-9]+]] +; CHECK-SPIRV: Store [[FOO_PTR_ALLOCA]] [[FOO_PTR]] +; CHECK-SPIRV: Load [[FOO_PTR_ID]] [[LOADED_FOO_PTR:[0-9]+]] [[FOO_PTR_ALLOCA]] +; CHECK-SPIRV: FunctionPointerCallINTEL 2 {{[0-9]+}} [[LOADED_FOO_PTR]] +; +; CHECK-LLVM: define spir_kernel void @test +; CHECK-LLVM: %fp = alloca i32 (i32) addrspace(9)* +; CHECK-LLVM: store i32 (i32) addrspace(9)* @foo, i32 (i32) addrspace(9)** %fp +; CHECK-LLVM: %0 = load i32 (i32) addrspace(9)*, i32 (i32) addrspace(9)** %fp +; CHECK-LLVM: %call = call spir_func addrspace(9) i32 %0(i32 %1) + +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" +target triple = "spir64-unknown-unknown" + +; Function Attrs: convergent noinline nounwind optnone +define spir_func i32 @foo(i32 %arg) #0 { +entry: + %arg.addr = alloca i32, align 4 + store i32 %arg, i32* %arg.addr, align 4 + %0 = load i32, i32* %arg.addr, align 4 + %add = add nsw i32 %0, 10 + ret i32 %add +} + +; Function Attrs: convergent noinline nounwind optnone +define spir_kernel void @test(i32 addrspace(1)* %data, i32 %input) #1 !kernel_arg_addr_space !1 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { +entry: + %data.addr = alloca i32 addrspace(1)*, align 8 + %input.addr = alloca i32, align 4 + %fp = alloca i32 (i32)*, align 8 + store i32 addrspace(1)* %data, i32 addrspace(1)** %data.addr, align 8 + store i32 %input, i32* %input.addr, align 4 + store i32 (i32)* @foo, i32 (i32)** %fp, align 8 + %0 = load i32 (i32)*, i32 (i32)** %fp, align 8 + %1 = load i32, i32* %input.addr, align 4 + %call = call spir_func i32 %0(i32 %1) #2 + %2 = load i32 addrspace(1)*, i32 addrspace(1)** %data.addr, align 8 + store i32 %call, i32 addrspace(1)* %2, align 4 + ret void +} + +attributes #0 = { convergent noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #1 = { convergent noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #2 = { convergent } + +!llvm.module.flags = !{!0} +!opencl.ocl.version = !{!1} +!opencl.spir.version = !{!2} +!opencl.used.extensions = !{!3} +!opencl.used.optional.core.features = !{!3} +!opencl.compiler.options = !{!3} +!llvm.ident = !{!4} + +!0 = !{i32 1, !"wchar_size", i32 4} +!1 = !{i32 1, i32 0} +!2 = !{i32 1, i32 2} +!3 = !{} +!4 = !{!"clang version 7.0.0 "} +!5 = !{!"none", !"none"} +!6 = !{!"int*", !"int"} +!7 = !{!"", !""} + diff --git a/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/gv-func-ptr.ll b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/gv-func-ptr.ll new file mode 100644 index 0000000000..31c36f98eb --- /dev/null +++ b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/gv-func-ptr.ll @@ -0,0 +1,40 @@ +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv %t.bc -spirv-ext=+SPV_INTEL_function_pointers -o %t.spv +; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.r.bc +; RUN: llvm-dis %t.r.bc -o %t.r.ll +; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM + +; ModuleID = 't.bc' +target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024" +target triple = "spir64-unknown-unknown" + +%structtype.3 = type { [1 x i8 addrspace(4)*] } + +; CHECK-LLVM: @A = addrspace(1) constant %structtype.3 { [1 x i8 addrspace(4)*] [i8 addrspace(4)* addrspacecast (i8 addrspace(9)* bitcast (void () addrspace(9)* @foo to i8 addrspace(9)*) to i8 addrspace(4)*)] }, align 8 + +@A = linkonce_odr addrspace(1) constant %structtype.3 { [1 x i8 addrspace(4)*] [i8 addrspace(4)* addrspacecast (i8* bitcast (void ()* @foo to i8*) to i8 addrspace(4)*)] }, align 8 + +; Function Attrs: nounwind +define linkonce_odr spir_func void @foo() #0 { +entry: +; CHECK-LLVM: %0 = getelementptr inbounds %structtype.3, %structtype.3 addrspace(1)* @A, i64 0, i32 0, i64 2 + %0 = getelementptr inbounds %structtype.3, %structtype.3 addrspace(1)* @A, i64 0, i32 0, i64 2 + ret void +} + +attributes #0 = { nounwind } + +!spirv.MemoryModel = !{!0} +!spirv.Source = !{!1} +!opencl.spir.version = !{!2} +!opencl.ocl.version = !{!2} +!opencl.used.extensions = !{!3} +!opencl.used.optional.core.features = !{!4} +!spirv.Generator = !{!5} + +!0 = !{i32 2, i32 2} +!1 = !{i32 4, i32 200000} +!2 = !{i32 2, i32 0} +!3 = !{!"cl_khr_int64_extended_atomics", !"cl_khr_subgroups"} +!4 = !{!"cl_doubles"} +!5 = !{i16 6, i16 14} diff --git a/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/non-uniform-function-pointer.ll b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/non-uniform-function-pointer.ll new file mode 100644 index 0000000000..b98fb602f4 --- /dev/null +++ b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/non-uniform-function-pointer.ll @@ -0,0 +1,139 @@ +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_INTEL_function_pointers -o %t.spt +; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV +; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -o %t.spv +; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.r.bc +; RUN: llvm-dis %t.r.bc -o %t.r.ll +; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM +; +; Generated from: +; int foo(int v) { +; return v + 1; +; } +; +; int bar(int v) { +; return v + 2; +; } +; +; __kernel void test(__global int *data, int control) { +; int (*fp)(int) = 0; +; +; if (get_global_id(0) % control == 0) +; fp = &foo; +; else +; fp = &bar; +; +; data[get_global_id(0)] = fp(data[get_global_id(0)]); +; } +; +; CHECK-SPIRV: Capability FunctionPointersINTEL +; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers" +; +; CHECK-SPIRV: Name [[KERNEL_ID:[0-9]+]] "test" +; CHECK-SPIRV: TypeInt [[TYPE_INT32_ID:[0-9+]]] 32 +; CHECK-SPIRV: TypeFunction [[FOO_TYPE_ID:[0-9]+]] [[TYPE_INT32_ID]] [[TYPE_INT32_ID]] +; CHECK-SPIRV: TypePointer [[FOO_PTR_TYPE_ID:[0-9]+]] {{[0-9]+}} [[FOO_TYPE_ID]] +; CHECK-SPIRV: TypePointer [[FOO_PTR_ALLOCA_TYPE_ID:[0-9]+]] 7 [[FOO_PTR_TYPE_ID]] +; CHECK-SPIRV: ConstantFunctionPointerINTEL [[FOO_PTR_TYPE_ID]] [[FOO_PTR_ID:[0-9]+]] [[FOO_ID:[0-9]+]] +; CHECK-SPIRV: ConstantFunctionPointerINTEL [[FOO_PTR_TYPE_ID]] [[BAR_PTR_ID:[0-9]+]] [[BAR_ID:[0-9]+]] +; +; CHECK-SPIRV: Function {{[0-9]+}} [[FOO_ID]] {{[0-9]+}} [[FOO_TYPE_ID]] +; CHECK-SPIRV: Function {{[0-9]+}} [[BAR_ID]] {{[0-9]+}} [[FOO_TYPE_ID]] +; +; CHECK-SPIRV: Function {{[0-9]+}} [[KERNEL_ID]] +; CHECK-SPIRV: Variable [[FOO_PTR_ALLOCA_TYPE_ID]] [[FOO_PTR_ALLOCA_ID:[0-9]+]] +; CHECK-SPIRV: Store [[FOO_PTR_ALLOCA_ID]] [[FOO_PTR_ID]] +; CHECK-SPIRV: Store [[FOO_PTR_ALLOCA_ID]] [[BAR_PTR_ID]] +; CHECK-SPIRV: Load [[FOO_PTR_TYPE_ID]] [[LOADED_FOO_PTR:[0-9]+]] [[FOO_PTR_ALLOCA_ID]] +; CHECK-SPIRV: FunctionPointerCallINTEL {{[0-9]+}} {{[0-9]+}} [[LOADED_FOO_PTR]] +; +; CHECK-LLVM: define spir_kernel void @test +; CHECK-LLVM: %fp = alloca i32 (i32) addrspace(9)* +; CHECK-LLVM: store i32 (i32) addrspace(9)* @foo, i32 (i32) addrspace(9)** %fp +; CHECK-LLVM: store i32 (i32) addrspace(9)* @bar, i32 (i32) addrspace(9)** %fp +; CHECK-LLVM: %[[FP:.*]] = load i32 (i32) addrspace(9)*, i32 (i32) addrspace(9)** %fp +; CHECK-LLVM: call spir_func addrspace(9) i32 %[[FP]](i32 %{{.*}}) + + +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" +target triple = "spir64-unknown-unknown" + +; Function Attrs: convergent noinline nounwind optnone +define spir_func i32 @foo(i32 %v) #0 { +entry: + %v.addr = alloca i32, align 4 + store i32 %v, i32* %v.addr, align 4 + %0 = load i32, i32* %v.addr, align 4 + %add = add nsw i32 %0, 1 + ret i32 %add +} + +; Function Attrs: convergent noinline nounwind optnone +define spir_func i32 @bar(i32 %v) #0 { +entry: + %v.addr = alloca i32, align 4 + store i32 %v, i32* %v.addr, align 4 + %0 = load i32, i32* %v.addr, align 4 + %add = add nsw i32 %0, 2 + ret i32 %add +} + +; Function Attrs: convergent noinline nounwind optnone +define spir_kernel void @test(i32 addrspace(1)* %data, i32 %control) #1 !kernel_arg_addr_space !1 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { +entry: + %data.addr = alloca i32 addrspace(1)*, align 8 + %control.addr = alloca i32, align 4 + %fp = alloca i32 (i32)*, align 8 + store i32 addrspace(1)* %data, i32 addrspace(1)** %data.addr, align 8 + store i32 %control, i32* %control.addr, align 4 + store i32 (i32)* null, i32 (i32)** %fp, align 8 + %call = call spir_func i64 @_Z13get_global_idj(i32 0) #3 + %0 = load i32, i32* %control.addr, align 4 + %conv = sext i32 %0 to i64 + %rem = urem i64 %call, %conv + %cmp = icmp eq i64 %rem, 0 + br i1 %cmp, label %if.then, label %if.else + +if.then: ; preds = %entry + store i32 (i32)* @foo, i32 (i32)** %fp, align 8 + br label %if.end + +if.else: ; preds = %entry + store i32 (i32)* @bar, i32 (i32)** %fp, align 8 + br label %if.end + +if.end: ; preds = %if.else, %if.then + %1 = load i32 (i32)*, i32 (i32)** %fp, align 8 + %2 = load i32 addrspace(1)*, i32 addrspace(1)** %data.addr, align 8 + %call2 = call spir_func i64 @_Z13get_global_idj(i32 0) #3 + %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %2, i64 %call2 + %3 = load i32, i32 addrspace(1)* %arrayidx, align 4 + %call3 = call spir_func i32 %1(i32 %3) #4 + %4 = load i32 addrspace(1)*, i32 addrspace(1)** %data.addr, align 8 + %call4 = call spir_func i64 @_Z13get_global_idj(i32 0) #3 + %arrayidx5 = getelementptr inbounds i32, i32 addrspace(1)* %4, i64 %call4 + store i32 %call3, i32 addrspace(1)* %arrayidx5, align 4 + ret void +} + +; Function Attrs: convergent nounwind readnone +declare spir_func i64 @_Z13get_global_idj(i32) #2 + +attributes #0 = { convergent noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #1 = { convergent noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #2 = { convergent nounwind readnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #3 = { convergent nounwind readnone } +attributes #4 = { convergent } + +!llvm.module.flags = !{!0} +!opencl.ocl.version = !{!1} +!opencl.spir.version = !{!2} +!llvm.ident = !{!3} + +!0 = !{i32 1, !"wchar_size", i32 4} +!1 = !{i32 1, i32 0} +!2 = !{i32 1, i32 2} +!3 = !{!"clang version 7.1.0 "} +!4 = !{!"none", !"none"} +!5 = !{!"int*", !"int"} +!6 = !{!"", !""} diff --git a/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/referenced-indirectly.ll b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/referenced-indirectly.ll new file mode 100644 index 0000000000..2d7d0dff76 --- /dev/null +++ b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/referenced-indirectly.ll @@ -0,0 +1,82 @@ +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_INTEL_function_pointers -o %t.spt +; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV +; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -o %t.spv +; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.r.bc +; RUN: llvm-dis %t.r.bc -o %t.r.ll +; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM +; +; Generated from: +; __attribute__((referenced_indirectly)) +; int foo(int arg) { +; return arg + 10; +; } +; +; void __kernel test(__global int *data, int input) { +; int (__constant *fp)(int) = &foo; +; +; *data = fp(input); +; } +; +; CHECK-SPIRV: Capability FunctionPointersINTEL +; CHECK-SPIRV: Capability IndirectReferencesINTEL +; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers" +; +; CHECK-SPIRV: Name [[FOO_ID:[0-9]+]] "foo" +; CHECK-SPIRV: Decorate [[FOO_ID]] ReferencedIndirectlyINTEL +; CHECK-SPIRV: Function {{[0-9]+}} [[FOO_ID]] +; +; CHECK-LLVM: define spir_func i32 @foo(i32 %arg) addrspace(9) #[[ATTRS:[0-9]+]] +; CHECK-LLVM: attributes #[[ATTRS]] = {{.*}} "referenced-indirectly" + +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" +target triple = "spir64-unknown-unknown" + +; Function Attrs: convergent noinline nounwind optnone +define spir_func i32 @foo(i32 %arg) #0 { +entry: + %arg.addr = alloca i32, align 4 + store i32 %arg, i32* %arg.addr, align 4 + %0 = load i32, i32* %arg.addr, align 4 + %add = add nsw i32 %0, 10 + ret i32 %add +} + +; Function Attrs: convergent noinline nounwind optnone +define spir_kernel void @test(i32 addrspace(1)* %data, i32 %input) #1 !kernel_arg_addr_space !1 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { +entry: + %data.addr = alloca i32 addrspace(1)*, align 8 + %input.addr = alloca i32, align 4 + %fp = alloca i32 (i32)*, align 8 + store i32 addrspace(1)* %data, i32 addrspace(1)** %data.addr, align 8 + store i32 %input, i32* %input.addr, align 4 + store i32 (i32)* @foo, i32 (i32)** %fp, align 8 + %0 = load i32 (i32)*, i32 (i32)** %fp, align 8 + %1 = load i32, i32* %input.addr, align 4 + %call = call spir_func i32 %0(i32 %1) #2 + %2 = load i32 addrspace(1)*, i32 addrspace(1)** %data.addr, align 8 + store i32 %call, i32 addrspace(1)* %2, align 4 + ret void +} + +attributes #0 = { convergent noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" "referenced-indirectly" } +attributes #1 = { convergent noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #2 = { convergent } + +!llvm.module.flags = !{!0} +!opencl.ocl.version = !{!1} +!opencl.spir.version = !{!2} +!opencl.used.extensions = !{!3} +!opencl.used.optional.core.features = !{!3} +!opencl.compiler.options = !{!3} +!llvm.ident = !{!4} + +!0 = !{i32 1, !"wchar_size", i32 4} +!1 = !{i32 1, i32 0} +!2 = !{i32 1, i32 2} +!3 = !{} +!4 = !{!"clang version 7.0.0 "} +!5 = !{!"none", !"none"} +!6 = !{!"int*", !"int"} +!7 = !{!"", !""} + diff --git a/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/select.ll b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/select.ll new file mode 100644 index 0000000000..020c93a996 --- /dev/null +++ b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/select.ll @@ -0,0 +1,147 @@ +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -o %t.spv +; RUN: llvm-spirv %t.spv -to-text -o %t.spt +; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV +; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.r.bc +; RUN: llvm-dis %t.r.bc -o %t.r.ll +; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM + +; CHECK-SPIRV: Name [[#KERNEL_ID:]] "_ZTS6kernel" +; CHECK-SPIRV-DAG: Name [[#BAR:]] "_Z3barii" +; CHECK-SPIRV-DAG: Name [[#BAZ:]] "_Z3bazii" +; CHECK-SPIRV: TypeInt [[#INT32:]] 32 +; CHECK-SPIRV: TypeFunction [[#FUNC_TYPE:]] [[#INT32]] [[#INT32]] +; CHECK-SPIRV: TypePointer [[#FUNC_PTR_TYPE:]] [[#]] [[#FUNC_TYPE]] +; CHECK-SPIRV: TypePointer [[#FUNC_PTR_ALLOCA_TYPE:]] [[#]] [[#FUNC_PTR_TYPE]] +; CHECK-SPIRV-DAG: ConstantFunctionPointerINTEL [[#FUNC_PTR_TYPE]] [[#BARPTR:]] [[#BAR]] +; CHECK-SPIRV-DAG: ConstantFunctionPointerINTEL [[#FUNC_PTR_TYPE]] [[#BAZPTR:]] [[#BAZ]] +; CHECK-SPIRV: Function [[#]] [[#KERNEL_ID]] +; CHECK-SPIRV: Variable [[#FUNC_PTR_ALLOCA_TYPE]] [[#FPTR:]] +; CHECK-SPIRV: Select [[#FUNC_PTR_TYPE]] [[#SELECT:]] [[#]] [[#BARPTR]] [[#BAZPTR]] +; CHECK-SPIRV: Store [[#FPTR]] [[#SELECT]] +; CHECK-SPIRV: Load [[#FUNC_PTR_TYPE]] [[#LOAD:]] [[#FPTR]] +; CHECK-SPIRV: FunctionPointerCallINTEL [[#]] [[#]] [[#LOAD]] + +; CHECK-LLVM: define spir_kernel void @_ZTS6kernel +; CHECK-LLVM: %[[FPTR_ALLOCA:.*]] = alloca i32 (i32, i32) addrspace(9)* +; CHECK-LLVM: %[[SELECT:.*]] = select i1 %{{.*}}, i32 (i32, i32) addrspace(9)* @_Z3barii, i32 (i32, i32) addrspace(9)* @_Z3bazii +; CHECK-LLVM: store i32 (i32, i32) addrspace(9)* %[[SELECT]], i32 (i32, i32) addrspace(9)** %[[FPTR_ALLOCA]] +; CHECK-LLVM: %[[FPTR:.*]] = load i32 (i32, i32) addrspace(9)*, i32 (i32, i32) addrspace(9)** %[[FPTR_ALLOCA]] +; CHECK-LLVM: call spir_func addrspace(9) i32 %[[FPTR]]( + +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" +target triple = "spir64-unknown-unknown" + +%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } +%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] } +%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } + +$_ZTS6kernel = comdat any + +@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32 + +; Function Attrs: norecurse +define weak_odr dso_local spir_kernel void @_ZTS6kernel(i32 addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3) local_unnamed_addr #0 comdat !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { +entry: + %fptr.alloca = alloca i32 (i32, i32)*, align 8 + %ref.tmp.i = alloca %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", align 8 + %agg.tmp2.i = alloca %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range", align 8 + %agg.tmp3.i = alloca %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range", align 8 + %agg.tmp6 = alloca %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", align 8 + %0 = bitcast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %agg.tmp2.i to i8* + call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %0) + %1 = bitcast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %agg.tmp3.i to i8* + call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %1) + %2 = addrspacecast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %agg.tmp2.i to %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* + %ptrint4.i = ptrtoint %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* %2 to i64 + %maskedptr5.i = and i64 %ptrint4.i, 7 + %maskcond6.i = icmp eq i64 %maskedptr5.i, 0 + %3 = addrspacecast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %agg.tmp3.i to %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* + %ptrint.i = ptrtoint %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* %3 to i64 + %maskedptr.i = and i64 %ptrint.i, 7 + %maskcond.i = icmp eq i64 %maskedptr.i, 0 + call void @llvm.lifetime.end.p0i8(i64 8, i8* nonnull %0) + call void @llvm.lifetime.end.p0i8(i64 8, i8* nonnull %1) + %4 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0 + %5 = load i64, i64* %4, align 8 + %add.ptr.i = getelementptr inbounds i32, i32 addrspace(1)* %_arg_, i64 %5 + %6 = addrspacecast %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %agg.tmp6 to %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* + %ptrint = ptrtoint %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* %6 to i64 + %maskedptr = and i64 %ptrint, 7 + %maskcond = icmp eq i64 %maskedptr, 0 + %7 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32, !noalias !8 + %8 = extractelement <3 x i64> %7, i64 0 + %arrayinit.begin.i.i.i.i.i = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* %6, i64 0, i32 0, i32 0, i64 0 + store i64 %8, i64 addrspace(4)* %arrayinit.begin.i.i.i.i.i, align 8, !tbaa !15, !alias.scope !8 + %9 = bitcast %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %ref.tmp.i to i8* + call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %9) #4 + %10 = addrspacecast %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %ref.tmp.i to %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* + %ptrint.i2 = ptrtoint %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* %10 to i64 + %maskedptr.i3 = and i64 %ptrint.i2, 7 + %maskcond.i4 = icmp eq i64 %maskedptr.i3, 0 + %rem.i.i = and i64 %8, 1 + %cmp.i.i = icmp eq i64 %rem.i.i, 0 + call void @llvm.lifetime.end.p0i8(i64 8, i8* nonnull %9) #4 + %_Z3barii._Z3bazii.i = select i1 %cmp.i.i, i32 (i32, i32)* @_Z3barii, i32 (i32, i32)* @_Z3bazii + store i32 (i32, i32)* %_Z3barii._Z3bazii.i, i32 (i32, i32)** %fptr.alloca, align 8 + %fptr = load i32 (i32, i32)*, i32 (i32, i32)** %fptr.alloca, align 8 + %call4.i = call spir_func i32 %fptr(i32 10, i32 10), !callees !19 + %arrayidx.i3.i = getelementptr inbounds i32, i32 addrspace(1)* %add.ptr.i, i64 %8 + %arrayidx.ascast.i.i = addrspacecast i32 addrspace(1)* %arrayidx.i3.i to i32 addrspace(4)* + store i32 %call4.i, i32 addrspace(4)* %arrayidx.ascast.i.i, align 4, !tbaa !20 + ret void +} + +; Function Attrs: argmemonly nounwind willreturn +declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1 + +; Function Attrs: argmemonly nounwind willreturn +declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #1 + +; Function Attrs: norecurse nounwind readnone +define dso_local spir_func i32 @_Z3barii(i32 %a, i32 %b) local_unnamed_addr #2 { +entry: + %add = add nsw i32 %b, %a + ret i32 %add +} + +; Function Attrs: norecurse nounwind readnone +define dso_local spir_func i32 @_Z3bazii(i32 %a, i32 %b) local_unnamed_addr #2 { +entry: + %sub = sub nsw i32 %a, %b + ret i32 %sub +} + +attributes #0 = { norecurse "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "sycl-module-id"="f.cpp" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #1 = { argmemonly nounwind willreturn } +attributes #2 = { norecurse nounwind readnone "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #3 = { nounwind willreturn } +attributes #4 = { nounwind } + +!llvm.module.flags = !{!0} +!opencl.spir.version = !{!1} +!spirv.Source = !{!2} +!llvm.ident = !{!3} + +!0 = !{i32 1, !"wchar_size", i32 4} +!1 = !{i32 1, i32 2} +!2 = !{i32 4, i32 100000} +!3 = !{!"clang version 11.0.0 "} +!4 = !{i32 1, i32 0, i32 0, i32 0} +!5 = !{!"none", !"none", !"none", !"none"} +!6 = !{!"int*", !"cl::sycl::range<1>", !"cl::sycl::range<1>", !"cl::sycl::id<1>"} +!7 = !{!"", !"", !"", !""} +!8 = !{!9, !11, !13} +!9 = distinct !{!9, !10, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"} +!10 = distinct !{!10, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"} +!11 = distinct !{!11, !12, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"} +!12 = distinct !{!12, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"} +!13 = distinct !{!13, !14, !"_ZN2cl4sycl6detail7Builder5getIdILi1EEEKNS0_2idIXT_EEEv: %agg.result"} +!14 = distinct !{!14, !"_ZN2cl4sycl6detail7Builder5getIdILi1EEEKNS0_2idIXT_EEEv"} +!15 = !{!16, !16, i64 0} +!16 = !{!"long", !17, i64 0} +!17 = !{!"omnipotent char", !18, i64 0} +!18 = !{!"Simple C++ TBAA"} +!19 = !{i32 (i32, i32)* @_Z3barii, i32 (i32, i32)* @_Z3bazii} +!20 = !{!21, !21, i64 0} +!21 = !{!"int", !17, i64 0} diff --git a/tools/llvm-spirv/llvm-spirv.cpp b/tools/llvm-spirv/llvm-spirv.cpp index 2a09c2f132..fadd3b9cec 100644 --- a/tools/llvm-spirv/llvm-spirv.cpp +++ b/tools/llvm-spirv/llvm-spirv.cpp @@ -161,6 +161,10 @@ static cl::opt "for the translation from SPIR-V."), cl::Hidden); +static cl::opt SPIRVEmitFunctionPtrAddrSpace( + "spirv-emit-function-ptr-addr-space", cl::init(false), + cl::desc("Emit and consume CodeSectionINTEL for function pointers")); + using SPIRV::ExtensionID; #ifdef _SPIRV_SUPPORT_TEXT_FMT @@ -767,6 +771,9 @@ int main(int Ac, char **Av) { if (PreserveOCLKernelArgTypeMetadataThroughString.getNumOccurrences() != 0) Opts.setPreserveOCLKernelArgTypeMetadataThroughString(true); + if (SPIRVEmitFunctionPtrAddrSpace.getNumOccurrences() != 0) + Opts.setEmitFunctionPtrAddrSpace(true); + #ifdef _SPIRV_SUPPORT_TEXT_FMT if (ToText && (ToBinary || IsReverse || IsRegularization)) { errs() << "Cannot use -to-text with -to-binary, -r, -s\n";