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
18 // This test is designed to stress changing kernel arguments between execute calls (that are asynchronous and thus
19 // potentially overlapping) to make sure each kernel gets the right arguments
20
21 // Note: put a delay loop in the kernel to make sure we have time to queue the next kernel before this one finishes
22 const char *inspect_image_kernel_source[] = {
23 "__kernel void sample_test(read_only image2d_t src, __global int *outDimensions )\n"
24 "{\n"
25 " int tid = get_global_id(0), i;\n"
26 " for( i = 0; i < 100000; i++ ); \n"
27 " outDimensions[tid * 2] = get_image_width(src) * tid;\n"
28 " outDimensions[tid * 2 + 1] = get_image_height(src) * tid;\n"
29 "\n"
30 "}\n" };
31
32 #define NUM_TRIES 100
33 #define NUM_THREADS 2048
34
test_kernel_arg_changes(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)35 int test_kernel_arg_changes(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
36 {
37 clProgramWrapper program;
38 clKernelWrapper kernel;
39 int error, i;
40 clMemWrapper images[ NUM_TRIES ];
41 size_t sizes[ NUM_TRIES ][ 2 ];
42 clMemWrapper results[ NUM_TRIES ];
43 cl_image_format imageFormat;
44 size_t maxWidth, maxHeight;
45 size_t threads[1], localThreads[1];
46 cl_int resultArray[ NUM_THREADS * 2 ];
47 char errStr[ 128 ];
48 RandomSeed seed( gRandomSeed );
49
50
51 PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
52
53 // Just get any ol format to test with
54 error = get_8_bit_image_format( context, CL_MEM_OBJECT_IMAGE2D, CL_MEM_READ_WRITE, 0, &imageFormat );
55 test_error( error, "Unable to obtain suitable image format to test with!" );
56
57 // Create our testing kernel
58 error = create_single_kernel_helper( context, &program, &kernel, 1, inspect_image_kernel_source, "sample_test" );
59 test_error( error, "Unable to create testing kernel" );
60
61 // Get max dimensions for each of our images
62 error = clGetDeviceInfo( device, CL_DEVICE_IMAGE2D_MAX_WIDTH, sizeof( maxWidth ), &maxWidth, NULL );
63 error |= clGetDeviceInfo( device, CL_DEVICE_IMAGE2D_MAX_HEIGHT, sizeof( maxHeight ), &maxHeight, NULL );
64 test_error( error, "Unable to get max image dimensions for device" );
65
66 // Get the number of threads we'll be able to run
67 threads[0] = NUM_THREADS;
68 error = get_max_common_work_group_size( context, kernel, threads[0], &localThreads[0] );
69 test_error( error, "Unable to get work group size for kernel" );
70
71 // Create a variety of images and output arrays
72 for( i = 0; i < NUM_TRIES; i++ )
73 {
74 sizes[ i ][ 0 ] = genrand_int32(seed) % (maxWidth/32) + 1;
75 sizes[ i ][ 1 ] = genrand_int32(seed) % (maxHeight/32) + 1;
76
77 images[i] = create_image_2d(context, CL_MEM_READ_ONLY, &imageFormat,
78 sizes[i][0], sizes[i][1], 0, NULL, &error);
79 if( images[i] == NULL )
80 {
81 log_error("Failed to create image %d of size %d x %d (%s).\n", i, (int)sizes[i][0], (int)sizes[i][1], IGetErrorString( error ));
82 return -1;
83 }
84 results[i] =
85 clCreateBuffer(context, CL_MEM_READ_WRITE,
86 sizeof(cl_int) * threads[0] * 2, NULL, &error);
87 if( results[i] == NULL)
88 {
89 log_error("Failed to create array %d of size %d.\n", i, (int)threads[0]*2);
90 return -1;
91 }
92 }
93
94 // Start setting arguments and executing kernels
95 for( i = 0; i < NUM_TRIES; i++ )
96 {
97 // Set the arguments for this try
98 error = clSetKernelArg( kernel, 0, sizeof( cl_mem ), &images[ i ] );
99 sprintf( errStr, "Unable to set argument 0 for kernel try %d", i );
100 test_error( error, errStr );
101
102 error = clSetKernelArg( kernel, 1, sizeof( cl_mem ), &results[ i ] );
103 sprintf( errStr, "Unable to set argument 1 for kernel try %d", i );
104 test_error( error, errStr );
105
106 // Queue up execution
107 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL );
108 sprintf( errStr, "Unable to execute kernel try %d", i );
109 test_error( error, errStr );
110 }
111
112 // Read the results back out, one at a time, and verify
113 for( i = 0; i < NUM_TRIES; i++ )
114 {
115 error = clEnqueueReadBuffer( queue, results[ i ], CL_TRUE, 0, sizeof( cl_int ) * threads[0] * 2, resultArray, 0, NULL, NULL );
116 sprintf( errStr, "Unable to read results for kernel try %d", i );
117 test_error( error, errStr );
118
119 // Verify. Each entry should be n * the (width/height) of image i
120 for( int j = 0; j < NUM_THREADS; j++ )
121 {
122 if( resultArray[ j * 2 + 0 ] != (int)sizes[ i ][ 0 ] * j )
123 {
124 log_error( "ERROR: Verficiation for kernel try %d, sample %d FAILED, expected a width of %d, got %d\n",
125 i, j, (int)sizes[ i ][ 0 ] * j, resultArray[ j * 2 + 0 ] );
126 return -1;
127 }
128 if( resultArray[ j * 2 + 1 ] != (int)sizes[ i ][ 1 ] * j )
129 {
130 log_error( "ERROR: Verficiation for kernel try %d, sample %d FAILED, expected a height of %d, got %d\n",
131 i, j, (int)sizes[ i ][ 1 ] * j, resultArray[ j * 2 + 1 ] );
132 return -1;
133 }
134 }
135 }
136
137 // If we got here, everything verified successfully
138 return 0;
139 }
140
141
142