| // |
| // Copyright (c) 2017 The Khronos Group Inc. |
| // |
| // Licensed under the Apache License, Version 2.0 (the "License"); |
| // you may not use this file except in compliance with the License. |
| // You may obtain a copy of the License at |
| // |
| // http://www.apache.org/licenses/LICENSE-2.0 |
| // |
| // Unless required by applicable law or agreed to in writing, software |
| // distributed under the License is distributed on an "AS IS" BASIS, |
| // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. |
| // See the License for the specific language governing permissions and |
| // limitations under the License. |
| // |
| #ifndef _KERNELS_H_ |
| #define _KERNELS_H_ |
| |
| static const char* pipe_readwrite_struct_kernel_code = { |
| "typedef struct{\n" |
| "char a;\n" |
| "int b;\n" |
| "}TestStruct;\n" |
| "__kernel void test_pipe_write_struct(__global TestStruct *src, __write_only pipe TestStruct out_pipe)\n" |
| "{\n" |
| " int gid = get_global_id(0);\n" |
| " reserve_id_t res_id; \n" |
| "\n" |
| " res_id = reserve_write_pipe(out_pipe, 1);\n" |
| " if(is_valid_reserve_id(res_id))\n" |
| " {\n" |
| " write_pipe(out_pipe, res_id, 0, &src[gid]);\n" |
| " commit_write_pipe(out_pipe, res_id);\n" |
| " }\n" |
| "}\n" |
| "\n" |
| "__kernel void test_pipe_read_struct(__read_only pipe TestStruct in_pipe, __global TestStruct *dst)\n" |
| "{\n" |
| " int gid = get_global_id(0);\n" |
| " reserve_id_t res_id; \n" |
| "\n" |
| " res_id = reserve_read_pipe(in_pipe, 1);\n" |
| " if(is_valid_reserve_id(res_id))\n" |
| " {\n" |
| " read_pipe(in_pipe, res_id, 0, &dst[gid]);\n" |
| " commit_read_pipe(in_pipe, res_id);\n" |
| " }\n" |
| "}\n" }; |
| |
| static const char* pipe_workgroup_readwrite_struct_kernel_code = { |
| "typedef struct{\n" |
| "char a;\n" |
| "int b;\n" |
| "}TestStruct;\n" |
| "__kernel void test_pipe_workgroup_write_struct(__global TestStruct *src, __write_only pipe TestStruct out_pipe)\n" |
| "{\n" |
| " int gid = get_global_id(0);\n" |
| " __local reserve_id_t res_id; \n" |
| "\n" |
| " res_id = work_group_reserve_write_pipe(out_pipe, get_local_size(0));\n" |
| " if(is_valid_reserve_id(res_id))\n" |
| " {\n" |
| " write_pipe(out_pipe, res_id, get_local_id(0), &src[gid]);\n" |
| " work_group_commit_write_pipe(out_pipe, res_id);\n" |
| " }\n" |
| "}\n" |
| "\n" |
| "__kernel void test_pipe_workgroup_read_struct(__read_only pipe TestStruct in_pipe, __global TestStruct *dst)\n" |
| "{\n" |
| " int gid = get_global_id(0);\n" |
| " __local reserve_id_t res_id; \n" |
| "\n" |
| " res_id = work_group_reserve_read_pipe(in_pipe, get_local_size(0));\n" |
| " if(is_valid_reserve_id(res_id))\n" |
| " {\n" |
| " read_pipe(in_pipe, res_id, get_local_id(0), &dst[gid]);\n" |
| " work_group_commit_read_pipe(in_pipe, res_id);\n" |
| " }\n" |
| "}\n" }; |
| |
| static const char* pipe_subgroup_readwrite_struct_kernel_code = { |
| "typedef struct{\n" |
| "char a;\n" |
| "int b;\n" |
| "}TestStruct;\n" |
| "#pragma OPENCL EXTENSION cl_khr_subgroups : enable\n" |
| "__kernel void test_pipe_subgroup_write_struct(__global TestStruct *src, __write_only pipe TestStruct out_pipe)\n" |
| "{\n" |
| " int gid = get_global_id(0);\n" |
| " reserve_id_t res_id; \n" |
| "\n" |
| " res_id = sub_group_reserve_write_pipe(out_pipe, get_sub_group_size());\n" |
| " if(is_valid_reserve_id(res_id))\n" |
| " {\n" |
| " write_pipe(out_pipe, res_id, get_sub_group_local_id(), &src[gid]);\n" |
| " sub_group_commit_write_pipe(out_pipe, res_id);\n" |
| " }\n" |
| "}\n" |
| "\n" |
| "__kernel void test_pipe_subgroup_read_struct(__read_only pipe TestStruct in_pipe, __global TestStruct *dst)\n" |
| "{\n" |
| " int gid = get_global_id(0);\n" |
| " reserve_id_t res_id; \n" |
| "\n" |
| " res_id = sub_group_reserve_read_pipe(in_pipe, get_sub_group_size());\n" |
| " if(is_valid_reserve_id(res_id))\n" |
| " {\n" |
| " read_pipe(in_pipe, res_id, get_sub_group_local_id(), &dst[gid]);\n" |
| " sub_group_commit_read_pipe(in_pipe, res_id);\n" |
| " }\n" |
| "}\n" }; |
| |
| static const char* pipe_convenience_readwrite_struct_kernel_code = { |
| "typedef struct{\n" |
| "char a;\n" |
| "int b;\n" |
| "}TestStruct;\n" |
| "__kernel void test_pipe_convenience_write_struct(__global TestStruct *src, __write_only pipe TestStruct out_pipe)\n" |
| "{\n" |
| " int gid = get_global_id(0);\n" |
| " write_pipe(out_pipe, &src[gid]);\n" |
| "}\n" |
| "\n" |
| "__kernel void test_pipe_convenience_read_struct(__read_only pipe TestStruct in_pipe, __global TestStruct *dst)\n" |
| "{\n" |
| " int gid = get_global_id(0);\n" |
| " read_pipe(in_pipe, &dst[gid]);\n" |
| "}\n" }; |
| |
| #endif //_KERNELS_H_ |