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 "harness/compat.h"
17
18 #include <stdio.h>
19 #include <stdlib.h>
20 #include <string.h>
21 #include <sys/types.h>
22 #include <sys/stat.h>
23
24
25 #include "procs.h"
26
27 static const char *image_dim_kernel_code =
28 "\n"
29 "__kernel void test_image_dim(read_only image2d_t srcimg, write_only image2d_t dstimg, sampler_t sampler)\n"
30 "{\n"
31 " int tid_x = get_global_id(0);\n"
32 " int tid_y = get_global_id(1);\n"
33 " float4 color;\n"
34 "\n"
35 " color = read_imagef(srcimg, sampler, (int2)(tid_x, tid_y));\n"
36 " write_imagef(dstimg, (int2)(tid_x, tid_y), color);\n"
37 "\n"
38 "}\n";
39
40
generate_8888_image(size_t w,size_t h,MTdata d)41 static unsigned char *generate_8888_image(size_t w, size_t h, MTdata d)
42 {
43 unsigned char *ptr = new unsigned char[4 * w * h];
44 size_t i;
45
46 for (i = 0; i < w * h * 4; i++)
47 {
48 ptr[i] = (unsigned char)genrand_int32(d);
49 }
50
51 return ptr;
52 }
53
verify_8888_image(unsigned char * image,unsigned char * outptr,size_t w,size_t h)54 static int verify_8888_image(unsigned char *image, unsigned char *outptr,
55 size_t w, size_t h)
56 {
57 size_t i;
58
59 for (i = 0; i < w * h; i++)
60 {
61 if (outptr[i] != image[i])
62 return -1;
63 }
64
65 return 0;
66 }
67
68
69 int
test_imagedim_pow2(cl_device_id device,cl_context context,cl_command_queue queue,int n_elems)70 test_imagedim_pow2(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems)
71 {
72 cl_mem streams[2];
73 cl_image_format img_format;
74 unsigned char *input_ptr, *output_ptr;
75 cl_program program;
76 cl_kernel kernel;
77 size_t threads[2];
78 cl_ulong max_mem_size;
79 size_t img_width, max_img_width;
80 size_t img_height, max_img_height;
81 size_t max_img_dim;
82 int i, j, i2, j2, err = 0;
83 size_t max_image2d_width, max_image2d_height;
84 int total_errors = 0;
85 MTdata d;
86
87 PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
88
89 err = create_single_kernel_helper( context, &program, &kernel, 1, &image_dim_kernel_code, "test_image_dim" );
90 if (err)
91 {
92 log_error("create_program_and_kernel_with_sources failed\n");
93 return -1;
94 }
95
96 err = clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE,sizeof(max_mem_size), &max_mem_size, NULL);
97 if (err)
98 {
99 log_error("clGetDeviceInfo for CL_DEVICE_GLOBAL_MEM_SIZE failed (%d)\n", err);
100 return -1;
101 }
102 err = clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_WIDTH, sizeof(max_image2d_width), &max_image2d_width, NULL);
103 if (err)
104 {
105 log_error("clGetDeviceInfo for CL_DEVICE_IMAGE2D_MAX_WIDTH failed (%d)\n", err);
106 return -1;
107 }
108 err = clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_HEIGHT, sizeof(max_image2d_width), &max_image2d_height, NULL);
109 if (err)
110 {
111 log_error("clGetDeviceInfo for CL_DEVICE_IMAGE2D_MAX_HEIGHT failed (%d)\n", err);
112 return -1;
113 }
114 log_info("Device reported max image sizes of %lu x %lu, and max mem size of %gMB.\n",
115 max_image2d_width, max_image2d_height, max_mem_size/(1024.0*1024.0));
116
117 if (max_mem_size > (cl_ulong)SIZE_MAX) {
118 max_mem_size = (cl_ulong)SIZE_MAX;
119 }
120
121 cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err);
122 test_error(err, "clCreateSampler failed");
123
124 max_img_width = max_image2d_width;
125 max_img_height = max_image2d_height;
126
127 // determine max image dim we can allocate - assume RGBA image, 4 bytes per pixel,
128 // and we want to consume 1/4 of global memory (this is the minimum required to be
129 // supported by the spec)
130 max_mem_size /= 4; // use 1/4
131 max_mem_size /= 4; // 4 bytes per pixel
132 max_img_dim = (size_t)sqrt((double)max_mem_size);
133 // convert to a power of 2
134 {
135 unsigned int n = (unsigned int)max_img_dim;
136 unsigned int m = 0x80000000;
137
138 // round-down to the nearest power of 2
139 while (m > n)
140 m >>= 1;
141
142 max_img_dim = m;
143 }
144
145 if (max_img_width > max_img_dim)
146 max_img_width = max_img_dim;
147 if (max_img_height > max_img_dim)
148 max_img_height = max_img_dim;
149
150 log_info("Adjusted maximum image size to test is %d x %d, which is a max mem size of %gMB.\n",
151 max_img_width, max_img_height, (max_img_width*max_img_height*4)/(1024.0*1024.0));
152
153 d = init_genrand( gRandomSeed );
154 input_ptr = generate_8888_image(max_img_width, max_img_height, d);
155
156 output_ptr = new unsigned char[4 * max_img_width * max_img_height];
157
158 // test power of 2 width, height starting at 1 to 4K
159 for (i = 1, i2 = 0; i <= max_img_height; i <<= 1, i2++)
160 {
161 img_height = (1 << i2);
162 for (j = 1, j2 = 0; j <= max_img_width; j <<= 1, j2++)
163 {
164 img_width = (1 << j2);
165
166 img_format.image_channel_order = CL_RGBA;
167 img_format.image_channel_data_type = CL_UNORM_INT8;
168 streams[0] =
169 create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
170 img_width, img_height, 0, NULL, NULL);
171 if (!streams[0])
172 {
173 log_error("create_image_2d failed. width = %d, height = %d\n", img_width, img_height);
174 delete[] input_ptr;
175 delete[] output_ptr;
176 free_mtdata(d);
177 return -1;
178 }
179 img_format.image_channel_order = CL_RGBA;
180 img_format.image_channel_data_type = CL_UNORM_INT8;
181 streams[1] =
182 create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
183 img_width, img_height, 0, NULL, NULL);
184 if (!streams[1])
185 {
186 log_error("create_image_2d failed. width = %d, height = %d\n", img_width, img_height);
187 clReleaseMemObject(streams[0]);
188 delete[] input_ptr;
189 delete[] output_ptr;
190 free_mtdata(d);
191 return -1;
192 }
193
194 size_t origin[3] = {0,0,0};
195 size_t region[3] = {img_width, img_height, 1};
196 err = clEnqueueWriteImage(queue, streams[0], CL_FALSE, origin, region, 0, 0, input_ptr, 0, NULL, NULL);
197 if (err != CL_SUCCESS)
198 {
199 log_error("clWriteImage failed\n");
200 clReleaseMemObject(streams[0]);
201 clReleaseMemObject(streams[1]);
202 delete[] input_ptr;
203 delete[] output_ptr;
204 free_mtdata(d);
205 return -1;
206 }
207
208 err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]);
209 err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]);
210 err |= clSetKernelArg(kernel, 2, sizeof sampler, &sampler);
211 if (err != CL_SUCCESS)
212 {
213 log_error("clSetKernelArgs failed\n");
214 clReleaseMemObject(streams[0]);
215 clReleaseMemObject(streams[1]);
216 delete[] input_ptr;
217 delete[] output_ptr;
218 free_mtdata(d);
219 return -1;
220 }
221
222 threads[0] = (size_t)img_width;
223 threads[1] = (size_t)img_height;
224 log_info("Testing image dimensions %d x %d with local threads NULL.\n", img_width, img_height);
225 err = clEnqueueNDRangeKernel( queue, kernel, 2, NULL, threads, NULL, 0, NULL, NULL );
226 if (err != CL_SUCCESS)
227 {
228 log_error("clEnqueueNDRangeKernel failed\n");
229 log_error("Image Dimension test failed. image width = %d, image height = %d, local NULL\n",
230 img_width, img_height);
231 clReleaseMemObject(streams[0]);
232 clReleaseMemObject(streams[1]);
233 delete[] input_ptr;
234 delete[] output_ptr;
235 free_mtdata(d);
236 return -1;
237 }
238 err = clEnqueueReadImage(queue, streams[1], CL_TRUE, origin, region, 0, 0, output_ptr, 0, NULL, NULL);
239 if (err != CL_SUCCESS)
240 {
241 log_error("clReadImage failed\n");
242 log_error("Image Dimension test failed. image width = %d, image height = %d, local NULL\n",
243 img_width, img_height);
244 clReleaseMemObject(streams[0]);
245 clReleaseMemObject(streams[1]);
246 delete[] input_ptr;
247 delete[] output_ptr;
248 free_mtdata(d);
249 return -1;
250 }
251 err = verify_8888_image(input_ptr, output_ptr, img_width, img_height);
252 if (err)
253 {
254 total_errors++;
255 log_error("Image Dimension test failed. image width = %d, image height = %d\n", img_width, img_height);
256 }
257
258 clReleaseMemObject(streams[0]);
259 clReleaseMemObject(streams[1]);
260 }
261 }
262
263 // cleanup
264 delete[] input_ptr;
265 delete[] output_ptr;
266 free_mtdata(d);
267 clReleaseSampler(sampler);
268 clReleaseKernel(kernel);
269 clReleaseProgram(program);
270
271 return total_errors;
272 }
273
274
275
276 int
test_imagedim_non_pow2(cl_device_id device,cl_context context,cl_command_queue queue,int n_elems)277 test_imagedim_non_pow2(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems)
278 {
279 cl_mem streams[2];
280 cl_image_format img_format;
281 unsigned char *input_ptr, *output_ptr;
282 cl_program program;
283 cl_kernel kernel;
284 size_t threads[2], local_threads[2];
285 cl_ulong max_mem_size;
286 size_t img_width, max_img_width;
287 size_t img_height, max_img_height;
288 size_t max_img_dim;
289 int i, j, i2, j2, err = 0;
290 size_t max_image2d_width, max_image2d_height;
291 int total_errors = 0;
292 size_t max_local_workgroup_size[3];
293 MTdata d;
294
295 PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
296
297 err = create_single_kernel_helper( context, &program, &kernel, 1, &image_dim_kernel_code, "test_image_dim" );
298 if (err)
299 {
300 log_error("create_program_and_kernel_with_sources failed\n");
301 return -1;
302 }
303
304 size_t work_group_size = 0;
305 err = clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(work_group_size), &work_group_size, NULL);
306 test_error(err, "clGetKerenlWorkgroupInfo failed for CL_KERNEL_WORK_GROUP_SIZE");
307
308 err = clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(max_local_workgroup_size), max_local_workgroup_size, NULL);
309 test_error(err, "clGetDeviceInfo failed for CL_DEVICE_MAX_WORK_ITEM_SIZES");
310
311 err = clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE,sizeof(max_mem_size), &max_mem_size, NULL);
312 if (err)
313 {
314 log_error("clGetDeviceInfo for CL_DEVICE_GLOBAL_MEM_SIZE failed (%d)\n", err);
315 return -1;
316 }
317 err = clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_WIDTH, sizeof(max_image2d_width), &max_image2d_width, NULL);
318 if (err)
319 {
320 log_error("clGetDeviceInfo for CL_DEVICE_IMAGE2D_MAX_WIDTH failed (%d)\n", err);
321 return -1;
322 }
323 err = clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_HEIGHT, sizeof(max_image2d_width), &max_image2d_height, NULL);
324 if (err)
325 {
326 log_error("clGetDeviceInfo for CL_DEVICE_IMAGE2D_MAX_HEIGHT failed (%d)\n", err);
327 return -1;
328 }
329 log_info("Device reported max image sizes of %lu x %lu, and max mem size of %gMB.\n",
330 max_image2d_width, max_image2d_height, max_mem_size/(1024.0*1024.0));
331
332 cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err);
333 test_error(err, "clCreateSampler failed");
334
335 max_img_width = (int)max_image2d_width;
336 max_img_height = (int)max_image2d_height;
337
338 if (max_mem_size > (cl_ulong)SIZE_MAX) {
339 max_mem_size = (cl_ulong)SIZE_MAX;
340 }
341
342 // determine max image dim we can allocate - assume RGBA image, 4 bytes per pixel,
343 // and we want to consume 1/4 of global memory (this is the minimum required to be
344 // supported by the spec)
345 max_mem_size /= 4; // use 1/4
346 max_mem_size /= 4; // 4 bytes per pixel
347 max_img_dim = (int)sqrt((double)max_mem_size);
348 // convert to a power of 2
349 {
350 unsigned int n = (unsigned int)max_img_dim;
351 unsigned int m = 0x80000000;
352
353 // round-down to the nearest power of 2
354 while (m > n)
355 m >>= 1;
356
357 max_img_dim = (int)m;
358 }
359
360 if (max_img_width > max_img_dim)
361 max_img_width = max_img_dim;
362 if (max_img_height > max_img_dim)
363 max_img_height = max_img_dim;
364
365 log_info("Adjusted maximum image size to test is %d x %d, which is a max mem size of %gMB.\n",
366 max_img_width, max_img_height, (max_img_width*max_img_height*4)/(1024.0*1024.0));
367
368 d = init_genrand( gRandomSeed );
369 input_ptr = generate_8888_image(max_img_width, max_img_height, d);
370 output_ptr = new unsigned char[4 * max_img_width * max_img_height];
371
372 int plus_minus;
373 for (plus_minus = 0; plus_minus < 3; plus_minus++)
374 {
375
376 // test power of 2 width, height starting at 1 to 4K
377 for (i=2,i2=1; i<=max_img_height; i<<=1,i2++)
378 {
379 img_height = (1 << i2);
380 for (j=2,j2=1; j<=max_img_width; j<<=1,j2++)
381 {
382 img_width = (1 << j2);
383
384 size_t effective_img_height = img_height;
385 size_t effective_img_width = img_width;
386
387 local_threads[0] = 1;
388 local_threads[1] = 1;
389
390 switch (plus_minus) {
391 case 0:
392 effective_img_height--;
393 local_threads[0] = work_group_size > max_local_workgroup_size[0] ? max_local_workgroup_size[0] : work_group_size;
394 while (img_width%local_threads[0] != 0)
395 local_threads[0]--;
396 break;
397 case 1:
398 effective_img_width--;
399 local_threads[1] = work_group_size > max_local_workgroup_size[1] ? max_local_workgroup_size[1] : work_group_size;
400 while (img_height%local_threads[1] != 0)
401 local_threads[1]--;
402 break;
403 case 2:
404 effective_img_width--;
405 effective_img_height--;
406 break;
407 default:
408 break;
409 }
410
411 img_format.image_channel_order = CL_RGBA;
412 img_format.image_channel_data_type = CL_UNORM_INT8;
413 streams[0] = create_image_2d(
414 context, CL_MEM_READ_WRITE, &img_format,
415 effective_img_width, effective_img_height, 0, NULL, NULL);
416 if (!streams[0])
417 {
418 log_error("create_image_2d failed. width = %d, height = %d\n", effective_img_width, effective_img_height);
419 delete[] input_ptr;
420 delete[] output_ptr;
421 free_mtdata(d);
422 return -1;
423 }
424 img_format.image_channel_order = CL_RGBA;
425 img_format.image_channel_data_type = CL_UNORM_INT8;
426 streams[1] = create_image_2d(
427 context, CL_MEM_READ_WRITE, &img_format,
428 effective_img_width, effective_img_height, 0, NULL, NULL);
429 if (!streams[1])
430 {
431 log_error("create_image_2d failed. width = %d, height = %d\n", effective_img_width, effective_img_height);
432 clReleaseMemObject(streams[0]);
433 delete[] input_ptr;
434 delete[] output_ptr;
435 free_mtdata(d);
436 return -1;
437 }
438
439 size_t origin[3] = {0,0,0};
440 size_t region[3] = {effective_img_width, effective_img_height, 1};
441 err = clEnqueueWriteImage(queue, streams[0], CL_FALSE, origin, region, 0, 0, input_ptr, 0, NULL, NULL);
442 if (err != CL_SUCCESS)
443 {
444 log_error("clWriteImage failed\n");
445 clReleaseMemObject(streams[0]);
446 clReleaseMemObject(streams[1]);
447 delete[] input_ptr;
448 delete[] output_ptr;
449 free_mtdata(d);
450 return -1;
451 }
452
453 err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]);
454 err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]);
455 err |= clSetKernelArg(kernel, 2, sizeof sampler, &sampler);
456 if (err != CL_SUCCESS)
457 {
458 log_error("clSetKernelArgs failed\n");
459 clReleaseMemObject(streams[0]);
460 clReleaseMemObject(streams[1]);
461 delete[] input_ptr;
462 delete[] output_ptr;
463 free_mtdata(d);
464 return -1;
465 }
466
467 threads[0] = (size_t)effective_img_width;
468 threads[1] = (size_t)effective_img_height;
469 log_info("Testing image dimensions %d x %d with local threads %d x %d.\n",
470 effective_img_width, effective_img_height, (int)local_threads[0], (int)local_threads[1]);
471 err = clEnqueueNDRangeKernel( queue, kernel, 2, NULL, threads, local_threads, 0, NULL, NULL );
472 if (err != CL_SUCCESS)
473 {
474 log_error("clEnqueueNDRangeKernel failed\n");
475 log_error("Image Dimension test failed. image width = %d, image height = %d, local %d x %d\n",
476 effective_img_width, effective_img_height, (int)local_threads[0], (int)local_threads[1]);
477 clReleaseMemObject(streams[0]);
478 clReleaseMemObject(streams[1]);
479 delete[] input_ptr;
480 delete[] output_ptr;
481 free_mtdata(d);
482 return -1;
483 }
484 err = clEnqueueReadImage(queue, streams[1], CL_TRUE, origin, region, 0, 0, output_ptr, 0, NULL, NULL);
485 if (err != CL_SUCCESS)
486 {
487 log_error("clReadImage failed\n");
488 log_error("Image Dimension test failed. image width = %d, image height = %d, local %d x %d\n",
489 effective_img_width, effective_img_height, (int)local_threads[0], (int)local_threads[1]);
490 clReleaseMemObject(streams[0]);
491 clReleaseMemObject(streams[1]);
492 delete[] input_ptr;
493 delete[] output_ptr;
494 free_mtdata(d);
495 return -1;
496 }
497 err = verify_8888_image(input_ptr, output_ptr, effective_img_width, effective_img_height);
498 if (err)
499 {
500 total_errors++;
501 log_error("Image Dimension test failed. image width = %d, image height = %d\n", effective_img_width, effective_img_height);
502 }
503
504 clReleaseMemObject(streams[0]);
505 clReleaseMemObject(streams[1]);
506 }
507 }
508
509 }
510
511 // cleanup
512 delete[] input_ptr;
513 delete[] output_ptr;
514 free_mtdata(d);
515 clReleaseSampler(sampler);
516 clReleaseKernel(kernel);
517 clReleaseProgram(program);
518
519 return total_errors;
520 }
521
522
523
524
525