Added first support for compilation from source. Quick and dirty (since uses system()) but works OK

21 files changed, 434 insertions, 54 deletions

diff --git a/backend/src/CMakeLists.txt b/backend/src/CMakeLists.txt
index 19a8cdce..26fd6cdf 100644
--- a/backend/src/CMakeLists.txt
+++ b/backend/src/CMakeLists.txt
@@ -1,9 +1,9 @@
 add_subdirectory(llvm)
 
-set (TO_STRINGIFY_FILES simulator sim_vector)
+macro (stringify TO_STRINGIFY_PATH TO_STRINGIFY_FILES)
 foreach (to_stringify_file ${TO_STRINGIFY_FILES})
-  set (input_file ${GBE_SOURCE_DIR}/src/backend/sim/${to_stringify_file}.h)
-  set (output_file ${GBE_SOURCE_DIR}/src/backend/sim/${to_stringify_file}_str.cpp)
+  set (input_file ${TO_STRINGIFY_PATH}/${to_stringify_file}.h)
+  set (output_file ${TO_STRINGIFY_PATH}/${to_stringify_file}_str.cpp)
   set (string_header "\\\"string\\\"")
   add_custom_command(
     OUTPUT ${output_file}
@@ -18,6 +18,12 @@ foreach (to_stringify_file ${TO_STRINGIFY_FILES})
     COMMAND echo "" >> ${output_file}
     MAIN_DEPENDENCY ${input_file})
 endforeach (to_stringify_file)
+endmacro (stringify)
+
+set (TO_STRINGIFY_FILES simulator sim_vector)
+stringify ("${GBE_SOURCE_DIR}/src/backend/sim/" "${TO_STRINGIFY_FILES}")
+set (TO_STRINGIFY_FILES stdlib)
+stringify ("${GBE_SOURCE_DIR}/src/llvm/" "${TO_STRINGIFY_FILES}")
 
 if (GBE_USE_BLOB)
   set (GBE_SRC blob.cpp)
@@ -63,6 +69,7 @@ else (GBE_USE_BLOB)
     ir/function.hpp
     ir/value.cpp
     ir/value.hpp
+    llvm/stdlib_str.cpp
     backend/context.cpp
     backend/context.hpp
     backend/program.cpp
diff --git a/backend/src/backend/context.cpp b/backend/src/backend/context.cpp
index 7f6ac89d..1a43ccbf 100644
--- a/backend/src/backend/context.cpp
+++ b/backend/src/backend/context.cpp
@@ -28,17 +28,21 @@
 #include "ir/profile.hpp"
 #include "ir/liveness.hpp"
 #include "ir/value.hpp"
+#include "sys/cvar.hpp"
 #include <algorithm>
 
 namespace gbe
 {
+
+  IVAR(OCL_SIMD_WIDTH, 8, 16, 32);
+
   Context::Context(const ir::Unit &unit, const std::string &name) :
     unit(unit), fn(*unit.getFunction(name)), name(name), liveness(NULL), dag(NULL)
   {
     GBE_ASSERT(unit.getPointerSize() == ir::POINTER_32_BITS);
     this->liveness = GBE_NEW(ir::Liveness, (ir::Function&) fn);
     this->dag = GBE_NEW(ir::FunctionDAG, *this->liveness);
-    this->simdWidth = 16; /* XXX environment variable for that to start with */
+    this->simdWidth = nextHighestPowerOf2(OCL_SIMD_WIDTH);
   }
   Context::~Context(void) {
     GBE_SAFE_DELETE(this->dag);
@@ -49,6 +53,7 @@ namespace gbe
     this->kernel = this->allocateKernel();
     this->buildPatchList();
     this->buildArgList();
+    this->buildUsedLabels();
     this->emitCode();
     return this->kernel;
   }
@@ -148,6 +153,16 @@ namespace gbe
     }
   }
 
+  void Context::buildUsedLabels(void) {
+    usedLabels.clear();
+    fn.foreachInstruction([this](const ir::Instruction &insn) {
+      using namespace ir;
+      if (insn.getOpcode() != OP_BRA) return;
+      const LabelIndex index = cast<BranchInstruction>(insn).getLabelIndex();
+      usedLabels.insert(index);
+    });
+  }
+
   bool Context::isScalarReg(const ir::Register &reg) const {
     GBE_ASSERT(fn.getProfile() == ir::Profile::PROFILE_OCL);
     if (fn.getInput(reg) != NULL)
diff --git a/backend/src/backend/context.hpp b/backend/src/backend/context.hpp
index cbb3c0cd..a88fd1d7 100644
--- a/backend/src/backend/context.hpp
+++ b/backend/src/backend/context.hpp
@@ -26,6 +26,8 @@
 #define __GBE_CONTEXT_HPP__
 
 #include "sys/platform.hpp"
+#include "sys/set.hpp"
+#include "ir/instruction.hpp"
 #include <string>
 
 namespace gbe {
@@ -58,24 +60,33 @@ namespace gbe
     ~Context(void);
     /*! Compile the code */
     Kernel *compileKernel(void);
+    /*! Tells if the labels is used */
+    INLINE bool isLabelUsed(ir::LabelIndex index) const {
+      return usedLabels.contains(index);
+    }
+    /*! Tells if the register is used */
+    bool isRegUsed(const ir::Register &reg) const;
   protected:
     /*! Build the curbe patch list for the given kernel */
     void buildPatchList(void);
     /*! Build the list of arguments to set to launch the kernel */
     void buildArgList(void);
+    /*! Build the sets of used labels */
+    void buildUsedLabels(void);
     /*! Indicate if a register is scalar or not */
     bool isScalarReg(const ir::Register &reg) const;
     /*! Build the instruction stream */
     virtual void emitCode(void) = 0;
     /*! Allocate a new empty kernel */
     virtual Kernel *allocateKernel(void) = 0;
-    const ir::Unit &unit;    //!< Unit that contains the kernel
-    const ir::Function &fn;  //!< Function to compile
-    std::string name;        //!< Name of the kernel to compile
-    Kernel *kernel;          //!< Kernel we are building
-    ir::Liveness *liveness;  //!< Liveness info for the variables
-    ir::FunctionDAG *dag;    //!< Complete DAG of values on the function
-    uint32_t simdWidth;      //!< Number of lanes per HW threads
+    const ir::Unit &unit;           //!< Unit that contains the kernel
+    const ir::Function &fn;         //!< Function to compile
+    std::string name;               //!< Name of the kernel to compile
+    Kernel *kernel;                 //!< Kernel we are building
+    ir::Liveness *liveness;         //!< Liveness info for the variables
+    ir::FunctionDAG *dag;           //!< Complete DAG of values on the function
+    set<ir::LabelIndex> usedLabels; //!< Set of all labels actually used
+    uint32_t simdWidth;             //!< Number of lanes per HW threads
   };
 
 } /* namespace gbe */
