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 <ctype.h>
18
19 static const char *linear_ids_source[1] = {
20 "__kernel void test_linear_ids(__global int2 *out)\n"
21 "{\n"
22 " size_t lid, gid;\n"
23 " uint d = get_work_dim();\n"
24 " if (d == 1U) {\n"
25 " gid = get_global_id(0) - get_global_offset(0);\n"
26 " lid = get_local_id(0);\n"
27 " } else if (d == 2U) {\n"
28 " gid = (get_global_id(1) - get_global_offset(1)) * get_global_size(0) +\n"
29 " (get_global_id(0) - get_global_offset(0));\n"
30 " lid = get_local_id(1) * get_local_size(0) + get_local_id(0);\n"
31 " } else {\n"
32 " gid = ((get_global_id(2) - get_global_offset(2)) * get_global_size(1) +\n"
33 " (get_global_id(1) - get_global_offset(1))) * get_global_size(0) +\n"
34 " (get_global_id(0) - get_global_offset(0));\n"
35 " lid = (get_local_id(2) * get_local_size(1) +\n"
36 " get_local_id(1)) * get_local_size(0) + get_local_id(0);\n"
37 " }\n"
38 " out[gid].x = gid == get_global_linear_id();\n"
39 " out[gid].y = lid == get_local_linear_id();\n"
40 "}\n"
41 };
42
43 #define NUM_ITER 12
44 #define MAX_1D 4096
45 #define MAX_2D 64
46 #define MAX_3D 16
47 #define MAX_OFFSET 100000
48
49 int
test_get_linear_ids(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)50 test_get_linear_ids(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
51 {
52 clProgramWrapper program;
53 clKernelWrapper kernel;
54 clMemWrapper outbuf;
55 int error, iter, i, j, k;
56 size_t lws[3], gws[3], gwo[3];
57 cl_uint dims;
58 cl_int outmem[2*MAX_1D], *om;
59
60
61 // Create the kernel
62 error = create_single_kernel_helper(context, &program, &kernel, 1,
63 linear_ids_source, "test_linear_ids");
64 if (error)
65 return error;
66
67 // Create the out buffer
68 outbuf = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(outmem), NULL, &error);
69 test_error(error, "failed to create result buffer\n");
70
71 // This will leak if there is an error, but this is what is done everywhere else
72 MTdata seed = init_genrand(gRandomSeed);
73
74 // Run some tests
75 for (iter=0; iter<NUM_ITER; ++iter) {
76 dims = iter % 3 + 1;
77
78 switch (dims) {
79 case 1:
80 gwo[0] = random_in_range(0, MAX_OFFSET, seed);
81 gws[0] = random_in_range(MAX_1D/8, MAX_1D/4, seed)*4;
82 error = get_max_common_work_group_size(context, kernel, gws[0], lws);
83 break;
84 case 2:
85 gwo[0] = random_in_range(0, MAX_OFFSET, seed);
86 gwo[1] = random_in_range(0, MAX_OFFSET, seed);
87 gws[0] = random_in_range(MAX_2D/8, MAX_2D/4, seed)*4;
88 gws[1] = random_in_range(MAX_2D/8, MAX_2D/4, seed)*4;
89 error = get_max_common_2D_work_group_size(context, kernel, gws, lws);
90 break;
91 case 3:
92 gwo[0] = random_in_range(0, MAX_OFFSET, seed);
93 gwo[1] = random_in_range(0, MAX_OFFSET, seed);
94 gwo[2] = random_in_range(0, MAX_OFFSET, seed);
95 gws[0] = random_in_range(MAX_3D/4, MAX_3D/2, seed)*2;
96 gws[1] = random_in_range(MAX_3D/4, MAX_3D/2, seed)*2;
97 gws[2] = random_in_range(MAX_3D/4, MAX_3D/2, seed)*2;
98 error = get_max_common_3D_work_group_size(context, kernel, gws, lws);
99 break;
100 }
101
102 test_error(error, "Failed to determine local work size\n");
103
104
105 switch (dims) {
106 case 1:
107 log_info(" testing offset=%zu global=%zu local=%zu...\n", gwo[0],
108 gws[0], lws[0]);
109 break;
110 case 2:
111 log_info(" testing offset=(%zu,%zu) global=(%zu,%zu) "
112 "local=(%zu,%zu)...\n",
113 gwo[0], gwo[1], gws[0], gws[1], lws[0], lws[1]);
114 break;
115 case 3:
116 log_info(" testing offset=(%zu,%zu,%zu) global=(%zu,%zu,%zu) "
117 "local=(%zu,%zu,%zu)...\n",
118 gwo[0], gwo[1], gwo[2], gws[0], gws[1], gws[2], lws[0],
119 lws[1], lws[2]);
120 break;
121 }
122
123 // Set up and run
124 memset(outmem, 0, sizeof(outmem));
125
126 error = clSetKernelArg(kernel, 0, sizeof(outbuf), (void *)&outbuf);
127 test_error(error, "clSetKernelArg failed\n");
128
129 error = clEnqueueWriteBuffer(queue, outbuf, CL_FALSE, 0, sizeof(outmem), (void *)outmem, 0, NULL, NULL);
130 test_error(error, "clEnqueueWriteBuffer failed\n");
131
132 error = clEnqueueNDRangeKernel(queue, kernel, dims, gwo, gws, lws, 0, NULL, NULL);
133 test_error(error, "clEnqueueNDRangeKernel failed\n");
134
135 error = clEnqueueReadBuffer(queue, outbuf, CL_FALSE, 0, sizeof(outmem), (void *)outmem, 0, NULL, NULL);
136 test_error(error, "clEnqueueReadBuffer failed\n");
137
138 error = clFinish(queue);
139 test_error(error, "clFinish failed\n");
140
141 // Check the return
142 switch (dims) {
143 case 1:
144 for (i=0, om=outmem; i<(int)gws[0]; ++i, om+=2) {
145 if (om[0] != 1) {
146 log_error("get_global_linear_id() failed at %d\n", i);
147 return -1;
148 }
149 if (om[1] != 1) {
150 log_error("get_local_linear_id() failed at (%d, %d)\n", i % (int)lws[0], i / (int)lws[0]);
151 return -1;
152 }
153 }
154 break;
155 case 2:
156 for (j=0, om=outmem; j<gws[1]; ++j) {
157 for (i=0; i<gws[0]; ++i, om+=2) {
158 if (om[0] != 1) {
159 log_error("get_global_linear_id() failed at (%d,%d)\n", i, j);
160 return -1;
161 }
162 if (om[1] != 1) {
163 log_error("get_local_linear_id() failed at (%d, %d), (%d, %d)\n",
164 i % (int)lws[0], j % (int)lws[1],
165 i / (int)lws[0], j / (int)lws[1]);
166 return -1;
167 }
168 }
169 }
170 break;
171 case 3:
172 for (k=0, om=outmem; k<gws[2]; ++k) {
173 for (j=0; j<gws[1]; ++j) {
174 for (i=0; i<gws[0]; ++i, om+=2) {
175 if (om[0] != 1) {
176 log_error("get_global_linear_id() failed at (%d,%d, %d)\n", i, j, k);
177 return -1;
178 }
179 if (om[1] != 1) {
180 log_error("get_local_linear_id() failed at (%d, %d), (%d, %d), (%d, %d)\n",
181 i % (int)lws[0], j % (int)lws[1], k % (int)lws[2],
182 i / (int)lws[0], j / (int)lws[1], k / (int)lws[2]);
183 return -1;
184 }
185 }
186 }
187 }
188 break;
189 }
190
191 }
192
193 free_mtdata(seed);
194 return 0;
195 }
196
197