Extended the llvm-to-gen translation pass

10 files changed, 92 insertions, 28 deletions

diff --git a/backend/kernels/get_global_id.cbe.c b/backend/kernels/get_global_id.cbe.c
index 4dbae41c..f88bd5c1 100644
--- a/backend/kernels/get_global_id.cbe.c
+++ b/backend/kernels/get_global_id.cbe.c
@@ -131,8 +131,9 @@ typedef union {
 double fmod(double, double);
 float fmodf(float, float);
 long double fmodl(long double, long double);
-void test_global_id(unsigned int *llvm_cbe_dst);
-unsigned int __gen_get_global_id0(void);
+void test_global_id(unsigned int *llvm_cbe_dst, unsigned int *llvm_cbe_p);
+unsigned int __gen_ocl_get_global_id0(void);
+unsigned int __gen_ocl_get_local_id0(void);
 void abort(void);
 
 
@@ -152,11 +153,14 @@ 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 test_global_id(unsigned int *llvm_cbe_dst) {
+void test_global_id(unsigned int *llvm_cbe_dst, unsigned int *llvm_cbe_p) {
   unsigned int llvm_cbe_call_2e_i;
+  unsigned int llvm_cbe_call_2e_i6;
 
-  llvm_cbe_call_2e_i =  /*tail*/ __gen_get_global_id0();
-  *((&llvm_cbe_dst[((signed int )llvm_cbe_call_2e_i)])) = 1u;
+  llvm_cbe_call_2e_i =  /*tail*/ __gen_ocl_get_local_id0();
+  llvm_cbe_call_2e_i6 =  /*tail*/ __gen_ocl_get_global_id0();
+  *((&llvm_cbe_dst[((signed int )llvm_cbe_call_2e_i6)])) = (((signed int )(((signed int )(llvm_cbe_call_2e_i << 16u)) >> ((signed int )16u))));
+  *((&llvm_cbe_p[((signed int )llvm_cbe_call_2e_i6)])) = llvm_cbe_call_2e_i;
   return;
 }
 
diff --git a/backend/kernels/get_global_id.cl b/backend/kernels/get_global_id.cl
index 86500ada..299d6c3d 100644
--- a/backend/kernels/get_global_id.cl
+++ b/backend/kernels/get_global_id.cl
@@ -1,17 +1,28 @@
-__attribute__((pure)) unsigned int __gen_get_global_id0(void);
-__attribute__((pure)) unsigned int __gen_get_global_id1(void);
-__attribute__((pure)) unsigned int __gen_get_global_id2(void);
+__attribute__((pure,const)) unsigned int __gen_ocl_get_global_id0(void);
+__attribute__((pure,const)) unsigned int __gen_ocl_get_global_id1(void);
+__attribute__((pure,const)) unsigned int __gen_ocl_get_global_id2(void);
+__attribute__((pure,const)) unsigned int __gen_ocl_get_local_id0(void);
+__attribute__((pure,const)) unsigned int __gen_ocl_get_local_id1(void);
+__attribute__((pure,const)) unsigned int __gen_ocl_get_local_id2(void);
 
 inline unsigned get_global_id(unsigned int dim) {
-  if (dim == 0) return __gen_get_global_id0();
-  else if (dim == 1) return __gen_get_global_id1();
-  else if (dim == 2) return __gen_get_global_id2();
+  if (dim == 0) return __gen_ocl_get_global_id0();
+  else if (dim == 1) return __gen_ocl_get_global_id1();
+  else if (dim == 2) return __gen_ocl_get_global_id2();
   else return 0;
 }
 
-__kernel void test_global_id(__global int *dst)
+inline unsigned get_local_id(unsigned int dim) {
+  if (dim == 0) return __gen_ocl_get_local_id0();
+  else if (dim == 1) return __gen_ocl_get_local_id1();
+  else if (dim == 2) return __gen_ocl_get_local_id2();
+  else return 0;
+}
+
+__kernel void test_global_id(__global int *dst, __global int *p)
 {
-  short hop = get_global_id(0);
+  short hop = get_local_id(0);
   dst[get_global_id(0)] = hop;
+  p[get_global_id(0)] = get_local_id(0);
 }
 
diff --git a/backend/kernels/get_global_id.ll b/backend/kernels/get_global_id.ll
index 965739a1..1df9fdff 100644
--- a/backend/kernels/get_global_id.ll
+++ b/backend/kernels/get_global_id.ll
@@ -2,21 +2,26 @@
 target datalayout = "e-p:32:32-i64:64:64-f64:64:64-n1:8:16:32:64"
 target triple = "ptx32--"
 
-define ptx_kernel void @test_global_id(i32* nocapture %dst) nounwind noinline {
-get_global_id.exit5:
-  %call.i = tail call ptx_device i32 @__gen_get_global_id0() nounwind readonly
+define ptx_kernel void @test_global_id(i32* nocapture %dst, i32* nocapture %p) nounwind noinline {
+get_global_id.exit13:
+  %call.i = tail call ptx_device i32 @__gen_ocl_get_local_id0() nounwind readnone
   %sext = shl i32 %call.i, 16
   %conv1 = ashr exact i32 %sext, 16
-  %arrayidx = getelementptr inbounds i32* %dst, i32 %call.i
+  %call.i6 = tail call ptx_device i32 @__gen_ocl_get_global_id0() nounwind readnone
+  %arrayidx = getelementptr inbounds i32* %dst, i32 %call.i6
   store i32 %conv1, i32* %arrayidx, align 4, !tbaa !1
+  %arrayidx5 = getelementptr inbounds i32* %p, i32 %call.i6
+  store i32 %call.i, i32* %arrayidx5, align 4, !tbaa !1
   ret void
 }
 
-declare ptx_device i32 @__gen_get_global_id0() nounwind readonly
+declare ptx_device i32 @__gen_ocl_get_global_id0() nounwind readnone
+
+declare ptx_device i32 @__gen_ocl_get_local_id0() nounwind readnone
 
 !opencl.kernels = !{!0}
 
-!0 = metadata !{void (i32*)* @test_global_id}
+!0 = metadata !{void (i32*, i32*)* @test_global_id}
 !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/get_global_id.o b/backend/kernels/get_global_id.o
index e21b2e1f..d1ddd393 100644
--- a/backend/kernels/get_global_id.o
+++ b/backend/kernels/get_global_id.o
diff --git a/backend/kernels/struct.cl b/backend/kernels/struct.cl
index 5915d311..af3b92d4 100644
--- a/backend/kernels/struct.cl
+++ b/backend/kernels/struct.cl
@@ -10,6 +10,6 @@ __kernel void struct_cl (struct my_struct s, int x, __global int *mem)
   __local int array[256];
   for (int i = 0; i < 256; ++i)
     array[i] = i;
-  mem[0] = array[x] + g[x];
+  mem[0] = s.a + array[x] + g[x];
 }
 
diff --git a/backend/kernels/struct.o.ll b/backend/kernels/struct.ll
index 517da9ef..acbb3fa2 100644
--- a/backend/kernels/struct.o.ll
+++ b/backend/kernels/struct.ll
@@ -12,20 +12,23 @@ entry:
   br label %for.body
 
 for.body:                                         ; preds = %for.body, %entry
-  %i.04 = phi i32 [ 0, %entry ], [ %inc, %for.body ]
-  %arrayidx = getelementptr inbounds [256 x i32] addrspace(4)* @struct_cl.array, i32 0, i32 %i.04
-  store i32 %i.04, i32 addrspace(4)* %arrayidx, align 4, !tbaa !1
-  %inc = add nsw i32 %i.04, 1
+  %i.05 = phi i32 [ 0, %entry ], [ %inc, %for.body ]
+  %arrayidx = getelementptr inbounds [256 x i32] addrspace(4)* @struct_cl.array, i32 0, i32 %i.05
+  store i32 %i.05, i32 addrspace(4)* %arrayidx, align 4, !tbaa !1
+  %inc = add nsw i32 %i.05, 1
   %exitcond = icmp eq i32 %inc, 256
   br i1 %exitcond, label %for.end, label %for.body
 
 for.end:                                          ; preds = %for.body
+  %a = getelementptr inbounds %struct.my_struct* %s, i32 0, i32 0
+  %0 = load i32* %a, align 4, !tbaa !1
   %arrayidx1 = getelementptr inbounds [256 x i32] addrspace(4)* @struct_cl.array, i32 0, i32 %x
-  %0 = load i32 addrspace(4)* %arrayidx1, align 4, !tbaa !1
+  %1 = load i32 addrspace(4)* %arrayidx1, align 4, !tbaa !1
   %arrayidx2 = getelementptr inbounds [4 x i32] addrspace(1)* @g, i32 0, i32 %x
-  %1 = load i32 addrspace(1)* %arrayidx2, align 4, !tbaa !1
-  %add = add nsw i32 %1, %0
-  store i32 %add, i32* %mem, align 4, !tbaa !1
+  %2 = load i32 addrspace(1)* %arrayidx2, align 4, !tbaa !1
+  %add = add i32 %1, %0
+  %add3 = add i32 %add, %2
+  store i32 %add3, i32* %mem, align 4, !tbaa !1
   ret void
 }
 
diff --git a/backend/kernels/struct.o b/backend/kernels/struct.o
index 48db4e33..4f6af9c7 100644
--- a/backend/kernels/struct.o
+++ b/backend/kernels/struct.o
diff --git a/backend/kernels/undefined.cl b/backend/kernels/undefined.cl
new file mode 100644
index 00000000..f9153ffb
--- /dev/null
+++ b/backend/kernels/undefined.cl
@@ -0,0 +1,9 @@
+__kernel void undefined(__global int *dst)
+{
+  int x;
+  if (x == 0)
+    dst[0] = 0;
+  else
+    dst[0] = 1;
+}
+
diff --git a/backend/kernels/undefined.ll b/backend/kernels/undefined.ll
new file mode 100644
index 00000000..a706e7ba
--- /dev/null
+++ b/backend/kernels/undefined.ll
@@ -0,0 +1,32 @@
+; ModuleID = 'undefined.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 @undefined(i32* %dst) nounwind noinline {
+entry:
+  %dst.addr = alloca i32*, align 4
+  %x = alloca i32, align 4
+  store i32* %dst, i32** %dst.addr, align 4
+  %0 = load i32* %x, align 4
+  %cmp = icmp eq i32 %0, 0
+  br i1 %cmp, label %if.then, label %if.else
+
+if.then:                                          ; preds = %entry
+  %1 = load i32** %dst.addr, align 4
+  %arrayidx = getelementptr inbounds i32* %1, i32 0
+  store i32 0, i32* %arrayidx
+  br label %if.end
+
+if.else:                                          ; preds = %entry
+  %2 = load i32** %dst.addr, align 4
+  %arrayidx1 = getelementptr inbounds i32* %2, i32 0
+  store i32 1, i32* %arrayidx1
+  br label %if.end
+
+if.end:                                           ; preds = %if.else, %if.then
+  ret void
+}
+
+!opencl.kernels = !{!0}
+
+!0 = metadata !{void (i32*)* @undefined}
diff --git a/backend/kernels/undefined.o b/backend/kernels/undefined.o
new file mode 100644
index 00000000..d20bc495
--- /dev/null
+++ b/backend/kernels/undefined.o