| // |
| // Copyright (c) 2017 The Khronos Group Inc. |
| // |
| // Licensed under the Apache License, Version 2.0 (the "License"); |
| // you may not use this file except in compliance with the License. |
| // You may obtain a copy of the License at |
| // |
| // http://www.apache.org/licenses/LICENSE-2.0 |
| // |
| // Unless required by applicable law or agreed to in writing, software |
| // distributed under the License is distributed on an "AS IS" BASIS, |
| // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. |
| // See the License for the specific language governing permissions and |
| // limitations under the License. |
| // |
| #include "testBase.h" |
| |
| // This test is designed to stress changing kernel arguments between execute calls (that are asynchronous and thus |
| // potentially overlapping) to make sure each kernel gets the right arguments |
| |
| // Note: put a delay loop in the kernel to make sure we have time to queue the next kernel before this one finishes |
| const char *inspect_image_kernel_source[] = { |
| "__kernel void sample_test(read_only image2d_t src, __global int *outDimensions )\n" |
| "{\n" |
| " int tid = get_global_id(0), i;\n" |
| " for( i = 0; i < 100000; i++ ); \n" |
| " outDimensions[tid * 2] = get_image_width(src) * tid;\n" |
| " outDimensions[tid * 2 + 1] = get_image_height(src) * tid;\n" |
| "\n" |
| "}\n" }; |
| |
| #define NUM_TRIES 100 |
| #define NUM_THREADS 2048 |
| |
| int test_kernel_arg_changes(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| clProgramWrapper program; |
| clKernelWrapper kernel; |
| int error, i; |
| clMemWrapper images[ NUM_TRIES ]; |
| size_t sizes[ NUM_TRIES ][ 2 ]; |
| clMemWrapper results[ NUM_TRIES ]; |
| cl_image_format imageFormat; |
| size_t maxWidth, maxHeight; |
| size_t threads[1], localThreads[1]; |
| cl_int resultArray[ NUM_THREADS * 2 ]; |
| char errStr[ 128 ]; |
| RandomSeed seed( gRandomSeed ); |
| |
| |
| PASSIVE_REQUIRE_IMAGE_SUPPORT( device ) |
| |
| // Just get any ol format to test with |
| error = get_8_bit_image_format( context, CL_MEM_OBJECT_IMAGE2D, CL_MEM_READ_WRITE, 0, &imageFormat ); |
| test_error( error, "Unable to obtain suitable image format to test with!" ); |
| |
| // Create our testing kernel |
| error = create_single_kernel_helper( context, &program, &kernel, 1, inspect_image_kernel_source, "sample_test" ); |
| test_error( error, "Unable to create testing kernel" ); |
| |
| // Get max dimensions for each of our images |
| error = clGetDeviceInfo( device, CL_DEVICE_IMAGE2D_MAX_WIDTH, sizeof( maxWidth ), &maxWidth, NULL ); |
| error |= clGetDeviceInfo( device, CL_DEVICE_IMAGE2D_MAX_HEIGHT, sizeof( maxHeight ), &maxHeight, NULL ); |
| test_error( error, "Unable to get max image dimensions for device" ); |
| |
| // Get the number of threads we'll be able to run |
| threads[0] = NUM_THREADS; |
| error = get_max_common_work_group_size( context, kernel, threads[0], &localThreads[0] ); |
| test_error( error, "Unable to get work group size for kernel" ); |
| |
| // Create a variety of images and output arrays |
| for( i = 0; i < NUM_TRIES; i++ ) |
| { |
| sizes[ i ][ 0 ] = genrand_int32(seed) % (maxWidth/32) + 1; |
| sizes[ i ][ 1 ] = genrand_int32(seed) % (maxHeight/32) + 1; |
| |
| images[ i ] = create_image_2d( context, (cl_mem_flags)(CL_MEM_READ_ONLY), |
| &imageFormat, sizes[ i ][ 0], sizes[ i ][ 1 ], 0, NULL, &error ); |
| if( images[i] == NULL ) |
| { |
| 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 )); |
| return -1; |
| } |
| results[ i ] = clCreateBuffer( context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof( cl_int ) * threads[0] * 2, NULL, &error ); |
| if( results[i] == NULL) |
| { |
| log_error("Failed to create array %d of size %d.\n", i, (int)threads[0]*2); |
| return -1; |
| } |
| } |
| |
| // Start setting arguments and executing kernels |
| for( i = 0; i < NUM_TRIES; i++ ) |
| { |
| // Set the arguments for this try |
| error = clSetKernelArg( kernel, 0, sizeof( cl_mem ), &images[ i ] ); |
| sprintf( errStr, "Unable to set argument 0 for kernel try %d", i ); |
| test_error( error, errStr ); |
| |
| error = clSetKernelArg( kernel, 1, sizeof( cl_mem ), &results[ i ] ); |
| sprintf( errStr, "Unable to set argument 1 for kernel try %d", i ); |
| test_error( error, errStr ); |
| |
| // Queue up execution |
| error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL ); |
| sprintf( errStr, "Unable to execute kernel try %d", i ); |
| test_error( error, errStr ); |
| } |
| |
| // Read the results back out, one at a time, and verify |
| for( i = 0; i < NUM_TRIES; i++ ) |
| { |
| error = clEnqueueReadBuffer( queue, results[ i ], CL_TRUE, 0, sizeof( cl_int ) * threads[0] * 2, resultArray, 0, NULL, NULL ); |
| sprintf( errStr, "Unable to read results for kernel try %d", i ); |
| test_error( error, errStr ); |
| |
| // Verify. Each entry should be n * the (width/height) of image i |
| for( int j = 0; j < NUM_THREADS; j++ ) |
| { |
| if( resultArray[ j * 2 + 0 ] != (int)sizes[ i ][ 0 ] * j ) |
| { |
| log_error( "ERROR: Verficiation for kernel try %d, sample %d FAILED, expected a width of %d, got %d\n", |
| i, j, (int)sizes[ i ][ 0 ] * j, resultArray[ j * 2 + 0 ] ); |
| return -1; |
| } |
| if( resultArray[ j * 2 + 1 ] != (int)sizes[ i ][ 1 ] * j ) |
| { |
| log_error( "ERROR: Verficiation for kernel try %d, sample %d FAILED, expected a height of %d, got %d\n", |
| i, j, (int)sizes[ i ][ 1 ] * j, resultArray[ j * 2 + 1 ] ); |
| return -1; |
| } |
| } |
| } |
| |
| // If we got here, everything verified successfully |
| return 0; |
| } |
| |
| |