From: Benjamin Segovia Date: Wed, 29 Feb 2012 18:09:08 +0000 (+0000) Subject: Finished a first very very limited LLVM to Gen-IR translation. It should be just... X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=ed6ae6a4408e2faa1349f105cd4967737e693de2;p=contrib%2Fbeignet.git Finished a first very very limited LLVM to Gen-IR translation. It should be just enough to write a first very simple kernel --- diff --git a/backend/kernels/load_store.cbe.c b/backend/kernels/load_store.cbe.c new file mode 100644 index 0000000..18768f9 --- /dev/null +++ b/backend/kernels/load_store.cbe.c @@ -0,0 +1,161 @@ +/* + * Copyright © 2012 Intel Corporation + * + * This library is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2 of the License, or (at your option) any later version. + * + * This library is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with this library. If not, see . + * + * Author: Benjamin Segovia + */ + +/* Provide Declarations */ +#include +#include +#include +/* get a declaration for alloca */ +#if defined(__CYGWIN__) || defined(__MINGW32__) +#define alloca(x) __builtin_alloca((x)) +#define _alloca(x) __builtin_alloca((x)) +#elif defined(__APPLE__) +extern void *__builtin_alloca(unsigned long); +#define alloca(x) __builtin_alloca(x) +#define longjmp _longjmp +#define setjmp _setjmp +#elif defined(__sun__) +#if defined(__sparcv9) +extern void *__builtin_alloca(unsigned long); +#else +extern void *__builtin_alloca(unsigned int); +#endif +#define alloca(x) __builtin_alloca(x) +#elif defined(__FreeBSD__) || defined(__NetBSD__) || defined(__OpenBSD__) || defined(__DragonFly__) || defined(__arm__) +#define alloca(x) __builtin_alloca(x) +#elif defined(_MSC_VER) +#define inline _inline +#define alloca(x) _alloca(x) +#else +#include +#endif + +#ifndef __GNUC__ /* Can only support "linkonce" vars with GCC */ +#define __attribute__(X) +#endif + +#if defined(__GNUC__) && defined(__APPLE_CC__) +#define __EXTERNAL_WEAK__ __attribute__((weak_import)) +#elif defined(__GNUC__) +#define __EXTERNAL_WEAK__ __attribute__((weak)) +#else +#define __EXTERNAL_WEAK__ +#endif + +#if defined(__GNUC__) && defined(__APPLE_CC__) +#define __ATTRIBUTE_WEAK__ +#elif defined(__GNUC__) +#define __ATTRIBUTE_WEAK__ __attribute__((weak)) +#else +#define __ATTRIBUTE_WEAK__ +#endif + +#if defined(__GNUC__) +#define __HIDDEN__ __attribute__((visibility("hidden"))) +#endif + +#ifdef __GNUC__ +#define LLVM_NAN(NanStr) __builtin_nan(NanStr) /* Double */ +#define LLVM_NANF(NanStr) __builtin_nanf(NanStr) /* Float */ +#define LLVM_NANS(NanStr) __builtin_nans(NanStr) /* Double */ +#define LLVM_NANSF(NanStr) __builtin_nansf(NanStr) /* Float */ +#define LLVM_INF __builtin_inf() /* Double */ +#define LLVM_INFF __builtin_inff() /* Float */ +#define LLVM_PREFETCH(addr,rw,locality) __builtin_prefetch(addr,rw,locality) +#define __ATTRIBUTE_CTOR__ __attribute__((constructor)) +#define __ATTRIBUTE_DTOR__ __attribute__((destructor)) +#define LLVM_ASM __asm__ +#else +#define LLVM_NAN(NanStr) ((double)0.0) /* Double */ +#define LLVM_NANF(NanStr) 0.0F /* Float */ +#define LLVM_NANS(NanStr) ((double)0.0) /* Double */ +#define LLVM_NANSF(NanStr) 0.0F /* Float */ +#define LLVM_INF ((double)0.0) /* Double */ +#define LLVM_INFF 0.0F /* Float */ +#define LLVM_PREFETCH(addr,rw,locality) /* PREFETCH */ +#define __ATTRIBUTE_CTOR__ +#define __ATTRIBUTE_DTOR__ +#define LLVM_ASM(X) +#endif + +#if __GNUC__ < 4 /* Old GCC's, or compilers not GCC */ +#define __builtin_stack_save() 0 /* not implemented */ +#define __builtin_stack_restore(X) /* noop */ +#endif + +#if __GNUC__ && __LP64__ /* 128-bit integer types */ +typedef int __attribute__((mode(TI))) llvmInt128; +typedef unsigned __attribute__((mode(TI))) llvmUInt128; +#endif + +#define CODE_FOR_MAIN() /* Any target-specific code for main()*/ + +#ifndef __cplusplus +typedef unsigned char bool; +#endif + + +/* Support for floating point constants */ +typedef unsigned long long ConstantDoubleTy; +typedef unsigned int ConstantFloatTy; +typedef struct { unsigned long long f1; unsigned short f2; unsigned short pad[3]; } ConstantFP80Ty; +typedef struct { unsigned long long f1; unsigned long long f2; } ConstantFP128Ty; + + +/* Global Declarations */ +/* Helper union for bitcasts */ +typedef union { + unsigned int Int32; + unsigned long long Int64; + float Float; + double Double; +} llvmBitCastUnion; + +/* Function Declarations */ +double fmod(double, double); +float fmodf(float, float); +long double fmodl(long double, long double); +void load_store(unsigned int *llvm_cbe_dst, unsigned int *llvm_cbe_src); +void abort(void); + + +/* Function Bodies */ +static inline int llvm_fcmp_ord(double X, double Y) { return X == X && Y == Y; } +static inline int llvm_fcmp_uno(double X, double Y) { return X != X || Y != Y; } +static inline int llvm_fcmp_ueq(double X, double Y) { return X == Y || llvm_fcmp_uno(X, Y); } +static inline int llvm_fcmp_une(double X, double Y) { return X != Y; } +static inline int llvm_fcmp_ult(double X, double Y) { return X < Y || llvm_fcmp_uno(X, Y); } +static inline int llvm_fcmp_ugt(double X, double Y) { return X > Y || llvm_fcmp_uno(X, Y); } +static inline int llvm_fcmp_ule(double X, double Y) { return X <= Y || llvm_fcmp_uno(X, Y); } +static inline int llvm_fcmp_uge(double X, double Y) { return X >= Y || llvm_fcmp_uno(X, Y); } +static inline int llvm_fcmp_oeq(double X, double Y) { return X == Y ; } +static inline int llvm_fcmp_one(double X, double Y) { return X != Y && llvm_fcmp_ord(X, Y); } +static inline int llvm_fcmp_olt(double X, double Y) { return X < Y ; } +static inline int llvm_fcmp_ogt(double X, double Y) { return X > Y ; } +static inline int llvm_fcmp_ole(double X, double Y) { return X <= Y ; } +static inline int llvm_fcmp_oge(double X, double Y) { return X >= Y ; } + +void load_store(unsigned int *llvm_cbe_dst, unsigned int *llvm_cbe_src) { + unsigned int llvm_cbe_tmp__1; + + llvm_cbe_tmp__1 = *llvm_cbe_src; + *llvm_cbe_dst = llvm_cbe_tmp__1; + return; +} + diff --git a/backend/kernels/load_store.cl b/backend/kernels/load_store.cl new file mode 100644 index 0000000..f88e4cc --- /dev/null +++ b/backend/kernels/load_store.cl @@ -0,0 +1,5 @@ +__kernel void load_store(__local int *dst, __local int *src) +{ + dst[0] = src[0]; +} + diff --git a/backend/kernels/load_store.ll b/backend/kernels/load_store.ll new file mode 100644 index 0000000..5329687 --- /dev/null +++ b/backend/kernels/load_store.ll @@ -0,0 +1,17 @@ +; ModuleID = 'load_store.o' +target datalayout = "e-p:32:32-i64:64:64-f64:64:64-n1:8:16:32:64" +target triple = "ptx32--" + +define ptx_kernel void @load_store(i32 addrspace(4)* nocapture %dst, i32 addrspace(4)* nocapture %src) nounwind noinline { +entry: + %0 = load i32 addrspace(4)* %src, align 4, !tbaa !1 + store i32 %0, i32 addrspace(4)* %dst, align 4, !tbaa !1 + ret void +} + +!opencl.kernels = !{!0} + +!0 = metadata !{void (i32 addrspace(4)*, i32 addrspace(4)*)* @load_store} +!1 = metadata !{metadata !"int", metadata !2} +!2 = metadata !{metadata !"omnipotent char", metadata !3} +!3 = metadata !{metadata !"Simple C/C++ TBAA", null} diff --git a/backend/kernels/load_store.o b/backend/kernels/load_store.o new file mode 100644 index 0000000..1ac5988 Binary files /dev/null and b/backend/kernels/load_store.o differ diff --git a/backend/src/ir/context.hpp b/backend/src/ir/context.hpp index 39cc727..bbdb8c6 100644 --- a/backend/src/ir/context.hpp +++ b/backend/src/ir/context.hpp @@ -85,22 +85,22 @@ namespace ir { /*! LOAD with the destinations directly specified */ template - void LOAD(Type type, Register offset, MemorySpace space, Args...values) + void LOAD(Type type, Register offset, MemorySpace space, bool dwAligned, Args...values) { const Tuple index = this->tuple(values...); const uint16_t valueNum = std::tuple_size>::value; GBE_ASSERT(valueNum > 0); - this->LOAD(type, index, offset, space, valueNum); + this->LOAD(type, index, offset, space, valueNum, dwAligned); } /*! STORE with the sources directly specified */ template - void STORE(Type type, Register offset, MemorySpace space, Args...values) + void STORE(Type type, Register offset, MemorySpace space, bool dwAligned, Args...values) { const Tuple index = this->tuple(values...); const uint16_t valueNum = std::tuple_size>::value; GBE_ASSERT(valueNum > 0); - this->STORE(type, index, offset, space, valueNum); + this->STORE(type, index, offset, space, valueNum, dwAligned); } private: diff --git a/backend/src/ir/instruction.cpp b/backend/src/ir/instruction.cpp index 629a6dc..fec684b 100644 --- a/backend/src/ir/instruction.cpp +++ b/backend/src/ir/instruction.cpp @@ -275,15 +275,17 @@ namespace ir { Tuple dstValues, Register offset, MemorySpace memSpace, - uint32_t valueNum) + uint32_t valueNum, + bool dwAligned) { - GBE_ASSERT(valueNum < 255); - this->opcode = OP_STORE; + GBE_ASSERT(valueNum < 128); + this->opcode = OP_LOAD; this->type = type; this->offset = offset; this->values = dstValues; this->memSpace = memSpace; this->valueNum = valueNum; + this->dwAligned = dwAligned ? 1 : 0; } INLINE Register getSrcIndex(const Function &fn, uint32_t ID) const { GBE_ASSERTM(ID == 0, "Only one source for the load instruction"); @@ -303,8 +305,9 @@ namespace ir { Type type; //!< Type to store Register offset; //!< First source is the offset where to store Tuple values; //!< Values to load - MemorySpace memSpace; //!< Where to store - uint8_t valueNum; //!< Number of values to store + MemorySpace memSpace; //!< Where to load + uint8_t valueNum:7; //!< Number of values to load + uint8_t dwAligned:1; //!< DWORD aligned is what matters with GEN }; class ALIGNED_INSTRUCTION StoreInstruction : @@ -315,7 +318,8 @@ namespace ir { Tuple values, Register offset, MemorySpace memSpace, - uint32_t valueNum) + uint32_t valueNum, + bool dwAligned) { GBE_ASSERT(valueNum < 255); this->opcode = OP_STORE; @@ -324,6 +328,7 @@ namespace ir { this->values = values; this->memSpace = memSpace; this->valueNum = valueNum; + this->dwAligned = dwAligned ? 1 : 0; } INLINE Register getSrcIndex(const Function &fn, uint32_t ID) const { GBE_ASSERTM(ID < valueNum + 1u, "Out-of-bound source register for store"); @@ -342,7 +347,8 @@ namespace ir { Register offset; //!< First source is the offset where to store Tuple values; //!< Values to store MemorySpace memSpace; //!< Where to store - uint8_t valueNum; //!< Number of values to store + uint8_t valueNum:7; //!< Number of values to store + uint8_t dwAligned:1; //!< DWORD aligned is what matters with GEN }; class ALIGNED_INSTRUCTION TextureInstruction : @@ -613,19 +619,20 @@ namespace ir { INLINE void LoadInstruction::out(std::ostream &out, const Function &fn) const { this->outOpcode(out); - out << "." << type << "." << memSpace << " {"; + out << "." << type << "." << memSpace << (dwAligned ? "." : ".un") << "aligned"; + out << " {"; for (uint32_t i = 0; i < valueNum; ++i) - out << this->getDstIndex(fn, i); + out << "%" << this->getDstIndex(fn, i) << (i != (valueNum-1) ? " " : ""); out << "}"; out << " %" << this->getSrcIndex(fn, 0); } INLINE void StoreInstruction::out(std::ostream &out, const Function &fn) const { this->outOpcode(out); - out << "." << type << "." << memSpace; + out << "." << type << "." << memSpace << (dwAligned ? "." : ".un") << "aligned"; out << " %" << this->getSrcIndex(fn, 0) << " {"; for (uint32_t i = 0; i < valueNum; ++i) - out << this->getSrcIndex(fn, i+1); + out << "%" << this->getSrcIndex(fn, i+1) << (i != (valueNum-1) ? " " : ""); out << "}"; } @@ -920,15 +927,16 @@ DECL_MEM_FN(BranchInstruction, LabelIndex, getLabelIndex(void), getLabelIndex()) } // LOAD and STORE -#define DECL_EMIT_FUNCTION(NAME, CLASS) \ - Instruction NAME(Type type, \ - Tuple tuple, \ - Register offset, \ - MemorySpace space, \ - uint16_t valueNum) \ - { \ - const internal::CLASS insn(type, tuple, offset, space, valueNum); \ - return insn.convert(); \ +#define DECL_EMIT_FUNCTION(NAME, CLASS) \ + Instruction NAME(Type type, \ + Tuple tuple, \ + Register offset, \ + MemorySpace space, \ + uint32_t valueNum, \ + bool dwAligned) \ + { \ + const internal::CLASS insn(type,tuple,offset,space,valueNum,dwAligned); \ + return insn.convert(); \ } DECL_EMIT_FUNCTION(LOAD, LoadInstruction) diff --git a/backend/src/ir/instruction.hpp b/backend/src/ir/instruction.hpp index 36ee60c..57853ff 100644 --- a/backend/src/ir/instruction.hpp +++ b/backend/src/ir/instruction.hpp @@ -189,6 +189,8 @@ namespace ir { uint32_t getValueNum(void) const; /*! Address space that is manipulated here */ MemorySpace getAddressSpace(void) const; + /*! DWORD aligned means untyped read for Gen. That is what matters */ + bool isDWORDAligned(void) const; /*! Return true if the given instruction is an instance of this class */ static bool isClassOf(const Instruction &insn); }; @@ -205,6 +207,8 @@ namespace ir { uint32_t getValueNum(void) const; /*! Address space that is manipulated here */ MemorySpace getAddressSpace(void) const; + /*! DWORD aligned means untyped read for Gen. That is what matters */ + bool isDWORDAligned(void) const; /*! Return true if the given instruction is an instance of this class */ static bool isClassOf(const Instruction &insn); }; @@ -370,9 +374,9 @@ namespace ir { /*! loadi.type dst value */ Instruction LOADI(Type type, Register dst, ImmediateIndex value); /*! load.type.space {dst1,...,dst_valueNum} offset value */ - Instruction LOAD(Type type, Tuple dst, Register offset, MemorySpace space, uint16_t valueNum); + Instruction LOAD(Type type, Tuple dst, Register offset, MemorySpace space, uint32_t valueNum, bool dwAligned); /*! store.type.space offset {src1,...,src_valueNum} value */ - Instruction STORE(Type type, Tuple src, Register offset, MemorySpace space, uint16_t valueNum); + Instruction STORE(Type type, Tuple src, Register offset, MemorySpace space, uint32_t valueNum, bool dwAligned); /*! fence.space */ Instruction FENCE(MemorySpace space); /*! label labelIndex */ diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp index b67ca08..ad902fc 100644 --- a/backend/src/llvm/llvm_gen_backend.cpp +++ b/backend/src/llvm/llvm_gen_backend.cpp @@ -209,9 +209,6 @@ namespace gbe void writeOperandWithCast(Value* Operand, const ICmpInst &I); bool writeInstructionCast(const Instruction &I); - void writeMemoryAccess(Value *Operand, Type *OperandType, - bool IsVolatile, unsigned Alignment); - private : std::string InterpretASMConstraint(InlineAsm::ConstraintInfo& c); @@ -247,7 +244,6 @@ namespace gbe INLINE ir::Type getType(const Type *type) const; void printBasicBlock(BasicBlock *BB); - void printLoop(Loop *L); void printCast(unsigned opcode, Type *SrcTy, Type *DstTy); void printConstant(Constant *CPV, bool Static); @@ -323,18 +319,16 @@ namespace gbe void visitReturnInst(ReturnInst &I); void visitBranchInst(BranchInst &I); - void visitSwitchInst(SwitchInst &I); - void visitIndirectBrInst(IndirectBrInst &I); - void visitInvokeInst(InvokeInst &I) { - llvm_unreachable("Lowerinvoke pass didn't work!"); - } - void visitUnwindInst(UnwindInst &I) { - llvm_unreachable("Lowerinvoke pass didn't work!"); - } - void visitResumeInst(ResumeInst &I) { - llvm_unreachable("DwarfEHPrepare pass didn't work!"); - } - void visitUnreachableInst(UnreachableInst &I); + + void visitVAArgInst(VAArgInst &I) {GBE_ASSERTM(false, "Not supported");} + void visitSwitchInst(SwitchInst &I) {GBE_ASSERTM(false, "Not supported");} + void visitInvokeInst(InvokeInst &I) {GBE_ASSERTM(false, "Not supported");} + void visitUnwindInst(UnwindInst &I) {GBE_ASSERTM(false, "Not supported");} + void visitResumeInst(ResumeInst &I) {GBE_ASSERTM(false, "Not supported");} + void visitInlineAsm(CallInst &I) {GBE_ASSERTM(false, "Not supported");} + void visitIndirectBrInst(IndirectBrInst &I) {GBE_ASSERTM(false, "Not supported");} + void visitUnreachableInst(UnreachableInst &I) {GBE_ASSERTM(false, "Not supported");} + void visitPHINode(PHINode &I); void visitBinaryOperator(Instruction &I); @@ -344,14 +338,13 @@ namespace gbe void visitCastInst (CastInst &I); void visitSelectInst(SelectInst &I); void visitCallInst (CallInst &I); - void visitInlineAsm(CallInst &I); bool visitBuiltinCall(CallInst &I, Intrinsic::ID ID, bool &WroteCallee); void visitAllocaInst(AllocaInst &I); + template void visitLoadOrStore(T &I); void visitLoadInst (LoadInst &I); void visitStoreInst (StoreInst &I); void visitGetElementPtrInst(GetElementPtrInst &I); - void visitVAArgInst (VAArgInst &I); void visitInsertElementInst(InsertElementInst &I); void visitExtractElementInst(ExtractElementInst &I); @@ -1921,97 +1914,9 @@ static std::string CBEMangle(const std::string &S) { // ... then, emit the code for all basic blocks for (Function::iterator BB = F.begin(), E = F.end(); BB != E; ++BB) emitBasicBlock(BB); -#if 0 - /// isStructReturn - Should this function actually return a struct by-value? - bool isStructReturn = F.hasStructRetAttr(); - - emitFunctionSignature(&F, false); - Out << " {\n"; - - // If this is a struct return function, handle the result with magic. - if (isStructReturn) { - Type *StructTy = - cast(F.arg_begin()->getType())->getElementType(); - Out << " "; - printType(Out, StructTy, false, "StructReturn"); - Out << "; /* Struct return temporary */\n"; - - Out << " "; - printType(Out, F.arg_begin()->getType(), false, - GetValueName(F.arg_begin())); - Out << " = &StructReturn;\n"; - } - - bool PrintedVar = false; - - // print local variable information for the function - for (inst_iterator I = inst_begin(&F), E = inst_end(&F); I != E; ++I) { - if (const AllocaInst *AI = isDirectAlloca(&*I)) { - GBE_ASSERT(0); - Out << " "; - printType(Out, AI->getAllocatedType(), false, GetValueName(AI)); - Out << "; /* Address-exposed local */\n"; - PrintedVar = true; - } else if (I->getType() != Type::getVoidTy(F.getContext()) && - !isInlinableInst(*I)) { - Out << " "; - printType(Out, I->getType(), false, GetValueName(&*I)); - Out << ";\n"; - - if (isa(*I)) { // Print out PHI node temporaries as well... - Out << " "; - printType(Out, I->getType(), false, - GetValueName(&*I)+"__PHI_TEMPORARY"); - Out << ";\n"; - } - PrintedVar = true; - } - // We need a temporary for the BitCast to use so it can pluck a value out - // of a union to do the BitCast. This is separate from the need for a - // variable to hold the result of the BitCast. - if (isFPIntBitCast(*I)) { - Out << " llvmBitCastUnion " << GetValueName(&*I) - << "__BITCAST_TEMPORARY;\n"; - PrintedVar = true; - } - } - - if (PrintedVar) - Out << '\n'; - - if (F.hasExternalLinkage() && F.getName() == "main") - Out << " CODE_FOR_MAIN();\n"; - - // print the basic blocks - for (Function::iterator BB = F.begin(), E = F.end(); BB != E; ++BB) { - if (Loop *L = LI->getLoopFor(BB)) { - if (L->getHeader() == BB && L->getParentLoop() == 0) - printLoop(L); - } else { - printBasicBlock(BB); - } - } - - Out << "}\n\n"; -#endif ctx.endFunction(); } - void GenWriter::printLoop(Loop *L) { - Out << " do { /* Syntactic loop '" << L->getHeader()->getName() - << "' to make GCC happy */\n"; - for (unsigned i = 0, e = L->getBlocks().size(); i != e; ++i) { - BasicBlock *BB = L->getBlocks()[i]; - Loop *BBLoop = LI->getLoopFor(BB); - if (BBLoop == L) - printBasicBlock(BB); - else if (BB == BBLoop->getHeader() && BBLoop->getParentLoop() == L) - printLoop(BBLoop); - } - Out << " } while (1); /* end of syntactic loop '" - << L->getHeader()->getName() << "' */\n"; - } - void GenWriter::printBasicBlock(BasicBlock *BB) { // Don't print the label for the basic block if there are no uses, or if @@ -2084,44 +1989,6 @@ static std::string CBEMangle(const std::string &S) { #endif } - void GenWriter::visitSwitchInst(SwitchInst &SI) { - - Value* Cond = SI.getCondition(); - - Out << " switch ("; - writeOperand(Cond); - Out << ") {\n default:\n"; - printPHICopiesForSuccessor (SI.getParent(), SI.getDefaultDest(), 2); - printBranchToBlock(SI.getParent(), SI.getDefaultDest(), 2); - Out << ";\n"; - - unsigned NumCases = SI.getNumCases(); - // Skip the first item since that's the default case. - for (unsigned i = 1; i < NumCases; ++i) { - ConstantInt* CaseVal = SI.getCaseValue(i); - BasicBlock* Succ = SI.getSuccessor(i); - Out << " case "; - writeOperand(CaseVal); - Out << ":\n"; - printPHICopiesForSuccessor (SI.getParent(), Succ, 2); - printBranchToBlock(SI.getParent(), Succ, 2); - if (Function::iterator(Succ) == llvm::next(Function::iterator(SI.getParent()))) - Out << " break;\n"; - } - - Out << " }\n"; - } - - void GenWriter::visitIndirectBrInst(IndirectBrInst &IBI) { - Out << " goto *(void*)("; - writeOperand(IBI.getOperand(0)); - Out << ");\n"; - } - - void GenWriter::visitUnreachableInst(UnreachableInst &I) { - Out << " /*UNREACHABLE*/;\n"; - } - bool GenWriter::isGotoCodeNecessary(BasicBlock *From, BasicBlock *To) { /// FIXME: This should be reenabled, but loop reordering safe!! return true; @@ -2850,146 +2717,6 @@ static std::string CBEMangle(const std::string &S) { return c.Codes[0]; } - //TODO: import logic from AsmPrinter.cpp - static std::string gccifyAsm(std::string asmstr) { - for (std::string::size_type i = 0; i != asmstr.size(); ++i) - if (asmstr[i] == '\n') - asmstr.replace(i, 1, "\\n"); - else if (asmstr[i] == '\t') - asmstr.replace(i, 1, "\\t"); - else if (asmstr[i] == '$') { - if (asmstr[i + 1] == '{') { - std::string::size_type a = asmstr.find_first_of(':', i + 1); - std::string::size_type b = asmstr.find_first_of('}', i + 1); - std::string n = "%" + - asmstr.substr(a + 1, b - a - 1) + - asmstr.substr(i + 2, a - i - 2); - asmstr.replace(i, b - i + 1, n); - i += n.size() - 1; - } else - asmstr.replace(i, 1, "%"); - } - else if (asmstr[i] == '%')//grr - { asmstr.replace(i, 1, "%%"); ++i;} - - return asmstr; - } - - //TODO: assumptions about what consume arguments from the call are likely wrong - // handle communitivity - void GenWriter::visitInlineAsm(CallInst &CI) { - InlineAsm* as = cast(CI.getCalledValue()); - InlineAsm::ConstraintInfoVector Constraints = as->ParseConstraints(); - - std::vector > ResultVals; - if (CI.getType() == Type::getVoidTy(CI.getContext())) - ; - else if (StructType *ST = dyn_cast(CI.getType())) { - for (unsigned i = 0, e = ST->getNumElements(); i != e; ++i) - ResultVals.push_back(std::make_pair(&CI, (int)i)); - } else { - ResultVals.push_back(std::make_pair(&CI, -1)); - } - - // Fix up the asm string for gcc and emit it. - Out << "__asm__ volatile (\"" << gccifyAsm(as->getAsmString()) << "\"\n"; - Out << " :"; - - unsigned ValueCount = 0; - bool IsFirst = true; - - // Convert over all the output constraints. - for (InlineAsm::ConstraintInfoVector::iterator I = Constraints.begin(), - E = Constraints.end(); I != E; ++I) { - - if (I->Type != InlineAsm::isOutput) { - ++ValueCount; - continue; // Ignore non-output constraints. - } - - assert(I->Codes.size() == 1 && "Too many asm constraint codes to handle"); - std::string C = InterpretASMConstraint(*I); - if (C.empty()) continue; - - if (!IsFirst) { - Out << ", "; - IsFirst = false; - } - - // Unpack the dest. - Value *DestVal; - int DestValNo = -1; - - if (ValueCount < ResultVals.size()) { - DestVal = ResultVals[ValueCount].first; - DestValNo = ResultVals[ValueCount].second; - } else - DestVal = CI.getArgOperand(ValueCount-ResultVals.size()); - - if (I->isEarlyClobber) - C = "&"+C; - - Out << "\"=" << C << "\"(" << GetValueName(DestVal); - if (DestValNo != -1) - Out << ".field" << DestValNo; // Multiple retvals. - Out << ")"; - ++ValueCount; - } - - - // Convert over all the input constraints. - Out << "\n :"; - IsFirst = true; - ValueCount = 0; - for (InlineAsm::ConstraintInfoVector::iterator I = Constraints.begin(), - E = Constraints.end(); I != E; ++I) { - if (I->Type != InlineAsm::isInput) { - ++ValueCount; - continue; // Ignore non-input constraints. - } - - assert(I->Codes.size() == 1 && "Too many asm constraint codes to handle"); - std::string C = InterpretASMConstraint(*I); - if (C.empty()) continue; - - if (!IsFirst) { - Out << ", "; - IsFirst = false; - } - - assert(ValueCount >= ResultVals.size() && "Input can't refer to result"); - Value *SrcVal = CI.getArgOperand(ValueCount-ResultVals.size()); - - Out << "\"" << C << "\"("; - if (!I->isIndirect) - writeOperand(SrcVal); - else - writeOperandDeref(SrcVal); - Out << ")"; - } - - // Convert over the clobber constraints. - IsFirst = true; - for (InlineAsm::ConstraintInfoVector::iterator I = Constraints.begin(), - E = Constraints.end(); I != E; ++I) { - if (I->Type != InlineAsm::isClobber) - continue; // Ignore non-input constraints. - - assert(I->Codes.size() == 1 && "Too many asm constraint codes to handle"); - std::string C = InterpretASMConstraint(*I); - if (C.empty()) continue; - - if (!IsFirst) { - Out << ", "; - IsFirst = false; - } - - Out << '\"' << C << '"'; - } - - Out << ")"; - } - void GenWriter::visitAllocaInst(AllocaInst &I) { Out << '('; printType(Out, I.getType()); @@ -3088,61 +2815,47 @@ static std::string CBEMangle(const std::string &S) { Out << ")"; } - void GenWriter::writeMemoryAccess(Value *Operand, Type *OperandType, - bool IsVolatile, unsigned Alignment) { - - bool IsUnaligned = Alignment && - Alignment < TD->getABITypeAlignment(OperandType); - - if (!IsUnaligned) - Out << '*'; - if (IsVolatile || IsUnaligned) { - Out << "(("; - if (IsUnaligned) - Out << "struct __attribute__ ((packed, aligned(" << Alignment << "))) {"; - printType(Out, OperandType, false, IsUnaligned ? "data" : "volatile*"); - if (IsUnaligned) { - Out << "; } "; - if (IsVolatile) Out << "volatile "; - Out << "*"; - } - Out << ")"; + static INLINE ir::MemorySpace addressSpaceLLVMToGen(unsigned llvmMemSpace) { + switch (llvmMemSpace) { + case 0: return ir::MEM_GLOBAL; + case 4: return ir::MEM_LOCAL; } + GBE_ASSERT(false); + return ir::MEM_GLOBAL; + } - writeOperand(Operand); + static INLINE Value *getLoadOrStoreValue(LoadInst &I) { + return &I; + } + static INLINE Value *getLoadOrStoreValue(StoreInst &I) { + return I.getValueOperand(); + } - if (IsVolatile || IsUnaligned) { - Out << ')'; - if (IsUnaligned) - Out << "->data"; - } + template + INLINE void GenWriter::visitLoadOrStore(T &I) + { + GBE_ASSERTM(I.isVolatile() == false, "Volatile pointer is not supported"); + unsigned int llvmSpace = I.getPointerAddressSpace(); + Value *llvmPtr = I.getPointerOperand(); + Value *llvmValues = getLoadOrStoreValue(I); + Type *llvmType = llvmValues->getType(); + const bool dwAligned = (I.getAlignment() % 4) == 0; + const ir::MemorySpace memSpace = addressSpaceLLVMToGen(llvmSpace); + const ir::Type type = getType(llvmType); + const ir::Register values = getRegister(llvmValues); + const ir::Register ptr = getRegister(llvmPtr); + if (isLoad) + ctx.LOAD(type, ptr, memSpace, dwAligned, values); + else + ctx.STORE(type, ptr, memSpace, dwAligned, values); } void GenWriter::visitLoadInst(LoadInst &I) { - writeMemoryAccess(I.getOperand(0), I.getType(), I.isVolatile(), - I.getAlignment()); - + this->visitLoadOrStore(I); } void GenWriter::visitStoreInst(StoreInst &I) { - writeMemoryAccess(I.getPointerOperand(), I.getOperand(0)->getType(), - I.isVolatile(), I.getAlignment()); - Out << " = "; - Value *Operand = I.getOperand(0); - Constant *BitMask = 0; - if (IntegerType* ITy = dyn_cast(Operand->getType())) - if (!ITy->isPowerOf2ByteWidth()) - // We have a bit width that doesn't match an even power-of-2 byte - // size. Consequently we must & the value with the type's bit mask - BitMask = ConstantInt::get(ITy, ITy->getBitMask()); - if (BitMask) - Out << "(("; - writeOperand(Operand); - if (BitMask) { - Out << ") & "; - printConstant(BitMask, false); - Out << ")"; - } + this->visitLoadOrStore(I); } void GenWriter::visitGetElementPtrInst(GetElementPtrInst &I) { @@ -3150,14 +2863,6 @@ static std::string CBEMangle(const std::string &S) { gep_type_end(I), false); } - void GenWriter::visitVAArgInst(VAArgInst &I) { - Out << "va_arg(*(va_list*)"; - writeOperand(I.getOperand(0)); - Out << ", "; - printType(Out, I.getType()); - Out << ");\n "; - } - void GenWriter::visitInsertElementInst(InsertElementInst &I) { Type *EltTy = I.getType()->getElementType(); writeOperand(I.getOperand(0)); diff --git a/backend/src/sys/assert.cpp b/backend/src/sys/assert.cpp index 0a0eccd..ea3d34c 100644 --- a/backend/src/sys/assert.cpp +++ b/backend/src/sys/assert.cpp @@ -39,6 +39,7 @@ namespace gbe + std::string(file) + ", function " + std::string(fn) + ", line " + std::string(lineString); + assert(0); throw Exception(str); } } /* namespace gbe */ diff --git a/backend/src/utest/utest_llvm.cpp b/backend/src/utest/utest_llvm.cpp index ee5fe64..f0d58e8 100644 --- a/backend/src/utest/utest_llvm.cpp +++ b/backend/src/utest/utest_llvm.cpp @@ -76,7 +76,8 @@ runTests: GBE_ASSERT(dummyKernel != NULL); fclose(dummyKernel); - UTEST_EXPECT_SUCCESS(utestLLVM2Gen("add.ll")); + //UTEST_EXPECT_SUCCESS(utestLLVM2Gen("add.ll")); + UTEST_EXPECT_SUCCESS(utestLLVM2Gen("load_store.ll")); //UTEST_EXPECT_SUCCESS(utestLLVM2Gen("add2.ll")); //UTEST_EXPECT_SUCCESS(utestLLVM2Gen("loop.ll")); //UTEST_EXPECT_SUCCESS(utestLLVM2Gen("void.ll"));