Implemented a first implementation of the function argument parsing from LLVM IR

10 files changed, 75 insertions, 0 deletions

diff --git a/backend/kernels/add.cl b/backend/kernels/add.cl
new file mode 100644
index 00000000..7786e875
--- /dev/null
+++ b/backend/kernels/add.cl
@@ -0,0 +1,5 @@
+__kernel unsigned int add(unsigned int x, unsigned int y)
+{
+  return x + y;
+}
+
diff --git a/backend/kernels/add.ll b/backend/kernels/add.ll
new file mode 100644
index 00000000..9b2c7413
--- /dev/null
+++ b/backend/kernels/add.ll
@@ -0,0 +1,13 @@
+; ModuleID = 'add.o'
+target datalayout = "e-p:32:32-i64:64:64-f64:64:64-n1:8:16:32:64"
+target triple = "ptx32--"
+
+define ptx_kernel i32 @add(i32 %x, i32 %y) nounwind readnone noinline {
+entry:
+  %add = add i32 %y, %x
+  ret i32 %add
+}
+
+!opencl.kernels = !{!0}
+
+!0 = metadata !{i32 (i32, i32)* @add}
diff --git a/backend/kernels/add.o b/backend/kernels/add.o
new file mode 100644
index 00000000..fcff924d
--- /dev/null
+++ b/backend/kernels/add.o
diff --git a/backend/kernels/add2.cl b/backend/kernels/add2.cl
new file mode 100644
index 00000000..c43e2c34
--- /dev/null
+++ b/backend/kernels/add2.cl
@@ -0,0 +1,12 @@
+struct big{
+  unsigned int a, b;
+};
+
+__kernel struct big add(unsigned int x, unsigned int y)
+{
+  struct big p;
+  p.a = x + y;
+  p.b = x - y;
+  return p;
+}
+
diff --git a/backend/kernels/add2.ll b/backend/kernels/add2.ll
new file mode 100644
index 00000000..37cf7a3e
--- /dev/null
+++ b/backend/kernels/add2.ll
@@ -0,0 +1,20 @@
+; ModuleID = 'add2.o'
+target datalayout = "e-p:32:32-i64:64:64-f64:64:64-n1:8:16:32:64"
+target triple = "ptx32--"
+
+%struct.big = type { i32, i32 }
+
+define ptx_kernel void @add(%struct.big* noalias nocapture sret %agg.result, i32 %x, i32 %y) nounwind noinline {
+entry:
+  %add = add i32 %y, %x
+  %sub = sub i32 %x, %y
+  %agg.result.0 = getelementptr inbounds %struct.big* %agg.result, i32 0, i32 0
+  store i32 %add, i32* %agg.result.0, align 4
+  %agg.result.1 = getelementptr inbounds %struct.big* %agg.result, i32 0, i32 1
+  store i32 %sub, i32* %agg.result.1, align 4
+  ret void
+}
+
+!opencl.kernels = !{!0}
+
+!0 = metadata !{void (%struct.big*, i32, i32)* @add}
diff --git a/backend/kernels/add2.o b/backend/kernels/add2.o
new file mode 100644
index 00000000..1feb0355
--- /dev/null
+++ b/backend/kernels/add2.o
diff --git a/backend/kernels/compile.sh b/backend/kernels/compile.sh
new file mode 100644
index 00000000..880da2c6
--- /dev/null
+++ b/backend/kernels/compile.sh
@@ -0,0 +1,4 @@
+clang -emit-llvm -O3 -ccc-host-triple ptx32 -c $1.cl -o $1.o
+llvm-dis $1.o
+mv $1.o.ll $1.ll
+
diff --git a/backend/kernels/store.cl b/backend/kernels/store.cl
new file mode 100644
index 00000000..b23b13cb
--- /dev/null
+++ b/backend/kernels/store.cl
@@ -0,0 +1,5 @@
+__kernel void store(__global int *dst, __local int *dst0, int x)
+{
+  dst[0] = 1;
+}
+
diff --git a/backend/kernels/store.ll b/backend/kernels/store.ll
new file mode 100644
index 00000000..c885d62e
--- /dev/null
+++ b/backend/kernels/store.ll
@@ -0,0 +1,16 @@
+; ModuleID = 'store.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 @store(i32* nocapture %dst, i32 addrspace(4)* nocapture %dst0, i32 %x) nounwind noinline {
+entry:
+  store i32 1, i32* %dst, align 4, !tbaa !1
+  ret void
+}
+
+!opencl.kernels = !{!0}
+
+!0 = metadata !{void (i32*, i32 addrspace(4)*, i32)* @store}
+!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/store.o b/backend/kernels/store.o
new file mode 100644
index 00000000..ea7b34af
--- /dev/null
+++ b/backend/kernels/store.o