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