| // |
| // 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" |
| #include "harness/typeWrappers.h" |
| #include "harness/testHarness.h" |
| #include "harness/conversions.h" |
| |
| const char *test_kernels[] = { "__kernel void kernelA(__global uint *dst)\n" |
| "{\n" |
| "\n" |
| " dst[get_global_id(0)]*=3;\n" |
| "\n" |
| "}\n" |
| "__kernel void kernelB(__global uint *dst)\n" |
| "{\n" |
| "\n" |
| " dst[get_global_id(0)]++;\n" |
| "\n" |
| "}\n" }; |
| |
| #define TEST_SIZE 512 |
| #define MAX_DEVICES 32 |
| #define MAX_QUEUES 1000 |
| |
| int test_device_set(size_t deviceCount, size_t queueCount, cl_device_id *devices, int num_elements) |
| { |
| int error; |
| clContextWrapper context; |
| clProgramWrapper program; |
| clKernelWrapper kernels[2]; |
| clMemWrapper stream; |
| clCommandQueueWrapper queues[MAX_QUEUES]; |
| size_t threads[1], localThreads[1]; |
| cl_uint data[TEST_SIZE]; |
| cl_uint outputData[TEST_SIZE]; |
| cl_uint expectedResults[TEST_SIZE]; |
| cl_uint expectedResultsOneDevice[MAX_DEVICES][TEST_SIZE]; |
| size_t i; |
| |
| memset(queues, 0, sizeof(queues)); |
| |
| RandomSeed seed( gRandomSeed ); |
| |
| if (deviceCount > MAX_DEVICES) { |
| log_error("Number of devices in set (%ld) is greater than the number for which the test was written (%d).", deviceCount, MAX_DEVICES); |
| return -1; |
| } |
| |
| if (queueCount > MAX_QUEUES) { |
| log_error("Number of queues (%ld) is greater than the number for which the test was written (%d).", queueCount, MAX_QUEUES); |
| return -1; |
| } |
| |
| log_info("Testing with %ld queues on %ld devices, %ld kernel executions.\n", queueCount, deviceCount, queueCount*num_elements/TEST_SIZE); |
| |
| for (i=0; i<deviceCount; i++) { |
| char deviceName[4096] = ""; |
| error = clGetDeviceInfo(devices[i], CL_DEVICE_NAME, sizeof(deviceName), deviceName, NULL); |
| test_error(error, "clGetDeviceInfo CL_DEVICE_NAME failed"); |
| log_info("Device %ld is \"%s\".\n", i, deviceName); |
| } |
| |
| /* Create a context */ |
| context = clCreateContext( NULL, (cl_uint)deviceCount, devices, notify_callback, NULL, &error ); |
| test_error( error, "Unable to create testing context" ); |
| |
| /* Create our kernels (they all have the same arguments so we don't need multiple ones for each device) */ |
| if( create_single_kernel_helper( context, &program, &kernels[0], 1, test_kernels, "kernelA" ) != 0 ) |
| { |
| return -1; |
| } |
| |
| kernels[1] = clCreateKernel(program, "kernelB", &error); |
| test_error(error, "clCreateKernel failed"); |
| |
| |
| /* Now create I/O streams */ |
| for( i = 0; i < TEST_SIZE; i++ ) |
| data[i] = genrand_int32(seed); |
| |
| stream = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, |
| sizeof(cl_uint) * TEST_SIZE, data, &error); |
| test_error(error, "Unable to create test array"); |
| |
| // Update the expected results |
| for( i = 0; i < TEST_SIZE; i++ ) { |
| expectedResults[i] = data[i]; |
| for (size_t j=0; j<deviceCount; j++) |
| expectedResultsOneDevice[j][i] = data[i]; |
| } |
| |
| |
| // Set the arguments |
| error = clSetKernelArg( kernels[0], 0, sizeof( stream ), &stream); |
| test_error( error, "Unable to set kernel arguments" ); |
| error = clSetKernelArg( kernels[1], 0, sizeof( stream ), &stream); |
| test_error( error, "Unable to set kernel arguments" ); |
| |
| /* Run the test */ |
| threads[0] = (size_t)TEST_SIZE; |
| |
| error = get_max_common_work_group_size( context, kernels[0], threads[0], &localThreads[ 0 ] ); |
| test_error( error, "Unable to calc work group size" ); |
| |
| /* Create work queues */ |
| for( i = 0; i < queueCount; i++ ) |
| { |
| queues[i] = clCreateCommandQueue( context, devices[ i % deviceCount ], 0, &error ); |
| if (error != CL_SUCCESS || queues[i] == NULL) { |
| log_info("Could not create queue[%d].\n", (int)i); |
| queueCount = i; |
| break; |
| } |
| } |
| log_info("Testing with %d queues.\n", (int)queueCount); |
| |
| /* Enqueue executions */ |
| for( int z = 0; z<num_elements/TEST_SIZE; z++) { |
| for( i = 0; i < queueCount; i++ ) |
| { |
| // Randomly choose a kernel to execute. |
| int kernel_selection = (int)get_random_float(0, 2, seed); |
| error = clEnqueueNDRangeKernel( queues[ i ], kernels[ kernel_selection ], 1, NULL, threads, localThreads, 0, NULL, NULL ); |
| test_error( error, "Kernel execution failed" ); |
| |
| // Update the expected results |
| for( int j = 0; j < TEST_SIZE; j++ ) { |
| expectedResults[j] = (kernel_selection) ? expectedResults[j]+1 : expectedResults[j]*3; |
| expectedResultsOneDevice[i % deviceCount][j] = (kernel_selection) ? expectedResultsOneDevice[i % deviceCount][j]+1 : expectedResultsOneDevice[i % deviceCount][j]*3; |
| } |
| |
| // Force the queue to finish so the next one will be in sync |
| error = clFinish(queues[i]); |
| test_error( error, "clFinish failed"); |
| } |
| } |
| |
| /* Read results */ |
| int errors = 0; |
| for (int q = 0; q<(int)queueCount; q++) { |
| error = clEnqueueReadBuffer( queues[ 0 ], stream, CL_TRUE, 0, sizeof(cl_int)*TEST_SIZE, (char *)outputData, 0, NULL, NULL ); |
| test_error( error, "Unable to get result data set" ); |
| |
| int errorsThisTime = 0; |
| /* Verify all of the data now */ |
| for( i = 0; i < TEST_SIZE; i++ ) |
| { |
| if( expectedResults[ i ] != outputData[ i ] ) |
| { |
| log_error( "ERROR: Sample data did not verify for queue %d on device %ld (sample %d, expected %d, got %d)\n", |
| q, q % deviceCount, (int)i, expectedResults[ i ], outputData[ i ] ); |
| for (size_t j=0; j<deviceCount; j++) { |
| if (expectedResultsOneDevice[j][i] == outputData[i]) |
| log_info("Sample consistent with only device %ld having modified the data.\n", j); |
| } |
| errorsThisTime++; |
| break; |
| } |
| } |
| if (errorsThisTime) |
| errors++; |
| } |
| |
| /* All done now! */ |
| if (errors) |
| return -1; |
| return 0; |
| } |
| |
| int test_two_devices(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| cl_platform_id platform; |
| cl_device_id devices[2]; |
| int err; |
| cl_uint numDevices; |
| |
| err = clGetPlatformIDs(1, &platform, NULL); |
| test_error( err, "Unable to get platform" ); |
| |
| /* Get some devices */ |
| err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 2, devices, &numDevices ); |
| test_error( err, "Unable to get 2 devices" ); |
| |
| if( numDevices < 2 ) |
| { |
| log_info( "WARNING: two device test unable to get two devices via CL_DEVICE_TYPE_ALL (got %d devices). Skipping test...\n", (int)numDevices ); |
| return 0; |
| } |
| else if (numDevices > 2) |
| { |
| log_info("Note: got %d devices, using just the first two.\n", (int)numDevices); |
| } |
| |
| /* Run test */ |
| return test_device_set( 2, 2, devices, num_elements ); |
| } |
| |
| int test_max_devices(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| cl_platform_id platform; |
| cl_device_id devices[MAX_DEVICES]; |
| cl_uint deviceCount; |
| int err; |
| |
| err = clGetPlatformIDs(1, &platform, NULL); |
| test_error( err, "Unable to get platform" ); |
| |
| /* Get some devices */ |
| err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, MAX_DEVICES, devices, &deviceCount ); |
| test_error( err, "Unable to get multiple devices" ); |
| |
| log_info("Testing with %d devices.", deviceCount); |
| |
| /* Run test */ |
| return test_device_set( deviceCount, deviceCount, devices, num_elements ); |
| } |
| |
| int test_hundred_queues(cl_device_id device, cl_context contextIgnore, cl_command_queue queueIgnore, int num_elements) |
| { |
| return test_device_set( 1, 100, &device, num_elements ); |
| } |
| |