From 9564e88c263551304ab037f2ef489ae215eae463 Mon Sep 17 00:00:00 2001 From: Benjamin Segovia Date: Tue, 28 Feb 2012 20:05:09 +0000 Subject: [PATCH] Added various helper functions in the LLVM to Gen translation Implemented binary instruction translation Implemented return instruction translation --- backend/kernels/compile.sh | 2 +- backend/kernels/loop.cl | 5 + backend/kernels/loop.ll | 29 +++++ backend/kernels/loop.o | Bin 0 -> 600 bytes backend/kernels/mad.cl | 16 +++ backend/kernels/mad.ll | 50 +++++++++ backend/kernels/mad.o | Bin 0 -> 884 bytes backend/src/ir/function.hpp | 4 + backend/src/ir/instruction.cpp | 45 ++++++-- backend/src/ir/instruction.hpp | 2 + backend/src/ir/instruction.hxx | 1 + backend/src/ir/type.hpp | 12 ++ backend/src/llvm/llvm_gen_backend.cpp | 202 ++++++++++++++++++++++++++++------ backend/src/utest/utest_llvm.cpp | 5 +- 14 files changed, 327 insertions(+), 46 deletions(-) mode change 100644 => 100755 backend/kernels/compile.sh create mode 100644 backend/kernels/loop.cl create mode 100644 backend/kernels/loop.ll create mode 100644 backend/kernels/loop.o create mode 100644 backend/kernels/mad.cl create mode 100644 backend/kernels/mad.ll create mode 100644 backend/kernels/mad.o diff --git a/backend/kernels/compile.sh b/backend/kernels/compile.sh old mode 100644 new mode 100755 index 880da2c..7545845 --- a/backend/kernels/compile.sh +++ b/backend/kernels/compile.sh @@ -1,4 +1,4 @@ -clang -emit-llvm -O3 -ccc-host-triple ptx32 -c $1.cl -o $1.o +clang -x cl -emit-llvm -O3 -ccc-host-triple ptx32 -c $1.cl -o $1.o llvm-dis $1.o mv $1.o.ll $1.ll diff --git a/backend/kernels/loop.cl b/backend/kernels/loop.cl new file mode 100644 index 0000000..e8ab630 --- /dev/null +++ b/backend/kernels/loop.cl @@ -0,0 +1,5 @@ +__kernel void add(__global int *dst, unsigned int x) +{ + for (int i = 0; i < x; ++i) dst[i]++; +} + diff --git a/backend/kernels/loop.ll b/backend/kernels/loop.ll new file mode 100644 index 0000000..c67faf7 --- /dev/null +++ b/backend/kernels/loop.ll @@ -0,0 +1,29 @@ +; ModuleID = 'loop.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 @add(i32* nocapture %dst, i32 %x) nounwind noinline { +entry: + %cmp2 = icmp eq i32 %x, 0 + br i1 %cmp2, label %for.end, label %for.body + +for.body: ; preds = %for.body, %entry + %i.03 = phi i32 [ %inc1, %for.body ], [ 0, %entry ] + %arrayidx = getelementptr inbounds i32* %dst, i32 %i.03 + %0 = load i32* %arrayidx, align 4, !tbaa !1 + %inc = add nsw i32 %0, 1 + store i32 %inc, i32* %arrayidx, align 4, !tbaa !1 + %inc1 = add nsw i32 %i.03, 1 + %exitcond = icmp eq i32 %inc1, %x + br i1 %exitcond, label %for.end, label %for.body + +for.end: ; preds = %for.body, %entry + ret void +} + +!opencl.kernels = !{!0} + +!0 = metadata !{void (i32*, i32)* @add} +!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/loop.o b/backend/kernels/loop.o new file mode 100644 index 0000000000000000000000000000000000000000..39a6897185b50d163951a6921a52fc404c1f0963 GIT binary patch literal 600 zcmZ>AK5$Qwhk;=-0|Nu200V;%kY;aGp6Gaj#gk3I$Z`^63y&%*hol@!cT$6ZE0fbC zZl@LjBM}P@2~RE~7Zne#35v=`f*N=v6^m5ZPjCn(tyr+a;n)X8pgEor3=Awln!AI6 z#ZiEz^FI*C9|&ORRA6A>1&SLfad0Z6HnlLQxp;7Ls<|BFbXL?aMV7AO?HoLKb?db)!4-43f71*l`*o!W(y-Z*) zS-@WAz+MGpTwr^K%6%umUe&-}CBR+=(E^k;a8dMQWMBdstiaDK1*Du=m?WG;oD6j& zE*+G8k;oA!;JWcaWqVmcdr3ok`GWSM0`@9~=0b<|f(h-4E{O~Wnli*LEtGxGD0_;L z$L%4nTO#i@2A&K@5hq0vH;eZS{7)b7JwCvfw}J1e0{^=JzBdN^!654r6&M&afV41+ zvkoxW!A=lSW>5t(84mLvVqn|M;}pOX!s2DX@R+mF@wh=7AOb;)r&WX_&s4i*Km`O$wjb1@o*~RVFll14hp>i^A7}CO!PN7%qZew>^ehK zfa&uMHG!CE1`53iY%Gs1Z`V~g((2Q685sYJh71hMK(nIRlsLBWxPhDm29>}#WCya0 pIT~b+Jd{ZjVYKaVwrDxr undef, float %conv, i32 0 + %splat = shufflevector <4 x float> %1, <4 x float> undef, <4 x i32> zeroinitializer + %call8 = tail call ptx_device <4 x float> @_Z3madDv4_fS_S_(<4 x float> %splat, <4 x float> , <4 x float> ) nounwind readonly + %conv9 = fptosi float %call5 to i32 + %add = add nsw i32 %conv9, %call2 + %conv10 = sitofp i32 %add to float + %2 = extractelement <4 x float> %call8, i32 0 + %add11 = fadd float %conv10, %2 + %3 = extractelement <4 x float> %call8, i32 1 + %add12 = fadd float %add11, %3 + %4 = extractelement <4 x float> %call8, i32 2 + %add13 = fadd float %add12, %4 + %conv14 = fptosi float %add13 to i32 + store i32 %conv14, i32* %arrayidx, align 4, !tbaa !1 + %inc = add nsw i32 %i.017, 1 + %exitcond = icmp eq i32 %inc, %x + br i1 %exitcond, label %for.end, label %for.body + +for.end: ; preds = %for.body, %entry + ret void +} + +declare ptx_device i32 @_Z3madiii(i32, i32, i32) nounwind readonly + +declare ptx_device float @_Z3madfff(float, float, float) nounwind readonly + +declare ptx_device <4 x float> @_Z3madDv4_fS_S_(<4 x float>, <4 x float>, <4 x float>) nounwind readonly + +!opencl.kernels = !{!0} + +!0 = metadata !{void (i32*, i32, float)* @add} +!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/mad.o b/backend/kernels/mad.o new file mode 100644 index 0000000000000000000000000000000000000000..26aa240908f640f28c87040101b904ddb207b35a GIT binary patch literal 884 zcmY*XZ%7ki7=P||-reTAd25&5O!KbULXtSA$^1}kw?9mk>YEie-^q;jrMxy-qEH!C!lT>grXX+Uz^ zZ&s0A>Z)KSL)O-->8qaHJc08x2sjtU!W2+mjym}VOL?YgoA^s-kkEZjiU2l@uH_-4 z*cOX}GY~Qwoq?RuCskGq^Y$fd!-$v)!bdiv-a{TqWwnr`uzihEPSB=hvYJ;aue6(~ zM^N$ zD&g69*{2rvDYyQ5zBk$3kvyy>kED}b9ZVFoPEWF9ASwCxKql5+{On>`zOU>dfr2yU zV2k-Bpmr~>NPN)!6^dr_miHOUflSTxsU`63r>)XQlR0k&YYQNAj@6f<6Y5Kg73 zhiXnB*qCUf_fR56p&$CKs>r@8zh+nKQ@NJ&^{KY)#W4qt0LYpj5cEx?$mUJdKCU#= zLwdt&f(S`YvM~I(q^>o>R!fmbu8H}`gsA-pIp_pu%hG|Nf*W(>cyBhx-M+L=R;#Ji z*JREh`0@iK2A8@@^JvyVJ5s|&HnH?^8no@Se*hm;zq^jNb>{f^?=J_zAfw_1lOcvZ|H)sV2LhV_ literal 0 HcmV?d00001 diff --git a/backend/src/ir/function.hpp b/backend/src/ir/function.hpp index 83d7823..6af43a4 100644 --- a/backend/src/ir/function.hpp +++ b/backend/src/ir/function.hpp @@ -79,6 +79,10 @@ namespace ir { const std::string &getName(void) const { return name; } /*! Extract the register from the register file */ INLINE RegisterData getRegisterData(Register ID) const { return file.get(ID); } + /*! Get the register family from the register itself */ + INLINE RegisterData::Family getRegisterFamiy(Register ID) const { + return this->getRegisterData(ID).family; + } /*! Get the register index from the tuple vector */ INLINE Register getRegister(Tuple ID, uint32_t which) const { return file.get(ID, which); diff --git a/backend/src/ir/instruction.cpp b/backend/src/ir/instruction.cpp index d60dd35..629a6dc 100644 --- a/backend/src/ir/instruction.cpp +++ b/backend/src/ir/instruction.cpp @@ -227,18 +227,31 @@ namespace ir { public BasePolicy, public NoDstPolicy { public: - INLINE BranchInstruction(LabelIndex labelIndex, Register predicate) { - this->opcode = OP_BRA; + INLINE BranchInstruction(Opcode op, LabelIndex labelIndex, Register predicate) { + GBE_ASSERT(op == OP_BRA); + this->opcode = op; this->predicate = predicate; this->labelIndex = labelIndex; this->hasPredicate = true; + this->hasLabel = true; } - INLINE BranchInstruction(LabelIndex labelIndex) { + INLINE BranchInstruction(Opcode op, LabelIndex labelIndex) { + GBE_ASSERT(op == OP_BRA); this->opcode = OP_BRA; this->labelIndex = labelIndex; this->hasPredicate = false; + this->hasLabel = true; + } + INLINE BranchInstruction(Opcode op) { + GBE_ASSERT(op == OP_RET); + this->opcode = OP_RET; + this->hasPredicate = false; + this->hasLabel = false; + } + INLINE LabelIndex getLabelIndex(void) const { + GBE_ASSERTM(hasLabel, "No target label for this branch instruction"); + return labelIndex; } - INLINE LabelIndex getLabelIndex(void) const { return labelIndex; } INLINE uint32_t getSrcNum(void) const { return hasPredicate ? 1 : 0; } INLINE Register getSrcIndex(const Function &fn, uint32_t ID) const { GBE_ASSERTM(hasPredicate, "No source for unpredicated branches"); @@ -250,7 +263,8 @@ namespace ir { INLINE void out(std::ostream &out, const Function &fn) const; Register predicate; //!< Predication means conditional branch LabelIndex labelIndex; //!< Index of the label the branch targets - bool hasPredicate; //!< Is it predicated? + bool hasPredicate:1; //!< Is it predicated? + bool hasLabel:1; //!< Is there any target label? }; class ALIGNED_INSTRUCTION LoadInstruction : @@ -556,10 +570,11 @@ namespace ir { // The label must exist and the register must of boolean family INLINE bool BranchInstruction::wellFormed(const Function &fn, std::string &whyNot) const { - if (UNLIKELY(labelIndex >= fn.labelNum())) { - whyNot = "Out-of-bound label index"; - return false; - } + if (hasLabel) + if (UNLIKELY(labelIndex >= fn.labelNum())) { + whyNot = "Out-of-bound label index"; + return false; + } if (hasPredicate) if (UNLIKELY(checkRegisterData(RegisterData::BOOL, predicate, fn, whyNot) == false)) return false; @@ -623,7 +638,7 @@ namespace ir { this->outOpcode(out); if (hasPredicate) out << "<%" << this->getSrcIndex(fn, 0) << ">"; - out << " -> label$" << labelIndex; + if (hasLabel) out << " -> label$" << labelIndex; } INLINE void LoadImmInstruction::out(std::ostream &out, const Function &fn) const { @@ -884,11 +899,17 @@ DECL_MEM_FN(BranchInstruction, LabelIndex, getLabelIndex(void), getLabelIndex()) // BRA Instruction BRA(LabelIndex labelIndex) { - const internal::BranchInstruction insn(labelIndex); + const internal::BranchInstruction insn(OP_BRA, labelIndex); return insn.convert(); } Instruction BRA(LabelIndex labelIndex, Register pred) { - const internal::BranchInstruction insn(labelIndex, pred); + const internal::BranchInstruction insn(OP_BRA, labelIndex, pred); + return insn.convert(); + } + + // RET + Instruction RET(void) { + const internal::BranchInstruction insn(OP_RET); return insn.convert(); } diff --git a/backend/src/ir/instruction.hpp b/backend/src/ir/instruction.hpp index 7c7aed7..36ee60c 100644 --- a/backend/src/ir/instruction.hpp +++ b/backend/src/ir/instruction.hpp @@ -365,6 +365,8 @@ namespace ir { Instruction BRA(LabelIndex labelIndex); /*! (pred) bra labelIndex */ Instruction BRA(LabelIndex labelIndex, Register pred); + /*! ret */ + Instruction RET(void); /*! loadi.type dst value */ Instruction LOADI(Type type, Register dst, ImmediateIndex value); /*! load.type.space {dst1,...,dst_valueNum} offset value */ diff --git a/backend/src/ir/instruction.hxx b/backend/src/ir/instruction.hxx index 7858f0d..2e0b5b2 100644 --- a/backend/src/ir/instruction.hxx +++ b/backend/src/ir/instruction.hxx @@ -55,6 +55,7 @@ DECL_INSN(GE, CompareInstruction) DECL_INSN(GT, CompareInstruction) DECL_INSN(CVT, ConvertInstruction) DECL_INSN(BRA, BranchInstruction) +DECL_INSN(RET, BranchInstruction) DECL_INSN(TEX, TextureInstruction) DECL_INSN(LOADI, LoadImmInstruction) DECL_INSN(LOAD, LoadInstruction) diff --git a/backend/src/ir/type.hpp b/backend/src/ir/type.hpp index ef8f2cf..56f24d2 100644 --- a/backend/src/ir/type.hpp +++ b/backend/src/ir/type.hpp @@ -76,6 +76,18 @@ namespace ir { return RegisterData::DWORD; } + /*! Return a type for each register family */ + INLINE Type getType(RegisterData::Family family) { + switch (family) { + case RegisterData::BOOL: return TYPE_BOOL; + case RegisterData::BYTE: return TYPE_U8; + case RegisterData::WORD: return TYPE_U16; + case RegisterData::DWORD: return TYPE_U32; + case RegisterData::QWORD: return TYPE_U64; + }; + return TYPE_U32; + } + } /* namespace ir */ } /* namespace gbe */ diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp index ab2d858..b67ca08 100644 --- a/backend/src/llvm/llvm_gen_backend.cpp +++ b/backend/src/llvm/llvm_gen_backend.cpp @@ -61,7 +61,7 @@ #include "ir/context.hpp" #include "ir/unit.hpp" -#include "sys/hash_map.hpp" +#include "sys/map.hpp" #include using namespace llvm; @@ -95,8 +95,11 @@ namespace gbe MCContext *TCtx; const TargetData* TD; - /*! Map value name to ir::Register*/ - hash_map registerMap; + /*! Map value to ir::Register*/ + map registerMap; + + /*! Map value to ir::LabelIndex */ + map labelMap; std::map FPConstantMap; std::set intrinsicPrototypesAlreadyGenerated; @@ -223,13 +226,26 @@ namespace gbe void printFloatingPointConstants(const Constant *C); void emitFunctionSignature(const Function *F, bool Prototype); + /*! Emit the complete function code and declaration */ + void emitFunction(Function &F); /*! Handle input and output function parameters */ void emitFunctionPrototype(const Function *F); + /*! Emit the code for a basic block */ + void emitBasicBlock(BasicBlock *BB); /*! Get the register family from the given type */ INLINE ir::RegisterData::Family getArgumentFamily(const Type*) const; + /*! Insert a new register when this is a scalar value */ + INLINE void newRegister(const Value *value); + /*! Return a valid register from an operand (can use LOADI to make one) */ + INLINE ir::Register getRegister(Value *value); + /*! Insert a new label index when this is a scalar value */ + INLINE void newLabelIndex(const Value *value); + /*! int / float / double / bool are scalars */ + INLINE bool isScalarType(const Type *type) const; + /*! Get the Gen IR type from the LLVM type */ + INLINE ir::Type getType(const Type *type) const; - void emitFunction(Function &); void printBasicBlock(BasicBlock *BB); void printLoop(Loop *L); @@ -367,6 +383,7 @@ namespace gbe }; char GenWriter::ID = 0; +#define PRINT_CODE 1 static std::string CBEMangle(const std::string &S) { std::string Result; @@ -1631,13 +1648,45 @@ static std::string CBEMangle(const std::string &S) { } } - INLINE ir::RegisterData::Family GenWriter::getArgumentFamily(const Type *type) const + INLINE bool GenWriter::isScalarType(const Type *type) const { - GBE_ASSERT(type->isFloatTy() || - type->isIntegerTy() || - type->isDoubleTy() || - type->isPointerTy()); + return type->isFloatTy() || + type->isIntegerTy() || + type->isDoubleTy() || + type->isPointerTy(); + } + INLINE ir::Type GenWriter::getType(const Type *type) const + { + GBE_ASSERT(this->isScalarType(type)); + if (type->isFloatTy() == true) + return ir::TYPE_FLOAT; + if (type->isDoubleTy() == true) + return ir::TYPE_DOUBLE; + if (type->isPointerTy() == true) { + if (ctx.getPointerSize() == ir::POINTER_32_BITS) + return ir::TYPE_U32; + else + return ir::TYPE_U64; + } + GBE_ASSERT(type->isIntegerTy() == true); + if (type == Type::getInt1Ty(type->getContext())) + return ir::TYPE_BOOL; + if (type == Type::getInt8Ty(type->getContext())) + return ir::TYPE_S8; + if (type == Type::getInt16Ty(type->getContext())) + return ir::TYPE_S16; + if (type == Type::getInt32Ty(type->getContext())) + return ir::TYPE_S32; + if (type == Type::getInt64Ty(type->getContext())) + return ir::TYPE_S64; + GBE_ASSERT(0); + return ir::TYPE_S64; + } + + INLINE ir::RegisterData::Family GenWriter::getArgumentFamily(const Type *type) const + { + GBE_ASSERT(this->isScalarType(type) == true); if (type == Type::getInt1Ty(type->getContext())) return ir::RegisterData::BOOL; if (type == Type::getInt8Ty(type->getContext())) @@ -1656,6 +1705,49 @@ static std::string CBEMangle(const std::string &S) { return ir::RegisterData::BOOL; } + void GenWriter::newRegister(const Value *value) { + if (registerMap.find(value) == registerMap.end()) { + const Type *type = value->getType(); + const ir::RegisterData::Family family = getArgumentFamily(type); + const ir::Register reg = ctx.reg(family); + ctx.input(reg); + registerMap[value] = reg; + } + } + + ir::Register GenWriter::getRegister(Value *value) { + Constant *CPV = dyn_cast(value); + if (CPV && !isa(CPV)) { + GBE_ASSERT(0); + // printConstant(CPV, Static); + } else { + GBE_ASSERT(this->registerMap.find(value) != this->registerMap.end()); + return this->registerMap[value]; + } + } + + void GenWriter::newLabelIndex(const Value *value) { + if (labelMap.find(value) == labelMap.end()) { + const ir::LabelIndex label = ctx.label(); + labelMap[value] = label; + } + } + + void GenWriter::emitBasicBlock(BasicBlock *BB) { + GBE_ASSERT(labelMap.find(BB) != labelMap.end()); + ctx.LABEL(labelMap[BB]); + for (auto II = BB->begin(), E = BB->end(); II != E; ++II) { + const Type *Ty = II->getType(); + GBE_ASSERT(!Ty->isIntegerTy() || + (Ty==Type::getInt1Ty(II->getContext()) || + Ty==Type::getInt8Ty(II->getContext()) || + Ty==Type::getInt16Ty(II->getContext()) || + Ty==Type::getInt32Ty(II->getContext()) || + Ty==Type::getInt64Ty(II->getContext()))); + visit(*II); + } + } + void GenWriter::emitFunctionPrototype(const Function *F) { const bool returnStruct = F->hasStructRetAttr(); @@ -1670,28 +1762,19 @@ static std::string CBEMangle(const std::string &S) { fn.setStructReturned(true); } - std::string ArgName; - for (; I != E; ++I) { - ArgName = GetValueName(I); - - // Insert a new register if we need to - if (registerMap.find(ArgName) == registerMap.end()) { - const Type *type = I->getType(); - const ir::RegisterData::Family family = getArgumentFamily(type); - const ir::Register reg = ctx.reg(family); - ctx.input(reg); - registerMap[ArgName] = reg; - } - } + // Insert a new register if we need to + for (; I != E; ++I) this->newRegister(I); } // When returning a structure, first input register is the pointer to the // structure if (!returnStruct) { const Type *type = F->getReturnType(); - const ir::RegisterData::Family family = getArgumentFamily(type); - const ir::Register reg = ctx.reg(family); - ctx.output(reg); + if (type->isVoidTy() == false) { + const ir::RegisterData::Family family = getArgumentFamily(type); + const ir::Register reg = ctx.reg(family); + ctx.output(reg); + } } #if GBE_DEBUG @@ -1823,8 +1906,22 @@ static std::string CBEMangle(const std::string &S) { { ctx.startFunction(GetValueName(&F)); this->registerMap.clear(); + this->labelMap.clear(); this->emitFunctionPrototype(&F); + // We create all the register variables + for (inst_iterator I = inst_begin(&F), E = inst_end(&F); I != E; ++I) + if (I->getType() != Type::getVoidTy(F.getContext())) + this->newRegister(&*I); + + // First create all the labels (one per block) + for (Function::iterator BB = F.begin(), E = F.end(); BB != E; ++BB) + this->newLabelIndex(BB); + + // ... 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(); @@ -1850,6 +1947,7 @@ static std::string CBEMangle(const std::string &S) { // 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"; @@ -1895,7 +1993,7 @@ static std::string CBEMangle(const std::string &S) { } Out << "}\n\n"; - +#endif ctx.endFunction(); } @@ -1954,8 +2052,17 @@ static std::string CBEMangle(const std::string &S) { // void GenWriter::visitReturnInst(ReturnInst &I) { // If this is a struct return function, return the temporary struct. - bool isStructReturn = I.getParent()->getParent()->hasStructRetAttr(); + const ir::Function &fn = ctx.getFunction(); + GBE_ASSERTM(fn.outputNum() <= 1, "no more than one value can be returned"); + if (fn.outputNum() == 1 && I.getNumOperands() > 0) { + const ir::Register dst = fn.getOutput(0); + const ir::Register src = this->getRegister(I.getOperand(0)); + const ir::RegisterData::Family family = fn.getRegisterFamiy(dst);; + ctx.MOV(ir::getType(family), dst, src); + } + ctx.RET(); + bool isStructReturn = I.getParent()->getParent()->hasStructRetAttr(); if (isStructReturn) { Out << " return StructReturn;\n"; return; @@ -1967,13 +2074,14 @@ static std::string CBEMangle(const std::string &S) { !I.getParent()->size() == 1) { return; } - +#if 0 Out << " return"; if (I.getNumOperands()) { Out << ' '; writeOperand(I.getOperand(0)); } Out << ";\n"; +#endif } void GenWriter::visitSwitchInst(SwitchInst &SI) { @@ -2099,10 +2207,41 @@ static std::string CBEMangle(const std::string &S) { } - void GenWriter::visitBinaryOperator(Instruction &I) { + void GenWriter::visitBinaryOperator(Instruction &I) + { + GBE_ASSERT(!I.getType()->isPointerTy()); + GBE_ASSERT(this->registerMap.find(&I) != this->registerMap.end()); + const ir::Register dst = this->registerMap[&I]; + const ir::Register src0 = this->getRegister(I.getOperand(0)); + const ir::Register src1 = this->getRegister(I.getOperand(1)); + const ir::Type type = this->getType(I.getType()); + + switch (I.getOpcode()) { + case Instruction::Add: + case Instruction::FAdd: ctx.ADD(type, dst, src0, src1); break; + case Instruction::Sub: + case Instruction::FSub: ctx.SUB(type, dst, src0, src1); break; + case Instruction::Mul: + case Instruction::FMul: ctx.MUL(type, dst, src0, src1); break; + case Instruction::URem: + case Instruction::SRem: + case Instruction::FRem: ctx.REM(type, dst, src0, src1); break; + case Instruction::UDiv: + case Instruction::SDiv: + case Instruction::FDiv: ctx.DIV(type, dst, src0, src1); break; + case Instruction::And: ctx.AND(type, dst, src0, src1); break; + case Instruction::Or: ctx.OR(type, dst, src0, src1); break; + case Instruction::Xor: ctx.XOR(type, dst, src0, src1); break; + case Instruction::Shl : ctx.SHL(type, dst, src0, src1); break; + case Instruction::LShr: ctx.SHR(type, dst, src0, src1); break; + case Instruction::AShr: ctx.ASR(type, dst, src0, src1); break; + default: + GBE_ASSERT(0); + }; + +#if 0 // binary instructions, shift instructions, setCond instructions. assert(!I.getType()->isPointerTy()); - // We must cast the results of binary operations which might be promoted. bool needsCast = false; if ((I.getType() == Type::getInt8Ty(I.getContext())) || @@ -2181,6 +2320,7 @@ static std::string CBEMangle(const std::string &S) { if (needsCast) { Out << "))"; } +#endif } void GenWriter::visitICmpInst(ICmpInst &I) { @@ -3076,7 +3216,7 @@ static std::string CBEMangle(const std::string &S) { void GenWriter::visitInsertValueInst(InsertValueInst &IVI) { // Start by copying the entire aggregate value into the result variable. - writeOperand(IVI.getOperand(0)); + writeOperand(IVI.getOperand(0)); Out << ";\n "; // Then do the insert to update the field. diff --git a/backend/src/utest/utest_llvm.cpp b/backend/src/utest/utest_llvm.cpp index df37db0..ee5fe64 100644 --- a/backend/src/utest/utest_llvm.cpp +++ b/backend/src/utest/utest_llvm.cpp @@ -76,8 +76,9 @@ runTests: GBE_ASSERT(dummyKernel != NULL); fclose(dummyKernel); - //UTEST_EXPECT_SUCCESS(utestLLVM2Gen("add.ll")); - UTEST_EXPECT_SUCCESS(utestLLVM2Gen("add2.ll")); + UTEST_EXPECT_SUCCESS(utestLLVM2Gen("add.ll")); + //UTEST_EXPECT_SUCCESS(utestLLVM2Gen("add2.ll")); + //UTEST_EXPECT_SUCCESS(utestLLVM2Gen("loop.ll")); //UTEST_EXPECT_SUCCESS(utestLLVM2Gen("void.ll")); } -- 2.7.4