1 files changed, 117 insertions, 0 deletions
diff --git a/kernels/compiler_pipe_builtin.cl b/kernels/compiler_pipe_builtin.cl
new file mode 100644
index 00000000..4e8dcc42
--- /dev/null
+++ b/kernels/compiler_pipe_builtin.cl
@@ -0,0 +1,117 @@
+typedef struct{
+  int a;
+  int b;
+}mystruct;
+
+__kernel void compiler_pipe_convenience_write_int(write_only pipe int p, __global int *src)
+{
+    int gid = get_global_id(0);
+    write_pipe(p, &src[gid]);
+}
+__kernel void compiler_pipe_convenience_read_int(read_only pipe int p, __global int *dst)
+{
+    int gid = get_global_id(0);
+    read_pipe(p, &dst[gid]);
+}
+__kernel void compiler_pipe_convenience_write_mystruct(write_only pipe mystruct p, __global mystruct *src)
+{
+    int gid = get_global_id(0);
+    write_pipe(p, &src[gid]);
+}
+__kernel void compiler_pipe_convenience_read_mystruct(read_only pipe mystruct p, __global mystruct *dst)
+{
+    int gid = get_global_id(0);
+    read_pipe(p, &dst[gid]);
+}
+
+__kernel void compiler_pipe_reserve_write_int(write_only pipe int p, __global int *src)
+{
+    int gid = get_global_id(0);
+    reserve_id_t res_id = reserve_write_pipe(p, 1);
+    if(is_valid_reserve_id(res_id))
+    {
+      write_pipe(p, res_id, 0, &src[gid]);
+      commit_write_pipe(p, res_id);
+    }
+}
+__kernel void compiler_pipe_reserve_read_int(read_only pipe int p, __global int *dst)
+{
+    int gid = get_global_id(0);
+    reserve_id_t res_id = reserve_read_pipe(p, 1);
+    if(is_valid_reserve_id(res_id))
+    {
+      read_pipe(p, res_id, 0, &dst[gid]);
+      commit_read_pipe(p, res_id);
+    }
+}
+__kernel void compiler_pipe_reserve_write_mystruct(write_only pipe mystruct p, __global mystruct *src)
+{
+    int gid = get_global_id(0);
+    reserve_id_t res_id = reserve_write_pipe(p, 1);
+    if(is_valid_reserve_id(res_id))
+    {
+      write_pipe(p, res_id, 0, &src[gid]);
+      commit_write_pipe(p, res_id);
+    }
+}
+__kernel void compiler_pipe_reserve_read_mystruct(read_only pipe mystruct p, __global mystruct *dst)
+{
+    int gid = get_global_id(0);
+    reserve_id_t res_id = reserve_read_pipe(p, 1);
+    if(is_valid_reserve_id(res_id))
+    {
+      read_pipe(p, res_id, 0, &dst[gid]);
+      commit_read_pipe(p, res_id);
+    }
+}
+
+__kernel void compiler_pipe_workgroup_write_int(write_only pipe int p, __global int *src)
+{
+    int gid = get_global_id(0);
+    reserve_id_t res_id = work_group_reserve_write_pipe(p, get_local_size(0));
+    if(is_valid_reserve_id(res_id))
+    {
+      write_pipe(p, res_id, get_local_id(0), &src[gid]);
+      work_group_commit_write_pipe(p, res_id);
+    }
+}
+__kernel void compiler_pipe_workgroup_read_int(read_only pipe int p, __global int *dst)
+{
+    int gid = get_global_id(0);
+    reserve_id_t res_id = work_group_reserve_read_pipe(p, get_local_size(0));
+    if(is_valid_reserve_id(res_id))
+    {
+      read_pipe(p, res_id, get_local_id(0), &dst[gid]);
+      work_group_commit_read_pipe(p, res_id);
+    }
+}
+__kernel void compiler_pipe_workgroup_write_mystruct(write_only pipe mystruct p, __global mystruct *src)
+{
+    int gid = get_global_id(0);
+    reserve_id_t res_id = work_group_reserve_write_pipe(p, get_local_size(0));
+    if(is_valid_reserve_id(res_id))
+    {
+      write_pipe(p, res_id, get_local_id(0), &src[gid]);
+      work_group_commit_write_pipe(p, res_id);
+    }
+}
+__kernel void compiler_pipe_workgroup_read_mystruct(read_only pipe mystruct p, __global mystruct *dst)
+{
+    int gid = get_global_id(0);
+    reserve_id_t res_id = work_group_reserve_read_pipe(p, get_local_size(0));
+    if(is_valid_reserve_id(res_id))
+    {
+      read_pipe(p, res_id, get_local_id(0), &dst[gid]);
+      work_group_commit_read_pipe(p, res_id);
+    }
+}
+
+__kernel void compiler_pipe_query(write_only pipe int p, __global uint *src)
+{
+    int gid = get_global_id(0);
+    write_pipe(p,&gid);
+    if(gid == 0) {
+      src[0] = get_pipe_num_packets(p);
+      src[1] = get_pipe_max_packets(p);
+    }
+}