1 //
2 // Copyright (c) 2021 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 #include "procs.h"
17 #include "subhelpers.h"
18 #include "subgroup_common_templates.h"
19 #include "harness/conversions.h"
20 #include "harness/typeWrappers.h"
21
22 namespace {
23
24 static const char* shuffle_down_source =
25 "__kernel void test_sub_group_shuffle_down(const __global Type *in, "
26 "__global int4 *xy, __global Type *out)\n"
27 "{\n"
28 " int gid = get_global_id(0);\n"
29 " XY(xy,gid);\n"
30 " Type x = in[gid];\n"
31 " out[gid] = sub_group_shuffle_down(x, xy[gid].z);"
32 "}\n";
33 static const char* shuffle_up_source =
34 "__kernel void test_sub_group_shuffle_up(const __global Type *in, __global "
35 "int4 *xy, __global Type *out)\n"
36 "{\n"
37 " int gid = get_global_id(0);\n"
38 " XY(xy,gid);\n"
39 " Type x = in[gid];\n"
40 " out[gid] = sub_group_shuffle_up(x, xy[gid].z);"
41 "}\n";
42
run_shuffle_relative_for_type(RunTestForType rft)43 template <typename T> int run_shuffle_relative_for_type(RunTestForType rft)
44 {
45 int error = rft.run_impl<T, SHF<T, ShuffleOp::shuffle_up>>(
46 "test_sub_group_shuffle_up", shuffle_up_source);
47 error |= rft.run_impl<T, SHF<T, ShuffleOp::shuffle_down>>(
48 "test_sub_group_shuffle_down", shuffle_down_source);
49 return error;
50 }
51
52 }
53
test_subgroup_functions_shuffle_relative(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)54 int test_subgroup_functions_shuffle_relative(cl_device_id device,
55 cl_context context,
56 cl_command_queue queue,
57 int num_elements)
58 {
59 std::vector<std::string> required_extensions = {
60 "cl_khr_subgroup_shuffle_relative"
61 };
62 constexpr size_t global_work_size = 2000;
63 constexpr size_t local_work_size = 200;
64 WorkGroupParams test_params(global_work_size, local_work_size,
65 required_extensions);
66 RunTestForType rft(device, context, queue, num_elements, test_params);
67
68 int error = run_shuffle_relative_for_type<cl_int>(rft);
69 error |= run_shuffle_relative_for_type<cl_uint>(rft);
70 error |= run_shuffle_relative_for_type<cl_long>(rft);
71 error |= run_shuffle_relative_for_type<cl_ulong>(rft);
72 error |= run_shuffle_relative_for_type<cl_short>(rft);
73 error |= run_shuffle_relative_for_type<cl_ushort>(rft);
74 error |= run_shuffle_relative_for_type<cl_char>(rft);
75 error |= run_shuffle_relative_for_type<cl_uchar>(rft);
76 error |= run_shuffle_relative_for_type<cl_float>(rft);
77 error |= run_shuffle_relative_for_type<cl_double>(rft);
78 error |= run_shuffle_relative_for_type<subgroups::cl_half>(rft);
79
80 return error;
81 }
82