Added more kernels for the very basic tests

diff --git a/backend/kernels/compile.sh b/backend/kernels/compile.sh
index 380017e..7f5ef36 100755 (executable)
--- a/backend/kernels/compile.sh
+++ b/backend/kernels/compile.sh
@@ -1,4 +1,4 @@
-clang -x cl -emit-llvm -O3 -ccc-host-triple ptx32 -c $1.cl -o $1.o
+clang -emit-llvm -O3 -ccc-host-triple ptx32 -c $1.cl -o $1.o
 llvm-dis $1.o
 rm $1.o
 mv $1.o.ll $1.ll

diff --git a/backend/kernels/insert.cl b/backend/kernels/insert.cl
new file mode 100644 (file)
index 0000000..5262266
--- /dev/null
+++ b/backend/kernels/insert.cl
@@ -0,0 +1,7 @@
+#include "stdlib.h"
+__kernel void insert(__global int4 *dst, __global int4 *src, int c)
+{
+  dst[0].x = dst[0][c];
+  dst[0].yzw = dst[1].xyz + src[0].xyz;
+}
+

diff --git a/backend/kernels/insert.ll b/backend/kernels/insert.ll
new file mode 100644 (file)
index 0000000..408102c
--- /dev/null
+++ b/backend/kernels/insert.ll
@@ -0,0 +1,27 @@
+; ModuleID = 'insert.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 @insert(<4 x i32>* nocapture %dst, <4 x i32>* nocapture %src, i32 %c) nounwind noinline {
+entry:
+  %0 = load <4 x i32>* %dst, align 16, !tbaa !1
+  %vecext = extractelement <4 x i32> %0, i32 %c
+  %1 = insertelement <4 x i32> %0, i32 %vecext, i32 0
+  store <4 x i32> %1, <4 x i32>* %dst, align 16
+  %arrayidx2 = getelementptr inbounds <4 x i32>* %dst, i32 1
+  %2 = load <4 x i32>* %arrayidx2, align 16
+  %3 = shufflevector <4 x i32> %2, <4 x i32> undef, <3 x i32> <i32 0, i32 1, i32 2>
+  %4 = load <4 x i32>* %src, align 16
+  %5 = shufflevector <4 x i32> %4, <4 x i32> undef, <3 x i32> <i32 0, i32 1, i32 2>
+  %add = add <3 x i32> %3, %5
+  %6 = shufflevector <3 x i32> %add, <3 x i32> undef, <4 x i32> <i32 0, i32 1, i32 2, i32 undef>
+  %7 = shufflevector <4 x i32> %1, <4 x i32> %6, <4 x i32> <i32 0, i32 4, i32 5, i32 6>
+  store <4 x i32> %7, <4 x i32>* %dst, align 16
+  ret void
+}
+
+!opencl.kernels = !{!0}
+
+!0 = metadata !{void (<4 x i32>*, <4 x i32>*, i32)* @insert}
+!1 = metadata !{metadata !"omnipotent char", metadata !2}
+!2 = metadata !{metadata !"Simple C/C++ TBAA", null}

diff --git a/backend/kernels/struct.cl b/backend/kernels/struct.cl
index a93d1e4..a4fa09c 100644 (file)
--- a/backend/kernels/struct.cl
+++ b/backend/kernels/struct.cl
@@ -7,13 +7,18 @@ const __constant int g[4] = {0,1,2,3};
 
 __kernel void struct_cl (struct my_struct s, int x, __global int *mem, int y)
 {
+  __local struct my_struct hop;
   __local struct my_struct array[256];
   for (int i = 0; i < 256; ++i) {
     array[i].a = i;
     array[i].b[0] = i;
     array[i].b[0] = i+1;
   }
-  array[0] = array[y];
+  if (y == 0)
+    hop = array[y];
+  else
+    hop = array[y+1];
+  array[0] = hop;
   mem[0] = s.a + array[x].a + array[x+1].b[0] + g[x] + g[3];
 }
 

diff --git a/backend/kernels/struct.ll b/backend/kernels/struct.ll
index ac4394d..d1273be 100644 (file)
--- a/backend/kernels/struct.ll
+++ b/backend/kernels/struct.ll
@@ -1,52 +1,121 @@
 ; ModuleID = 'struct.o'
-target datalayout = "e-p:32:32-i64:64:64-f64:64:64-n1:8:16:32:64"
-target triple = "ptx32--"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v64:64:64-v128:128:128-a0:0:64-s0:64:64-f80:128:128-n8:16:32:64-S128"
+target triple = "x86_64-unknown-linux-gnu"
 
 %struct.my_struct = type { i32, [2 x i32] }
 
