From 46c9eb9ea9b4c6f4f0c6bbc1f2f66a9fbc7141b1 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Fri, 27 Sep 2024 10:49:56 +0200 Subject: [PATCH] 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 | 51 ++++- lib/SPIRV/SPIRVReader.h | 2 +- lib/SPIRV/SPIRVWriter.cpp | 8 + lib/SPIRV/libSPIRV/SPIRVInstruction.cpp | 5 + lib/SPIRV/libSPIRV/SPIRVModule.cpp | 7 + lib/SPIRV/libSPIRV/SPIRVModule.h | 5 + .../CodeSectionINTEL/alias.ll | 49 +++++ .../CodeSectionINTEL/bitcast.ll | 53 ++++++ .../const-function-pointer.ll | 67 +++++++ .../decor-func-ptr-arg-attr.ll | 67 +++++++ .../CodeSectionINTEL/fp-from-host.ll | 69 +++++++ .../function-pointer-as-function-arg.ll | 177 ++++++++++++++++++ .../function-pointer-dedicated-as.ll | 107 +++++++++++ .../CodeSectionINTEL/function-pointer.ll | 92 +++++++++ .../global-function-pointer.ll | 25 +++ .../CodeSectionINTEL/global_ctor_dtor.ll | 77 ++++++++ .../global_ctor_dtor_addrspace.ll | 34 ++++ .../CodeSectionINTEL/gv-func-ptr.ll | 40 ++++ .../non-uniform-function-pointer.ll | 138 ++++++++++++++ .../CodeSectionINTEL/referenced-indirectly.ll | 82 ++++++++ .../CodeSectionINTEL/select.ll | 142 ++++++++++++++ tools/llvm-spirv/llvm-spirv.cpp | 7 + 24 files changed, 1311 insertions(+), 9 deletions(-) create mode 100644 test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/alias.ll create mode 100644 test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/bitcast.ll create mode 100644 test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/const-function-pointer.ll create mode 100644 test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/decor-func-ptr-arg-attr.ll create mode 100644 test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/fp-from-host.ll create mode 100644 test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer-as-function-arg.ll create mode 100644 test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer-dedicated-as.ll create mode 100644 test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer.ll create mode 100644 test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/global-function-pointer.ll create mode 100644 test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/global_ctor_dtor.ll create mode 100644 test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/global_ctor_dtor_addrspace.ll create mode 100644 test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/gv-func-ptr.ll create mode 100644 test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/non-uniform-function-pointer.ll create mode 100644 test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/referenced-indirectly.ll create mode 100644 test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/select.ll diff --git a/include/LLVMSPIRVOpts.h b/include/LLVMSPIRVOpts.h index 73af0e97fa..08ef7b0e9e 100644 --- a/include/LLVMSPIRVOpts.h +++ b/include/LLVMSPIRVOpts.h @@ -236,6 +236,14 @@ class TranslatorOpts { PreserveOCLKernelArgTypeMetadataThroughString = Value; } + bool shouldEmitFunctionPtrAddrSpace() const noexcept { + return EmitFunctionPtrAddrSpace; + } + + void setEmitFunctionPtrAddrSpace(bool Value) noexcept { + EmitFunctionPtrAddrSpace = Value; + } + void setBuiltinFormat(BuiltinFormat Value) noexcept { SPIRVBuiltinFormat = Value; } @@ -287,6 +295,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; BuiltinFormat SPIRVBuiltinFormat = BuiltinFormat::Function; diff --git a/lib/SPIRV/SPIRVInternal.h b/lib/SPIRV/SPIRVInternal.h index 8c313f0ef9..427d947ce3 100644 --- a/lib/SPIRV/SPIRVInternal.h +++ b/lib/SPIRV/SPIRVInternal.h @@ -189,6 +189,7 @@ enum SPIRAddressSpace { SPIRAS_GlobalHost, SPIRAS_Input, SPIRAS_Output, + SPIRAS_CodeSectionINTEL, SPIRAS_Count, }; @@ -199,6 +200,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"); } @@ -215,6 +218,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 903ff95d43..90a358df05 100644 --- a/lib/SPIRV/SPIRVReader.cpp +++ b/lib/SPIRV/SPIRVReader.cpp @@ -348,16 +348,21 @@ Type *SPIRVToLLVM::transType(SPIRVType *T, bool UseTPT) { case internal::OpTypeTokenINTEL: return mapType(T, Type::getTokenTy(*Context)); case OpTypePointer: { - const unsigned AS = - SPIRSPIRVAddrSpaceMap::rmap(T->getPointerStorageClass()); + 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; Type *ElementTy = transType(T->getPointerElementType(), UseTPT); if (UseTPT) return TypedPointerType::get(ElementTy, AS); return mapType(T, PointerType::get(ElementTy, AS)); } case OpTypeUntypedPointerKHR: { - const unsigned AS = - SPIRSPIRVAddrSpaceMap::rmap(T->getPointerStorageClass()); + unsigned AS = SPIRSPIRVAddrSpaceMap::rmap(T->getPointerStorageClass()); + if (AS == SPIRAS_CodeSectionINTEL && !BM->shouldEmitFunctionPtrAddrSpace()) + AS = SPIRAS_Private; return mapType(T, PointerType::get(*Context, AS)); } case OpTypeVector: @@ -1469,6 +1474,17 @@ Value *SPIRVToLLVM::transValueWithoutDecoration(SPIRVValue *BV, Function *F, case OpTypeMatrix: 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: { @@ -1485,7 +1501,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)); } } @@ -1521,7 +1542,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: @@ -3046,7 +3070,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; @@ -3094,7 +3118,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); @@ -3501,6 +3525,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 19a5f52d1d..c1b74874c7 100644 --- a/lib/SPIRV/SPIRVReader.h +++ b/lib/SPIRV/SPIRVReader.h @@ -103,7 +103,7 @@ class SPIRVToLLVM : private BuiltinCallHelper { 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 ab84eb2b8c..30cd445e2c 100644 --- a/lib/SPIRV/SPIRVWriter.cpp +++ b/lib/SPIRV/SPIRVWriter.cpp @@ -655,6 +655,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(); @@ -750,6 +755,9 @@ SPIRVType *LLVMToSPIRVBase::transPointerType(SPIRVType *ET, unsigned AddrSpc) { return Loc->second; SPIRVType *TranslatedTy = nullptr; + if (AddrSpc == SPIRAS_CodeSectionINTEL && + !BM->shouldEmitFunctionPtrAddrSpace()) + return transPointerType(ET, SPIRAS_Private); if (BM->isAllowedToUseExtension(ExtensionID::SPV_KHR_untyped_pointers) && !(ET->isTypeArray() || ET->isTypeVector() || ET->isTypeStruct() || ET->isTypeImage() || ET->isTypeSampler() || ET->isTypePipe())) { diff --git a/lib/SPIRV/libSPIRV/SPIRVInstruction.cpp b/lib/SPIRV/libSPIRV/SPIRVInstruction.cpp index 67d8b1ca24..048722eb1f 100644 --- a/lib/SPIRV/libSPIRV/SPIRVInstruction.cpp +++ b/lib/SPIRV/libSPIRV/SPIRVInstruction.cpp @@ -282,6 +282,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 7261f9b3a7..94d0818528 100644 --- a/lib/SPIRV/libSPIRV/SPIRVModule.cpp +++ b/lib/SPIRV/libSPIRV/SPIRVModule.cpp @@ -142,6 +142,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 8915eb7ecc..9a9958ecdd 100644 --- a/lib/SPIRV/libSPIRV/SPIRVModule.h +++ b/lib/SPIRV/libSPIRV/SPIRVModule.h @@ -142,6 +142,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; @@ -554,6 +555,10 @@ class SPIRVModule { .shouldPreserveOCLKernelArgTypeMetadataThroughString(); } + bool shouldEmitFunctionPtrAddrSpace() const noexcept { + return TranslationOpts.shouldEmitFunctionPtrAddrSpace(); + } + bool preserveAuxData() const noexcept { return TranslationOpts.preserveAuxData(); } diff --git a/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/alias.ll b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/alias.ll new file mode 100644 index 0000000000..108b04ef58 --- /dev/null +++ b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/alias.ll @@ -0,0 +1,49 @@ +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv -spirv-ext=+SPV_INTEL_function_pointers -spirv-text %t.bc -o - | FileCheck %s --check-prefix=CHECK-SPIRV +; RUN: llvm-spirv -spirv-ext=+SPV_INTEL_function_pointers %t.bc -o %t.spv +; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o - | llvm-dis -o - | FileCheck %s --check-prefix=CHECK-LLVM + +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" + +; Check that aliases are dereferenced and translated to their aliasee values +; when used since they can't be translated directly. + +; CHECK-SPIRV-DAG: Name [[#FOO:]] "foo" +; CHECK-SPIRV-DAG: Name [[#BAR:]] "bar" +; CHECK-SPIRV-DAG: Name [[#Y:]] "y" +; CHECK-SPIRV-DAG: Name [[#FOOPTR:]] "foo.alias" +; CHECK-SPIRV-DAG: Decorate [[#FOO]] LinkageAttributes "foo" Export +; CHECK-SPIRV-DAG: Decorate [[#BAR]] LinkageAttributes "bar" Export +; CHECK-SPIRV-DAG: TypeInt [[#I32:]] 32 0 +; CHECK-SPIRV-DAG: TypeInt [[#I64:]] 64 0 +; CHECK-SPIRV-DAG: TypeFunction [[#FOO_TYPE:]] [[#I32]] [[#I32]] +; CHECK-SPIRV-DAG: TypeVoid [[#VOID:]] +; CHECK-SPIRV-DAG: TypePointer [[#I64PTR:]] 7 [[#I64]] +; CHECK-SPIRV-DAG: TypeFunction [[#BAR_TYPE:]] [[#VOID]] [[#I64PTR]] +; CHECK-SPIRV-DAG: TypePointer [[#FOOPTR_TYPE:]] 7 [[#FOO_TYPE]] +; CHECK-SPIRV-DAG: ConstantFunctionPointerINTEL [[#FOOPTR_TYPE]] [[#FOOPTR]] [[#FOO]] + +; CHECK-SPIRV: Function [[#I32]] [[#FOO]] 0 [[#FOO_TYPE]] + +; CHECK-SPIRV: Function [[#VOID]] [[#BAR]] 0 [[#BAR_TYPE]] +; CHECK-SPIRV: FunctionParameter [[#I64PTR]] [[#Y]] +; CHECK-SPIRV: ConvertPtrToU [[#I64]] [[#PTRTOINT:]] [[#FOOPTR]] +; CHECK-SPIRV: Store [[#Y]] [[#PTRTOINT]] 2 8 + +; CHECK-LLVM: define spir_func i32 @foo(i32 %x) addrspace(9) + +; CHECK-LLVM: define spir_kernel void @bar(ptr %y) +; CHECK-LLVM: [[PTRTOINT:%.*]] = ptrtoint ptr addrspace(9) @foo to i64 +; CHECK-LLVM: store i64 [[PTRTOINT]], ptr %y, align 8 + +define spir_func i32 @foo(i32 %x) { + ret i32 %x +} + +@foo.alias = internal alias i32 (i32), ptr @foo + +define spir_kernel void @bar(ptr %y) { + store i64 ptrtoint (ptr @foo.alias to i64), ptr %y + ret void +} diff --git a/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/bitcast.ll b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/bitcast.ll new file mode 100644 index 0000000000..0d71b99cc6 --- /dev/null +++ b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/bitcast.ll @@ -0,0 +1,53 @@ +; OpenCL C source: +; char foo(char a) { +; return a; +; } +; void bar() { +; int (*fun_ptr)(int) = &foo; +; fun_ptr(0); +; } + +; 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-DAG: TypeInt [[#I8:]] 8 +; CHECK-SPIRV-DAG: TypeInt [[#I32:]] 32 +; CHECK-SPIRV-DAG: TypeFunction [[#FOO_TY:]] [[#I8]] [[#I8]] +; CHECK-SPIRV-DAG: TypeFunction [[#DEST_TY:]] [[#I32]] [[#I32]] +; CHECK-SPIRV-DAG: TypePointer [[#DEST_TY_PTR:]] [[#]] [[#DEST_TY]] +; CHECK-SPIRV-DAG: 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 ptr addrspace(9) @foo to ptr 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 = call i32 @foo(i32 0) + 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/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/const-function-pointer.ll b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/const-function-pointer.ll new file mode 100644 index 0000000000..07d46b43dc --- /dev/null +++ b/test/extensions/INTEL/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-DAG: TypeInt [[Int32:[0-9]+]] 32 +; CHECK-SPIRV-DAG: 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 addrspace(1) constant [2 x ptr addrspace(9)] [ptr addrspace(9) @f1, ptr addrspace(9) @f2], align 16 +; CHECK-LLVM: %[[Idx:[a-z0-9]+]] = getelementptr inbounds [2 x ptr addrspace(9)], ptr addrspace(1) @__const.main.funcs, i64 0, i64 %{{[a-z0-9]+}} +; CHECK-LLVM: %[[FuncPtr:[a-z0-9]+]] = load ptr addrspace(9), ptr addrspace(1) %[[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 addrspace(1) constant [2 x ptr] [ptr @f1, ptr @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 ptr], ptr addrspace(1) @__const.main.funcs, i64 0, i64 %idxprom + %0 = load ptr, ptr addrspace(1) %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/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/decor-func-ptr-arg-attr.ll b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/decor-func-ptr-arg-attr.ll new file mode 100644 index 0000000000..9e07baf650 --- /dev/null +++ b/test/extensions/INTEL/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(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 { ptr } +%"range" = type { %"array" } +%"array" = type { [1 x i64] } +%wrapper_class = type { ptr addrspace(1) } +%wrapper_class.0 = type { ptr addrspace(1) } + +$RoundedRangeKernel = comdat any + +; Function Attrs: nounwind +define spir_func void @inc_function(ptr byval(%"multi_ptr") noalias nocapture %ptr) #0 { +entry: + ret void +} + + +; Function Attrs: convergent norecurse +define weak_odr dso_local spir_kernel void @RoundedRangeKernel(ptr byval(%"range") align 8 %_arg_NumWorkItems, i1 zeroext %_arg_, ptr byval(%wrapper_class) align 8 %_arg_1, ptr 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_, ptr @inc_function, ptr null + call spir_func void %cond.i.i(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 = !{ptr @inc_function} diff --git a/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/fp-from-host.ll b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/fp-from-host.ll new file mode 100644 index 0000000000..aacdcc4fbc --- /dev/null +++ b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/fp-from-host.ll @@ -0,0 +1,69 @@ +; 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 %t.spv -spirv-emit-function-ptr-addr-space -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(ptr 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(ptr addrspace(1) %fp, ptr 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, ptr addrspace(1) %data, i64 1 + %0 = load i32, ptr addrspace(1) %arrayidx, align 4, !tbaa !8 + %1 = load i32, ptr addrspace(1) %fp, align 4, !tbaa !8 + %2 = inttoptr i32 %1 to ptr + %call = call spir_func i32 %2(i32 %0) #1 + store i32 %call, ptr addrspace(1) %data, 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/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer-as-function-arg.ll b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer-as-function-arg.ll new file mode 100644 index 0000000000..a933712f4d --- /dev/null +++ b/test/extensions/INTEL/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(ptr addrspace(9) %[[F:.*]], +; CHECK-LLVM: %[[F_ADDR:.*]] = alloca ptr addrspace(9) +; CHECK-LLVM: store ptr addrspace(9) %[[F]], ptr %[[F_ADDR]] +; CHECK-LLVM: %[[F_LOADED:.*]] = load ptr addrspace(9), ptr %[[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: %{{.*}} = alloca ptr +; CHECK-LLVM: %[[FP:.*]] = alloca ptr addrspace(9) +; CHECK-LLVM: store ptr addrspace(9) @foo, ptr %[[FP]] +; CHECK-LLVM: store ptr addrspace(9) @bar, ptr %[[FP]] +; CHECK-LLVM: %[[FP_LOADED:.*]] = load ptr addrspace(9), ptr %[[FP]] +; CHECK-LLVM: call spir_func i32 @helper(ptr 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(ptr %f, i32 %arg) #0 { +entry: + %f.addr = alloca ptr, align 8 + %arg.addr = alloca i32, align 4 + store ptr %f, ptr %f.addr, align 8 + store i32 %arg, ptr %arg.addr, align 4 + %0 = load ptr, ptr %f.addr, align 8 + %1 = load i32, ptr %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, ptr %v.addr, align 4 + %0 = load i32, ptr %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, ptr %v.addr, align 4 + %0 = load i32, ptr %v.addr, align 4 + %add = add nsw i32 %0, 2 + ret i32 %add +} + +; Function Attrs: convergent noinline nounwind optnone +define spir_kernel void @test(ptr 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 ptr addrspace(1), align 8 + %control.addr = alloca i32, align 4 + %fp = alloca ptr, align 8 + store ptr addrspace(1) %data, ptr %data.addr, align 8 + store i32 %control, ptr %control.addr, align 4 + %call = call spir_func i64 @_Z13get_global_idj(i32 0) #4 + %0 = load i32, ptr %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 ptr @foo, ptr %fp, align 8 + br label %if.end + +if.else: ; preds = %entry + store ptr @bar, ptr %fp, align 8 + br label %if.end + +if.end: ; preds = %if.else, %if.then + %1 = load ptr, ptr %fp, align 8 + %2 = load ptr addrspace(1), ptr %data.addr, align 8 + %call2 = call spir_func i64 @_Z13get_global_idj(i32 0) #4 + %arrayidx = getelementptr inbounds i32, ptr addrspace(1) %2, i64 %call2 + %3 = load i32, ptr addrspace(1) %arrayidx, align 4 + %call3 = call spir_func i32 @helper(ptr %1, i32 %3) #3 + %4 = load ptr addrspace(1), ptr %data.addr, align 8 + %call4 = call spir_func i64 @_Z13get_global_idj(i32 0) #4 + %arrayidx5 = getelementptr inbounds i32, ptr addrspace(1) %4, i64 %call4 + store i32 %call3, ptr 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/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer-dedicated-as.ll b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer-dedicated-as.ll new file mode 100644 index 0000000000..d728d91d19 --- /dev/null +++ b/test/extensions/INTEL/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, ptr %arg.addr, align 4 + %0 = load i32, ptr %arg.addr, align 4 + %add = add nsw i32 %0, 10 + ret i32 %add +} + +; Function Attrs: noinline nounwind +define spir_kernel void @test(ptr 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 ptr addrspace(1), align 8 + %input.addr = alloca i32, align 4 + %fp = alloca ptr addrspace(9), align 8 + store ptr addrspace(1) %data, ptr %data.addr, align 8 + store i32 %input, ptr %input.addr, align 4 + store ptr addrspace(9) @foo, ptr %fp, align 8 + %0 = load ptr addrspace(9), ptr %fp, align 8 + %1 = load i32, ptr %input.addr, align 4 + %call = call spir_func addrspace(9) i32 %0(i32 %1) + %2 = load ptr addrspace(1), ptr %data.addr, align 8 + store i32 %call, ptr 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/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer.ll b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer.ll new file mode 100644 index 0000000000..bd2ceb32d4 --- /dev/null +++ b/test/extensions/INTEL/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 ptr addrspace(9) +; CHECK-LLVM: store ptr addrspace(9) @foo, ptr %fp +; CHECK-LLVM: %0 = load ptr addrspace(9), ptr %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, ptr %arg.addr, align 4 + %0 = load i32, ptr %arg.addr, align 4 + %add = add nsw i32 %0, 10 + ret i32 %add +} + +; Function Attrs: convergent noinline nounwind optnone +define spir_kernel void @test(ptr 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 ptr addrspace(1), align 8 + %input.addr = alloca i32, align 4 + %fp = alloca ptr, align 8 + store ptr addrspace(1) %data, ptr %data.addr, align 8 + store i32 %input, ptr %input.addr, align 4 + store ptr @foo, ptr %fp, align 8 + %0 = load ptr, ptr %fp, align 8 + %1 = load i32, ptr %input.addr, align 4 + %call = call spir_func i32 %0(i32 %1) #2 + %2 = load ptr addrspace(1), ptr %data.addr, align 8 + store i32 %call, ptr 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/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/global-function-pointer.ll b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/global-function-pointer.ll new file mode 100644 index 0000000000..59e6b238d3 --- /dev/null +++ b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/global-function-pointer.ll @@ -0,0 +1,25 @@ +; RUN: llvm-as < %s | llvm-spirv -spirv-ext=+SPV_INTEL_function_pointers -o %t.spv +; RUN: llvm-spirv %t.spv -spirv-ext=+SPV_INTEL_function_pointers -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV +; RUN: llvm-spirv -r %t.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM + +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" + + +; CHECK-SPIRV: Capability FunctionPointersINTEL +; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers" +; CHECK-SPIRV: TypeFunction [[#FOO_TY:]] [[#]] [[#]] +; CHECK-SPIRV: TypePointer [[#FOO_TY_PTR:]] [[#]] [[#FOO_TY]] +; CHECK-SPIRV: ConstantFunctionPointerINTEL [[#FOO_TY_PTR]] [[#FOO_PTR:]] [[#FOO:]] +; CHECK-SPIRV: Function [[#]] [[#]] [[#]] [[#FOO_TY]] + +; CHECK-LLVM: @two = internal addrspace(1) global ptr @_Z4barrii +; CHECK-LLVM: define spir_func i32 @_Z4barrii(i32 %[[#]], i32 %[[#]]) + +@two = internal addrspace(1) global ptr @_Z4barrii, align 8 + +; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn writeonly +define protected spir_func noundef i32 @_Z4barrii(i32 %0, i32 %1) { +entry: + ret i32 1 +} diff --git a/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/global_ctor_dtor.ll b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/global_ctor_dtor.ll new file mode 100644 index 0000000000..26c8864a12 --- /dev/null +++ b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/global_ctor_dtor.ll @@ -0,0 +1,77 @@ +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -spirv-text -o - | FileCheck %s +; 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 + + +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" + +@llvm.global_ctors = appending addrspace(1) global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 65535, ptr @_GLOBAL__sub_I_test.cpp.ctor, ptr null }] +@llvm.global_dtors = appending addrspace(1) global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 65535, ptr @_GLOBAL__sub_I_test.cpp.dtor, ptr null }] +@_ZL15DIVERGENCE = internal addrspace(1) global i32 0, align 4 + +; CHECK: Name [[NameCtor:[0-9]+]] "_GLOBAL__sub_I_test.cpp.ctor" +; CHECK: Name [[Name1:[0-9]+]] "llvm.global_ctors" +; CHECK: Name [[NameDtor:[0-9]+]] "_GLOBAL__sub_I_test.cpp.dtor" +; CHECK: Name [[Name2:[0-9]+]] "llvm.global_dtors" +; CHECK: Name [[NameInit:[0-9]+]] "__cxx_global_var_init" + +; CHECK-DAG: Decorate {{[0-9]+}} LinkageAttributes "llvm.global_ctors" Export +; CHECK-DAG: Decorate {{[0-9]+}} LinkageAttributes "llvm.global_dtors" Export + +; CHECK: TypeFunction {{[0-9]+}} [[TF:[0-9]+]] + +; CHECK: TypePointer [[TP:[0-9]+]] +; CHECK: ConstantFunctionPointerINTEL [[TP]] [[FPCtor:[0-9]+]] [[NameCtor]] +; CHECK: SpecConstantOp {{[0-9]+}} [[FPCtorI8:[0-9]+]] 124 [[FPCtor]] +; CHECK: ConstantComposite {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} [[FPCtorI8]] +; CHECK: ConstantFunctionPointerINTEL [[TP]] [[FPDtor:[0-9]+]] [[NameDtor]] +; CHECK: SpecConstantOp {{[0-9]+}} [[FPDtorI8:[0-9]+]] 124 [[FPDtor]] +; CHECK: ConstantComposite {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} [[FPDtorI8]] + +; CHECK: 5 Function [[TF]] [[NameCtor]] 0 +; CHECK-EMPTY: +; CHECK-NEXT: Label {{[0-9]+}} +; CHECK-NEXT: FunctionCall {{[0-9]+}} {{[0-9]+}} [[NameInit]] +; CHECK-NEXT: Return +; CHECK-EMPTY: +; CHECK-NEXT: FunctionEnd + +; Function Attrs: nounwind sspstrong +define internal void @_GLOBAL__sub_I_test.cpp.ctor() #0 { + call void @__cxx_global_var_init() + ret void +} + +; Function Attrs: nounwind sspstrong +define internal void @__cxx_global_var_init() #0 { + store i32 0, ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL15DIVERGENCE to ptr addrspace(4)), align 4 + ret void +} + +; CHECK: 5 Function [[TF]] [[NameDtor]] 0 +; CHECK-EMPTY: +; CHECK-NEXT: Label {{[0-9]+}} +; CHECK-NEXT: Return +; CHECK-EMPTY: +; CHECK-NEXT: FunctionEnd + +; Function Attrs: nounwind sspstrong +define internal void @_GLOBAL__sub_I_test.cpp.dtor() #0 { + ret void +} + +; Ctor/dtor functions should not be serialized twice. +; CHECK-NOT: 5 Function [[TF]] [[NameCtor]] 0 +; CHECK-NOT: 5 Function [[TF]] [[NameDtor]] 0 + +attributes #0 = { nounwind sspstrong "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" } + +!opencl.spir.version = !{!0} +!opencl.used.extensions = !{!1} +!opencl.used.optional.core.features = !{!1} +!opencl.compiler.options = !{!1} + +!0 = !{i32 1, i32 2} +!1 = !{} diff --git a/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/global_ctor_dtor_addrspace.ll b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/global_ctor_dtor_addrspace.ll new file mode 100644 index 0000000000..24cdc7278f --- /dev/null +++ b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/global_ctor_dtor_addrspace.ll @@ -0,0 +1,34 @@ +; +; This test case checks that LLVM -> SPIR-V -> LLVM translation +; produces valid LLVM IR, where intrinsic global variables +; llvm.global_ctors and llvm.global_dtors, defined with non-default +; address space have correct (appending) linkage. +; +; No additional checks are needed in addition to simple translation +; to and from SPIR-V. In case of an error newly produced LLVM module +; validation would fail with the message: +; +; "Fails to verify module: invalid linkage for intrinsic global variable". +; +; +; 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 + +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" + +@llvm.global_ctors = appending addrspace(1) global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 65535, ptr @_GLOBAL__sub_I_test.cpp.ctor, ptr null }] +@llvm.global_dtors = appending addrspace(1) global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 65535, ptr @_GLOBAL__sub_I_test.cpp.dtor, ptr null }] + +; Function Attrs: nounwind sspstrong +define internal void @_GLOBAL__sub_I_test.cpp.ctor() #0 { + ret void +} + +; Function Attrs: nounwind sspstrong +define internal void @_GLOBAL__sub_I_test.cpp.dtor() #0 { + ret void +} + +attributes #0 = { nounwind sspstrong } diff --git a/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/gv-func-ptr.ll b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/gv-func-ptr.ll new file mode 100644 index 0000000000..8d5b3b8e9b --- /dev/null +++ b/test/extensions/INTEL/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 ptr addrspace(4)] } + +; CHECK-LLVM: @A = addrspace(1) constant %structtype.3 { [1 x ptr addrspace(4)] [ptr addrspace(4) addrspacecast (ptr addrspace(9) @foo to ptr addrspace(4))] }, align 8 + +@A = linkonce_odr addrspace(1) constant %structtype.3 { [1 x ptr addrspace(4)] [ptr addrspace(4) addrspacecast (ptr @foo to ptr addrspace(4))] }, align 8 + +; Function Attrs: nounwind +define linkonce_odr spir_func void @foo() #0 { +entry: +; CHECK-LLVM: %0 = getelementptr inbounds %structtype.3, ptr addrspace(1) @A, i64 0, i32 0, i64 2 + %0 = getelementptr inbounds %structtype.3, ptr 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/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/non-uniform-function-pointer.ll b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/non-uniform-function-pointer.ll new file mode 100644 index 0000000000..f4e46456f9 --- /dev/null +++ b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/non-uniform-function-pointer.ll @@ -0,0 +1,138 @@ +; 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 ptr addrspace(9) +; CHECK-LLVM: store ptr addrspace(9) @foo, ptr %fp +; CHECK-LLVM: store ptr addrspace(9) @bar, ptr %fp +; CHECK-LLVM: %[[FP:.*]] = load ptr addrspace(9), ptr %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, ptr %v.addr, align 4 + %0 = load i32, ptr %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, ptr %v.addr, align 4 + %0 = load i32, ptr %v.addr, align 4 + %add = add nsw i32 %0, 2 + ret i32 %add +} + +; Function Attrs: convergent noinline nounwind optnone +define spir_kernel void @test(ptr 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 ptr addrspace(1), align 8 + %control.addr = alloca i32, align 4 + %fp = alloca ptr, align 8 + store ptr addrspace(1) %data, ptr %data.addr, align 8 + store i32 %control, ptr %control.addr, align 4 + %call = call spir_func i64 @_Z13get_global_idj(i32 0) #3 + %0 = load i32, ptr %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 ptr @foo, ptr %fp, align 8 + br label %if.end + +if.else: ; preds = %entry + store ptr @bar, ptr %fp, align 8 + br label %if.end + +if.end: ; preds = %if.else, %if.then + %1 = load ptr, ptr %fp, align 8 + %2 = load ptr addrspace(1), ptr %data.addr, align 8 + %call2 = call spir_func i64 @_Z13get_global_idj(i32 0) #3 + %arrayidx = getelementptr inbounds i32, ptr addrspace(1) %2, i64 %call2 + %3 = load i32, ptr addrspace(1) %arrayidx, align 4 + %call3 = call spir_func i32 %1(i32 %3) #4 + %4 = load ptr addrspace(1), ptr %data.addr, align 8 + %call4 = call spir_func i64 @_Z13get_global_idj(i32 0) #3 + %arrayidx5 = getelementptr inbounds i32, ptr addrspace(1) %4, i64 %call4 + store i32 %call3, ptr 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/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/referenced-indirectly.ll b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/referenced-indirectly.ll new file mode 100644 index 0000000000..40730bdd53 --- /dev/null +++ b/test/extensions/INTEL/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, ptr %arg.addr, align 4 + %0 = load i32, ptr %arg.addr, align 4 + %add = add nsw i32 %0, 10 + ret i32 %add +} + +; Function Attrs: convergent noinline nounwind optnone +define spir_kernel void @test(ptr 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 ptr addrspace(1), align 8 + %input.addr = alloca i32, align 4 + %fp = alloca ptr, align 8 + store ptr addrspace(1) %data, ptr %data.addr, align 8 + store i32 %input, ptr %input.addr, align 4 + store ptr @foo, ptr %fp, align 8 + %0 = load ptr, ptr %fp, align 8 + %1 = load i32, ptr %input.addr, align 4 + %call = call spir_func i32 %0(i32 %1) #2 + %2 = load ptr addrspace(1), ptr %data.addr, align 8 + store i32 %call, ptr 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/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/select.ll b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/select.ll new file mode 100644 index 0000000000..67eebd988e --- /dev/null +++ b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/select.ll @@ -0,0 +1,142 @@ +; 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 ptr addrspace(9) +; CHECK-LLVM: %[[SELECT:.*]] = select i1 %{{.*}}, ptr addrspace(9) @_Z3barii, ptr addrspace(9) @_Z3bazii +; CHECK-LLVM: store ptr addrspace(9) %[[SELECT]], ptr %[[FPTR_ALLOCA]] +; CHECK-LLVM: %[[FPTR:.*]] = load ptr addrspace(9), ptr %[[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(ptr addrspace(1) %_arg_, ptr byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, ptr byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, ptr 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 ptr, 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 + call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %agg.tmp2.i) + call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %agg.tmp3.i) + %0 = addrspacecast ptr %agg.tmp2.i to ptr addrspace(4) + %ptrint4.i = ptrtoint ptr addrspace(4) %0 to i64 + %maskedptr5.i = and i64 %ptrint4.i, 7 + %maskcond6.i = icmp eq i64 %maskedptr5.i, 0 + %1 = addrspacecast ptr %agg.tmp3.i to ptr addrspace(4) + %ptrint.i = ptrtoint ptr addrspace(4) %1 to i64 + %maskedptr.i = and i64 %ptrint.i, 7 + %maskcond.i = icmp eq i64 %maskedptr.i, 0 + call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %agg.tmp2.i) + call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %agg.tmp3.i) + %2 = load i64, ptr %_arg_3, align 8 + %add.ptr.i = getelementptr inbounds i32, ptr addrspace(1) %_arg_, i64 %2 + %3 = addrspacecast ptr %agg.tmp6 to ptr addrspace(4) + %ptrint = ptrtoint ptr addrspace(4) %3 to i64 + %maskedptr = and i64 %ptrint, 7 + %maskcond = icmp eq i64 %maskedptr, 0 + %4 = load <3 x i64>, ptr addrspace(4) addrspacecast (ptr addrspace(1) @__spirv_BuiltInGlobalInvocationId to ptr addrspace(4)), align 32, !noalias !8 + %5 = extractelement <3 x i64> %4, i64 0 + store i64 %5, ptr addrspace(4) %3, align 8, !tbaa !15, !alias.scope !8 + call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %ref.tmp.i) #4 + %6 = addrspacecast ptr %ref.tmp.i to ptr addrspace(4) + %ptrint.i2 = ptrtoint ptr addrspace(4) %6 to i64 + %maskedptr.i3 = and i64 %ptrint.i2, 7 + %maskcond.i4 = icmp eq i64 %maskedptr.i3, 0 + %rem.i.i = and i64 %5, 1 + %cmp.i.i = icmp eq i64 %rem.i.i, 0 + call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %ref.tmp.i) #4 + %_Z3barii._Z3bazii.i = select i1 %cmp.i.i, ptr @_Z3barii, ptr @_Z3bazii + store ptr %_Z3barii._Z3bazii.i, ptr %fptr.alloca, align 8 + %fptr = load ptr, ptr %fptr.alloca, align 8 + %call4.i = call spir_func i32 %fptr(i32 10, i32 10), !callees !19 + %arrayidx.i3.i = getelementptr inbounds i32, ptr addrspace(1) %add.ptr.i, i64 %5 + %arrayidx.ascast.i.i = addrspacecast ptr addrspace(1) %arrayidx.i3.i to ptr addrspace(4) + store i32 %call4.i, ptr addrspace(4) %arrayidx.ascast.i.i, align 4, !tbaa !20 + ret void +} + +; Function Attrs: argmemonly nounwind willreturn +declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 + +; Function Attrs: argmemonly nounwind willreturn +declare void @llvm.lifetime.end.p0(i64 immarg, ptr 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 = !{ptr @_Z3barii, ptr @_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 828169d60f..d174aafb68 100644 --- a/tools/llvm-spirv/llvm-spirv.cpp +++ b/tools/llvm-spirv/llvm-spirv.cpp @@ -161,6 +161,10 @@ static cl::opt SPIRVToolsDis("spirv-tools-dis", cl::init(false), cl::desc("Emit textual assembly using SPIRV-Tools")); +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 @@ -845,6 +849,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";