Hello community, here is the log from the commit of package beignet for openSUSE:Factory checked in at 2017-05-17 10:54:32 ++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ Comparing /work/SRC/openSUSE:Factory/beignet (Old) and /work/SRC/openSUSE:Factory/.beignet.new (New) ++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++
Package is "beignet" Wed May 17 10:54:32 2017 rev:11 rq:495212 version:1.3.1 Changes: -------- --- /work/SRC/openSUSE:Factory/beignet/beignet.changes 2017-03-20 17:09:34.441775042 +0100 +++ /work/SRC/openSUSE:Factory/.beignet.new/beignet.changes 2017-05-17 10:54:34.109018441 +0200 @@ -1,0 +2,12 @@ +Thu May 4 19:47:06 UTC 2017 - [email protected] + +- Add patches for LLVM 4.0 support + * 0001-Backend-Remove-old-llvm-support-code.patch + * 0002-Backend-Fix-an-include-file-error-problem.patch + * 0003-Backend-Refine-GEP-lowering-code.patch + * 0004-Backend-Refine-LLVM-version-check-macro.patch + * 0005-Backend-Refine-FCmp-one-and-une.patch + * 0006-utest-fix-image-qualifier-of-compiler_fill_gl_image-.patch + * 0007-Backend-Add-LLVM40-support.patch + +------------------------------------------------------------------- New: ---- 0001-Backend-Remove-old-llvm-support-code.patch 0002-Backend-Fix-an-include-file-error-problem.patch 0003-Backend-Refine-GEP-lowering-code.patch 0004-Backend-Refine-LLVM-version-check-macro.patch 0005-Backend-Refine-FCmp-one-and-une.patch 0006-utest-fix-image-qualifier-of-compiler_fill_gl_image-.patch 0007-Backend-Add-LLVM40-support.patch ++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ Other differences: ------------------ ++++++ beignet.spec ++++++ --- /var/tmp/diff_new_pack.br5NDD/_old 2017-05-17 10:54:34.992894098 +0200 +++ /var/tmp/diff_new_pack.br5NDD/_new 2017-05-17 10:54:34.996893536 +0200 @@ -25,6 +25,13 @@ Url: https://01.org/beignet/ Source0: https://01.org/sites/default/files/%{name}-%{version}-source.tar.gz Source99: beignet-rpmlintrc +Patch0: 0001-Backend-Remove-old-llvm-support-code.patch +Patch1: 0002-Backend-Fix-an-include-file-error-problem.patch +Patch2: 0003-Backend-Refine-GEP-lowering-code.patch +Patch3: 0004-Backend-Refine-LLVM-version-check-macro.patch +Patch4: 0005-Backend-Refine-FCmp-one-and-une.patch +Patch5: 0006-utest-fix-image-qualifier-of-compiler_fill_gl_image-.patch +Patch6: 0007-Backend-Add-LLVM40-support.patch BuildRequires: cmake BuildRequires: gcc-c++ BuildRequires: llvm >= 3.3 @@ -61,6 +68,13 @@ %prep %setup -q -n Beignet-%{version}-Source +%patch0 -p1 +%patch1 -p1 +%patch2 -p1 +%patch3 -p1 +%patch4 -p1 +%patch5 -p1 +%patch6 -p1 %build %cmake \ ++++++ 0001-Backend-Remove-old-llvm-support-code.patch ++++++ >From c6497b19e7b32bcb0d1f4aceb1988488dfec639d Mon Sep 17 00:00:00 2001 From: Pan Xiuli <[email protected]> Date: Fri, 17 Mar 2017 14:15:58 +0800 Subject: [PATCH 1/7] Backend: Remove old llvm support code. LLVM 3.3 or older is not supportted by Beignet now, and we need delete these codes. Signed-off-by: Pan Xiuli <[email protected]> Reviewed-by: Yang Rong <[email protected]> --- backend/src/backend/gen_program.cpp | 6 ------ backend/src/backend/program.cpp | 30 ------------------------------ backend/src/llvm/llvm_gen_backend.cpp | 22 ---------------------- backend/src/llvm/llvm_printf_parser.cpp | 6 ------ backend/src/llvm/llvm_profiling.cpp | 20 -------------------- backend/src/llvm/llvm_scalarize.cpp | 6 ------ 6 files changed, 90 deletions(-) diff --git a/backend/src/backend/gen_program.cpp b/backend/src/backend/gen_program.cpp index 073ede64..376342b8 100644 --- a/backend/src/backend/gen_program.cpp +++ b/backend/src/backend/gen_program.cpp @@ -24,15 +24,9 @@ #ifdef GBE_COMPILER_AVAILABLE #include "llvm/Config/llvm-config.h" -#if LLVM_VERSION_MAJOR == 3 && LLVM_VERSION_MINOR <= 2 -#include "llvm/LLVMContext.h" -#include "llvm/Module.h" -#include "llvm/DataLayout.h" -#else #include "llvm/IR/LLVMContext.h" #include "llvm/IR/Module.h" #include "llvm/IR/DataLayout.h" -#endif /* LLVM_VERSION_MINOR <= 2 */ #include "llvm-c/Linker.h" #include "llvm/Transforms/Utils/Cloning.h" #include "llvm/Bitcode/ReaderWriter.h" diff --git a/backend/src/backend/program.cpp b/backend/src/backend/program.cpp index 8b8abc4d..6b5fce06 100644 --- a/backend/src/backend/program.cpp +++ b/backend/src/backend/program.cpp @@ -52,33 +52,16 @@ #include <mutex> #ifdef GBE_COMPILER_AVAILABLE -/* Not defined for LLVM 3.0 */ -#if !defined(LLVM_VERSION_MAJOR) -#define LLVM_VERSION_MAJOR 3 -#endif /* !defined(LLVM_VERSION_MAJOR) */ - -/* Not defined for LLVM 3.0 */ -#if !defined(LLVM_VERSION_MINOR) -#define LLVM_VERSION_MINOR 0 -#endif /* !defined(LLVM_VERSION_MINOR) */ #include <clang/CodeGen/CodeGenAction.h> #include <clang/Frontend/CompilerInstance.h> #include <clang/Frontend/CompilerInvocation.h> -#if LLVM_VERSION_MINOR <= 1 -#include <clang/Frontend/DiagnosticOptions.h> -#else #include <clang/Basic/DiagnosticOptions.h> -#endif /* LLVM_VERSION_MINOR <= 1 */ #include <clang/Frontend/TextDiagnosticPrinter.h> #include <clang/Basic/TargetInfo.h> #include <clang/Basic/TargetOptions.h> #include <llvm/ADT/IntrusiveRefCntPtr.h> -#if LLVM_VERSION_MINOR <= 2 -#include <llvm/Module.h> -#else #include <llvm/IR/Module.h> -#endif /* LLVM_VERSION_MINOR <= 2 */ #include <llvm/Bitcode/ReaderWriter.h> #include <llvm/Support/raw_ostream.h> #endif @@ -686,10 +669,6 @@ namespace gbe { args.push_back("-disable-llvm-optzns"); if(bFastMath) args.push_back("-D __FAST_RELAXED_MATH__=1"); -#if LLVM_VERSION_MINOR <= 2 - args.push_back("-triple"); - args.push_back("nvptx"); -#else args.push_back("-x"); args.push_back("cl"); args.push_back("-triple"); @@ -698,7 +677,6 @@ namespace gbe { args.push_back("-fblocks"); } else args.push_back("spir"); -#endif /* LLVM_VERSION_MINOR <= 2 */ args.push_back("stringInput.cl"); args.push_back("-ffp-contract=on"); if(OCL_DEBUGINFO) args.push_back("-g"); @@ -791,11 +769,7 @@ namespace gbe { std::string err; llvm::raw_fd_ostream ostream (dumpLLVMFileName.c_str(), err, - #if LLVM_VERSION_MINOR == 3 - 0 - #else llvm::sys::fs::F_None - #endif ); if (err.empty()) { @@ -807,11 +781,7 @@ namespace gbe { std::string err; llvm::raw_fd_ostream ostream (dumpSPIRBinaryName.c_str(), err, - #if LLVM_VERSION_MINOR == 3 - 0 - #else llvm::sys::fs::F_None - #endif ); if (err.empty()) llvm::WriteBitcodeToFile(*out_module, ostream); diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp index d7dabe3c..6fa1ae02 100644 --- a/backend/src/llvm/llvm_gen_backend.cpp +++ b/backend/src/llvm/llvm_gen_backend.cpp @@ -746,9 +746,6 @@ namespace gbe void visitVAArgInst(VAArgInst &I) {NOT_SUPPORTED;} void visitSwitchInst(SwitchInst &I) {NOT_SUPPORTED;} void visitInvokeInst(InvokeInst &I) {NOT_SUPPORTED;} -#if LLVM_VERSION_MINOR == 0 - void visitUnwindInst(UnwindInst &I) {NOT_SUPPORTED;} -#endif /* __LLVM_30__ */ void visitResumeInst(ResumeInst &I) {NOT_SUPPORTED;} void visitInlineAsm(CallInst &I) {NOT_SUPPORTED;} void visitIndirectBrInst(IndirectBrInst &I) {NOT_SUPPORTED;} @@ -1750,7 +1747,6 @@ namespace gbe { GBE_ASSERT(dyn_cast<ConstantExpr>(CPV) == NULL); -#if LLVM_VERSION_MINOR > 0 ConstantDataSequential *seq = dyn_cast<ConstantDataSequential>(CPV); if (seq) { @@ -1773,7 +1769,6 @@ namespace gbe GBE_ASSERTM(0, "Const data array never be half float\n"); } } else -#endif /* LLVM_VERSION_MINOR > 0 */ if (dyn_cast<ConstantAggregateZero>(CPV)) { Type* Ty = CPV->getType(); @@ -2344,9 +2339,6 @@ namespace gbe Function::arg_iterator I = F.arg_begin(), E = F.arg_end(); // Insert a new register for each function argument -#if LLVM_VERSION_MINOR <= 1 - const AttrListPtr &PAL = F.getAttributes(); -#endif /* LLVM_VERSION_MINOR <= 1 */ for (; I != E; ++I, ++argID) { uint32_t opID = argID; #if LLVM_VERSION_MAJOR == 3 && LLVM_VERSION_MINOR < 9 @@ -2436,11 +2428,7 @@ namespace gbe continue; Type *pointed = pointerType->getElementType(); // By value structure -#if LLVM_VERSION_MINOR <= 1 - if (PAL.paramHasAttr(argID+1, Attribute::ByVal)) { -#else if (I->hasByValAttr()) { -#endif /* LLVM_VERSION_MINOR <= 1 */ const size_t structSize = getTypeByteSize(unit, pointed); ctx.input(argName, ir::FunctionArgument::STRUCTURE, reg, llvmInfo, structSize, getAlignmentByte(unit, type), 0); } @@ -3164,15 +3152,9 @@ namespace gbe void GenWriter::emitFunction(Function &F) { switch (F.getCallingConv()) { -#if LLVM_VERSION_MINOR <= 2 - case CallingConv::PTX_Device: // we do not emit device function - return; - case CallingConv::PTX_Kernel: -#else case CallingConv::C: case CallingConv::Fast: case CallingConv::SPIR_KERNEL: -#endif break; default: GBE_ASSERTM(false, "Unsupported calling convention"); @@ -3789,14 +3771,12 @@ namespace gbe break; case Intrinsic::stackrestore: break; -#if LLVM_VERSION_MINOR >= 2 case Intrinsic::lifetime_start: case Intrinsic::lifetime_end: break; case Intrinsic::fmuladd: this->newRegister(&I); break; -#endif /* LLVM_VERSION_MINOR >= 2 */ case Intrinsic::debugtrap: case Intrinsic::trap: case Intrinsic::dbg_value: @@ -4644,11 +4624,9 @@ namespace gbe ctx.MOV(ir::getType(family), dst, src); } break; -#if LLVM_VERSION_MINOR >= 2 case Intrinsic::lifetime_start: case Intrinsic::lifetime_end: break; -#endif /* LLVM_VERSION_MINOR >= 2 */ case Intrinsic::debugtrap: case Intrinsic::trap: case Intrinsic::dbg_value: diff --git a/backend/src/llvm/llvm_printf_parser.cpp b/backend/src/llvm/llvm_printf_parser.cpp index 800f343f..d64fc60c 100644 --- a/backend/src/llvm/llvm_printf_parser.cpp +++ b/backend/src/llvm/llvm_printf_parser.cpp @@ -389,15 +389,9 @@ error: { bool hasPrintf = false; switch (F.getCallingConv()) { -#if LLVM_VERSION_MAJOR == 3 && LLVM_VERSION_MINOR <= 2 - case CallingConv::PTX_Device: - return false; - case CallingConv::PTX_Kernel: -#else case CallingConv::C: case CallingConv::Fast: case CallingConv::SPIR_KERNEL: -#endif break; default: GBE_ASSERTM(false, "Unsupported calling convention"); diff --git a/backend/src/llvm/llvm_profiling.cpp b/backend/src/llvm/llvm_profiling.cpp index 96c95eeb..734c69d9 100644 --- a/backend/src/llvm/llvm_profiling.cpp +++ b/backend/src/llvm/llvm_profiling.cpp @@ -26,27 +26,13 @@ #include <stdlib.h> #include "llvm/Config/llvm-config.h" -#if LLVM_VERSION_MINOR <= 2 -#include "llvm/Function.h" -#include "llvm/InstrTypes.h" -#include "llvm/Instructions.h" -#include "llvm/IntrinsicInst.h" -#include "llvm/Module.h" -#else #include "llvm/IR/Function.h" #include "llvm/IR/InstrTypes.h" #include "llvm/IR/Instructions.h" #include "llvm/IR/IntrinsicInst.h" #include "llvm/IR/Module.h" -#endif /* LLVM_VERSION_MINOR <= 2 */ #include "llvm/Pass.h" -#if LLVM_VERSION_MINOR <= 1 -#include "llvm/Support/IRBuilder.h" -#elif LLVM_VERSION_MINOR == 2 -#include "llvm/IRBuilder.h" -#else #include "llvm/IR/IRBuilder.h" -#endif /* LLVM_VERSION_MINOR <= 1 */ #if LLVM_VERSION_MINOR >= 5 #include "llvm/IR/CallSite.h" @@ -111,15 +97,9 @@ namespace gbe int pointNum = 0; switch (F.getCallingConv()) { -#if LLVM_VERSION_MAJOR == 3 && LLVM_VERSION_MINOR <= 2 - case CallingConv::PTX_Device: - return false; - case CallingConv::PTX_Kernel: -#else case CallingConv::C: case CallingConv::Fast: case CallingConv::SPIR_KERNEL: -#endif break; default: GBE_ASSERTM(false, "Unsupported calling convention"); diff --git a/backend/src/llvm/llvm_scalarize.cpp b/backend/src/llvm/llvm_scalarize.cpp index 6f46c9da..329e3a91 100644 --- a/backend/src/llvm/llvm_scalarize.cpp +++ b/backend/src/llvm/llvm_scalarize.cpp @@ -884,15 +884,9 @@ namespace gbe { bool Scalarize::runOnFunction(Function& F) { switch (F.getCallingConv()) { -#if LLVM_VERSION_MAJOR == 3 && LLVM_VERSION_MINOR <= 2 - case CallingConv::PTX_Device: - return false; - case CallingConv::PTX_Kernel: -#else case CallingConv::C: case CallingConv::Fast: case CallingConv::SPIR_KERNEL: -#endif break; default: GBE_ASSERTM(false, "Unsupported calling convention"); -- 2.12.0 ++++++ 0002-Backend-Fix-an-include-file-error-problem.patch ++++++ >From 1d4fc66b580e409cb2d0a9df0e7c0ef7ebd5bdb5 Mon Sep 17 00:00:00 2001 From: Pan Xiuli <[email protected]> Date: Fri, 17 Mar 2017 14:15:59 +0800 Subject: [PATCH 2/7] Backend: Fix an include file error problem We should not include any llvm header in ir unit, and we need add missing headers for proliling after deleting llvm headers. Signed-off-by: Pan Xiuli <[email protected]> Reviewed-by: Yang Rong <[email protected]> --- backend/src/ir/profiling.cpp | 1 + backend/src/ir/unit.hpp | 4 +--- backend/src/llvm/llvm_gen_backend.cpp | 2 +- backend/src/llvm/llvm_printf_parser.cpp | 2 +- 4 files changed, 4 insertions(+), 5 deletions(-) diff --git a/backend/src/ir/profiling.cpp b/backend/src/ir/profiling.cpp index ac61e9b2..3289e769 100644 --- a/backend/src/ir/profiling.cpp +++ b/backend/src/ir/profiling.cpp @@ -24,6 +24,7 @@ #include <stdlib.h> #include "ir/profiling.hpp" #include "src/cl_device_data.h" +#include <inttypes.h> namespace gbe { diff --git a/backend/src/ir/unit.hpp b/backend/src/ir/unit.hpp index 46d7be79..d7a2a672 100644 --- a/backend/src/ir/unit.hpp +++ b/backend/src/ir/unit.hpp @@ -32,8 +32,6 @@ #include "sys/map.hpp" #include <string.h> -#include "llvm/IR/Instructions.h" - namespace gbe { namespace ir { @@ -46,7 +44,7 @@ namespace ir { public: typedef map<std::string, Function*> FunctionSet; /*! Moved from printf pass */ - map<llvm::CallInst*, PrintfSet::PrintfFmt*> printfs; + map<void *, PrintfSet::PrintfFmt*> printfs; vector<std::string> blockFuncs; /*! Create an empty unit */ Unit(PointerSize pointerSize = POINTER_32_BITS); diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp index 6fa1ae02..8c046644 100644 --- a/backend/src/llvm/llvm_gen_backend.cpp +++ b/backend/src/llvm/llvm_gen_backend.cpp @@ -764,7 +764,7 @@ namespace gbe void emitUnalignedDQLoadStore(ir::Register ptr, Value *llvmValues, ir::AddressSpace addrSpace, ir::Register bti, bool isLoad, bool dwAligned, bool fixedBTI); void visitInstruction(Instruction &I) {NOT_SUPPORTED;} ir::PrintfSet::PrintfFmt* getPrintfInfo(CallInst* inst) { - if (unit.printfs.find(inst) == unit.printfs.end()) + if (unit.printfs.find((void *)inst) == unit.printfs.end()) return NULL; return unit.printfs[inst]; } diff --git a/backend/src/llvm/llvm_printf_parser.cpp b/backend/src/llvm/llvm_printf_parser.cpp index d64fc60c..a1b1c2c9 100644 --- a/backend/src/llvm/llvm_printf_parser.cpp +++ b/backend/src/llvm/llvm_printf_parser.cpp @@ -381,7 +381,7 @@ error: } GBE_ASSERT(unit.printfs.find(call) == unit.printfs.end()); - unit.printfs.insert(std::pair<llvm::CallInst*, PrintfSet::PrintfFmt*>(call, printf_fmt)); + unit.printfs.insert(std::pair<void *, PrintfSet::PrintfFmt*>((void *)call, printf_fmt)); return true; } -- 2.12.0 ++++++ 0003-Backend-Refine-GEP-lowering-code.patch ++++++ >From c2b2d76e73c6c7751e3671021fdfc5510cd145f0 Mon Sep 17 00:00:00 2001 From: Pan Xiuli <[email protected]> Date: Fri, 17 Mar 2017 14:16:00 +0800 Subject: [PATCH 3/7] Backend: Refine GEP lowering code Pointer is not as like as array or vector, we should handle it in a standalone path to fit furture change about PointerType inheritance. Signed-off-by: Pan Xiuli <[email protected]> Reviewed-by: Yang Rong <[email protected]> --- backend/src/llvm/llvm_gen_backend.cpp | 6 +++--- backend/src/llvm/llvm_gen_backend.hpp | 5 ++++- backend/src/llvm/llvm_passes.cpp | 35 +++++++++++++++++++++++------------ 3 files changed, 30 insertions(+), 16 deletions(-) diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp index 8c046644..742c9476 100644 --- a/backend/src/llvm/llvm_gen_backend.cpp +++ b/backend/src/llvm/llvm_gen_backend.cpp @@ -1503,15 +1503,15 @@ namespace gbe Value *pointer = expr->getOperand(0); if (expr->getOpcode() == Instruction::GetElementPtr) { uint32_t constantOffset = 0; - CompositeType* CompTy = cast<CompositeType>(pointer->getType()); + Type* EltTy = pointer->getType(); for(uint32_t op=1; op<expr->getNumOperands(); ++op) { int32_t TypeIndex; ConstantInt* ConstOP = dyn_cast<ConstantInt>(expr->getOperand(op)); GBE_ASSERTM(ConstOP != NULL, "must be constant index"); TypeIndex = ConstOP->getZExtValue(); GBE_ASSERT(TypeIndex >= 0); - constantOffset += getGEPConstOffset(unit, CompTy, TypeIndex); - CompTy = dyn_cast<CompositeType>(CompTy->getTypeAtIndex(TypeIndex)); + constantOffset += getGEPConstOffset(unit, pointer->getType(), TypeIndex); + EltTy = getEltType(EltTy, TypeIndex); } ir::Constant cc = unit.getConstantSet().getConstant(pointer->getName()); diff --git a/backend/src/llvm/llvm_gen_backend.hpp b/backend/src/llvm/llvm_gen_backend.hpp index ae486c5e..f1c791ea 100644 --- a/backend/src/llvm/llvm_gen_backend.hpp +++ b/backend/src/llvm/llvm_gen_backend.hpp @@ -118,7 +118,10 @@ namespace gbe uint32_t getTypeByteSize(const ir::Unit &unit, llvm::Type* Ty); /*! Get GEP constant offset for the specified operand.*/ - int32_t getGEPConstOffset(const ir::Unit &unit, llvm::CompositeType *CompTy, int32_t TypeIndex); + int32_t getGEPConstOffset(const ir::Unit &unit, llvm::Type *eltTy, int32_t TypeIndex); + + /*! Get element type for a type.*/ + llvm::Type* getEltType(llvm::Type *eltTy, uint32_t index = 0); /*! whether this is a kernel function */ bool isKernelFunction(const llvm::Function &f); diff --git a/backend/src/llvm/llvm_passes.cpp b/backend/src/llvm/llvm_passes.cpp index c5f3ffe4..8f5bcc9f 100644 --- a/backend/src/llvm/llvm_passes.cpp +++ b/backend/src/llvm/llvm_passes.cpp @@ -180,12 +180,23 @@ namespace gbe return size_bit/8; } - int32_t getGEPConstOffset(const ir::Unit &unit, CompositeType *CompTy, int32_t TypeIndex) { + Type* getEltType(Type* eltTy, uint32_t index) { + Type *elementType = NULL; + if (PointerType* ptrType = dyn_cast<PointerType>(eltTy)) + elementType = ptrType->getElementType(); + else if(SequentialType * seqType = dyn_cast<SequentialType>(eltTy)) + elementType = seqType->getElementType(); + else if(CompositeType * compTy= dyn_cast<CompositeType>(eltTy)) + elementType = compTy->getTypeAtIndex(index); + GBE_ASSERT(elementType); + return elementType; + } + + int32_t getGEPConstOffset(const ir::Unit &unit, Type *eltTy, int32_t TypeIndex) { int32_t offset = 0; - SequentialType * seqType = dyn_cast<SequentialType>(CompTy); - if (seqType != NULL) { + if (!eltTy->isStructTy()) { if (TypeIndex != 0) { - Type *elementType = seqType->getElementType(); + Type *elementType = getEltType(eltTy); uint32_t elementSize = getTypeByteSize(unit, elementType); uint32_t align = getAlignmentByte(unit, elementType); elementSize += getPadding(elementSize, align); @@ -193,17 +204,16 @@ namespace gbe } } else { int32_t step = TypeIndex > 0 ? 1 : -1; - GBE_ASSERT(CompTy->isStructTy()); for(int32_t ty_i=0; ty_i != TypeIndex; ty_i += step) { - Type* elementType = CompTy->getTypeAtIndex(ty_i); + Type* elementType = getEltType(eltTy, ty_i); uint32_t align = getAlignmentByte(unit, elementType); offset += getPadding(offset, align * step); offset += getTypeByteSize(unit, elementType) * step; } //add getPaddingding for accessed type - const uint32_t align = getAlignmentByte(unit, CompTy->getTypeAtIndex(TypeIndex)); + const uint32_t align = getAlignmentByte(unit, getEltType(eltTy ,TypeIndex)); offset += getPadding(offset, align * step); } return offset; @@ -247,8 +257,8 @@ namespace gbe { const uint32_t ptrSize = unit.getPointerSize(); Value* parentPointer = GEPInst->getOperand(0); - CompositeType* CompTy = parentPointer ? cast<CompositeType>(parentPointer->getType()) : NULL; - if(!CompTy) + Type* eltTy = parentPointer ? parentPointer->getType() : NULL; + if(!eltTy) return false; Value* currentAddrInst = @@ -262,14 +272,15 @@ namespace gbe ConstantInt* ConstOP = dyn_cast<ConstantInt>(GEPInst->getOperand(op)); if (ConstOP != NULL) { TypeIndex = ConstOP->getZExtValue(); - constantOffset += getGEPConstOffset(unit, CompTy, TypeIndex); + constantOffset += getGEPConstOffset(unit, eltTy, TypeIndex); } else { // we only have array/vectors here, // therefore all elements have the same size TypeIndex = 0; - Type* elementType = CompTy->getTypeAtIndex(TypeIndex); + Type* elementType = getEltType(eltTy); + uint32_t size = getTypeByteSize(unit, elementType); //add padding @@ -326,7 +337,7 @@ namespace gbe } //step down in type hirachy - CompTy = dyn_cast<CompositeType>(CompTy->getTypeAtIndex(TypeIndex)); + eltTy = getEltType(eltTy, TypeIndex); } //insert addition of new offset before GEPInst when it is not zero -- 2.12.0 ++++++ 0004-Backend-Refine-LLVM-version-check-macro.patch ++++++ ++++ 940 lines (skipped) ++++++ 0005-Backend-Refine-FCmp-one-and-une.patch ++++++ >From 90e7415816a44576cd410a6f8d748c039d025f20 Mon Sep 17 00:00:00 2001 From: Pan Xiuli <[email protected]> Date: Fri, 17 Mar 2017 14:16:02 +0800 Subject: [PATCH 5/7] Backend: Refine FCmp one and une llvm will merge: %1 = fcmp olt %a, %b %2 = fcmp ogt %a, %b %dst = or %1, %2 into %dst = fcmp one %a, %b And own CMP.NE is actually une so refine Fcmp one into CMP.LT and CMP.GT and OR Signed-off-by: Pan Xiuli <[email protected]> Reviewed-by: Yang Rong <[email protected]> --- backend/src/llvm/llvm_gen_backend.cpp | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp index 4621f6d3..fa45a420 100644 --- a/backend/src/llvm/llvm_gen_backend.cpp +++ b/backend/src/llvm/llvm_gen_backend.cpp @@ -3384,11 +3384,12 @@ namespace gbe const ir::Register src0 = this->getRegister(I.getOperand(0)); const ir::Register src1 = this->getRegister(I.getOperand(1)); const ir::Register tmp = ctx.reg(getFamily(ctx, I.getType())); + const ir::Register tmp1 = ctx.reg(getFamily(ctx, I.getType())); Value *cv = ConstantInt::get(I.getType(), 1); switch (I.getPredicate()) { case ICmpInst::FCMP_OEQ: ctx.EQ(type, dst, src0, src1); break; - case ICmpInst::FCMP_ONE: ctx.NE(type, dst, src0, src1); break; + case ICmpInst::FCMP_UNE: ctx.NE(type, dst, src0, src1); break; case ICmpInst::FCMP_OLE: ctx.LE(type, dst, src0, src1); break; case ICmpInst::FCMP_OGE: ctx.GE(type, dst, src0, src1); break; case ICmpInst::FCMP_OLT: ctx.LT(type, dst, src0, src1); break; @@ -3434,9 +3435,10 @@ namespace gbe ctx.GT(type, tmp, src0, src1); ctx.XOR(insnType, dst, tmp, getRegister(cv)); break; - case ICmpInst::FCMP_UNE: - ctx.EQ(type, tmp, src0, src1); - ctx.XOR(insnType, dst, tmp, getRegister(cv)); + case ICmpInst::FCMP_ONE: + ctx.LT(type, tmp, src0, src1); + ctx.GT(type, tmp1, src0, src1); + ctx.OR(insnType, dst, tmp, tmp1); break; case ICmpInst::FCMP_TRUE: ctx.MOV(insnType, dst, getRegister(cv)); -- 2.12.0 ++++++ 0006-utest-fix-image-qualifier-of-compiler_fill_gl_image-.patch ++++++ >From 8e119eb32e01066c23c3d96bd2b42032e03f7628 Mon Sep 17 00:00:00 2001 From: Yang Rong <[email protected]> Date: Thu, 13 Apr 2017 19:32:51 +0800 Subject: [PATCH 6/7] utest: fix image qualifier of compiler_fill_gl_image test. After clang check the image qualifier, can't use default qualifier to write_image. Signed-off-by: Yang Rong <[email protected]> Reviewed-by: Pan Xiuli <[email protected]> --- kernels/test_fill_gl_image.cl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/kernels/test_fill_gl_image.cl b/kernels/test_fill_gl_image.cl index 4250a571..7b5dce78 100644 --- a/kernels/test_fill_gl_image.cl +++ b/kernels/test_fill_gl_image.cl @@ -1,5 +1,5 @@ __kernel void -test_fill_gl_image(image2d_t img, int color) +test_fill_gl_image(write_only image2d_t img, int color) { int2 coord; float4 color_v4; -- 2.12.0 ++++++ 0007-Backend-Add-LLVM40-support.patch ++++++ ++++ 715 lines (skipped)
