1*6467f958SSadaf Ebrahimi // 2*6467f958SSadaf Ebrahimi // Copyright (c) 2017 The Khronos Group Inc. 3*6467f958SSadaf Ebrahimi // 4*6467f958SSadaf Ebrahimi // Licensed under the Apache License, Version 2.0 (the "License"); 5*6467f958SSadaf Ebrahimi // you may not use this file except in compliance with the License. 6*6467f958SSadaf Ebrahimi // You may obtain a copy of the License at 7*6467f958SSadaf Ebrahimi // 8*6467f958SSadaf Ebrahimi // http://www.apache.org/licenses/LICENSE-2.0 9*6467f958SSadaf Ebrahimi // 10*6467f958SSadaf Ebrahimi // Unless required by applicable law or agreed to in writing, software 11*6467f958SSadaf Ebrahimi // distributed under the License is distributed on an "AS IS" BASIS, 12*6467f958SSadaf Ebrahimi // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. 13*6467f958SSadaf Ebrahimi // See the License for the specific language governing permissions and 14*6467f958SSadaf Ebrahimi // limitations under the License. 15*6467f958SSadaf Ebrahimi // 16*6467f958SSadaf Ebrahimi #ifndef KERNELS_H_ 17*6467f958SSadaf Ebrahimi #define KERNELS_H_ 18*6467f958SSadaf Ebrahimi 19*6467f958SSadaf Ebrahimi static const char* pipe_readwrite_struct_kernel_code = { 20*6467f958SSadaf Ebrahimi "typedef struct{\n" 21*6467f958SSadaf Ebrahimi "char a;\n" 22*6467f958SSadaf Ebrahimi "int b;\n" 23*6467f958SSadaf Ebrahimi "}TestStruct;\n" 24*6467f958SSadaf Ebrahimi "__kernel void test_pipe_write_struct(__global TestStruct *src, __write_only pipe TestStruct out_pipe)\n" 25*6467f958SSadaf Ebrahimi "{\n" 26*6467f958SSadaf Ebrahimi " int gid = get_global_id(0);\n" 27*6467f958SSadaf Ebrahimi " reserve_id_t res_id; \n" 28*6467f958SSadaf Ebrahimi "\n" 29*6467f958SSadaf Ebrahimi " res_id = reserve_write_pipe(out_pipe, 1);\n" 30*6467f958SSadaf Ebrahimi " if(is_valid_reserve_id(res_id))\n" 31*6467f958SSadaf Ebrahimi " {\n" 32*6467f958SSadaf Ebrahimi " write_pipe(out_pipe, res_id, 0, &src[gid]);\n" 33*6467f958SSadaf Ebrahimi " commit_write_pipe(out_pipe, res_id);\n" 34*6467f958SSadaf Ebrahimi " }\n" 35*6467f958SSadaf Ebrahimi "}\n" 36*6467f958SSadaf Ebrahimi "\n" 37*6467f958SSadaf Ebrahimi "__kernel void test_pipe_read_struct(__read_only pipe TestStruct in_pipe, __global TestStruct *dst)\n" 38*6467f958SSadaf Ebrahimi "{\n" 39*6467f958SSadaf Ebrahimi " int gid = get_global_id(0);\n" 40*6467f958SSadaf Ebrahimi " reserve_id_t res_id; \n" 41*6467f958SSadaf Ebrahimi "\n" 42*6467f958SSadaf Ebrahimi " res_id = reserve_read_pipe(in_pipe, 1);\n" 43*6467f958SSadaf Ebrahimi " if(is_valid_reserve_id(res_id))\n" 44*6467f958SSadaf Ebrahimi " {\n" 45*6467f958SSadaf Ebrahimi " read_pipe(in_pipe, res_id, 0, &dst[gid]);\n" 46*6467f958SSadaf Ebrahimi " commit_read_pipe(in_pipe, res_id);\n" 47*6467f958SSadaf Ebrahimi " }\n" 48*6467f958SSadaf Ebrahimi "}\n" }; 49*6467f958SSadaf Ebrahimi 50*6467f958SSadaf Ebrahimi static const char* pipe_workgroup_readwrite_struct_kernel_code = { 51*6467f958SSadaf Ebrahimi "typedef struct{\n" 52*6467f958SSadaf Ebrahimi "char a;\n" 53*6467f958SSadaf Ebrahimi "int b;\n" 54*6467f958SSadaf Ebrahimi "}TestStruct;\n" 55*6467f958SSadaf Ebrahimi "__kernel void test_pipe_workgroup_write_struct(__global TestStruct *src, __write_only pipe TestStruct out_pipe)\n" 56*6467f958SSadaf Ebrahimi "{\n" 57*6467f958SSadaf Ebrahimi " int gid = get_global_id(0);\n" 58*6467f958SSadaf Ebrahimi " __local reserve_id_t res_id; \n" 59*6467f958SSadaf Ebrahimi "\n" 60*6467f958SSadaf Ebrahimi " res_id = work_group_reserve_write_pipe(out_pipe, get_local_size(0));\n" 61*6467f958SSadaf Ebrahimi " if(is_valid_reserve_id(res_id))\n" 62*6467f958SSadaf Ebrahimi " {\n" 63*6467f958SSadaf Ebrahimi " write_pipe(out_pipe, res_id, get_local_id(0), &src[gid]);\n" 64*6467f958SSadaf Ebrahimi " work_group_commit_write_pipe(out_pipe, res_id);\n" 65*6467f958SSadaf Ebrahimi " }\n" 66*6467f958SSadaf Ebrahimi "}\n" 67*6467f958SSadaf Ebrahimi "\n" 68*6467f958SSadaf Ebrahimi "__kernel void test_pipe_workgroup_read_struct(__read_only pipe TestStruct in_pipe, __global TestStruct *dst)\n" 69*6467f958SSadaf Ebrahimi "{\n" 70*6467f958SSadaf Ebrahimi " int gid = get_global_id(0);\n" 71*6467f958SSadaf Ebrahimi " __local reserve_id_t res_id; \n" 72*6467f958SSadaf Ebrahimi "\n" 73*6467f958SSadaf Ebrahimi " res_id = work_group_reserve_read_pipe(in_pipe, get_local_size(0));\n" 74*6467f958SSadaf Ebrahimi " if(is_valid_reserve_id(res_id))\n" 75*6467f958SSadaf Ebrahimi " {\n" 76*6467f958SSadaf Ebrahimi " read_pipe(in_pipe, res_id, get_local_id(0), &dst[gid]);\n" 77*6467f958SSadaf Ebrahimi " work_group_commit_read_pipe(in_pipe, res_id);\n" 78*6467f958SSadaf Ebrahimi " }\n" 79*6467f958SSadaf Ebrahimi "}\n" }; 80*6467f958SSadaf Ebrahimi 81*6467f958SSadaf Ebrahimi static const char* pipe_subgroup_readwrite_struct_kernel_code = { 82*6467f958SSadaf Ebrahimi "typedef struct{\n" 83*6467f958SSadaf Ebrahimi "char a;\n" 84*6467f958SSadaf Ebrahimi "int b;\n" 85*6467f958SSadaf Ebrahimi "}TestStruct;\n" 86*6467f958SSadaf Ebrahimi "#pragma OPENCL EXTENSION cl_khr_subgroups : enable\n" 87*6467f958SSadaf Ebrahimi "__kernel void test_pipe_subgroup_write_struct(__global TestStruct *src, __write_only pipe TestStruct out_pipe)\n" 88*6467f958SSadaf Ebrahimi "{\n" 89*6467f958SSadaf Ebrahimi " int gid = get_global_id(0);\n" 90*6467f958SSadaf Ebrahimi " reserve_id_t res_id; \n" 91*6467f958SSadaf Ebrahimi "\n" 92*6467f958SSadaf Ebrahimi " res_id = sub_group_reserve_write_pipe(out_pipe, get_sub_group_size());\n" 93*6467f958SSadaf Ebrahimi " if(is_valid_reserve_id(res_id))\n" 94*6467f958SSadaf Ebrahimi " {\n" 95*6467f958SSadaf Ebrahimi " write_pipe(out_pipe, res_id, get_sub_group_local_id(), &src[gid]);\n" 96*6467f958SSadaf Ebrahimi " sub_group_commit_write_pipe(out_pipe, res_id);\n" 97*6467f958SSadaf Ebrahimi " }\n" 98*6467f958SSadaf Ebrahimi "}\n" 99*6467f958SSadaf Ebrahimi "\n" 100*6467f958SSadaf Ebrahimi "__kernel void test_pipe_subgroup_read_struct(__read_only pipe TestStruct in_pipe, __global TestStruct *dst)\n" 101*6467f958SSadaf Ebrahimi "{\n" 102*6467f958SSadaf Ebrahimi " int gid = get_global_id(0);\n" 103*6467f958SSadaf Ebrahimi " reserve_id_t res_id; \n" 104*6467f958SSadaf Ebrahimi "\n" 105*6467f958SSadaf Ebrahimi " res_id = sub_group_reserve_read_pipe(in_pipe, get_sub_group_size());\n" 106*6467f958SSadaf Ebrahimi " if(is_valid_reserve_id(res_id))\n" 107*6467f958SSadaf Ebrahimi " {\n" 108*6467f958SSadaf Ebrahimi " read_pipe(in_pipe, res_id, get_sub_group_local_id(), &dst[gid]);\n" 109*6467f958SSadaf Ebrahimi " sub_group_commit_read_pipe(in_pipe, res_id);\n" 110*6467f958SSadaf Ebrahimi " }\n" 111*6467f958SSadaf Ebrahimi "}\n" }; 112*6467f958SSadaf Ebrahimi 113*6467f958SSadaf Ebrahimi static const char* pipe_convenience_readwrite_struct_kernel_code = { 114*6467f958SSadaf Ebrahimi "typedef struct{\n" 115*6467f958SSadaf Ebrahimi "char a;\n" 116*6467f958SSadaf Ebrahimi "int b;\n" 117*6467f958SSadaf Ebrahimi "}TestStruct;\n" 118*6467f958SSadaf Ebrahimi "__kernel void test_pipe_convenience_write_struct(__global TestStruct *src, __write_only pipe TestStruct out_pipe)\n" 119*6467f958SSadaf Ebrahimi "{\n" 120*6467f958SSadaf Ebrahimi " int gid = get_global_id(0);\n" 121*6467f958SSadaf Ebrahimi " write_pipe(out_pipe, &src[gid]);\n" 122*6467f958SSadaf Ebrahimi "}\n" 123*6467f958SSadaf Ebrahimi "\n" 124*6467f958SSadaf Ebrahimi "__kernel void test_pipe_convenience_read_struct(__read_only pipe TestStruct in_pipe, __global TestStruct *dst)\n" 125*6467f958SSadaf Ebrahimi "{\n" 126*6467f958SSadaf Ebrahimi " int gid = get_global_id(0);\n" 127*6467f958SSadaf Ebrahimi " read_pipe(in_pipe, &dst[gid]);\n" 128*6467f958SSadaf Ebrahimi "}\n" }; 129*6467f958SSadaf Ebrahimi 130*6467f958SSadaf Ebrahimi #endif // KERNELS_H_ 131