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=%u global=%u local=%u...\n", gwo[0], gws[0], lws[0]);
108 break;
109 case 2:
110 log_info(" testing offset=(%u,%u) global=(%u,%u) local=(%u,%u)...\n",
111 gwo[0], gwo[1], gws[0], gws[1], lws[0], lws[1]);
112 break;
113 case 3:
114 log_info(" testing offset=(%u,%u,%u) global=(%u,%u,%u) local=(%u,%u,%u)...\n",
115 gwo[0], gwo[1], gwo[2], gws[0], gws[1], gws[2], lws[0], lws[1], lws[2]);
116 break;
117 }
118
119 // Set up and run
120 memset(outmem, 0, sizeof(outmem));
121
122 error = clSetKernelArg(kernel, 0, sizeof(outbuf), (void *)&outbuf);
123 test_error(error, "clSetKernelArg failed\n");
124
125 error = clEnqueueWriteBuffer(queue, outbuf, CL_FALSE, 0, sizeof(outmem), (void *)outmem, 0, NULL, NULL);
126 test_error(error, "clEnqueueWriteBuffer failed\n");
127
128 error = clEnqueueNDRangeKernel(queue, kernel, dims, gwo, gws, lws, 0, NULL, NULL);
129 test_error(error, "clEnqueueNDRangeKernel failed\n");
130
131 error = clEnqueueReadBuffer(queue, outbuf, CL_FALSE, 0, sizeof(outmem), (void *)outmem, 0, NULL, NULL);
132 test_error(error, "clEnqueueReadBuffer failed\n");
133
134 error = clFinish(queue);
135 test_error(error, "clFinish failed\n");
136
137 // Check the return
138 switch (dims) {
139 case 1:
140 for (i=0, om=outmem; i<(int)gws[0]; ++i, om+=2) {
141 if (om[0] != 1) {
142 log_error("get_global_linear_id() failed at %d\n", i);
143 return -1;
144 }
145 if (om[1] != 1) {
146 log_error("get_local_linear_id() failed at (%d, %d)\n", i % (int)lws[0], i / (int)lws[0]);
147 return -1;
148 }
149 }
150 break;
151 case 2:
152 for (j=0, om=outmem; j<gws[1]; ++j) {
153 for (i=0; i<gws[0]; ++i, om+=2) {
154 if (om[0] != 1) {
155 log_error("get_global_linear_id() failed at (%d,%d)\n", i, j);
156 return -1;
157 }
158 if (om[1] != 1) {
159 log_error("get_local_linear_id() failed at (%d, %d), (%d, %d)\n",
160 i % (int)lws[0], j % (int)lws[1],
161 i / (int)lws[0], j / (int)lws[1]);
162 return -1;
163 }
164 }
165 }
166 break;
167 case 3:
168 for (k=0, om=outmem; k<gws[2]; ++k) {
169 for (j=0; j<gws[1]; ++j) {
170 for (i=0; i<gws[0]; ++i, om+=2) {
171 if (om[0] != 1) {
172 log_error("get_global_linear_id() failed at (%d,%d, %d)\n", i, j, k);
173 return -1;
174 }
175 if (om[1] != 1) {
176 log_error("get_local_linear_id() failed at (%d, %d), (%d, %d), (%d, %d)\n",
177 i % (int)lws[0], j % (int)lws[1], k % (int)lws[2],
178 i / (int)lws[0], j / (int)lws[1], k / (int)lws[2]);
179 return -1;
180 }
181 }
182 }
183 }
184 break;
185 }
186
187 }
188
189 free_mtdata(seed);
190 return 0;
191 }
192
193