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/testHarness.h"
18
19
20 const char *sample_single_kernel[] = {
21 "__kernel void sample_test(__global float *src, __global int *dst)\n"
22 "{\n"
23 " int tid = get_global_id(0);\n"
24 "\n"
25 " dst[tid] = (int)src[tid];\n"
26 "\n"
27 "}\n" };
28
29 size_t sample_single_kernel_lengths[1];
30
31 const char *sample_two_kernels[] = {
32 "__kernel void sample_test(__global float *src, __global int *dst)\n"
33 "{\n"
34 " int tid = get_global_id(0);\n"
35 "\n"
36 " dst[tid] = (int)src[tid];\n"
37 "\n"
38 "}\n",
39 "__kernel void sample_test2(__global int *src, __global float *dst)\n"
40 "{\n"
41 " int tid = get_global_id(0);\n"
42 "\n"
43 " dst[tid] = (float)src[tid];\n"
44 "\n"
45 "}\n" };
46
47 size_t sample_two_kernel_lengths[2];
48
49 const char *sample_two_kernels_in_1[] = {
50 "__kernel void sample_test(__global float *src, __global int *dst)\n"
51 "{\n"
52 " int tid = get_global_id(0);\n"
53 "\n"
54 " dst[tid] = (int)src[tid];\n"
55 "\n"
56 "}\n"
57 "__kernel void sample_test2(__global int *src, __global float *dst)\n"
58 "{\n"
59 " int tid = get_global_id(0);\n"
60 "\n"
61 " dst[tid] = (float)src[tid];\n"
62 "\n"
63 "}\n" };
64
65 size_t sample_two_kernels_in_1_lengths[1];
66
67
68 const char *repeate_test_kernel =
69 "__kernel void test_kernel(__global int *src, __global int *dst)\n"
70 "{\n"
71 " dst[get_global_id(0)] = src[get_global_id(0)]+1;\n"
72 "}\n";
73
74
75
test_load_single_kernel(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)76 int test_load_single_kernel(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
77 {
78 int error;
79 clProgramWrapper program;
80 cl_program testProgram;
81 clKernelWrapper kernel;
82 cl_context testContext;
83 unsigned int numKernels;
84 cl_char testName[512];
85 cl_uint testArgCount;
86 size_t realSize;
87
88
89 error = create_single_kernel_helper(context, &program, NULL, 1, sample_single_kernel, NULL);
90 test_error( error, "Unable to build test program" );
91
92 error = clCreateKernelsInProgram(program, 1, &kernel, &numKernels);
93 test_error( error, "Unable to create single kernel program" );
94
95 /* Check program and context pointers */
96 error = clGetKernelInfo( kernel, CL_KERNEL_PROGRAM, sizeof( cl_program ), &testProgram, &realSize );
97 test_error( error, "Unable to get kernel's program" );
98 if( (cl_program)testProgram != (cl_program)program )
99 {
100 log_error( "ERROR: Returned kernel's program does not match program used to create it! (Got %p, expected %p)\n", (cl_program)testProgram, (cl_program)program );
101 return -1;
102 }
103 if( realSize != sizeof( cl_program ) )
104 {
105 log_error( "ERROR: Returned size of kernel's program does not match expected size (expected %d, got %d)\n", (int)sizeof( cl_program ), (int)realSize );
106 return -1;
107 }
108
109 error = clGetKernelInfo( kernel, CL_KERNEL_CONTEXT, sizeof( cl_context ), &testContext, &realSize );
110 test_error( error, "Unable to get kernel's context" );
111 if( (cl_context)testContext != (cl_context)context )
112 {
113 log_error( "ERROR: Returned kernel's context does not match program used to create it! (Got %p, expected %p)\n", (cl_context)testContext, (cl_context)context );
114 return -1;
115 }
116 if( realSize != sizeof( cl_context ) )
117 {
118 log_error( "ERROR: Returned size of kernel's context does not match expected size (expected %d, got %d)\n", (int)sizeof( cl_context ), (int)realSize );
119 return -1;
120 }
121
122 /* Test arg count */
123 error = clGetKernelInfo( kernel, CL_KERNEL_NUM_ARGS, 0, NULL, &realSize );
124 test_error( error, "Unable to get size of arg count info from kernel" );
125
126 if( realSize != sizeof( testArgCount ) )
127 {
128 log_error( "ERROR: size of arg count not valid! %d\n", (int)realSize );
129 return -1;
130 }
131
132 error = clGetKernelInfo( kernel, CL_KERNEL_NUM_ARGS, sizeof( testArgCount ), &testArgCount, NULL );
133 test_error( error, "Unable to get arg count from kernel" );
134
135 if( testArgCount != 2 )
136 {
137 log_error( "ERROR: Kernel arg count does not match!\n" );
138 return -1;
139 }
140
141
142 /* Test function name */
143 error = clGetKernelInfo( kernel, CL_KERNEL_FUNCTION_NAME, sizeof( testName ), testName, &realSize );
144 test_error( error, "Unable to get name from kernel" );
145
146 if( strcmp( (char *)testName, "sample_test" ) != 0 )
147 {
148 log_error( "ERROR: Kernel names do not match!\n" );
149 return -1;
150 }
151 if( realSize != strlen( (char *)testName ) + 1 )
152 {
153 log_error( "ERROR: Length of kernel name returned does not validate (expected %d, got %d)\n", (int)strlen( (char *)testName ) + 1, (int)realSize );
154 return -1;
155 }
156
157 /* All done */
158
159 return 0;
160 }
161
test_load_two_kernels(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)162 int test_load_two_kernels(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
163 {
164 int error;
165 clProgramWrapper program;
166 clKernelWrapper kernel[2];
167 unsigned int numKernels;
168 cl_char testName[ 512 ];
169 cl_uint testArgCount;
170
171
172 error = create_single_kernel_helper(context, &program, NULL, 2, sample_two_kernels, NULL);
173 test_error( error, "Unable to build test program" );
174
175 error = clCreateKernelsInProgram(program, 2, &kernel[0], &numKernels);
176 test_error( error, "Unable to create dual kernel program" );
177
178 if( numKernels != 2 )
179 {
180 log_error( "ERROR: wrong # of kernels! (%d)\n", numKernels );
181 return -1;
182 }
183
184 /* Check first kernel */
185 error = clGetKernelInfo( kernel[0], CL_KERNEL_FUNCTION_NAME, sizeof( testName ), testName, NULL );
186 test_error( error, "Unable to get function name from kernel" );
187
188 int found_kernel1 = 0, found_kernel2 = 0;
189
190 if( strcmp( (char *)testName, "sample_test" ) == 0 ) {
191 found_kernel1 = 1;
192 } else if( strcmp( (char *)testName, "sample_test2" ) == 0 ) {
193 found_kernel2 = 1;
194 } else {
195 log_error( "ERROR: Invalid kernel name returned: \"%s\" expected \"%s\" or \"%s\".\n", testName, "sample_test", "sample_test2");
196 return -1;
197 }
198
199 error = clGetKernelInfo( kernel[1], CL_KERNEL_FUNCTION_NAME, sizeof( testName ), testName, NULL );
200 test_error( error, "Unable to get function name from second kernel" );
201
202 if( strcmp( (char *)testName, "sample_test" ) == 0 ) {
203 if (found_kernel1) {
204 log_error("Kernel \"%s\" returned twice.\n", (char *)testName);
205 return -1;
206 }
207 found_kernel1 = 1;
208 } else if( strcmp( (char *)testName, "sample_test2" ) == 0 ) {
209 if (found_kernel2) {
210 log_error("Kernel \"%s\" returned twice.\n", (char *)testName);
211 return -1;
212 }
213 found_kernel2 = 1;
214 } else {
215 log_error( "ERROR: Invalid kernel name returned: \"%s\" expected \"%s\" or \"%s\".\n", testName, "sample_test", "sample_test2");
216 return -1;
217 }
218
219 if( !found_kernel1 || !found_kernel2 )
220 {
221 log_error( "ERROR: Kernel names do not match.\n" );
222 if (!found_kernel1)
223 log_error("Kernel \"%s\" not returned.\n", "sample_test");
224 if (!found_kernel2)
225 log_error("Kernel \"%s\" not returned.\n", "sample_test");
226 return -1;
227 }
228
229 error = clGetKernelInfo( kernel[0], CL_KERNEL_NUM_ARGS, sizeof( testArgCount ), &testArgCount, NULL );
230 test_error( error, "Unable to get arg count from kernel" );
231
232 if( testArgCount != 2 )
233 {
234 log_error( "ERROR: wrong # of args for kernel\n" );
235 return -1;
236 }
237
238 /* All done */
239 return 0;
240 }
241
test_load_two_kernels_in_one(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)242 int test_load_two_kernels_in_one(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
243 {
244 int error;
245 clProgramWrapper program;
246 clKernelWrapper kernel[2];
247 unsigned int numKernels;
248 cl_char testName[512];
249 cl_uint testArgCount;
250
251
252 error = create_single_kernel_helper(context, &program, NULL, 1, sample_two_kernels_in_1, NULL);
253 test_error( error, "Unable to build test program" );
254
255 error = clCreateKernelsInProgram(program, 2, &kernel[0], &numKernels);
256 test_error( error, "Unable to create dual kernel program" );
257
258 if( numKernels != 2 )
259 {
260 log_error( "ERROR: wrong # of kernels! (%d)\n", numKernels );
261 return -1;
262 }
263
264 /* Check first kernel */
265 error = clGetKernelInfo( kernel[0], CL_KERNEL_FUNCTION_NAME, sizeof( testName ), testName, NULL );
266 test_error( error, "Unable to get function name from kernel" );
267
268 int found_kernel1 = 0, found_kernel2 = 0;
269
270 if( strcmp( (char *)testName, "sample_test" ) == 0 ) {
271 found_kernel1 = 1;
272 } else if( strcmp( (char *)testName, "sample_test2" ) == 0 ) {
273 found_kernel2 = 1;
274 } else {
275 log_error( "ERROR: Invalid kernel name returned: \"%s\" expected \"%s\" or \"%s\".\n", testName, "sample_test", "sample_test2");
276 return -1;
277 }
278
279 error = clGetKernelInfo( kernel[0], CL_KERNEL_NUM_ARGS, sizeof( testArgCount ), &testArgCount, NULL );
280 test_error( error, "Unable to get arg count from kernel" );
281
282 if( testArgCount != 2 )
283 {
284 log_error( "ERROR: wrong # of args for kernel\n" );
285 return -1;
286 }
287
288 /* Check second kernel */
289 error = clGetKernelInfo( kernel[1], CL_KERNEL_FUNCTION_NAME, sizeof( testName ), testName, NULL );
290 test_error( error, "Unable to get function name from kernel" );
291
292 if( strcmp( (char *)testName, "sample_test" ) == 0 ) {
293 if (found_kernel1) {
294 log_error("Kernel \"%s\" returned twice.\n", (char *)testName);
295 return -1;
296 }
297 found_kernel1 = 1;
298 } else if( strcmp( (char *)testName, "sample_test2" ) == 0 ) {
299 if (found_kernel2) {
300 log_error("Kernel \"%s\" returned twice.\n", (char *)testName);
301 return -1;
302 }
303 found_kernel2 = 1;
304 } else {
305 log_error( "ERROR: Invalid kernel name returned: \"%s\" expected \"%s\" or \"%s\".\n", testName, "sample_test", "sample_test2");
306 return -1;
307 }
308
309 if( !found_kernel1 || !found_kernel2 )
310 {
311 log_error( "ERROR: Kernel names do not match.\n" );
312 if (!found_kernel1)
313 log_error("Kernel \"%s\" not returned.\n", "sample_test");
314 if (!found_kernel2)
315 log_error("Kernel \"%s\" not returned.\n", "sample_test");
316 return -1;
317 }
318
319 /* All done */
320 return 0;
321 }
322
test_load_two_kernels_manually(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)323 int test_load_two_kernels_manually( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
324 {
325 clProgramWrapper program;
326 clKernelWrapper kernel1, kernel2;
327 int error;
328
329
330 /* Now create a test program */
331 error = create_single_kernel_helper(context, &program, NULL, 1, sample_two_kernels_in_1, NULL);
332 test_error( error, "Unable to build test program" );
333
334 /* Try manually creating kernels (backwards just in case) */
335 kernel1 = clCreateKernel( program, "sample_test2", &error );
336
337 if( kernel1 == NULL || error != CL_SUCCESS )
338 {
339 print_error( error, "Could not get kernel 1" );
340 return -1;
341 }
342
343 kernel2 = clCreateKernel( program, "sample_test", &error );
344
345 if( kernel2 == NULL )
346 {
347 print_error( error, "Could not get kernel 2" );
348 return -1;
349 }
350
351 return 0;
352 }
353
test_get_program_info_kernel_names(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)354 int test_get_program_info_kernel_names( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
355 {
356 clProgramWrapper program;
357 clKernelWrapper kernel1, kernel2;
358 int error;
359 size_t i;
360
361 /* Now create a test program */
362 error = create_single_kernel_helper(context, &program, NULL, 1, sample_two_kernels_in_1, NULL);
363 test_error( error, "Unable to build test program" );
364
365 /* Lookup the number of kernels in the program. */
366 size_t total_kernels = 0;
367 error = clGetProgramInfo(program, CL_PROGRAM_NUM_KERNELS, sizeof(size_t),&total_kernels,NULL);
368 test_error( error, "Unable to get program info num kernels");
369
370 if (total_kernels != 2)
371 {
372 print_error( error, "Program did not contain two kernels" );
373 return -1;
374 }
375
376 /* Lookup the kernel names. */
377 const char* actual_names[] = { "sample_test;sample_test2", "sample_test2;sample_test"} ;
378
379 size_t kernel_names_len = 0;
380 error = clGetProgramInfo(program,CL_PROGRAM_KERNEL_NAMES,0,NULL,&kernel_names_len);
381 test_error( error, "Unable to get length of kernel names list." );
382
383 if (kernel_names_len != (strlen(actual_names[0])+1))
384 {
385 print_error( error, "Kernel names length did not match");
386 return -1;
387 }
388
389 const size_t len = (kernel_names_len+1)*sizeof(char);
390 char* kernel_names = (char*)malloc(len);
391 error = clGetProgramInfo(program,CL_PROGRAM_KERNEL_NAMES,len,kernel_names,&kernel_names_len);
392 test_error( error, "Unable to get kernel names list." );
393
394 /* Check to see if the kernel name array is null terminated. */
395 if (kernel_names[kernel_names_len-1] != '\0')
396 {
397 free(kernel_names);
398 print_error( error, "Kernel name list was not null terminated");
399 return -1;
400 }
401
402 /* Check to see if the correct kernel name string was returned. */
403 for( i = 0; i < sizeof( actual_names ) / sizeof( actual_names[0] ); i++ )
404 if( 0 == strcmp(actual_names[i],kernel_names) )
405 break;
406
407 if (i == sizeof( actual_names ) / sizeof( actual_names[0] ) )
408 {
409 free(kernel_names);
410 log_error( "Kernel names \"%s\" did not match:\n", kernel_names );
411 for( i = 0; i < sizeof( actual_names ) / sizeof( actual_names[0] ); i++ )
412 log_error( "\t\t\"%s\"\n", actual_names[0] );
413 return -1;
414 }
415 free(kernel_names);
416
417 /* Try manually creating kernels (backwards just in case) */
418 kernel1 = clCreateKernel( program, "sample_test", &error );
419 if( kernel1 == NULL || error != CL_SUCCESS )
420 {
421 print_error( error, "Could not get kernel 1" );
422 return -1;
423 }
424
425 kernel2 = clCreateKernel( program, "sample_test2", &error );
426 if( kernel2 == NULL )
427 {
428 print_error( error, "Could not get kernel 2" );
429 return -1;
430 }
431
432 return 0;
433 }
434
435 static const char *single_task_kernel[] = {
436 "__kernel void sample_test(__global int *dst, int count)\n"
437 "{\n"
438 " int tid = get_global_id(0);\n"
439 "\n"
440 " for( int i = 0; i < count; i++ )\n"
441 " dst[i] = tid + i;\n"
442 "\n"
443 "}\n" };
444
test_enqueue_task(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)445 int test_enqueue_task(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
446 {
447 int error;
448 clProgramWrapper program;
449 clKernelWrapper kernel;
450 clMemWrapper output;
451 cl_int count;
452
453
454 if( create_single_kernel_helper( context, &program, &kernel, 1, single_task_kernel, "sample_test" ) )
455 return -1;
456
457 // Create args
458 count = 100;
459 output = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_int) * count,
460 NULL, &error);
461 test_error( error, "Unable to create output buffer" );
462
463 error = clSetKernelArg( kernel, 0, sizeof( cl_mem ), &output );
464 test_error( error, "Unable to set kernel argument" );
465 error = clSetKernelArg( kernel, 1, sizeof( cl_int ), &count );
466 test_error( error, "Unable to set kernel argument" );
467
468 // Run task
469 error = clEnqueueTask( queue, kernel, 0, NULL, NULL );
470 test_error( error, "Unable to run task" );
471
472 // Read results
473 cl_int *results = (cl_int*)malloc(sizeof(cl_int)*count);
474 error = clEnqueueReadBuffer( queue, output, CL_TRUE, 0, sizeof( cl_int ) * count, results, 0, NULL, NULL );
475 test_error( error, "Unable to read results" );
476
477 // Validate
478 for( cl_int i = 0; i < count; i++ )
479 {
480 if( results[ i ] != i )
481 {
482 log_error( "ERROR: Task result value %d did not validate! Expected %d, got %d\n", (int)i, (int)i, (int)results[ i ] );
483 free(results);
484 return -1;
485 }
486 }
487
488 /* All done */
489 free(results);
490 return 0;
491 }
492
493
494
495 #define TEST_SIZE 1000
test_repeated_setup_cleanup(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)496 int test_repeated_setup_cleanup(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
497 {
498
499 cl_context local_context;
500 cl_command_queue local_queue;
501 cl_program local_program;
502 cl_kernel local_kernel;
503 cl_mem local_mem_in, local_mem_out;
504 cl_event local_event;
505 size_t global_dim[3];
506 int i, j, error;
507 global_dim[0] = TEST_SIZE;
508 global_dim[1] = 1; global_dim[2] = 1;
509 cl_int *inData, *outData;
510 cl_int status;
511
512 inData = (cl_int*)malloc(sizeof(cl_int)*TEST_SIZE);
513 outData = (cl_int*)malloc(sizeof(cl_int)*TEST_SIZE);
514 for (i=0; i<TEST_SIZE; i++) {
515 inData[i] = i;
516 }
517
518
519 for (i=0; i<100; i++) {
520 memset(outData, 0, sizeof(cl_int)*TEST_SIZE);
521
522 local_context = clCreateContext(NULL, 1, &deviceID, notify_callback, NULL, &error);
523 test_error( error, "clCreateContext failed");
524
525 local_queue = clCreateCommandQueue(local_context, deviceID, 0, &error);
526 test_error( error, "clCreateCommandQueue failed");
527
528 error = create_single_kernel_helper(
529 local_context, &local_program, &local_kernel, 1,
530 &repeate_test_kernel, "test_kernel");
531 test_error(error, "Unable to create kernel");
532
533 local_mem_in = clCreateBuffer(local_context, CL_MEM_READ_ONLY, TEST_SIZE*sizeof(cl_int), NULL, &error);
534 test_error( error, "clCreateBuffer failed");
535
536 local_mem_out = clCreateBuffer(local_context, CL_MEM_WRITE_ONLY, TEST_SIZE*sizeof(cl_int), NULL, &error);
537 test_error( error, "clCreateBuffer failed");
538
539 error = clEnqueueWriteBuffer(local_queue, local_mem_in, CL_TRUE, 0, TEST_SIZE*sizeof(cl_int), inData, 0, NULL, NULL);
540 test_error( error, "clEnqueueWriteBuffer failed");
541
542 error = clEnqueueWriteBuffer(local_queue, local_mem_out, CL_TRUE, 0, TEST_SIZE*sizeof(cl_int), outData, 0, NULL, NULL);
543 test_error( error, "clEnqueueWriteBuffer failed");
544
545 error = clSetKernelArg(local_kernel, 0, sizeof(local_mem_in), &local_mem_in);
546 test_error( error, "clSetKernelArg failed");
547
548 error = clSetKernelArg(local_kernel, 1, sizeof(local_mem_out), &local_mem_out);
549 test_error( error, "clSetKernelArg failed");
550
551 error = clEnqueueNDRangeKernel(local_queue, local_kernel, 1, NULL, global_dim, NULL, 0, NULL, &local_event);
552 test_error( error, "clEnqueueNDRangeKernel failed");
553
554 error = clWaitForEvents(1, &local_event);
555 test_error( error, "clWaitForEvents failed");
556
557 error = clGetEventInfo(local_event, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(status), &status, NULL);
558 test_error( error, "clGetEventInfo failed");
559
560 if (status != CL_COMPLETE) {
561 log_error( "Kernel execution not complete: status %d.\n", status);
562 free(inData);
563 free(outData);
564 return -1;
565 }
566
567 error = clEnqueueReadBuffer(local_queue, local_mem_out, CL_TRUE, 0, TEST_SIZE*sizeof(cl_int), outData, 0, NULL, NULL);
568 test_error( error, "clEnqueueReadBuffer failed");
569
570 clReleaseEvent(local_event);
571 clReleaseMemObject(local_mem_in);
572 clReleaseMemObject(local_mem_out);
573 clReleaseKernel(local_kernel);
574 clReleaseProgram(local_program);
575 clReleaseCommandQueue(local_queue);
576 clReleaseContext(local_context);
577
578 for (j=0; j<TEST_SIZE; j++) {
579 if (outData[j] != inData[j] + 1) {
580 log_error("Results failed to validate at iteration %d. %d != %d.\n", i, outData[j], inData[j] + 1);
581 free(inData);
582 free(outData);
583 return -1;
584 }
585 }
586 }
587
588 free(inData);
589 free(outData);
590
591 return 0;
592 }
593
594
595
596