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 #include "testBase.h"
17 #include "harness/typeWrappers.h"
18 #include "harness/conversions.h"
19
20 const char *subgroup_dispatch_kernel[] = {
21 "#pragma OPENCL EXTENSION cl_khr_subgroups : enable\n"
22 "__kernel void subgroup_dispatch_kernel(__global int *output)\n"
23 "{\n"
24 " size_t size = get_num_sub_groups ();\n"
25 "\n"
26 " output[0] = size;\n"
27 "\n"
28 "}\n" };
29
flatten_ndrange(size_t * ndrange,size_t dim)30 size_t flatten_ndrange(size_t* ndrange, size_t dim)
31 {
32 switch(dim)
33 {
34 case 1:
35 return *ndrange;
36 case 2:
37 return ndrange[0] * ndrange[1];
38 case 3:
39 return ndrange[0] * ndrange[1] * ndrange[2];
40 default:
41 log_error("ERROR: bad ndrange value");
42 return 0;
43 }
44 }
45
get_sub_group_num(cl_command_queue queue,cl_kernel kernel,clMemWrapper & out,size_t & size,size_t local_size,size_t dim)46 cl_int get_sub_group_num(cl_command_queue queue, cl_kernel kernel, clMemWrapper& out, size_t& size, size_t local_size, size_t dim)
47 {
48 size_t ndrange[3] = {local_size, 1, 1};
49 cl_int error = CL_SUCCESS;
50 size = 0;
51 error = clSetKernelArg(kernel, 0, sizeof(out), &out);
52 error += clEnqueueNDRangeKernel(queue, kernel, dim, NULL, ndrange, ndrange, 0, NULL, NULL);
53 error += clEnqueueReadBuffer(queue, out, CL_TRUE, 0, 4, &size, 0, NULL, NULL);
54 return error;
55 }
56
test_sub_group_dispatch(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)57 int test_sub_group_dispatch(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
58 {
59 int error;
60 size_t realSize;
61 size_t kernel_max_subgroup_size, kernel_subgroup_count;
62 size_t max_local;
63
64 cl_platform_id platform;
65 clProgramWrapper program;
66 clKernelWrapper kernel;
67 clMemWrapper out;
68
69 size_t ret_ndrange1d;
70 size_t ret_ndrange2d[2];
71 size_t ret_ndrange3d[3];
72
73 size_t ret_ndrange2d_flattened;
74 size_t ret_ndrange3d_flattened;
75
76 if (get_device_cl_version(deviceID) >= Version(3, 0))
77 {
78 int error;
79 cl_uint max_num_sub_groups;
80
81 error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_NUM_SUB_GROUPS,
82 sizeof(max_num_sub_groups), &max_num_sub_groups,
83 NULL);
84 if (error != CL_SUCCESS)
85 {
86 print_error(error, "Unable to get max num subgroups");
87 return error;
88 }
89
90 if (max_num_sub_groups == 0)
91 {
92 return TEST_SKIPPED_ITSELF;
93 }
94 }
95
96 error = create_single_kernel_helper(context, &program, &kernel, 1,
97 subgroup_dispatch_kernel,
98 "subgroup_dispatch_kernel");
99 if (error != 0)
100 return error;
101
102 out = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, sizeof(size_t), NULL, &error);
103 test_error(error, "clCreateBuffer failed");
104
105 error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t), &max_local, NULL);
106 test_error(error, "clGetDeviceInfo failed");
107
108
109 error = clGetDeviceInfo(deviceID, CL_DEVICE_PLATFORM, sizeof(platform), (void *)&platform, NULL);
110 test_error(error, "clDeviceInfo failed for CL_DEVICE_PLATFORM");
111
112 // Get the max subgroup size
113 error = clGetKernelSubGroupInfo(kernel, deviceID, CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE,
114 sizeof(max_local), &max_local, sizeof(kernel_max_subgroup_size), (void *)&kernel_max_subgroup_size, &realSize);
115 test_error(error, "clGetKernelSubGroupInfo failed for CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE");
116 log_info("The CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE for the kernel is %d.\n", (int)kernel_max_subgroup_size);
117
118 if (realSize != sizeof(kernel_max_subgroup_size)) {
119 log_error( "ERROR: Returned size of max sub group size not valid! (Expected %d, got %d)\n", (int)sizeof(kernel_max_subgroup_size), (int)realSize );
120 return -1;
121 }
122
123 // Get the number of subgroup for max local size
124 error = clGetKernelSubGroupInfo(kernel, deviceID, CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE,
125 sizeof(max_local), &max_local, sizeof(kernel_subgroup_count), (void *)&kernel_subgroup_count, &realSize);
126 test_error(error, "clGetKernelSubGroupInfo failed for CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE");
127 log_info("The CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE for the kernel is %d.\n", (int)kernel_subgroup_count);
128
129 if (realSize != sizeof(kernel_subgroup_count)) {
130 log_error( "ERROR: Returned size of sub group count not valid! (Expected %d, got %d)\n", (int)sizeof(kernel_subgroup_count), (int)realSize );
131 return -1;
132 }
133
134 // test CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT
135 for (size_t i = kernel_subgroup_count; i > 0; --i)
136 {
137 // test all 3 different dimention of requested local size
138 size_t expect_size = kernel_max_subgroup_size * i;
139 size_t kernel_ret_size = 0;
140 error = clGetKernelSubGroupInfo(kernel, deviceID, CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT, sizeof(i), &i, sizeof(ret_ndrange1d), &ret_ndrange1d, &realSize);
141 test_error(error, "clGetKernelSubGroupInfo failed for CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT");
142 if (realSize != sizeof(ret_ndrange1d)) {
143 log_error( "ERROR: Returned size of sub group count not valid! (Expected %d, got %d)\n", (int)sizeof(kernel_subgroup_count), (int)realSize );
144 return -1;
145 }
146
147 if (ret_ndrange1d != expect_size)
148 {
149 log_error( "ERROR: Incorrect value returned for CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT! (Expected %d, got %d)\n", (int)expect_size, (int)ret_ndrange1d );
150 return -1;
151 }
152
153 error = get_sub_group_num(queue, kernel, out, kernel_ret_size, ret_ndrange1d, 1);
154 test_error(error, "Failed to query number of subgroups from kernel");
155 if (i != kernel_ret_size)
156 {
157 log_error( "ERROR: Mismatch between requested number of subgroups and what get_num_sub_groups() in kernel returned! (Expected %d, got %d)\n", (int)i, (int)kernel_ret_size );
158 return -1;
159 }
160
161 error = clGetKernelSubGroupInfo(kernel, deviceID, CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT, sizeof(i), &i, sizeof(ret_ndrange2d), ret_ndrange2d, &realSize);
162 test_error(error, "clGetKernelSubGroupInfo failed for CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT");
163 if (realSize != sizeof(ret_ndrange2d)) {
164 log_error( "ERROR: Returned size of sub group count not valid! (Expected %d, got %d)\n", (int)sizeof(kernel_subgroup_count), (int)realSize );
165 return -1;
166 }
167
168 ret_ndrange2d_flattened = flatten_ndrange(ret_ndrange2d, 2);
169 if (ret_ndrange2d_flattened != expect_size ||
170 ret_ndrange2d[1] != 1)
171 {
172 log_error( "ERROR: Incorrect value returned for CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT! (Expected %d, got %d)\n", (int)expect_size, (int)ret_ndrange2d_flattened );
173 return -1;
174 }
175
176 error = get_sub_group_num(queue, kernel, out, kernel_ret_size, ret_ndrange2d_flattened, 2);
177 test_error(error, "Failed to query number of subgroups from kernel");
178 if (i != kernel_ret_size)
179 {
180 log_error( "ERROR: Mismatch between requested number of subgroups and what get_num_sub_groups() in kernel returned! (Expected %d, got %d)\n", (int)i, (int)kernel_ret_size );
181 return -1;
182 }
183
184 error = clGetKernelSubGroupInfo(kernel, deviceID, CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT, sizeof(i), &i, sizeof(ret_ndrange3d), ret_ndrange3d, &realSize);
185 test_error(error, "clGetKernelSubGroupInfo failed for CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT");
186 if (realSize != sizeof(ret_ndrange3d)) {
187 log_error( "ERROR: Returned size of sub group count not valid! (Expected %d, got %d)\n", (int)sizeof(kernel_subgroup_count), (int)realSize );
188 return -1;
189 }
190
191 ret_ndrange3d_flattened = flatten_ndrange(ret_ndrange3d, 3);
192 if (ret_ndrange3d_flattened != expect_size ||
193 ret_ndrange3d[1] != 1 ||
194 ret_ndrange3d[2] != 1)
195 {
196 log_error( "ERROR: Incorrect value returned for CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT! (Expected %d, got %d)\n", (int)expect_size, (int)ret_ndrange3d_flattened );
197 return -1;
198 }
199
200 error = get_sub_group_num(queue, kernel, out, kernel_ret_size, ret_ndrange3d_flattened, 3);
201 test_error(error, "Failed to query number of subgroups from kernel");
202 if (i != kernel_ret_size)
203 {
204 log_error( "ERROR: Mismatch between requested number of subgroups and what get_num_sub_groups() in kernel returned! (Expected %d, got %d)\n", (int)i, (int)kernel_ret_size );
205 return -1;
206 }
207 }
208
209 // test when input subgroup count exceeds max wg size
210 size_t large_sg_size = kernel_subgroup_count + 1;
211 error = clGetKernelSubGroupInfo(kernel, deviceID, CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT, sizeof(size_t), &large_sg_size, sizeof(ret_ndrange1d), &ret_ndrange1d, &realSize);
212 test_error(error, "clGetKernelSubGroupInfo failed for CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT");
213 if (ret_ndrange1d != 0)
214 {
215 log_error( "ERROR: Incorrect value returned for CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT! (Expected %d, got %d)\n", 0, (int)ret_ndrange1d );
216 return -1;
217 }
218
219 error = clGetKernelSubGroupInfo(kernel, deviceID, CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT, sizeof(size_t), &large_sg_size, sizeof(ret_ndrange2d), ret_ndrange2d, &realSize);
220 test_error(error, "clGetKernelSubGroupInfo failed for CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT");
221 if (ret_ndrange2d[0] != 0 ||
222 ret_ndrange2d[1] != 0)
223 {
224 log_error( "ERROR: Incorrect value returned for CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT!" );
225 return -1;
226 }
227
228 error = clGetKernelSubGroupInfo(kernel, deviceID, CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT, sizeof(size_t), &large_sg_size, sizeof(ret_ndrange3d), ret_ndrange3d, &realSize);
229 test_error(error, "clGetKernelSubGroupInfo failed for CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT");
230 if (ret_ndrange3d[0] != 0 ||
231 ret_ndrange3d[1] != 0 ||
232 ret_ndrange3d[2] != 0)
233 {
234 log_error( "ERROR: Incorrect value returned for CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT!" );
235 return -1;
236 }
237
238 return 0;
239 }
240