-@g = addrspace(1) constant [4 x i32] [i32 0, i32 1, i32 2, i32 3], align 4
-@struct_cl.array = internal addrspace(4) global [256 x %struct.my_struct] zeroinitializer, align 4
+@g = constant [4 x i32] [i32 0, i32 1, i32 2, i32 3], align 16
+@struct_cl.hop = internal global %struct.my_struct zeroinitializer, align 4
+@struct_cl.array = internal global [256 x %struct.my_struct] zeroinitializer, align 16
 
-define ptx_kernel void @struct_cl(%struct.my_struct* nocapture byval %s, i32 %x, i32* nocapture %mem, i32 %y) nounwind noinline {
+define void @struct_cl(i64 %s.coerce0, i32 %s.coerce1, i32 %x, i32* %mem, i32 %y) nounwind uwtable {
 entry:
-  br label %for.body
-
-for.body:                                         ; preds = %for.body, %entry
-  %i.020 = phi i32 [ 0, %entry ], [ %add, %for.body ]
-  %a = getelementptr inbounds [256 x %struct.my_struct] addrspace(4)* @struct_cl.array, i32 0, i32 %i.020, i32 0
-  store i32 %i.020, i32 addrspace(4)* %a, align 4, !tbaa !1
-  %arrayidx2 = getelementptr inbounds [256 x %struct.my_struct] addrspace(4)* @struct_cl.array, i32 0, i32 %i.020, i32 1, i32 0
-  %add = add nsw i32 %i.020, 1
-  store i32 %add, i32 addrspace(4)* %arrayidx2, align 4, !tbaa !1
-  %exitcond = icmp eq i32 %add, 256
-  br i1 %exitcond, label %for.end, label %for.body
-
-for.end:                                          ; preds = %for.body
-  %arrayidx6 = getelementptr inbounds [256 x %struct.my_struct] addrspace(4)* @struct_cl.array, i32 0, i32 %y
-  %0 = bitcast %struct.my_struct addrspace(4)* %arrayidx6 to i8 addrspace(4)*
-  tail call void @llvm.memcpy.p4i8.p4i8.i32(i8 addrspace(4)* bitcast ([256 x %struct.my_struct] addrspace(4)* @struct_cl.array to i8 addrspace(4)*), i8 addrspace(4)* %0, i32 12, i32 4, i1 false)
-  %a7 = getelementptr inbounds %struct.my_struct* %s, i32 0, i32 0
-  %1 = load i32* %a7, align 4, !tbaa !1
-  %a9 = getelementptr inbounds [256 x %struct.my_struct] addrspace(4)* @struct_cl.array, i32 0, i32 %x, i32 0
-  %2 = load i32 addrspace(4)* %a9, align 4, !tbaa !1
-  %add11 = add nsw i32 %x, 1
-  %arrayidx14 = getelementptr inbounds [256 x %struct.my_struct] addrspace(4)* @struct_cl.array, i32 0, i32 %add11, i32 1, i32 0
-  %3 = load i32 addrspace(4)* %arrayidx14, align 4, !tbaa !1
-  %arrayidx16 = getelementptr inbounds [4 x i32] addrspace(1)* @g, i32 0, i32 %x
-  %4 = load i32 addrspace(1)* %arrayidx16, align 4, !tbaa !1
-  %add10 = add i32 %1, 3
-  %add15 = add i32 %add10, %2
-  %add17 = add i32 %add15, %3
-  %add18 = add i32 %add17, %4
-  store i32 %add18, i32* %mem, align 4, !tbaa !1
+  %s = alloca %struct.my_struct, align 8
+  %x.addr = alloca i32, align 4
+  %mem.addr = alloca i32*, align 8
+  %y.addr = alloca i32, align 4
+  %i = alloca i32, align 4
+  %0 = bitcast %struct.my_struct* %s to { i64, i32 }*
+  %1 = getelementptr { i64, i32 }* %0, i32 0, i32 0
+  store i64 %s.coerce0, i64* %1
+  %2 = getelementptr { i64, i32 }* %0, i32 0, i32 1
+  store i32 %s.coerce1, i32* %2
+  store i32 %x, i32* %x.addr, align 4
+  store i32* %mem, i32** %mem.addr, align 8
+  store i32 %y, i32* %y.addr, align 4
+  store i32 0, i32* %i, align 4
+  br label %for.cond
+
+for.cond:                                         ; preds = %for.inc, %entry
+  %3 = load i32* %i, align 4
+  %cmp = icmp slt i32 %3, 256
+  br i1 %cmp, label %for.body, label %for.end
+
+for.body:                                         ; preds = %for.cond
+  %4 = load i32* %i, align 4
+  %5 = load i32* %i, align 4
+  %idxprom = sext i32 %5 to i64
+  %arrayidx = getelementptr inbounds [256 x %struct.my_struct]* @struct_cl.array, i32 0, i64 %idxprom
+  %a = getelementptr inbounds %struct.my_struct* %arrayidx, i32 0, i32 0
+  store i32 %4, i32* %a, align 4
+  %6 = load i32* %i, align 4
+  %7 = load i32* %i, align 4
+  %idxprom1 = sext i32 %7 to i64
+  %arrayidx2 = getelementptr inbounds [256 x %struct.my_struct]* @struct_cl.array, i32 0, i64 %idxprom1
+  %b = getelementptr inbounds %struct.my_struct* %arrayidx2, i32 0, i32 1
+  %arrayidx3 = getelementptr inbounds [2 x i32]* %b, i32 0, i64 0
+  store i32 %6, i32* %arrayidx3, align 4
+  %8 = load i32* %i, align 4
+  %add = add nsw i32 %8, 1
+  %9 = load i32* %i, align 4
+  %idxprom4 = sext i32 %9 to i64
+  %arrayidx5 = getelementptr inbounds [256 x %struct.my_struct]* @struct_cl.array, i32 0, i64 %idxprom4
+  %b6 = getelementptr inbounds %struct.my_struct* %arrayidx5, i32 0, i32 1
+  %arrayidx7 = getelementptr inbounds [2 x i32]* %b6, i32 0, i64 0
+  store i32 %add, i32* %arrayidx7, align 4
+  br label %for.inc
+
+for.inc:                                          ; preds = %for.body
+  %10 = load i32* %i, align 4
+  %inc = add nsw i32 %10, 1
+  store i32 %inc, i32* %i, align 4
+  br label %for.cond
+
+for.end:                                          ; preds = %for.cond
+  %11 = load i32* %y.addr, align 4
+  %cmp8 = icmp eq i32 %11, 0
+  br i1 %cmp8, label %if.then, label %if.else
+
+if.then:                                          ; preds = %for.end
+  %12 = load i32* %y.addr, align 4
+  %idxprom9 = sext i32 %12 to i64
+  %arrayidx10 = getelementptr inbounds [256 x %struct.my_struct]* @struct_cl.array, i32 0, i64 %idxprom9
+  %13 = bitcast %struct.my_struct* %arrayidx10 to i8*
+  call void @llvm.memcpy.p0i8.p0i8.i64(i8* bitcast (%struct.my_struct* @struct_cl.hop to i8*), i8* %13, i64 12, i32 4, i1 false)
+  br label %if.end
+
+if.else:                                          ; preds = %for.end
+  %14 = load i32* %y.addr, align 4
+  %add11 = add nsw i32 %14, 1
+  %idxprom12 = sext i32 %add11 to i64
+  %arrayidx13 = getelementptr inbounds [256 x %struct.my_struct]* @struct_cl.array, i32 0, i64 %idxprom12
+  %15 = bitcast %struct.my_struct* %arrayidx13 to i8*
+  call void @llvm.memcpy.p0i8.p0i8.i64(i8* bitcast (%struct.my_struct* @struct_cl.hop to i8*), i8* %15, i64 12, i32 4, i1 false)
+  br label %if.end
+
+if.end:                                           ; preds = %if.else, %if.then
+  call void @llvm.memcpy.p0i8.p0i8.i64(i8* bitcast ([256 x %struct.my_struct]* @struct_cl.array to i8*), i8* bitcast (%struct.my_struct* @struct_cl.hop to i8*), i64 12, i32 4, i1 false)
+  %a14 = getelementptr inbounds %struct.my_struct* %s, i32 0, i32 0
+  %16 = load i32* %a14, align 4
+  %17 = load i32* %x.addr, align 4
+  %idxprom15 = sext i32 %17 to i64
+  %arrayidx16 = getelementptr inbounds [256 x %struct.my_struct]* @struct_cl.array, i32 0, i64 %idxprom15
+  %a17 = getelementptr inbounds %struct.my_struct* %arrayidx16, i32 0, i32 0
+  %18 = load i32* %a17, align 4
+  %add18 = add nsw i32 %16, %18
+  %19 = load i32* %x.addr, align 4
+  %add19 = add nsw i32 %19, 1
+  %idxprom20 = sext i32 %add19 to i64
+  %arrayidx21 = getelementptr inbounds [256 x %struct.my_struct]* @struct_cl.array, i32 0, i64 %idxprom20
+  %b22 = getelementptr inbounds %struct.my_struct* %arrayidx21, i32 0, i32 1
+  %arrayidx23 = getelementptr inbounds [2 x i32]* %b22, i32 0, i64 0
+  %20 = load i32* %arrayidx23, align 4
+  %add24 = add nsw i32 %add18, %20
+  %21 = load i32* %x.addr, align 4
+  %idxprom25 = sext i32 %21 to i64
+  %arrayidx26 = getelementptr inbounds [4 x i32]* @g, i32 0, i64 %idxprom25
+  %22 = load i32* %arrayidx26, align 4
+  %add27 = add nsw i32 %add24, %22
+  %23 = load i32* getelementptr inbounds ([4 x i32]* @g, i32 0, i64 3), align 4
+  %add28 = add nsw i32 %add27, %23
+  %24 = load i32** %mem.addr, align 8
+  %arrayidx29 = getelementptr inbounds i32* %24, i64 0
+  store i32 %add28, i32* %arrayidx29
   ret void
 }
 
-declare void @llvm.memcpy.p4i8.p4i8.i32(i8 addrspace(4)* nocapture, i8 addrspace(4)* nocapture, i32, i32, i1) nounwind
+declare void @llvm.memcpy.p0i8.p0i8.i64(i8* nocapture, i8* nocapture, i64, i32, i1) nounwind
 
 !opencl.kernels = !{!0}
 
-!0 = metadata !{void (%struct.my_struct*, i32, i32*, i32)* @struct_cl}
-!1 = metadata !{metadata !"int", metadata !2}
-!2 = metadata !{metadata !"omnipotent char", metadata !3}
-!3 = metadata !{metadata !"Simple C/C++ TBAA", null}
+!0 = metadata !{void (i64, i32, i32, i32*, i32)* @struct_cl}

diff --git a/backend/kernels/struct2.cl b/backend/kernels/struct2.cl
new file mode 100644 (file)
index 0000000..5c3aa41
--- /dev/null
+++ b/backend/kernels/struct2.cl
@@ -0,0 +1,20 @@
+struct my_struct {
+  int a;
+  int b[2];
+};
+
+const __constant int g[4] = {0,1,2,3};
+
+__kernel void struct_cl (struct my_struct s, int x, __global struct my_struct *mem, int y)
+{
+  struct my_struct hop;
+  if (y == 0) {
+    hop.a = 1;
+    hop.b[0] = 2;
+    hop.b[1] = 2;
+  } else {
+    hop = s;
+  }
+  mem[0] = hop;
+}
+

diff --git a/backend/kernels/struct2.ll b/backend/kernels/struct2.ll
new file mode 100644 (file)
index 0000000..ef1f3ab
--- /dev/null
+++ b/backend/kernels/struct2.ll
@@ -0,0 +1,38 @@
+; ModuleID = 'struct2.o'
+target datalayout = "e-p:32:32-i64:64:64-f64:64:64-n1:8:16:32:64"
+target triple = "ptx32--"
+
+%struct.my_struct = type { i32, [2 x i32] }
+
+@g = addrspace(1) constant [4 x i32] [i32 0, i32 1, i32 2, i32 3], align 4
+
+define ptx_kernel void @struct_cl(%struct.my_struct* nocapture byval %s, i32 %x, %struct.my_struct* nocapture %mem, i32 %y) nounwind noinline {
+entry:
+  %cmp = icmp eq i32 %y, 0
+  br i1 %cmp, label %if.end, label %if.else
+
+if.else:                                          ; preds = %entry
+  %s.0 = getelementptr inbounds %struct.my_struct* %s, i32 0, i32 0
+  %tmp4 = load i32* %s.0, align 4
+  %s.1.0 = getelementptr inbounds %struct.my_struct* %s, i32 0, i32 1, i32 0
+  %tmp5 = load i32* %s.1.0, align 4
+  %s.1.1 = getelementptr inbounds %struct.my_struct* %s, i32 0, i32 1, i32 1
+  %tmp6 = load i32* %s.1.1, align 4
+  br label %if.end
+
+if.end:                                           ; preds = %if.else, %entry
+  %hop.1.1.0 = phi i32 [ %tmp6, %if.else ], [ 2, %entry ]
+  %hop.1.0.0 = phi i32 [ %tmp5, %if.else ], [ 2, %entry ]
+  %hop.0.0 = phi i32 [ %tmp4, %if.else ], [ 1, %entry ]
+  %mem.0 = getelementptr inbounds %struct.my_struct* %mem, i32 0, i32 0
+  store i32 %hop.0.0, i32* %mem.0, align 4
+  %mem.1.0 = getelementptr inbounds %struct.my_struct* %mem, i32 0, i32 1, i32 0
+  store i32 %hop.1.0.0, i32* %mem.1.0, align 4
+  %mem.1.1 = getelementptr inbounds %struct.my_struct* %mem, i32 0, i32 1, i32 1
+  store i32 %hop.1.1.0, i32* %mem.1.1, align 4
+  ret void
+}
+
+!opencl.kernels = !{!0}
+
+!0 = metadata !{void (%struct.my_struct*, i32, %struct.my_struct*, i32)* @struct_cl}