diff --git a/backend/src/backend/gen_program.cpp b/backend/src/backend/gen_program.cpp
index 5029f349..dafacea6 100644
--- a/backend/src/backend/gen_program.cpp
+++ b/backend/src/backend/gen_program.cpp
@@ -47,11 +47,6 @@ namespace gbe {
     return ker;
   }
 
-  static gbe_program genProgramNewFromSource(const char *source) {
-    NOT_IMPLEMENTED;
-    return NULL;
-  }
-
   static gbe_program genProgramNewFromBinary(const char *binary, size_t size) {
     NOT_IMPLEMENTED;
     return NULL;
@@ -65,7 +60,6 @@ namespace gbe {
     using namespace gbe;
     GenProgram *program = GBE_NEW(GenProgram);
     std::string error;
-
     // Try to compile the program
     if (program->buildFromLLVMFile(fileName, error) == false) {
       if (err != NULL && errSize != NULL && stringSize > 0u) {
@@ -83,7 +77,6 @@ namespace gbe {
 
 void genSetupCallBacks(void)
 {
-  gbe_program_new_from_source = gbe::genProgramNewFromSource;
   gbe_program_new_from_binary = gbe::genProgramNewFromBinary;
   gbe_program_new_from_llvm = gbe::genProgramNewFromLLVM;
 }
diff --git a/backend/src/backend/program.cpp b/backend/src/backend/program.cpp
index 6ffd1583..08dac6ef 100644
--- a/backend/src/backend/program.cpp
+++ b/backend/src/backend/program.cpp
@@ -59,21 +59,9 @@ namespace gbe {
   }
 
   BVAR(OCL_OUTPUT_GEN_IR, false);
-  BVAR(OCL_OUTPUT_LLVM, false);
 
   bool Program::buildFromLLVMFile(const char *fileName, std::string &error) {
     ir::Unit unit;
-    if (OCL_OUTPUT_LLVM) {
-      std::ifstream llvmFile;
-      llvmFile.open(fileName);
-      if (llvmFile.is_open() == true) {
-        std::string line;
-        while (llvmFile.good() == true) {
-          std::getline(llvmFile ,line);
-          std::cout << line << std::endl;
-        }
-      }
-    }
     if (llvmToGen(unit, fileName) == false) {
       error = std::string(fileName) + " not found";
       return false;
@@ -100,6 +88,35 @@ namespace gbe {
     GBE_SAFE_DELETE(program);
   }
 
+  extern std::string stdlib_str;
+  static gbe_program programNewFromSource(const char *source,
+                                          size_t stringSize,
+                                          char *err,
+                                          size_t *errSize)
+  {
+    char clStr[L_tmpnam+1], llStr[L_tmpnam+1];
+    const std::string clName = std::string(tmpnam_r(clStr)) + ".cl"; /* unsafe! */
+    const std::string llName = std::string(tmpnam_r(llStr)) + ".ll"; /* unsafe! */
+
+    // Write the source to the cl file
+    FILE *clFile = fopen(clName.c_str(), "w");
+    FATAL_IF(clFile == NULL, "Failed to open temporary file");
+    fwrite(stdlib_str.c_str(), strlen(stdlib_str.c_str()), 1, clFile);
+    fwrite(source, strlen(source), 1, clFile);
+    fclose(clFile);
+
+    // Now compile the code to llvm using clang
+    // XXX use popen and stuff instead of that
+    std::string compileCmd = "clang -emit-llvm -O3 -ccc-host-triple ptx32 -c ";
+    compileCmd += clName;
+    compileCmd += " -o ";
+    compileCmd += llName;
+    if (UNLIKELY(system(compileCmd.c_str()) != 0)) return NULL;
+
+    // Now build the program from llvm
+    return gbe_program_new_from_llvm(llName.c_str(), stringSize, err, errSize);
+  }
+
   static uint32_t programGetKernelNum(gbe_program gbeProgram) {
     if (gbeProgram == NULL) return 0;
     const gbe::Program *program = (const gbe::Program*) gbeProgram;
@@ -199,6 +216,7 @@ namespace gbe
   struct CallBackInitializer
   {
     CallBackInitializer(void) {
+      gbe_program_new_from_source = gbe::programNewFromSource;
       gbe_program_delete = gbe::programDelete;
       gbe_program_get_kernel_num = gbe::programGetKernelNum;
       gbe_program_get_kernel_by_name = gbe::programGetKernelByName;
diff --git a/backend/src/backend/program.h b/backend/src/backend/program.h
index 387fba66..56de46a3 100644
--- a/backend/src/backend/program.h
+++ b/backend/src/backend/program.h
@@ -75,7 +75,10 @@ enum gbe_curbe_type {
 };
 
 /*! Create a new program from the given source code (zero terminated string) */
-typedef gbe_program (gbe_program_new_from_source_cb)(const char *source);
+typedef gbe_program (gbe_program_new_from_source_cb)(const char *source,
+                                                     size_t stringSize,
+                                                     char *err,
+                                                     size_t *err_size);
 extern gbe_program_new_from_source_cb *gbe_program_new_from_source;
 
 /*! Create a new program from the given blob */
diff --git a/backend/src/backend/sim/sim_vector.h b/backend/src/backend/sim/sim_vector.h
index e8970d93..34cec348 100644
--- a/backend/src/backend/sim/sim_vector.h
+++ b/backend/src/backend/sim/sim_vector.h
@@ -342,13 +342,13 @@ template <uint32_t vectorNum>
 INLINE void SCATTER(const simd_dw<vectorNum> &offset,
                     const scalar_dw &value,
                     char *base_address) {
-  SCATTER(simd_dw<vectorNum>(value), offset, base_address);
+  SCATTER(offset, simd_dw<vectorNum>(value), base_address);
 }
 template <uint32_t vectorNum>
 INLINE void SCATTER(const scalar_dw &offset,
                     const simd_dw<vectorNum> &value,
                     char *base_address) {
-  SCATTER(value, simd_dw<vectorNum>(offset), base_address);
+  SCATTER(simd_dw<vectorNum>(offset), value, base_address);
 }
 #include <cstdio>
 /* Gather */
@@ -418,7 +418,7 @@ INLINE void GT_U32(scalar_m &dst, scalar_dw v0, scalar_dw v1) { dst.u = (v0.u >
 INLINE void LOAD(scalar_dw &dst, const char *ptr) { dst.u = *(const uint32_t *) ptr; }
 INLINE void STORE(scalar_dw src, char *ptr) { *(uint32_t *) ptr = src.u; }
 INLINE void LOADI(scalar_dw &dst, uint32_t u) { dst.u = u; }
-INLINE void SCATTER(scalar_dw value, scalar_dw offset, char *base) { *(uint32_t*)(base + offset.u) = value.u; }
+INLINE void SCATTER(scalar_dw offset, scalar_dw value, char *base) { *(uint32_t*)(base + offset.u) = value.u; }
 INLINE void GATHER(scalar_dw &dst, scalar_dw offset, const char *base) { dst.u = *(const uint32_t*)(base + offset.u); }
 
 //////////////////////////////////////////////////////////////////////////////
diff --git a/backend/src/backend/sim/sim_vector_str.cpp b/backend/src/backend/sim/sim_vector_str.cpp
index 7699de47..f900c3fc 100644
--- a/backend/src/backend/sim/sim_vector_str.cpp
+++ b/backend/src/backend/sim/sim_vector_str.cpp
@@ -368,13 +368,13 @@ std::string sim_vector_str =
 "INLINE void SCATTER(const simd_dw<vectorNum> &offset,\n"
 "                    const scalar_dw &value,\n"
 "                    char *base_address) {\n"
-"  SCATTER(simd_dw<vectorNum>(value), offset, base_address);\n"
+"  SCATTER(offset, simd_dw<vectorNum>(value), base_address);\n"
 "}\n"
 "template <uint32_t vectorNum>\n"
 "INLINE void SCATTER(const scalar_dw &offset,\n"
 "                    const simd_dw<vectorNum> &value,\n"
 "                    char *base_address) {\n"
-"  SCATTER(value, simd_dw<vectorNum>(offset), base_address);\n"
+"  SCATTER(simd_dw<vectorNum>(offset), value, base_address);\n"
 "}\n"
 "#include <cstdio>\n"
 "/* Gather */\n"
@@ -444,7 +444,7 @@ std::string sim_vector_str =
 "INLINE void LOAD(scalar_dw &dst, const char *ptr) { dst.u = *(const uint32_t *) ptr; }\n"
 "INLINE void STORE(scalar_dw src, char *ptr) { *(uint32_t *) ptr = src.u; }\n"
 "INLINE void LOADI(scalar_dw &dst, uint32_t u) { dst.u = u; }\n"
-"INLINE void SCATTER(scalar_dw value, scalar_dw offset, char *base) { *(uint32_t*)(base + offset.u) = value.u; }\n"
+"INLINE void SCATTER(scalar_dw offset, scalar_dw value, char *base) { *(uint32_t*)(base + offset.u) = value.u; }\n"
 "INLINE void GATHER(scalar_dw &dst, scalar_dw offset, const char *base) { dst.u = *(const uint32_t*)(base + offset.u); }\n"
 "\n"
 "//////////////////////////////////////////////////////////////////////////////\n"
diff --git a/backend/src/backend/sim_context.cpp b/backend/src/backend/sim_context.cpp
index 280d5712..b01ca7b9 100644
--- a/backend/src/backend/sim_context.cpp
+++ b/backend/src/backend/sim_context.cpp
@@ -44,10 +44,22 @@ namespace gbe
 
   void SimContext::emitRegisters(void) {
     GBE_ASSERT(fn.getProfile() == ir::PROFILE_OCL);
+
+    // First we build the set of all used registers
+    set<ir::Register> usedRegs;
+    fn.foreachInstruction([&usedRegs](const ir::Instruction &insn) {
+      const uint32_t srcNum = insn.getSrcNum(), dstNum = insn.getDstNum();
+      for (uint32_t srcID = 0; srcID < srcNum; ++srcID)
+        usedRegs.insert(insn.getSrc(srcID));
+      for (uint32_t dstID = 0; dstID < dstNum; ++dstID)
+        usedRegs.insert(insn.getDst(dstID));
+    });
+
     const uint32_t regNum = fn.regNum();
     bool lid0 = false, lid1 = false, lid2 = false; // for local id registers
     for (uint32_t regID = 0; regID < regNum; ++regID) {
       const ir::Register reg(regID);
+      if (usedRegs.contains(reg) == false) continue;
       if (reg == ir::ocl::groupid0 ||
           reg == ir::ocl::groupid1 ||
           reg == ir::ocl::groupid2)
@@ -153,13 +165,15 @@ namespace gbe
 #undef DECL_INSN
       }
       if (opcode == OP_LABEL) {
-        o << "label" << cast<LabelInstruction>(insn).getLabelIndex() << ":\n";
+        const LabelInstruction labelInsn = cast<LabelInstruction>(insn);
+        const LabelIndex index = labelInsn.getLabelIndex();
+        if (usedLabels.contains(index) == true)
+          o << "label" << index << ":\n";
         return;
       } else if (opcode == OP_BRA) {
         NOT_IMPLEMENTED;
         return;
       } else if (opcode == OP_RET) {
-        std::cout << "BE AWARE OF RET: ONLY ONE RET AT THE END OF THE FUNCTION SHOULD BE OUTPUTTED!";
         o << "return;\n";
         return;
       }
diff --git a/backend/src/backend/sim_program.cpp b/backend/src/backend/sim_program.cpp
index 4df627d3..5c18af7a 100644
--- a/backend/src/backend/sim_program.cpp
+++ b/backend/src/backend/sim_program.cpp
@@ -45,11 +45,6 @@ namespace gbe {
     return ker;
   }
 
-  static gbe_program simProgramNewFromSource(const char *source) {
-    NOT_IMPLEMENTED;
-    return NULL;
-  }
-
   static gbe_program simProgramNewFromBinary(const char *binary, size_t size) {
     NOT_IMPLEMENTED;
     return NULL;
@@ -63,7 +58,7 @@ namespace gbe {
     using namespace gbe;
     SimProgram *program = GBE_NEW(SimProgram);
     std::string error;
-    /* Try to compile the program */
+    // Try to compile the program
     if (program->buildFromLLVMFile(fileName, error) == false) {
       if (err != NULL && errSize != NULL && stringSize > 0u) {
         const size_t msgSize = std::min(error.size(), stringSize-1u);
@@ -73,7 +68,7 @@ namespace gbe {
       GBE_DELETE(program);
       return NULL;
     }
-    /* Everything run fine */
+    // Everything run fine
     return (gbe_program) program;
   }
 
@@ -81,7 +76,6 @@ namespace gbe {
 
 void simSetupCallBacks(void)
 {
-  gbe_program_new_from_source = gbe::simProgramNewFromSource;
   gbe_program_new_from_binary = gbe::simProgramNewFromBinary;
   gbe_program_new_from_llvm = gbe::simProgramNewFromLLVM;
 }
diff --git a/backend/src/ir/context.cpp b/backend/src/ir/context.cpp
index 838ce9ec..7de50ec1 100644
--- a/backend/src/ir/context.cpp
+++ b/backend/src/ir/context.cpp
@@ -51,6 +51,7 @@ namespace ir {
     // Check first that all branch instructions point to valid labels
     for (auto it = usedLabels->begin(); it != usedLabels->end(); ++it)
       GBE_ASSERTM(*it != LABEL_IS_POINTED, "A label is used and not defined");
+    fn->sortLabels();
     fn->computeCFG();
     GBE_DELETE(usedLabels);
     const StackElem elem = fnStack.back();
diff --git a/backend/src/ir/function.cpp b/backend/src/ir/function.cpp
index a91cec36..e2a233fc 100644
--- a/backend/src/ir/function.cpp
+++ b/backend/src/ir/function.cpp
@@ -23,6 +23,7 @@
  */
 #include "ir/function.hpp"
 #include "sys/string.hpp"
+#include "sys/map.hpp"
 
 namespace gbe {
 namespace ir {
@@ -37,6 +38,44 @@ namespace ir {
       GBE_DELETE(*it);
   }
 
+  void Function::sortLabels(void) {
+    uint32_t last = 0;
+
+    // Compute the new labels and patch the label instruction
+    map<LabelIndex, LabelIndex> labelMap;
+    foreachInstruction([&](Instruction &insn) {
+      if (insn.getOpcode() != OP_LABEL) return;
+
+      // Create the new label
+      Instruction *newLabel = newInstruction();
+      *newLabel = LABEL(LabelIndex(last));
+
+      // Replace the previous label instruction
+      LabelInstruction &label = cast<LabelInstruction>(insn);
+      const LabelIndex index = label.getLabelIndex();
+      labelMap.insert(std::make_pair(index, LabelIndex(last++)));
+      newLabel->replace(&insn);
+    });
+
+    // Patch all branch instructions with the new labels
+    foreachInstruction([&](Instruction &insn) {
+      if (insn.getOpcode() != OP_BRA) return;
+
+      // Get the current branch instruction
+      BranchInstruction &bra = cast<BranchInstruction>(insn);
+      const LabelIndex index = bra.getLabelIndex();
+      const LabelIndex newIndex = labelMap.find(index)->second;
+
+      // Insert the patched branch instruction
+      Instruction *newBra = newInstruction();
+      if (bra.isPredicated() == true)
+        *newBra = BRA(newIndex, bra.getPredicateIndex());
+      else
+        *newBra = BRA(newIndex);
+      newBra->replace(&insn);
+    });
+  }
+
   LabelIndex Function::newLabel(void) {
     GBE_ASSERTM(labels.size() < 0xffff,
                 "Too many labels are defined (65536 only are supported)");
diff --git a/backend/src/ir/function.hpp b/backend/src/ir/function.hpp
index 165bc6f7..2ae8427d 100644
--- a/backend/src/ir/function.hpp
+++ b/backend/src/ir/function.hpp
@@ -66,8 +66,10 @@ namespace ir {
     INLINE void foreach(const T &functor) const {
       Instruction *curr = first;
       while (curr) {
+        // Be aware the current instruction can be destroyed in functor
+        Instruction *succ = curr->getSuccessor();
         functor(*curr);
-        curr = curr->getSuccessor();
+        curr = succ;
       }
     }
     /*! Apply the given functor on all instructions (reverse order) */
@@ -75,8 +77,10 @@ namespace ir {
     INLINE void rforeach(const T &functor) const {
       Instruction *curr = last;
       while (curr) {
+        // Be aware the current instruction can be destroyed in functor
+        Instruction *pred = curr->getPredecessor();
         functor(*curr);
-        curr = curr->getPredecessor();
+        curr = pred;
       }
     }
     /*! Get the parent function */
@@ -85,9 +89,11 @@ namespace ir {
     /*! Get the next and previous allocated block */
     BasicBlock *getNextBlock(void) const { return this->nextBlock; }
     BasicBlock *getPrevBlock(void) const { return this->prevBlock; }
-    /*! Get the first and last instructions */
+    /*! Get / set the first and last instructions */
     Instruction *getFirstInstruction(void) const { return this->first; }
     Instruction *getLastInstruction(void) const { return this->last; }
+    void setFirstInstruction(Instruction *insn) { this->first = insn; }
+    void setLastInstruction(Instruction *insn) { this->last = insn; }
     /*! Get successors and predecessors */
     const BlockSet &getSuccessorSet(void) const { return successors; }
     const BlockSet &getPredecessorSet(void) const { return predecessors; }
@@ -214,6 +220,10 @@ namespace ir {
     LabelIndex newLabel(void);
     /*! Create the control flow graph */
     void computeCFG(void);
+    /*! Sort the labels in increasing orders (ie top block has the smallest
+     *  labels)
+     */
+    void sortLabels(void);
     /*! Number of registers in the register file */
     INLINE uint32_t regNum(void) const { return file.regNum(); }
     /*! Number of register tuples in the register file */
diff --git a/backend/src/ir/instruction.cpp b/backend/src/ir/instruction.cpp
index 6941ac5c..7b04bd25 100644
--- a/backend/src/ir/instruction.cpp
+++ b/backend/src/ir/instruction.cpp
@@ -884,6 +884,24 @@ END_FUNCTION(Instruction, Register)
     GBE_ASSERT(bb != NULL);
     return bb->getParent();
   }
+  Function &Instruction::getFunction(void) {
+    BasicBlock *bb = this->getParent();
+    GBE_ASSERT(bb != NULL);
+    return bb->getParent();
+  }
+
+  void Instruction::replace(Instruction *other) {
+    Function &fn = other->getFunction();
+    BasicBlock *bb = other->getParent();
+    if (bb->getFirstInstruction() == other) bb->setFirstInstruction(this);
+    if (bb->getLastInstruction() == other) bb->setLastInstruction(this);
+    if (other->predecessor) other->predecessor->successor = this;
+    if (other->successor) other->successor->predecessor = this;
+    this->parent = other->parent;
+    this->predecessor = other->predecessor;
+    this->successor = other->successor;
+    fn.deleteInstruction(other);
+  }
 
 #define DECL_MEM_FN(CLASS, RET, PROTOTYPE, CALL)                  \
   RET CLASS::PROTOTYPE const {                                    \
diff --git a/backend/src/ir/instruction.hpp b/backend/src/ir/instruction.hpp
index 31acd176..5a485aba 100644
--- a/backend/src/ir/instruction.hpp
+++ b/backend/src/ir/instruction.hpp
@@ -109,11 +109,14 @@ namespace ir {
     void setParent(BasicBlock *block) { this->parent = block; }
     /*! Get the function from the parent basic block */
     const Function &getFunction(void) const;
+    Function &getFunction(void);
     /*! Check that the instruction is well formed (type properly match,
      *  registers not of bound and so on). If not well formed, provide a reason
      *  in string why
      */
     bool wellFormed(const Function &fn, std::string &why) const;
+    /*! Replace other by this instruction */
+    void replace(Instruction *other);
     /*! Indicates if the instruction belongs to instruction type T. Typically, T
      *  can be BinaryInstruction, UnaryInstruction, LoadInstruction and so on
      */
diff --git a/backend/src/ir/value.hpp b/backend/src/ir/value.hpp
index 673a2ab2..9d856f94 100644
--- a/backend/src/ir/value.hpp
+++ b/backend/src/ir/value.hpp
@@ -199,8 +199,12 @@ namespace ir {
     const ValueDef *getDefAddress(const Register &reg) const;
     /*! Get the pointer to the use *as stored in the DAG* */
     const ValueUse *getUseAddress(const Instruction *insn, uint32_t srcID) const;
+    /*! Get the set of all uses for the register */
+    const UseSet *getRegUse(const Register &reg) const;
+    /*! Get the set of all definitions for the register */
+    const DefSet *getRegDef(const Register &reg) const;
     /*! Get the function we have the graph for */
-    const Function &getFunction(void) const { return fn; }
+    INLINE const Function &getFunction(void) const { return fn; }
     /*! The DefSet for each definition use */
     typedef map<ValueUse, DefSet*> UDGraph;
     /*! The UseSet for each definition */
diff --git a/backend/src/llvm/llvm_to_gen.cpp b/backend/src/llvm/llvm_to_gen.cpp
index 802aed2b..9cdbbecf 100644
--- a/backend/src/llvm/llvm_to_gen.cpp
+++ b/backend/src/llvm/llvm_to_gen.cpp
@@ -27,19 +27,27 @@
 #include "llvm/PassManager.h"
 #include "llvm/Pass.h"
 #include "llvm/Support/IRReader.h"
+#include "llvm/Support/raw_ostream.h"
 #include "llvm/Transforms/Scalar.h"
+#include "llvm/Assembly/PrintModulePass.h"
 
 #include "llvm/llvm_gen_backend.hpp"
 #include "llvm/llvm_to_gen.hpp"
+#include "sys/cvar.hpp"
 #include "sys/platform.hpp"
 
 namespace gbe
 {
+  BVAR(OCL_OUTPUT_LLVM, false);
+  BVAR(OCL_OUTPUT_LLVM_BEFORE_EXTRA_PASS, false);
+
   bool llvmToGen(ir::Unit &unit, const char *fileName)
   {
     using namespace llvm;
     // Get the global LLVM context
     llvm::LLVMContext& c = llvm::getGlobalContext();
+    std::string errInfo;
+    llvm::raw_fd_ostream o("-", errInfo);
 
     // Get the module from its file
     SMDiagnostic Err;
@@ -49,14 +57,22 @@ namespace gbe
     Module &mod = *M.get();
 
     llvm::PassManager passes;
+
+    // Print the code before further optimizations
+    if (OCL_OUTPUT_LLVM_BEFORE_EXTRA_PASS)
+      passes.add(createPrintModulePass(&o));
     passes.add(createScalarReplAggregatesPass()); // Break up allocas
     passes.add(createRemoveGEPPass(unit));
     passes.add(createConstantPropagationPass());
-    passes.add(createDeadInstEliminationPass()); // remove simplified instructions
+    passes.add(createDeadInstEliminationPass());  // Remove simplified instructions
     passes.add(createLowerSwitchPass());
     passes.add(createPromoteMemoryToRegisterPass());
     passes.add(createGVNPass());                  // Remove redundancies
     passes.add(createGenPass(unit));
+
+    // Print the code extra optimization passes
+    if (OCL_OUTPUT_LLVM)
+      passes.add(createPrintModulePass(&o));
     passes.run(mod);
     return true;
   }
diff --git a/backend/src/llvm/stdlib.h b/backend/src/llvm/stdlib.h
new file mode 100644
index 00000000..eaf4b171
--- /dev/null
+++ b/backend/src/llvm/stdlib.h
@@ -0,0 +1,114 @@
+/* 
+ * 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 <http://www.gnu.org/licenses/>.
+ *
+ * Author: Benjamin Segovia <benjamin.segovia@intel.com>
+ */
+
+#define DECL_INTERNAL_WORK_ITEM_FN(NAME)                             \
+__attribute__((pure,const)) unsigned int __gen_ocl_##NAME##0(void);  \
+__attribute__((pure,const)) unsigned int __gen_ocl_##NAME##1(void);  \
+__attribute__((pure,const)) unsigned int __gen_ocl_##NAME##2(void);
+DECL_INTERNAL_WORK_ITEM_FN(get_group_id)
+DECL_INTERNAL_WORK_ITEM_FN(get_local_id)
+DECL_INTERNAL_WORK_ITEM_FN(get_local_size)
+DECL_INTERNAL_WORK_ITEM_FN(get_global_size)
+DECL_INTERNAL_WORK_ITEM_FN(get_num_groups)
+#undef DECL_INTERNAL_WORK_ITEM_FN
+
+#define DECL_PUBLIC_WORK_ITEM_FN(NAME)              \
+inline unsigned NAME(unsigned int dim) {            \
+  if (dim == 0) return __gen_ocl_##NAME##0();       \
+  else if (dim == 1) return __gen_ocl_##NAME##1();  \
+  else if (dim == 2) return __gen_ocl_##NAME##2();  \
+  else return 0;                                    \
+}
+DECL_PUBLIC_WORK_ITEM_FN(get_group_id)
+DECL_PUBLIC_WORK_ITEM_FN(get_local_id)
+DECL_PUBLIC_WORK_ITEM_FN(get_local_size)
+DECL_PUBLIC_WORK_ITEM_FN(get_global_size)
+DECL_PUBLIC_WORK_ITEM_FN(get_num_groups)
+#undef DECL_PUBLIC_WORK_ITEM_FN
+
+inline unsigned int get_global_id(unsigned int dim) {
+  return get_local_id(dim) + get_local_size(dim) * get_group_id(dim);
+}
+
+__attribute__ ((pure,const,overloadable)) float mad(float a, float b, float c);
+__attribute__((overloadable)) inline unsigned select(unsigned src0, unsigned src1, unsigned cond) {
+  return cond ? src0 : src1;
+}
+__attribute__((overloadable)) inline int select(int src0, int src1, int cond) {
+  return cond ? src0 : src1;
+}
+
+typedef float float2 __attribute__((ext_vector_type(2)));
+typedef float float3 __attribute__((ext_vector_type(3)));
+typedef float float4 __attribute__((ext_vector_type(4)));
+typedef int int2 __attribute__((ext_vector_type(2)));
+typedef int int3 __attribute__((ext_vector_type(3)));
+typedef int int4 __attribute__((ext_vector_type(4)));
+typedef int uint2 __attribute__((ext_vector_type(2)));
+typedef unsigned uint3 __attribute__((ext_vector_type(3)));
+typedef unsigned uint4 __attribute__((ext_vector_type(4)));
+typedef bool bool2 __attribute__((ext_vector_type(2)));
+typedef bool bool3 __attribute__((ext_vector_type(3)));
+typedef bool bool4 __attribute__((ext_vector_type(4)));
+
+// This will be optimized out by LLVM and will output LLVM select instructions
+#define DECL_SELECT4(TYPE4, TYPE, COND_TYPE4, MASK)                       \
+__attribute__((overloadable))                                             \
+inline TYPE4 select(TYPE4 src0, TYPE4 src1, COND_TYPE4 cond) {            \
+  TYPE4 dst;                                                              \
+  const TYPE x0 = src0.x; /* Fix performance issue with CLANG */          \
+  const TYPE x1 = src1.x;                                                 \
+  const TYPE y0 = src0.y;                                                 \
+  const TYPE y1 = src1.y;                                                 \
+  const TYPE z0 = src0.z;                                                 \
+  const TYPE z1 = src1.z;                                                 \
+  const TYPE w0 = src0.w;                                                 \
+  const TYPE w1 = src1.w;                                                 \
+                                                                          \
+  dst.x = (cond.x & MASK) ? x1 : x0;                                      \
+  dst.y = (cond.y & MASK) ? y1 : y0;                                      \
+  dst.z = (cond.z & MASK) ? z1 : z0;                                      \
+  dst.w = (cond.w & MASK) ? w1 : w0;                                      \
+  return dst;                                                             \
+}
+DECL_SELECT4(int4, int, int4, 0x80000000)
+DECL_SELECT4(float4, float, int4, 0x80000000)
+#undef DECL_SELECT4
+
+__attribute__((overloadable,always_inline)) inline float2 mad(float2 a, float2 b, float2 c) {
+  return (float2)(mad(a.x,b.x,c.x), mad(a.y,b.y,c.y));
+}
+__attribute__((overloadable,always_inline)) inline float3 mad(float3 a, float3 b, float3 c) {
+  return (float3)(mad(a.x,b.x,c.x), mad(a.y,b.y,c.y), mad(a.z,b.z,c.z));
+}
+__attribute__((overloadable,always_inline)) inline float4 mad(float4 a, float4 b, float4 c) {
+  return (float4)(mad(a.x,b.x,c.x), mad(a.y,b.y,c.y),
+                  mad(a.z,b.z,c.z), mad(a.w,b.w,c.w));
+}
+
+#define __private __attribute__((address_space(0)))
+#define __global __attribute__((address_space(1)))
+#define __constant __attribute__((address_space(2)))
+//#define __local __attribute__((address_space(3)))
+#define global __global
+//#define local __local
+#define constant __constant
+#define private __private
+
+#define NULL ((void*)0)
diff --git a/backend/src/llvm/stdlib_str.cpp b/backend/src/llvm/stdlib_str.cpp
new file mode 100644
index 00000000..41ce7fe6
--- /dev/null
+++ b/backend/src/llvm/stdlib_str.cpp
@@ -0,0 +1,120 @@
+/* 
+ * 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 <http://www.gnu.org/licenses/>.
+ *
+ * Author: Benjamin Segovia <benjamin.segovia@intel.com>
+ */
+
+#include "string"
+namespace gbe {
+std::string stdlib_str = 
+"#define DECL_INTERNAL_WORK_ITEM_FN(NAME)                             \\\n"
+"__attribute__((pure,const)) unsigned int __gen_ocl_##NAME##0(void);  \\\n"
+"__attribute__((pure,const)) unsigned int __gen_ocl_##NAME##1(void);  \\\n"
+"__attribute__((pure,const)) unsigned int __gen_ocl_##NAME##2(void);\n"
+"DECL_INTERNAL_WORK_ITEM_FN(get_group_id)\n"
+"DECL_INTERNAL_WORK_ITEM_FN(get_local_id)\n"
+"DECL_INTERNAL_WORK_ITEM_FN(get_local_size)\n"
+"DECL_INTERNAL_WORK_ITEM_FN(get_global_size)\n"
+"DECL_INTERNAL_WORK_ITEM_FN(get_num_groups)\n"
+"#undef DECL_INTERNAL_WORK_ITEM_FN\n"
+"\n"
+"#define DECL_PUBLIC_WORK_ITEM_FN(NAME)              \\\n"
+"inline unsigned NAME(unsigned int dim) {            \\\n"
+"  if (dim == 0) return __gen_ocl_##NAME##0();       \\\n"
+"  else if (dim == 1) return __gen_ocl_##NAME##1();  \\\n"
+"  else if (dim == 2) return __gen_ocl_##NAME##2();  \\\n"
+"  else return 0;                                    \\\n"
+"}\n"
+"DECL_PUBLIC_WORK_ITEM_FN(get_group_id)\n"
+"DECL_PUBLIC_WORK_ITEM_FN(get_local_id)\n"
+"DECL_PUBLIC_WORK_ITEM_FN(get_local_size)\n"
+"DECL_PUBLIC_WORK_ITEM_FN(get_global_size)\n"
+"DECL_PUBLIC_WORK_ITEM_FN(get_num_groups)\n"
+"#undef DECL_PUBLIC_WORK_ITEM_FN\n"
+"\n"
+"inline unsigned int get_global_id(unsigned int dim) {\n"
+"  return get_local_id(dim) + get_local_size(dim) * get_group_id(dim);\n"
+"}\n"
+"\n"
+"__attribute__ ((pure,const,overloadable)) float mad(float a, float b, float c);\n"
+"__attribute__((overloadable)) inline unsigned select(unsigned src0, unsigned src1, unsigned cond) {\n"
+"  return cond ? src0 : src1;\n"
+"}\n"
+"__attribute__((overloadable)) inline int select(int src0, int src1, int cond) {\n"
+"  return cond ? src0 : src1;\n"
+"}\n"
+"\n"
+"typedef float float2 __attribute__((ext_vector_type(2)));\n"
+"typedef float float3 __attribute__((ext_vector_type(3)));\n"
+"typedef float float4 __attribute__((ext_vector_type(4)));\n"
+"typedef int int2 __attribute__((ext_vector_type(2)));\n"
+"typedef int int3 __attribute__((ext_vector_type(3)));\n"
+"typedef int int4 __attribute__((ext_vector_type(4)));\n"
+"typedef int uint2 __attribute__((ext_vector_type(2)));\n"
+"typedef unsigned uint3 __attribute__((ext_vector_type(3)));\n"
+"typedef unsigned uint4 __attribute__((ext_vector_type(4)));\n"
+"typedef bool bool2 __attribute__((ext_vector_type(2)));\n"
+"typedef bool bool3 __attribute__((ext_vector_type(3)));\n"
+"typedef bool bool4 __attribute__((ext_vector_type(4)));\n"
+"\n"
+"// This will be optimized out by LLVM and will output LLVM select instructions\n"
+"#define DECL_SELECT4(TYPE4, TYPE, COND_TYPE4, MASK)                       \\\n"
+"__attribute__((overloadable))                                             \\\n"
+"inline TYPE4 select(TYPE4 src0, TYPE4 src1, COND_TYPE4 cond) {            \\\n"
+"  TYPE4 dst;                                                              \\\n"
+"  const TYPE x0 = src0.x; /* Fix performance issue with CLANG */          \\\n"
+"  const TYPE x1 = src1.x;                                                 \\\n"
+"  const TYPE y0 = src0.y;                                                 \\\n"
+"  const TYPE y1 = src1.y;                                                 \\\n"
+"  const TYPE z0 = src0.z;                                                 \\\n"
+"  const TYPE z1 = src1.z;                                                 \\\n"
+"  const TYPE w0 = src0.w;                                                 \\\n"
+"  const TYPE w1 = src1.w;                                                 \\\n"
+"                                                                          \\\n"
+"  dst.x = (cond.x & MASK) ? x1 : x0;                                      \\\n"
+"  dst.y = (cond.y & MASK) ? y1 : y0;                                      \\\n"
+"  dst.z = (cond.z & MASK) ? z1 : z0;                                      \\\n"
+"  dst.w = (cond.w & MASK) ? w1 : w0;                                      \\\n"
+"  return dst;                                                             \\\n"
+"}\n"
+"DECL_SELECT4(int4, int, int4, 0x80000000)\n"
+"DECL_SELECT4(float4, float, int4, 0x80000000)\n"
+"#undef DECL_SELECT4\n"
+"\n"
+"__attribute__((overloadable,always_inline)) inline float2 mad(float2 a, float2 b, float2 c) {\n"
+"  return (float2)(mad(a.x,b.x,c.x), mad(a.y,b.y,c.y));\n"
+"}\n"
+"__attribute__((overloadable,always_inline)) inline float3 mad(float3 a, float3 b, float3 c) {\n"
+"  return (float3)(mad(a.x,b.x,c.x), mad(a.y,b.y,c.y), mad(a.z,b.z,c.z));\n"
+"}\n"
+"__attribute__((overloadable,always_inline)) inline float4 mad(float4 a, float4 b, float4 c) {\n"
+"  return (float4)(mad(a.x,b.x,c.x), mad(a.y,b.y,c.y),\n"
+"                  mad(a.z,b.z,c.z), mad(a.w,b.w,c.w));\n"
+"}\n"
+"\n"
+"#define __private __attribute__((address_space(0)))\n"
+"#define __global __attribute__((address_space(1)))\n"
+"#define __constant __attribute__((address_space(2)))\n"
+"//#define __local __attribute__((address_space(3)))\n"
+"#define global __global\n"
+"//#define local __local\n"
+"#define constant __constant\n"
+"#define private __private\n"
+"\n"
+"#define NULL ((void*)0)\n"
+;
+}
+
diff --git a/backend/src/sys/assert.cpp b/backend/src/sys/assert.cpp
index d13017cf..ea3d34c2 100644
--- a/backend/src/sys/assert.cpp
+++ b/backend/src/sys/assert.cpp
@@ -39,7 +39,7 @@ namespace gbe
                           + std::string(file)
                           + ", function " + std::string(fn)
                           + ", line " + std::string(lineString);
-    // assert(0);
+    assert(0);
     throw Exception(str);
   }
 } /* namespace gbe */
diff --git a/backend/src/utest/utest_vector.cpp b/backend/src/utest/utest_vector.cpp
index 26ef64b8..6d9e77d9 100644
--- a/backend/src/utest/utest_vector.cpp
+++ b/backend/src/utest/utest_vector.cpp
@@ -353,7 +353,7 @@ static void utestScatterGather(void)
     LOAD(_##INDEX##g, (const char *) (gatherOffsets+index##INDEX));\
     LOAD(_##INDEX##s, (const char *) (scatterOffsets+index##INDEX));\
     GATHER(_##INDEX, _##INDEX##g, (const char *) data);\
-    SCATTER(_##INDEX, _##INDEX##s, (char *) dst);\
+    SCATTER(_##INDEX##s, _##INDEX, (char *) dst);\
     for (uint32_t i = 0; i < elemNum(_##INDEX); ++i)\
       GBE_ASSERT(data[gatherOffsets[index##INDEX+i] / sizeof(uint32_t)] ==\
                  dst[scatterOffsets[index##INDEX+i] / sizeof(uint32_t)]);