Started to implement liveness analysis

diff --git a/backend/kernels/function.cl b/backend/kernels/function.cl
new file mode 100644 (file)
index 0000000..0cc6873
--- /dev/null
+++ b/backend/kernels/function.cl
@@ -0,0 +1,12 @@
+#include <stdlib.h>
+
+void write(__global int *dst)
+{
+  dst[0] = 1;
+}
+
+__kernel void write2(__global int *dst, int x)
+{
+  write(dst);
+  dst[x] = 1;
+}

diff --git a/backend/kernels/function.ll b/backend/kernels/function.ll
new file mode 100644 (file)
index 0000000..e428448
--- /dev/null
+++ b/backend/kernels/function.ll
@@ -0,0 +1,24 @@
+; ModuleID = 'function.o'
+target datalayout = "e-p:32:32-i64:64:64-f64:64:64-n1:8:16:32:64"
+target triple = "ptx32--"
+
+define ptx_device void @write(i32 addrspace(1)* nocapture %dst) nounwind {
+entry:
+  store i32 1, i32 addrspace(1)* %dst, align 4, !tbaa !1
+  ret void
+}
+
+define ptx_kernel void @write2(i32 addrspace(1)* nocapture %dst, i32 %x) nounwind noinline {
+entry:
+  store i32 1, i32 addrspace(1)* %dst, align 4, !tbaa !1
+  %arrayidx = getelementptr inbounds i32 addrspace(1)* %dst, i32 %x
+  store i32 1, i32 addrspace(1)* %arrayidx, align 4, !tbaa !1
+  ret void
+}
+
+!opencl.kernels = !{!0}
+
+!0 = metadata !{void (i32 addrspace(1)*, i32)* @write2}
+!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/src/ir/function.cpp b/backend/src/ir/function.cpp
index e2f06da..9fb1c92 100644 (file)
--- a/backend/src/ir/function.cpp
+++ b/backend/src/ir/function.cpp
@@ -62,7 +62,7 @@ namespace ir {
       case TYPE_HALF: out << "half(" << imm.data.u16 << ")"; break;
       case TYPE_FLOAT: out << imm.data.f32; break;
       case TYPE_DOUBLE: out << imm.data.f64; break;
-    };
+    }
   }
 
   std::ostream &operator<< (std::ostream &out, const Function &fn)

diff --git a/backend/src/ir/function.hpp b/backend/src/ir/function.hpp
index c4aaf04..94b49ed 100644 (file)
--- a/backend/src/ir/function.hpp
+++ b/backend/src/ir/function.hpp
@@ -161,6 +161,12 @@ namespace ir {
     INLINE uint32_t blockNum(void) const { return blocks.size(); }
     /*! Output an immediate value in a stream */
     void outImmediate(std::ostream &out, ImmediateIndex index) const;
+    /*! Apply the given functor on all basic blocks */
+    template <typename T>
+    INLINE void apply(const T &functor) const {
+      for (auto it = blocks.begin; it != blocks.end(); ++it)
+        functor(*it);
+    }
   private:
     friend class Context;         //!< Can freely modify a function
     std::string name;             //!< Function name

diff --git a/backend/src/ir/liveness.cpp b/backend/src/ir/liveness.cpp
new file mode 100644 (file)
index 0000000..7e55fac
--- /dev/null
+++ b/backend/src/ir/liveness.cpp
@@ -0,0 +1,58 @@
+/* 
+ * 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>
+ */
+
+/**
+ * \file liveness.cpp
+ * \author Benjamin Segovia <benjamin.segovia@intel.com>
+ */
+#include <sys/map.hpp>
+#include <sys/set.hpp>
+
+namespace gbe
+{
+  /*! Compute liveness of each register */
+  class LivenessInfo
+  {
+  public:
+    LivenessInfo(Function &fn) : fn(fn) {}
+    /*! Set of variables used upwards in the block (before a definition) */
+    typedef set<Register> UsedVar;
+    /*! Set of variables alive at the exit of the block */
+    typedef set<Register> LiveOut;
+    /*! Set of variables actually killed in each block */
+    typedef set<Register> Kill;
+    /*! Per-block info */
+    struct BlockInfo {
+      UEVar upwardUsed;
+      LiveOut liveOut;
+      Kill kill;
+    };
+    /*! Gives for each block the variables alive at entry / exit */
+    typedef map<BasicBlock*, BlockInfo> BlockLiveness;
+    /*! Compute the liveness for this function */
+    Function &fn;
+  };
+
+  LivenessInfo::LivenessInfo(Function &fn) : fn(fn) {
+    
+
+  }
+
+} /* namespace gbe */
+