You can not select more than 25 topics
Topics must start with a letter or number, can include dashes ('-') and can be up to 35 characters long.
131 lines
4.7 KiB
131 lines
4.7 KiB
//
|
|
// 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_
|