| // |
| // 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" |
| |
| #include <vector> |
| |
| typedef long long int lld; |
| typedef long long unsigned llu; |
| |
| const char *test_kernels[] = { |
| "__kernel void kernelA(__global int *dst)\n" |
| "{\n" |
| "\n" |
| " dst[get_global_id(0)]*=3;\n" |
| "\n" |
| "}\n" |
| "__kernel void kernelB(__global int *dst)\n" |
| "{\n" |
| "\n" |
| " dst[get_global_id(0)]++;\n" |
| "\n" |
| "}\n" |
| }; |
| |
| #define TEST_SIZE 512 |
| #define MAX_QUEUES 1000 |
| |
| const char *printPartition(cl_device_partition_property partition) |
| { |
| switch (partition) { |
| case (0): return "<NONE>"; |
| case (CL_DEVICE_PARTITION_EQUALLY): return "CL_DEVICE_PARTITION_EQUALLY"; |
| case (CL_DEVICE_PARTITION_BY_COUNTS): return "CL_DEVICE_PARTITION_BY_COUNTS"; |
| case (CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN): return "CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN"; |
| default: return "<unknown>"; |
| } // switch |
| } |
| |
| const char *printAffinity(cl_device_affinity_domain affinity) |
| { |
| switch (affinity) { |
| case (0): return "<NONE>"; |
| case (CL_DEVICE_AFFINITY_DOMAIN_NUMA): return "CL_DEVICE_AFFINITY_DOMAIN_NUMA"; |
| case (CL_DEVICE_AFFINITY_DOMAIN_L4_CACHE): return "CL_DEVICE_AFFINITY_DOMAIN_L4_CACHE"; |
| case (CL_DEVICE_AFFINITY_DOMAIN_L3_CACHE): return "CL_DEVICE_AFFINITY_DOMAIN_L3_CACHE"; |
| case (CL_DEVICE_AFFINITY_DOMAIN_L2_CACHE): return "CL_DEVICE_AFFINITY_DOMAIN_L2_CACHE"; |
| case (CL_DEVICE_AFFINITY_DOMAIN_L1_CACHE): return "CL_DEVICE_AFFINITY_DOMAIN_L1_CACHE"; |
| case (CL_DEVICE_AFFINITY_DOMAIN_NEXT_PARTITIONABLE): return "CL_DEVICE_AFFINITY_DOMAIN_NEXT_PARTITIONABLE"; |
| default: return "<unknown>"; |
| } // switch |
| } |
| int create_single_kernel_helper( cl_context context, cl_program *outProgram, cl_kernel *outKernel, unsigned int numKernelLines, const char **kernelProgram, const char *kernelName, const cl_device_id *parentDevice ) |
| { |
| int error = CL_SUCCESS; |
| |
| /* Create the program object from source */ |
| error = create_single_kernel_helper_create_program(context, outProgram, numKernelLines, kernelProgram); |
| if( *outProgram == NULL || error != CL_SUCCESS) |
| { |
| print_error( error, "clCreateProgramWithSource failed" ); |
| return error; |
| } |
| |
| /* Compile the program */ |
| int buildProgramFailed = 0; |
| int printedSource = 0; |
| error = clBuildProgram( *outProgram, ((parentDevice == NULL) ? 0 : 1), parentDevice, NULL, NULL, NULL ); |
| if (error != CL_SUCCESS) |
| { |
| unsigned int i; |
| print_error(error, "clBuildProgram failed"); |
| buildProgramFailed = 1; |
| printedSource = 1; |
| log_error( "Original source is: ------------\n" ); |
| for( i = 0; i < numKernelLines; i++ ) |
| log_error( "%s", kernelProgram[ i ] ); |
| } |
| |
| // Verify the build status on all devices |
| cl_uint deviceCount = 0; |
| error = clGetProgramInfo( *outProgram, CL_PROGRAM_NUM_DEVICES, sizeof( deviceCount ), &deviceCount, NULL ); |
| if (error != CL_SUCCESS) { |
| print_error(error, "clGetProgramInfo CL_PROGRAM_NUM_DEVICES failed"); |
| return error; |
| } |
| |
| if (deviceCount == 0) { |
| log_error("No devices found for program.\n"); |
| return -1; |
| } |
| |
| cl_device_id *devices = (cl_device_id*) malloc( deviceCount * sizeof( cl_device_id ) ); |
| if( NULL == devices ) |
| return -1; |
| memset( devices, 0, deviceCount * sizeof( cl_device_id )); |
| error = clGetProgramInfo( *outProgram, CL_PROGRAM_DEVICES, sizeof( cl_device_id ) * deviceCount, devices, NULL ); |
| if (error != CL_SUCCESS) { |
| print_error(error, "clGetProgramInfo CL_PROGRAM_DEVICES failed"); |
| free( devices ); |
| return error; |
| } |
| |
| cl_uint z; |
| for( z = 0; z < deviceCount; z++ ) |
| { |
| char deviceName[4096] = ""; |
| error = clGetDeviceInfo(devices[z], CL_DEVICE_NAME, sizeof( deviceName), deviceName, NULL); |
| if (error != CL_SUCCESS || deviceName[0] == '\0') { |
| log_error("Device \"%d\" failed to return a name\n", z); |
| print_error(error, "clGetDeviceInfo CL_DEVICE_NAME failed"); |
| } |
| |
| cl_build_status buildStatus; |
| error = clGetProgramBuildInfo(*outProgram, devices[z], CL_PROGRAM_BUILD_STATUS, sizeof(buildStatus), &buildStatus, NULL); |
| if (error != CL_SUCCESS) { |
| print_error(error, "clGetProgramBuildInfo CL_PROGRAM_BUILD_STATUS failed"); |
| free( devices ); |
| return error; |
| } |
| |
| if (buildStatus != CL_BUILD_SUCCESS || buildProgramFailed) { |
| char log[10240] = ""; |
| if (buildStatus == CL_BUILD_SUCCESS && buildProgramFailed) log_error("clBuildProgram returned an error, but buildStatus is marked as CL_BUILD_SUCCESS.\n"); |
| |
| char statusString[64] = ""; |
| if (buildStatus == (cl_build_status)CL_BUILD_SUCCESS) |
| sprintf(statusString, "CL_BUILD_SUCCESS"); |
| else if (buildStatus == (cl_build_status)CL_BUILD_NONE) |
| sprintf(statusString, "CL_BUILD_NONE"); |
| else if (buildStatus == (cl_build_status)CL_BUILD_ERROR) |
| sprintf(statusString, "CL_BUILD_ERROR"); |
| else if (buildStatus == (cl_build_status)CL_BUILD_IN_PROGRESS) |
| sprintf(statusString, "CL_BUILD_IN_PROGRESS"); |
| else |
| sprintf(statusString, "UNKNOWN (%d)", buildStatus); |
| |
| if (buildStatus != CL_BUILD_SUCCESS) log_error("Build not successful for device \"%s\", status: %s\n", deviceName, statusString); |
| error = clGetProgramBuildInfo( *outProgram, devices[z], CL_PROGRAM_BUILD_LOG, sizeof(log), log, NULL ); |
| if (error != CL_SUCCESS || log[0]=='\0'){ |
| log_error("Device %d (%s) failed to return a build log\n", z, deviceName); |
| if (error) { |
| print_error(error, "clGetProgramBuildInfo CL_PROGRAM_BUILD_LOG failed"); |
| free( devices ); |
| return error; |
| } else { |
| log_error("clGetProgramBuildInfo returned an empty log.\n"); |
| free( devices ); |
| return -1; |
| } |
| } |
| // In this case we've already printed out the code above. |
| if (!printedSource) |
| { |
| unsigned int i; |
| log_error( "Original source is: ------------\n" ); |
| for( i = 0; i < numKernelLines; i++ ) |
| log_error( "%s", kernelProgram[ i ] ); |
| printedSource = 1; |
| } |
| log_error( "Build log for device \"%s\" is: ------------\n", deviceName ); |
| log_error( "%s\n", log ); |
| log_error( "\n----------\n" ); |
| free( devices ); |
| return -1; |
| } |
| } |
| |
| /* And create a kernel from it */ |
| *outKernel = clCreateKernel( *outProgram, kernelName, &error ); |
| if( *outKernel == NULL || error != CL_SUCCESS) |
| { |
| print_error( error, "Unable to create kernel" ); |
| free( devices ); |
| return error; |
| } |
| |
| free( devices ); |
| return 0; |
| } |
| |
| template<class T> |
| class AutoDestructArray |
| { |
| public: |
| AutoDestructArray(T* arr) : m_arr(arr) {} |
| ~AutoDestructArray() { if (m_arr) delete [] m_arr; } |
| |
| private: |
| T* m_arr; |
| }; |
| |
| int test_device_set(size_t deviceCount, size_t queueCount, cl_device_id *devices, int num_elements, cl_device_id *parentDevice = NULL) |
| { |
| int error; |
| clContextWrapper context; |
| clProgramWrapper program; |
| clKernelWrapper kernels[2]; |
| clMemWrapper stream; |
| clCommandQueueWrapper queues[MAX_QUEUES]; |
| size_t threads[1], localThreads[1]; |
| int data[TEST_SIZE]; |
| int outputData[TEST_SIZE]; |
| int expectedResults[TEST_SIZE]; |
| int *expectedResultsOneDeviceArray = new int[deviceCount * TEST_SIZE]; |
| int **expectedResultsOneDevice = (int**)alloca(sizeof(int**) * deviceCount); |
| size_t i; |
| AutoDestructArray<int> autoDestruct(expectedResultsOneDeviceArray); |
| |
| for (i=0; i<deviceCount; i++) { |
| expectedResultsOneDevice[i] = expectedResultsOneDeviceArray + (i * TEST_SIZE); |
| } |
| |
| memset(queues, 0, sizeof(queues)); |
| |
| RandomSeed seed( gRandomSeed ); |
| |
| 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++) { |
| size_t deviceNameSize; |
| error = clGetDeviceInfo(devices[i], CL_DEVICE_NAME, 0, NULL, &deviceNameSize); |
| test_error(error, "clGetDeviceInfo CL_DEVICE_NAME failed"); |
| char *deviceName = (char *)alloca(deviceNameSize * (sizeof(char))); |
| error = clGetDeviceInfo(devices[i], CL_DEVICE_NAME, deviceNameSize, 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", parentDevice ) != 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_int) * 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] = clCreateCommandQueueWithProperties( 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[ q ], 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 init_device_partition_test(cl_device_id parentDevice, cl_uint &maxComputeUnits, cl_uint &maxSubDevices) |
| { |
| int err = clGetDeviceInfo(parentDevice, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(maxComputeUnits), &maxComputeUnits, NULL); |
| test_error( err, "Unable to get maximal number of compute units" ); |
| err = clGetDeviceInfo(parentDevice, CL_DEVICE_PARTITION_MAX_SUB_DEVICES, sizeof(maxSubDevices), &maxSubDevices, NULL); |
| test_error( err, "Unable to get maximal number of sub-devices" ); |
| |
| log_info("Maximal number of sub-devices on device %p is %d.\n", parentDevice, maxSubDevices ); |
| return 0; |
| } |
| |
| int test_device_partition_type_support(cl_device_id parentDevice, const cl_device_partition_property partitionType, const cl_device_affinity_domain affinityDomain) |
| { |
| typedef std::vector< cl_device_partition_property > properties_t; |
| properties_t supportedProps( 3 ); // only 3 types defined in the spec (but implementation can define more) |
| size_t const propSize = sizeof( cl_device_partition_property ); // Size of one property in bytes. |
| size_t size; // size of all properties in bytes. |
| cl_int err; |
| size = 0; |
| err = clGetDeviceInfo( parentDevice, CL_DEVICE_PARTITION_PROPERTIES, 0, NULL, & size ); |
| if ( err == CL_SUCCESS ) { |
| if ( size % propSize != 0 ) { |
| log_error( "ERROR: clGetDeviceInfo: Bad size of returned partition properties (%llu), it must me a multiply of partition property size (%llu)\n", llu( size ), llu( propSize ) ); |
| return -1; |
| } |
| supportedProps.resize( size / propSize ); |
| size = 0; |
| err = clGetDeviceInfo( parentDevice, CL_DEVICE_PARTITION_PROPERTIES, supportedProps.size() * propSize, & supportedProps.front(), & size ); |
| test_error_ret( err, "Unable to get device partition properties (2)", -1 ); |
| } else if ( err == CL_INVALID_VALUE ) { |
| log_error( "ERROR: clGetDeviceInfo: CL_DEVICE_PARTITION_PROPERTIES is not supported.\n" ); |
| return -1; |
| } else { |
| test_error_ret( err, "Unable to get device partition properties (1)", -1 ); |
| }; |
| for ( int i = 0; i < supportedProps.size(); i++) |
| { |
| if (supportedProps[i] == partitionType) |
| { |
| if (partitionType == CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN) |
| { |
| cl_device_affinity_domain supportedAffinityDomain; |
| err = clGetDeviceInfo(parentDevice, CL_DEVICE_PARTITION_AFFINITY_DOMAIN, sizeof(supportedAffinityDomain), &supportedAffinityDomain, NULL); |
| test_error( err, "Unable to get supported affinity domains" ); |
| if (supportedAffinityDomain & affinityDomain) |
| return 0; |
| } |
| else |
| return 0; |
| } |
| } |
| |
| return -1; |
| } |
| |
| int test_partition_of_device(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, cl_device_partition_property *partition_type, |
| cl_uint starting_property, cl_uint ending_property) |
| { |
| cl_uint maxComputeUnits; |
| cl_uint maxSubDevices; // maximal number of sub-devices that can be created in one call to clCreateSubDevices |
| int err = 0; |
| |
| if (init_device_partition_test(deviceID, maxComputeUnits, maxSubDevices) != 0) |
| return -1; |
| |
| if (maxComputeUnits <= 1) |
| return 0; |
| // confirm that this devices reports how it was partitioned |
| if (partition_type != NULL) |
| { // if we're not the root device |
| size_t psize; |
| err = clGetDeviceInfo(deviceID, CL_DEVICE_PARTITION_TYPE, 0, NULL, &psize); |
| test_error( err, "Unable to get CL_DEVICE_PARTITION_TYPE" ); |
| cl_device_partition_property *properties_returned = (cl_device_partition_property *)alloca(psize); |
| err = clGetDeviceInfo(deviceID, CL_DEVICE_PARTITION_TYPE, psize, (void *) properties_returned, NULL); |
| test_error( err, "Unable to get CL_DEVICE_PARTITION_TYPE" ); |
| |
| // test returned type |
| for (cl_uint i = 0;i < psize / sizeof(cl_device_partition_property);i++) { |
| if (properties_returned[i] != partition_type[i]) { |
| if (!(partition_type[0] == CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN && |
| i == 1 && partition_type[1] == CL_DEVICE_AFFINITY_DOMAIN_NEXT_PARTITIONABLE && |
| (properties_returned[1] == CL_DEVICE_AFFINITY_DOMAIN_NUMA || |
| properties_returned[1] == CL_DEVICE_AFFINITY_DOMAIN_L4_CACHE || |
| properties_returned[1] == CL_DEVICE_AFFINITY_DOMAIN_L3_CACHE || |
| properties_returned[1] == CL_DEVICE_AFFINITY_DOMAIN_L2_CACHE || |
| properties_returned[1] == CL_DEVICE_AFFINITY_DOMAIN_L1_CACHE))) { |
| log_error("properties_returned[%d] 0x%x != 0x%x partition_type[%d].", i, properties_returned[i], partition_type[i], i); |
| return -1; |
| } |
| } |
| } // for |
| } |
| |
| #define PROPERTY_TYPES 8 |
| cl_device_partition_property partitionProp[PROPERTY_TYPES][5] = { |
| { CL_DEVICE_PARTITION_EQUALLY, maxComputeUnits / 2, 0, 0, 0 } , |
| { CL_DEVICE_PARTITION_BY_COUNTS, 1, maxComputeUnits - 1, CL_DEVICE_PARTITION_BY_COUNTS_LIST_END, 0 } , |
| { CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN, CL_DEVICE_AFFINITY_DOMAIN_NUMA, 0, 0, 0 } , |
| { CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN, CL_DEVICE_AFFINITY_DOMAIN_L4_CACHE, 0, 0, 0 } , |
| { CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN, CL_DEVICE_AFFINITY_DOMAIN_L3_CACHE, 0, 0, 0 } , |
| { CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN, CL_DEVICE_AFFINITY_DOMAIN_L2_CACHE, 0, 0, 0 } , |
| { CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN, CL_DEVICE_AFFINITY_DOMAIN_L1_CACHE, 0, 0, 0 } , |
| { CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN, CL_DEVICE_AFFINITY_DOMAIN_NEXT_PARTITIONABLE, 0, 0, 0 } |
| }; |
| |
| // loop thru each type, creating sub-devices for each type |
| for (cl_uint i = starting_property;i < ending_property;i++) { |
| |
| if (test_device_partition_type_support(deviceID, partitionProp[i][0], partitionProp[i][1]) != 0) |
| { |
| if (partitionProp[i][0] == CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN) |
| { |
| log_info( "Device partition type \"%s\" \"%s\" is not supported on device %p. Skipping test...\n", |
| printPartition(partitionProp[i][0]), |
| printAffinity(partitionProp[i][1]), deviceID); |
| } |
| else |
| { |
| log_info( "Device partition type \"%s\" is not supported on device %p. Skipping test...\n", |
| printPartition(partitionProp[i][0]), deviceID); |
| } |
| continue; |
| } |
| |
| if (partitionProp[i][0] == CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN) |
| { |
| log_info("Testing on device %p partition type \"%s\" \"%s\"\n", deviceID, printPartition(partitionProp[i][0]), |
| printAffinity(partitionProp[i][1])); |
| } |
| else |
| { |
| log_info("Testing on device %p partition type \"%s\" (%d,%d)\n", deviceID, printPartition(partitionProp[i][0]), |
| partitionProp[i][1], partitionProp[i][2]); |
| } |
| |
| cl_uint deviceCount; |
| |
| // how many sub-devices can we create? |
| err = clCreateSubDevices(deviceID, partitionProp[i], 0, NULL, &deviceCount); |
| if ( err == CL_DEVICE_PARTITION_FAILED ) { |
| log_info( "The device %p could not be further partitioned.\n", deviceID ); |
| continue; |
| } |
| test_error( err, "Failed to get number of sub-devices" ); |
| |
| // get the list of subDevices |
| // create room for 1 more device_id, so that we can put the parent device in there. |
| cl_device_id *subDevices = (cl_device_id*)alloca(sizeof(cl_device_id) * (deviceCount + 1)); |
| err = clCreateSubDevices(deviceID, partitionProp[i], deviceCount, subDevices, &deviceCount); |
| test_error( err, "Actual creation of sub-devices failed" ); |
| |
| log_info("Testing on all devices in context\n"); |
| err = test_device_set(deviceCount, deviceCount, subDevices, num_elements); |
| if (err == 0) |
| { |
| log_info("Testing on a parent device for context\n"); |
| |
| // add the parent device |
| subDevices[deviceCount] = deviceID; |
| err = test_device_set(deviceCount + 1, deviceCount, subDevices, num_elements, &deviceID); |
| } |
| if (err != 0) |
| { |
| printf("error! returning %d\n",err); |
| return err; |
| } |
| |
| // now, recurse and test the FIRST of these sub-devices, to make sure it can be further partitioned |
| err = test_partition_of_device(subDevices[0], context, queue, num_elements, partitionProp[i], starting_property, ending_property); |
| if (err != 0) |
| { |
| printf("error! returning %d\n",err); |
| return err; |
| } |
| |
| for (cl_uint j=0;j < deviceCount;j++) |
| { |
| err = clReleaseDevice(subDevices[j]); |
| test_error( err, "\n Releasing sub-device failed \n" ); |
| } |
| |
| } // for |
| |
| log_info("Testing on all device %p finished\n", deviceID); |
| return 0; |
| } |
| |
| |
| int test_partition_equally(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return test_partition_of_device(deviceID, context, queue, num_elements, NULL, 0, 1); |
| } |
| |
| int test_partition_by_counts(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return test_partition_of_device(deviceID, context, queue, num_elements, NULL, 1, 2); |
| } |
| |
| int test_partition_by_affinity_domain_numa(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return test_partition_of_device(deviceID, context, queue, num_elements, NULL, 2, 3); |
| } |
| |
| int test_partition_by_affinity_domain_l4_cache(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return test_partition_of_device(deviceID, context, queue, num_elements, NULL, 3, 4); |
| } |
| |
| int test_partition_by_affinity_domain_l3_cache(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return test_partition_of_device(deviceID, context, queue, num_elements, NULL, 4, 5); |
| } |
| |
| int test_partition_by_affinity_domain_l2_cache(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return test_partition_of_device(deviceID, context, queue, num_elements, NULL, 5, 6); |
| } |
| |
| int test_partition_by_affinity_domain_l1_cache(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return test_partition_of_device(deviceID, context, queue, num_elements, NULL, 6, 7); |
| } |
| |
| int test_partition_by_affinity_domain_next_partitionable(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return test_partition_of_device(deviceID, context, queue, num_elements, NULL, 7, 8); |
| } |
| |
| int test_partition_all(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return test_partition_of_device(deviceID, context, queue, num_elements, NULL, 0, 8); |
| } |
