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 "harness/typeWrappers.h"
19 #include "subgroup_common_templates.h"
20
21 namespace {
22
23 std::string sub_group_non_uniform_arithmetic_source = R"(
24 __kernel void test_%s(const __global Type *in, __global int4 *xy, __global Type *out, uint4 work_item_mask_vector) {
25 int gid = get_global_id(0);
26 XY(xy,gid);
27 uint subgroup_local_id = get_sub_group_local_id();
28 uint elect_work_item = 1 << (subgroup_local_id % 32);
29 uint work_item_mask;
30 if(subgroup_local_id < 32) {
31 work_item_mask = work_item_mask_vector.x;
32 } else if(subgroup_local_id < 64) {
33 work_item_mask = work_item_mask_vector.y;
34 } else if(subgroup_local_id < 96) {
35 work_item_mask = work_item_mask_vector.z;
36 } else if(subgroup_local_id < 128) {
37 work_item_mask = work_item_mask_vector.w;
38 }
39 if (elect_work_item & work_item_mask){
40 out[gid] = %s(in[gid]);
41 }
42 }
43 )";
44
45 template <typename T>
run_functions_add_mul_max_min_for_type(RunTestForType rft)46 int run_functions_add_mul_max_min_for_type(RunTestForType rft)
47 {
48 int error = rft.run_impl<T, SCIN_NU<T, ArithmeticOp::add_>>(
49 "sub_group_non_uniform_scan_inclusive_add");
50 error |= rft.run_impl<T, SCIN_NU<T, ArithmeticOp::mul_>>(
51 "sub_group_non_uniform_scan_inclusive_mul");
52 error |= rft.run_impl<T, SCIN_NU<T, ArithmeticOp::max_>>(
53 "sub_group_non_uniform_scan_inclusive_max");
54 error |= rft.run_impl<T, SCIN_NU<T, ArithmeticOp::min_>>(
55 "sub_group_non_uniform_scan_inclusive_min");
56 error |= rft.run_impl<T, SCEX_NU<T, ArithmeticOp::add_>>(
57 "sub_group_non_uniform_scan_exclusive_add");
58 error |= rft.run_impl<T, SCEX_NU<T, ArithmeticOp::mul_>>(
59 "sub_group_non_uniform_scan_exclusive_mul");
60 error |= rft.run_impl<T, SCEX_NU<T, ArithmeticOp::max_>>(
61 "sub_group_non_uniform_scan_exclusive_max");
62 error |= rft.run_impl<T, SCEX_NU<T, ArithmeticOp::min_>>(
63 "sub_group_non_uniform_scan_exclusive_min");
64 error |= rft.run_impl<T, RED_NU<T, ArithmeticOp::add_>>(
65 "sub_group_non_uniform_reduce_add");
66 error |= rft.run_impl<T, RED_NU<T, ArithmeticOp::mul_>>(
67 "sub_group_non_uniform_reduce_mul");
68 error |= rft.run_impl<T, RED_NU<T, ArithmeticOp::max_>>(
69 "sub_group_non_uniform_reduce_max");
70 error |= rft.run_impl<T, RED_NU<T, ArithmeticOp::min_>>(
71 "sub_group_non_uniform_reduce_min");
72 return error;
73 }
74
run_functions_and_or_xor_for_type(RunTestForType rft)75 template <typename T> int run_functions_and_or_xor_for_type(RunTestForType rft)
76 {
77 int error = rft.run_impl<T, SCIN_NU<T, ArithmeticOp::and_>>(
78 "sub_group_non_uniform_scan_inclusive_and");
79 error |= rft.run_impl<T, SCIN_NU<T, ArithmeticOp::or_>>(
80 "sub_group_non_uniform_scan_inclusive_or");
81 error |= rft.run_impl<T, SCIN_NU<T, ArithmeticOp::xor_>>(
82 "sub_group_non_uniform_scan_inclusive_xor");
83 error |= rft.run_impl<T, SCEX_NU<T, ArithmeticOp::and_>>(
84 "sub_group_non_uniform_scan_exclusive_and");
85 error |= rft.run_impl<T, SCEX_NU<T, ArithmeticOp::or_>>(
86 "sub_group_non_uniform_scan_exclusive_or");
87 error |= rft.run_impl<T, SCEX_NU<T, ArithmeticOp::xor_>>(
88 "sub_group_non_uniform_scan_exclusive_xor");
89 error |= rft.run_impl<T, RED_NU<T, ArithmeticOp::and_>>(
90 "sub_group_non_uniform_reduce_and");
91 error |= rft.run_impl<T, RED_NU<T, ArithmeticOp::or_>>(
92 "sub_group_non_uniform_reduce_or");
93 error |= rft.run_impl<T, RED_NU<T, ArithmeticOp::xor_>>(
94 "sub_group_non_uniform_reduce_xor");
95 return error;
96 }
97
98 template <typename T>
run_functions_logical_and_or_xor_for_type(RunTestForType rft)99 int run_functions_logical_and_or_xor_for_type(RunTestForType rft)
100 {
101 int error = rft.run_impl<T, SCIN_NU<T, ArithmeticOp::logical_and>>(
102 "sub_group_non_uniform_scan_inclusive_logical_and");
103 error |= rft.run_impl<T, SCIN_NU<T, ArithmeticOp::logical_or>>(
104 "sub_group_non_uniform_scan_inclusive_logical_or");
105 error |= rft.run_impl<T, SCIN_NU<T, ArithmeticOp::logical_xor>>(
106 "sub_group_non_uniform_scan_inclusive_logical_xor");
107 error |= rft.run_impl<T, SCEX_NU<T, ArithmeticOp::logical_and>>(
108 "sub_group_non_uniform_scan_exclusive_logical_and");
109 error |= rft.run_impl<T, SCEX_NU<T, ArithmeticOp::logical_or>>(
110 "sub_group_non_uniform_scan_exclusive_logical_or");
111 error |= rft.run_impl<T, SCEX_NU<T, ArithmeticOp::logical_xor>>(
112 "sub_group_non_uniform_scan_exclusive_logical_xor");
113 error |= rft.run_impl<T, RED_NU<T, ArithmeticOp::logical_and>>(
114 "sub_group_non_uniform_reduce_logical_and");
115 error |= rft.run_impl<T, RED_NU<T, ArithmeticOp::logical_or>>(
116 "sub_group_non_uniform_reduce_logical_or");
117 error |= rft.run_impl<T, RED_NU<T, ArithmeticOp::logical_xor>>(
118 "sub_group_non_uniform_reduce_logical_xor");
119 return error;
120 }
121
122 }
123
test_subgroup_functions_non_uniform_arithmetic(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)124 int test_subgroup_functions_non_uniform_arithmetic(cl_device_id device,
125 cl_context context,
126 cl_command_queue queue,
127 int num_elements)
128 {
129 if (!is_extension_available(device,
130 "cl_khr_subgroup_non_uniform_arithmetic"))
131 {
132 log_info("cl_khr_subgroup_non_uniform_arithmetic is not supported on "
133 "this device, skipping test.\n");
134 return TEST_SKIPPED_ITSELF;
135 }
136
137 constexpr size_t global_work_size = 2000;
138 constexpr size_t local_work_size = 200;
139 WorkGroupParams test_params(global_work_size, local_work_size, 3);
140 test_params.save_kernel_source(sub_group_non_uniform_arithmetic_source);
141 RunTestForType rft(device, context, queue, num_elements, test_params);
142
143 int error = run_functions_add_mul_max_min_for_type<cl_int>(rft);
144 error |= run_functions_add_mul_max_min_for_type<cl_uint>(rft);
145 error |= run_functions_add_mul_max_min_for_type<cl_long>(rft);
146 error |= run_functions_add_mul_max_min_for_type<cl_ulong>(rft);
147 error |= run_functions_add_mul_max_min_for_type<cl_short>(rft);
148 error |= run_functions_add_mul_max_min_for_type<cl_ushort>(rft);
149 error |= run_functions_add_mul_max_min_for_type<cl_char>(rft);
150 error |= run_functions_add_mul_max_min_for_type<cl_uchar>(rft);
151 error |= run_functions_add_mul_max_min_for_type<cl_float>(rft);
152 error |= run_functions_add_mul_max_min_for_type<cl_double>(rft);
153 error |= run_functions_add_mul_max_min_for_type<subgroups::cl_half>(rft);
154
155 error |= run_functions_and_or_xor_for_type<cl_int>(rft);
156 error |= run_functions_and_or_xor_for_type<cl_uint>(rft);
157 error |= run_functions_and_or_xor_for_type<cl_long>(rft);
158 error |= run_functions_and_or_xor_for_type<cl_ulong>(rft);
159 error |= run_functions_and_or_xor_for_type<cl_short>(rft);
160 error |= run_functions_and_or_xor_for_type<cl_ushort>(rft);
161 error |= run_functions_and_or_xor_for_type<cl_char>(rft);
162 error |= run_functions_and_or_xor_for_type<cl_uchar>(rft);
163
164 error |= run_functions_logical_and_or_xor_for_type<cl_int>(rft);
165 return error;
166 }
167