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 "procs.h"
17 #include "subhelpers.h"
18 #include "harness/conversions.h"
19 #include "harness/typeWrappers.h"
20
21 static const char *lbar_source =
22 "__kernel void test_lbar(const __global Type *in, __global int2 *xy, "
23 "__global Type *out)\n"
24 "{\n"
25 " __local int tmp[200];\n"
26 " int gid = get_global_id(0);\n"
27 " int nid = get_sub_group_size();\n"
28 " int lid = get_sub_group_local_id();\n"
29 " xy[gid].x = lid;\n"
30 " xy[gid].y = get_sub_group_id();\n"
31 " if (get_sub_group_id() == 0) {\n"
32 " tmp[lid] = in[gid];\n"
33 " sub_group_barrier(CLK_LOCAL_MEM_FENCE);\n"
34 " out[gid] = tmp[nid-1-lid];\n"
35 " } else {\n"
36 " out[gid] = -in[gid];\n"
37 " }\n"
38 "}\n";
39
40 static const char *gbar_source =
41 "__kernel void test_gbar(const __global Type *in, __global int2 *xy, "
42 "__global Type *out, __global Type *tmp)\n"
43 "{\n"
44 " int gid = get_global_id(0);\n"
45 " int nid = get_sub_group_size();\n"
46 " int lid = get_sub_group_local_id();\n"
47 " int tof = get_group_id(0)*get_max_sub_group_size();\n"
48 " xy[gid].x = lid;\n"
49 " xy[gid].y = get_sub_group_id();\n"
50 " if (get_sub_group_id() == 0) {\n"
51 " tmp[tof+lid] = in[gid];\n"
52 " sub_group_barrier(CLK_GLOBAL_MEM_FENCE);\n"
53 " out[gid] = tmp[tof+nid-1-lid];\n"
54 " } else {\n"
55 " out[gid] = -in[gid];\n"
56 " }\n"
57 "}\n";
58
59 // barrier test functions
60 template <int Which> struct BAR
61 {
genBAR62 static void gen(cl_int *x, cl_int *t, cl_int *m, int ns, int nw, int ng)
63 {
64 int i, ii, j, k, n;
65 int nj = (nw + ns - 1) / ns;
66 int e;
67
68 ii = 0;
69 for (k = 0; k < ng; ++k)
70 {
71 for (j = 0; j < nj; ++j)
72 {
73 ii = j * ns;
74 n = ii + ns > nw ? nw - ii : ns;
75
76 for (i = 0; i < n; ++i) t[ii + i] = genrand_int32(gMTdata);
77 }
78
79 // Now map into work group using map from device
80 for (j = 0; j < nw; ++j)
81 {
82 i = m[2 * j + 1] * ns + m[2 * j];
83 x[j] = t[i];
84 }
85
86 x += nw;
87 m += 2 * nw;
88 }
89 }
90
chkBAR91 static int chk(cl_int *x, cl_int *y, cl_int *mx, cl_int *my, cl_int *m,
92 int ns, int nw, int ng)
93 {
94 int ii, i, j, k, n;
95 int nj = (nw + ns - 1) / ns;
96 cl_int tr, rr;
97
98 if (Which == 0)
99 log_info(" sub_group_barrier(CLK_LOCAL_MEM_FENCE)...\n");
100 else
101 log_info(" sub_group_barrier(CLK_GLOBAL_MEM_FENCE)...\n");
102
103 for (k = 0; k < ng; ++k)
104 {
105 // Map to array indexed to array indexed by local ID and sub group
106 for (j = 0; j < nw; ++j)
107 {
108 i = m[2 * j + 1] * ns + m[2 * j];
109 mx[i] = x[j];
110 my[i] = y[j];
111 }
112
113 for (j = 0; j < nj; ++j)
114 {
115 ii = j * ns;
116 n = ii + ns > nw ? nw - ii : ns;
117
118 for (i = 0; i < n; ++i)
119 {
120 tr = j == 0 ? mx[ii + n - 1 - i] : -mx[ii + i];
121 rr = my[ii + i];
122
123 if (tr != rr)
124 {
125 log_error("ERROR: sub_group_barrier mismatch for local "
126 "id %d in sub group %d in group %d\n",
127 i, j, k);
128 return -1;
129 }
130 }
131 }
132
133 x += nw;
134 y += nw;
135 m += 2 * nw;
136 }
137
138 return 0;
139 }
140 };
141
142 // Entry point from main
test_barrier_functions(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements,bool useCoreSubgroups)143 int test_barrier_functions(cl_device_id device, cl_context context,
144 cl_command_queue queue, int num_elements,
145 bool useCoreSubgroups)
146 {
147 int error;
148
149 // Adjust these individually below if desired/needed
150 #define G 2000
151 #define L 200
152
153 error = test<cl_int, BAR<0>, G, L>::run(device, context, queue,
154 num_elements, "test_lbar",
155 lbar_source, 0, useCoreSubgroups);
156 error = test<cl_int, BAR<1>, G, L, G>::run(
157 device, context, queue, num_elements, "test_gbar", gbar_source, 0,
158 useCoreSubgroups);
159
160 return error;
161 }
162
test_barrier_functions_core(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)163 int test_barrier_functions_core(cl_device_id device, cl_context context,
164 cl_command_queue queue, int num_elements)
165 {
166 return test_barrier_functions(device, context, queue, num_elements, true);
167 }
168
test_barrier_functions_ext(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)169 int test_barrier_functions_ext(cl_device_id device, cl_context context,
170 cl_command_queue queue, int num_elements)
171 {
172 bool hasExtension = is_extension_available(device, "cl_khr_subgroups");
173
174 if (!hasExtension)
175 {
176 log_info(
177 "Device does not support 'cl_khr_subgroups'. Skipping the test.\n");
178 return TEST_SKIPPED_ITSELF;
179 }
180
181 return test_barrier_functions(device, context, queue, num_elements, false);
182 }