Started to fix problem with blocks that do not end with BRA instructions

diff --git a/backend/kernels/loop2.cl b/backend/kernels/loop2.cl
new file mode 100644 (file)
index 0000000..04997ae
--- /dev/null
+++ b/backend/kernels/loop2.cl
@@ -0,0 +1,14 @@
+#include "stdlib.h"
+
+struct big { int x[10]; };
+
+__kernel void add(__global int *dst, unsigned int x, struct big b)
+{
+  for (int i = 0; i < x; ++i) {
+    if (dst[i+1] > 0)
+      dst[i]++;
+    else
+      dst[i] += 2;
+  }
+}
+

diff --git a/backend/kernels/loop2.ll b/backend/kernels/loop2.ll
new file mode 100644 (file)
index 0000000..abb16d7
--- /dev/null
+++ b/backend/kernels/loop2.ll
@@ -0,0 +1,39 @@
+; ModuleID = 'loop2.o'
+target datalayout = "e-p:32:32-i64:64:64-f64:64:64-n1:8:16:32:64"
+target triple = "ptx32--"
+
+%struct.big = type { [10 x i32] }
+
+define ptx_kernel void @add(i32 addrspace(1)* nocapture %dst, i32 %x, %struct.big* nocapture byval %b) nounwind noinline {
+entry:
+  %cmp6 = icmp eq i32 %x, 0
+  br i1 %cmp6, label %for.end, label %for.body.lr.ph
+
+for.body.lr.ph:                                   ; preds = %entry
+  %.pre = load i32 addrspace(1)* %dst, align 4, !tbaa !1
+  br label %for.body
+
+for.body:                                         ; preds = %for.body, %for.body.lr.ph
+  %0 = phi i32 [ %.pre, %for.body.lr.ph ], [ %1, %for.body ]
+  %i.07 = phi i32 [ 0, %for.body.lr.ph ], [ %add, %for.body ]
+  %add = add nsw i32 %i.07, 1
+  %arrayidx = getelementptr inbounds i32 addrspace(1)* %dst, i32 %add
+  %1 = load i32 addrspace(1)* %arrayidx, align 4, !tbaa !1
+  %cmp1 = icmp sgt i32 %1, 0
+  %arrayidx2 = getelementptr inbounds i32 addrspace(1)* %dst, i32 %i.07
+  %storemerge.v = select i1 %cmp1, i32 1, i32 2
+  %storemerge = add i32 %storemerge.v, %0
+  store i32 %storemerge, i32 addrspace(1)* %arrayidx2, align 4
+  %exitcond = icmp eq i32 %add, %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 addrspace(1)*, i32, %struct.big*)* @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/src/ir/context.cpp b/backend/src/ir/context.cpp
index 838ce9e..1652b2d 100644 (file)
--- 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");
+    // std::cout << *fn << std::endl;
     fn->computeCFG();
     GBE_DELETE(usedLabels);
     const StackElem elem = fnStack.back();

diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp
index 31b4dfc..e666dad 100644 (file)
--- a/backend/src/llvm/llvm_gen_backend.cpp
+++ b/backend/src/llvm/llvm_gen_backend.cpp
@@ -1106,12 +1106,13 @@ namespace gbe
     // Emit MOVs if required
     BasicBlock *bb = I.getParent();
     this->emitMovForPHI(bb, I.getSuccessor(0));
-    this->emitMovForPHI(bb, I.getSuccessor(1));
+    if (I.isConditional())
+      this->emitMovForPHI(bb, I.getSuccessor(1));
 
     // Inconditional branch. Just check that we jump to a block which is not our
     // successor
     if (I.isConditional() == false) {
-    BasicBlock *target = I.getSuccessor(0);
+      BasicBlock *target = I.getSuccessor(0);
       if (llvm::next(Function::iterator(bb)) != Function::iterator(target)) {
         GBE_ASSERT(labelMap.find(target) != labelMap.end());
         const ir::LabelIndex labelIndex = labelMap[bb];

diff --git a/backend/src/utest/utest_llvm.cpp b/backend/src/utest/utest_llvm.cpp
index d368c35..172c73d 100644 (file)
--- a/backend/src/utest/utest_llvm.cpp
+++ b/backend/src/utest/utest_llvm.cpp
@@ -84,7 +84,7 @@ runTests:
   GBE_ASSERT(dummyKernel != NULL);
   fclose(dummyKernel);
 
-  UTEST_EXPECT_SUCCESS(utestLLVM2Gen("loop.ll"));
+  UTEST_EXPECT_SUCCESS(utestLLVM2Gen("loop2.ll"));
   //UTEST_EXPECT_SUCCESS(utestLLVM2Gen("function_param.ll"));
   //UTEST_EXPECT_SUCCESS(utestLLVM2Gen("function.ll"));
   //UTEST_EXPECT_SUCCESS(utestLLVM2Gen("mad.ll"));