// // Copyright (c) 2020 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/testHarness.h" #include "harness/deviceInfo.h" static const char* test_kernel = R"CLC( __kernel void test(__global int* dst) { dst[0] = 0; } )CLC"; int test_consistency_svm(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { // clGetDeviceInfo, passing CL_DEVICE_SVM_CAPABILITIES: // May return 0, indicating that device does not support Shared Virtual // Memory. cl_int error; const size_t allocSize = 16; clMemWrapper mem; clProgramWrapper program; clKernelWrapper kernel; cl_device_svm_capabilities svmCaps = 0; error = clGetDeviceInfo(deviceID, CL_DEVICE_SVM_CAPABILITIES, sizeof(svmCaps), &svmCaps, NULL); test_error(error, "Unable to query CL_DEVICE_SVM_CAPABILITIES"); if (svmCaps == 0) { // Test setup: mem = clCreateBuffer(context, CL_MEM_READ_WRITE, allocSize, NULL, &error); test_error(error, "Unable to create test buffer"); error = create_single_kernel_helper(context, &program, &kernel, 1, &test_kernel, "test"); test_error(error, "Unable to create test kernel"); // clGetMemObjectInfo, passing CL_MEM_USES_SVM_POINTER // Returns CL_FALSE if no devices in the context associated with // memobj support Shared Virtual Memory. cl_bool usesSVMPointer; error = clGetMemObjectInfo(mem, CL_MEM_USES_SVM_POINTER, sizeof(usesSVMPointer), &usesSVMPointer, NULL); test_error(error, "Unable to query CL_MEM_USES_SVM_POINTER"); test_assert_error(usesSVMPointer == CL_FALSE, "CL_DEVICE_SVM_CAPABILITIES returned 0 but " "CL_MEM_USES_SVM_POINTER did not return CL_FALSE"); // Check that the SVM APIs can be called. // Returns NULL if no devices in context support Shared Virtual Memory. void* ptr0 = clSVMAlloc(context, CL_MEM_READ_WRITE, allocSize, 0); void* ptr1 = clSVMAlloc(context, CL_MEM_READ_WRITE, allocSize, 0); test_assert_error(ptr0 == NULL && ptr1 == NULL, "CL_DEVICE_SVM_CAPABILITIES returned 0 but " "clSVMAlloc returned a non-NULL value"); // clEnqueueSVMFree, clEnqueueSVMMemcpy, clEnqueueSVMMemFill, // clEnqueueSVMMap, clEnqueueSVMUnmap, clEnqueueSVMMigrateMem Returns // CL_INVALID_OPERATION if the device associated with command_queue does // not support Shared Virtual Memory. // These calls purposefully pass bogus pointers to the functions to // better test that they are a NOP when SVM is not supported. void* bogus0 = (void*)0xDEADBEEF; void* bogus1 = (void*)0xDEADDEAD; cl_uint pattern = 0xAAAAAAAA; error = clEnqueueSVMMemFill(queue, bogus0, &pattern, sizeof(pattern), allocSize, 0, NULL, NULL); test_failure_error( error, CL_INVALID_OPERATION, "CL_DEVICE_SVM_CAPABILITIES returned 0 but clEnqueueSVMMemFill did " "not return CL_INVALID_OPERATION"); error = clEnqueueSVMMemcpy(queue, CL_TRUE, bogus1, bogus0, allocSize, 0, NULL, NULL); test_failure_error( error, CL_INVALID_OPERATION, "CL_DEVICE_SVM_CAPABILITIES returned 0 but " "clEnqueueSVMMemcpy did not return CL_INVALID_OPERATION"); error = clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_READ, bogus1, allocSize, 0, NULL, NULL); test_failure_error( error, CL_INVALID_OPERATION, "CL_DEVICE_SVM_CAPABILITIES returned 0 but " "clEnqueueSVMMap did not return CL_INVALID_OPERATION"); error = clEnqueueSVMUnmap(queue, bogus1, 0, NULL, NULL); test_failure_error( error, CL_INVALID_OPERATION, "CL_DEVICE_SVM_CAPABILITIES returned 0 but " "clEnqueueSVMUnmap did not return CL_INVALID_OPERATION"); error = clEnqueueSVMMigrateMem(queue, 1, (const void**)&bogus1, NULL, 0, 0, NULL, NULL); test_failure_error( error, CL_INVALID_OPERATION, "CL_DEVICE_SVM_CAPABILITIES returned 0 but " "clEnqueueSVMMigrateMem did not return CL_INVALID_OPERATION"); // If the enqueue calls above did not return errors, a clFinish would be // needed here to ensure the SVM operations are complete before freeing // the SVM pointers. clSVMFree(context, bogus0); error = clEnqueueSVMFree(queue, 1, &bogus0, NULL, NULL, 0, NULL, NULL); test_failure_error( error, CL_INVALID_OPERATION, "CL_DEVICE_SVM_CAPABILITIES returned 0 but " "clEnqueueSVMFree did not return CL_INVALID_OPERATION"); // If the enqueue calls above did not return errors, a clFinish should // be included here to ensure the enqueued SVM free is complete. // clSetKernelArgSVMPointer, clSetKernelExecInfo // Returns CL_INVALID_OPERATION if no devices in the context associated // with kernel support Shared Virtual Memory. error = clSetKernelArgSVMPointer(kernel, 0, NULL); test_failure_error( error, CL_INVALID_OPERATION, "CL_DEVICE_SVM_CAPABILITIES returned 0 but " "clSetKernelArgSVMPointer did not return CL_INVALID_OPERATION"); error = clSetKernelExecInfo(kernel, CL_KERNEL_EXEC_INFO_SVM_PTRS, 0, NULL); test_failure_error( error, CL_INVALID_OPERATION, "CL_DEVICE_SVM_CAPABILITIES returned 0 but " "clSetKernelExecInfo did not return CL_INVALID_OPERATION"); } return TEST_PASS; } static int check_atomic_capabilities(cl_device_atomic_capabilities atomicCaps, cl_device_atomic_capabilities requiredCaps) { if ((atomicCaps & requiredCaps) != requiredCaps) { log_error("Atomic capabilities %llx is missing support for at least " "one required capability %llx!\n", atomicCaps, requiredCaps); return TEST_FAIL; } if ((atomicCaps & CL_DEVICE_ATOMIC_SCOPE_ALL_DEVICES) != 0 && (atomicCaps & CL_DEVICE_ATOMIC_SCOPE_DEVICE) == 0) { log_error("Support for CL_DEVICE_ATOMIC_SCOPE_ALL_DEVICES requires " "support for CL_DEVICE_ATOMIC_SCOPE_DEVICE!\n"); return TEST_FAIL; } if ((atomicCaps & CL_DEVICE_ATOMIC_SCOPE_DEVICE) != 0 && (atomicCaps & CL_DEVICE_ATOMIC_SCOPE_WORK_GROUP) == 0) { log_error("Support for CL_DEVICE_ATOMIC_SCOPE_DEVICE requires " "support for CL_DEVICE_ATOMIC_SCOPE_WORK_GROUP!\n"); return TEST_FAIL; } if ((atomicCaps & CL_DEVICE_ATOMIC_ORDER_SEQ_CST) != 0 && (atomicCaps & CL_DEVICE_ATOMIC_ORDER_ACQ_REL) == 0) { log_error("Support for CL_DEVICE_ATOMIC_ORDER_SEQ_CST requires " "support for CL_DEVICE_ATOMIC_ORDER_ACQ_REL!\n"); return TEST_FAIL; } if ((atomicCaps & CL_DEVICE_ATOMIC_ORDER_ACQ_REL) != 0 && (atomicCaps & CL_DEVICE_ATOMIC_ORDER_RELAXED) == 0) { log_error("Support for CL_DEVICE_ATOMIC_ORDER_ACQ_REL requires " "support for CL_DEVICE_ATOMIC_ORDER_RELAXED!\n"); return TEST_FAIL; } return TEST_PASS; } int test_consistency_memory_model(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { cl_int error; cl_device_atomic_capabilities atomicCaps = 0; error = clGetDeviceInfo(deviceID, CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES, sizeof(atomicCaps), &atomicCaps, NULL); test_error(error, "Unable to query CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES"); error = check_atomic_capabilities(atomicCaps, CL_DEVICE_ATOMIC_ORDER_RELAXED | CL_DEVICE_ATOMIC_SCOPE_WORK_GROUP); if (error == TEST_FAIL) { log_error("Checks failed for CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES\n"); return error; } error = clGetDeviceInfo(deviceID, CL_DEVICE_ATOMIC_FENCE_CAPABILITIES, sizeof(atomicCaps), &atomicCaps, NULL); test_error(error, "Unable to query CL_DEVICE_ATOMIC_FENCE_CAPABILITIES"); error = check_atomic_capabilities(atomicCaps, CL_DEVICE_ATOMIC_ORDER_RELAXED | CL_DEVICE_ATOMIC_ORDER_ACQ_REL | CL_DEVICE_ATOMIC_SCOPE_WORK_GROUP); if (error == TEST_FAIL) { log_error("Checks failed for CL_DEVICE_ATOMIC_FENCE_CAPABILITIES\n"); return error; } return TEST_PASS; } int test_consistency_device_enqueue(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { // clGetDeviceInfo, passing CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES // May return 0, indicating that device does not support Device-Side Enqueue // and On-Device Queues. cl_int error; cl_device_device_enqueue_capabilities dseCaps = 0; error = clGetDeviceInfo(deviceID, CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES, sizeof(dseCaps), &dseCaps, NULL); test_error(error, "Unable to query CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES"); if (dseCaps == 0) { // clGetDeviceInfo, passing CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES // Returns 0 if device does not support Device-Side Enqueue and // On-Device Queues. cl_command_queue_properties devQueueProps = 0; error = clGetDeviceInfo(deviceID, CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES, sizeof(devQueueProps), &devQueueProps, NULL); test_error(error, "Unable to query CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES"); test_assert_error( devQueueProps == 0, "CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 but " "CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES returned a non-zero value"); // clGetDeviceInfo, passing // CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE, // CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE, // CL_DEVICE_MAX_ON_DEVICE_QUEUES, or // CL_DEVICE_MAX_ON_DEVICE_EVENTS // Returns 0 if device does not support Device-Side Enqueue and // On-Device Queues. cl_uint u = 0; error = clGetDeviceInfo(deviceID, CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE, sizeof(u), &u, NULL); test_error(error, "Unable to query CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE"); test_assert_error(u == 0, "CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 " "but CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE " "returned a non-zero value"); error = clGetDeviceInfo(deviceID, CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE, sizeof(u), &u, NULL); test_error(error, "Unable to query CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE"); test_assert_error( u == 0, "CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 but " "CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE returned a non-zero value"); error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_ON_DEVICE_QUEUES, sizeof(u), &u, NULL); test_error(error, "Unable to query CL_DEVICE_MAX_ON_DEVICE_QUEUES"); test_assert_error( u == 0, "CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 but " "CL_DEVICE_MAX_ON_DEVICE_QUEUES returned a non-zero value"); error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_ON_DEVICE_EVENTS, sizeof(u), &u, NULL); test_error(error, "Unable to query CL_DEVICE_MAX_ON_DEVICE_EVENTS"); test_assert_error( u == 0, "CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 but " "CL_DEVICE_MAX_ON_DEVICE_EVENTS returned a non-zero value"); // clGetCommandQueueInfo, passing CL_QUEUE_SIZE // Returns CL_INVALID_COMMAND_QUEUE since command_queue cannot be a // valid device command-queue. error = clGetCommandQueueInfo(queue, CL_QUEUE_SIZE, sizeof(u), &u, NULL); test_failure_error( error, CL_INVALID_COMMAND_QUEUE, "CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 but " "CL_QUEUE_SIZE did not return CL_INVALID_COMMAND_QUEUE"); cl_command_queue q = NULL; error = clGetCommandQueueInfo(queue, CL_QUEUE_DEVICE_DEFAULT, sizeof(q), &q, NULL); test_error(error, "Unable to query CL_QUEUE_DEVICE_DEFAULT"); test_assert_error( q == NULL, "CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 but " "CL_QUEUE_DEVICE_DEFAULT returned a non-NULL value"); // clSetDefaultDeviceCommandQueue // Returns CL_INVALID_OPERATION if device does not support On-Device // Queues. error = clSetDefaultDeviceCommandQueue(context, deviceID, NULL); test_failure_error(error, CL_INVALID_OPERATION, "CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 " "but clSetDefaultDeviceCommandQueue did not return " "CL_INVALID_OPERATION"); } else { if ((dseCaps & CL_DEVICE_QUEUE_REPLACEABLE_DEFAULT) == 0) { // clSetDefaultDeviceCommandQueue // Returns CL_INVALID_OPERATION if device does not support a // replaceable default On-Device Queue. error = clSetDefaultDeviceCommandQueue(context, deviceID, NULL); test_failure_error( error, CL_INVALID_OPERATION, "CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES did not " "include CL_DEVICE_QUEUE_REPLACEABLE_DEFAULT but " "clSetDefaultDeviceCommandQueue did not return " "CL_INVALID_OPERATION"); } // If CL_DEVICE_QUEUE_REPLACEABLE_DEFAULT is set, // CL_DEVICE_QUEUE_SUPPORTED must also be set. if ((dseCaps & CL_DEVICE_QUEUE_REPLACEABLE_DEFAULT) != 0 && (dseCaps & CL_DEVICE_QUEUE_SUPPORTED) == 0) { log_error("DEVICE_QUEUE_REPLACEABLE_DEFAULT is set but " "DEVICE_QUEUE_SUPPORTED is not set\n"); return TEST_FAIL; } // Devices that set CL_DEVICE_QUEUE_SUPPORTED must also return CL_TRUE // for CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT. if ((dseCaps & CL_DEVICE_QUEUE_SUPPORTED) != 0) { cl_bool b; error = clGetDeviceInfo(deviceID, CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT, sizeof(b), &b, NULL); test_error( error, "Unable to query CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT"); test_assert_error( b == CL_TRUE, "DEVICE_QUEUE_SUPPORTED is set but " "CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT returned CL_FALSE"); } } return TEST_PASS; } int test_consistency_pipes(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { // clGetDeviceInfo, passing CL_DEVICE_PIPE_SUPPORT // May return CL_FALSE, indicating that device does not support Pipes. cl_int error; cl_bool pipeSupport = CL_FALSE; error = clGetDeviceInfo(deviceID, CL_DEVICE_PIPE_SUPPORT, sizeof(pipeSupport), &pipeSupport, NULL); test_error(error, "Unable to query CL_DEVICE_PIPE_SUPPORT"); if (pipeSupport == CL_FALSE) { // clGetDeviceInfo, passing // CL_DEVICE_MAX_PIPE_ARGS, // CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS, or // CL_DEVICE_PIPE_MAX_PACKET_SIZE // Returns 0 if device does not support Pipes. cl_uint u = 0; error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_PIPE_ARGS, sizeof(u), &u, NULL); test_error(error, "Unable to query CL_DEVICE_MAX_PIPE_ARGS"); test_assert_error(u == 0, "CL_DEVICE_PIPE_SUPPORT returned CL_FALSE, but " "CL_DEVICE_MAX_PIPE_ARGS returned a non-zero value"); error = clGetDeviceInfo(deviceID, CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS, sizeof(u), &u, NULL); test_error(error, "Unable to query CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS"); test_assert_error(u == 0, "CL_DEVICE_PIPE_SUPPORT returned CL_FALSE, but " "CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS returned " "a non-zero value"); error = clGetDeviceInfo(deviceID, CL_DEVICE_PIPE_MAX_PACKET_SIZE, sizeof(u), &u, NULL); test_error(error, "Unable to query CL_DEVICE_PIPE_MAX_PACKET_SIZE"); test_assert_error( u == 0, "CL_DEVICE_PIPE_SUPPORT returned CL_FALSE, but " "CL_DEVICE_PIPE_MAX_PACKET_SIZE returned a non-zero value"); // clCreatePipe // Returns CL_INVALID_OPERATION if no devices in context support Pipes. clMemWrapper mem = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, 4, 4, NULL, &error); test_failure_error(error, CL_INVALID_OPERATION, "CL_DEVICE_PIPE_SUPPORT returned CL_FALSE but " "clCreatePipe did not return CL_INVALID_OPERATION"); // clGetPipeInfo // Returns CL_INVALID_MEM_OBJECT since pipe cannot be a valid pipe // object. clMemWrapper not_a_pipe = clCreateBuffer(context, CL_MEM_READ_WRITE, 4, NULL, &error); test_error(error, "Unable to create non-pipe buffer"); error = clGetPipeInfo(not_a_pipe, CL_PIPE_PACKET_SIZE, sizeof(u), &u, NULL); test_failure_error( error, CL_INVALID_MEM_OBJECT, "CL_DEVICE_PIPE_SUPPORT returned CL_FALSE but clGetPipeInfo did " "not return CL_INVALID_MEM_OBJECT"); } else { // Devices that support pipes must also return CL_TRUE // for CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT. cl_bool b; error = clGetDeviceInfo(deviceID, CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT, sizeof(b), &b, NULL); test_error(error, "Unable to query CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT"); test_assert_error( b == CL_TRUE, "CL_DEVICE_PIPE_SUPPORT returned CL_TRUE but " "CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT returned CL_FALSE"); } return TEST_PASS; } int test_consistency_progvar(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { // clGetDeviceInfo, passing CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE // May return 0, indicating that device does not support Program Scope // Global Variables. cl_int error; clProgramWrapper program; clKernelWrapper kernel; size_t maxGlobalVariableSize = 0; error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE, sizeof(maxGlobalVariableSize), &maxGlobalVariableSize, NULL); test_error(error, "Unable to query CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE"); if (maxGlobalVariableSize == 0) { // Test setup: error = create_single_kernel_helper(context, &program, &kernel, 1, &test_kernel, "test"); test_error(error, "Unable to create test kernel"); size_t sz = SIZE_MAX; // clGetDeviceInfo, passing // CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE // Returns 0 if device does not support Program Scope Global Variables. error = clGetDeviceInfo(deviceID, CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE, sizeof(sz), &sz, NULL); test_error( error, "Unable to query CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE"); test_assert_error( sz == 0, "CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE returned 0 but " "CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE returned a " "non-zero value"); // clGetProgramBuildInfo, passing // CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE // Returns 0 if device does not support Program Scope Global Variables. error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE, sizeof(sz), &sz, NULL); test_error( error, "Unable to query CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE"); test_assert_error(sz == 0, "CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE returned 0 " "but CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE " "returned a non-zero value"); } return TEST_PASS; } int test_consistency_non_uniform_work_group(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { // clGetDeviceInfo, passing CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT: // May return CL_FALSE, indicating that device does not support Non-Uniform // Work Groups. cl_int error; const size_t allocSize = 16; clMemWrapper mem; clProgramWrapper program; clKernelWrapper kernel; cl_bool nonUniformWorkGroupSupport = CL_FALSE; error = clGetDeviceInfo(deviceID, CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT, sizeof(nonUniformWorkGroupSupport), &nonUniformWorkGroupSupport, NULL); test_error(error, "Unable to query CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT"); if (nonUniformWorkGroupSupport == CL_FALSE) { // Test setup: mem = clCreateBuffer(context, CL_MEM_READ_WRITE, allocSize, NULL, &error); test_error(error, "Unable to create test buffer"); error = create_single_kernel_helper(context, &program, &kernel, 1, &test_kernel, "test"); test_error(error, "Unable to create test kernel"); error = clSetKernelArg(kernel, 0, sizeof(mem), &mem); // clEnqueueNDRangeKernel // Behaves as though Non-Uniform Work Groups were not enabled for // kernel, if the device associated with command_queue does not support // Non-Uniform Work Groups. size_t global_work_size[] = { 3, 3, 3 }; size_t local_work_size[] = { 2, 2, 2 }; // First, check that a NULL local work size succeeds. error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size, NULL, 0, NULL, NULL); test_error(error, "Unable to enqueue kernel with a NULL local work size"); error = clFinish(queue); test_error(error, "Error calling clFinish after NULL local work size"); // 1D non-uniform work group: error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL); test_failure_error( error, CL_INVALID_WORK_GROUP_SIZE, "CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT returned CL_FALSE but 1D " "clEnqueueNDRangeKernel did not return CL_INVALID_WORK_GROUP_SIZE"); // 2D non-uniform work group: global_work_size[0] = local_work_size[0]; error = clEnqueueNDRangeKernel(queue, kernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL); test_failure_error( error, CL_INVALID_WORK_GROUP_SIZE, "CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT returned CL_FALSE but 2D " "clEnqueueNDRangeKernel did not return CL_INVALID_WORK_GROUP_SIZE"); // 3D non-uniform work group: global_work_size[1] = local_work_size[1]; error = clEnqueueNDRangeKernel(queue, kernel, 3, NULL, global_work_size, local_work_size, 0, NULL, NULL); test_failure_error( error, CL_INVALID_WORK_GROUP_SIZE, "CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT returned CL_FALSE but 3D " "clEnqueueNDRangeKernel did not return CL_INVALID_WORK_GROUP_SIZE"); } return TEST_PASS; } int test_consistency_read_write_images(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { // clGetDeviceInfo, passing // CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS May return 0, // indicating that device does not support Read-Write Images. cl_int error; cl_uint maxReadWriteImageArgs = 0; error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS, sizeof(maxReadWriteImageArgs), &maxReadWriteImageArgs, NULL); test_error(error, "Unable to query " "CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS"); // clGetSupportedImageFormats, passing // CL_MEM_KERNEL_READ_AND_WRITE // Returns an empty set (such as num_image_formats equal to 0), indicating // that no image formats are supported for reading and writing in the same // kernel, if no devices in context support Read-Write Images. cl_uint totalReadWriteImageFormats = 0; const cl_mem_object_type image_types[] = { CL_MEM_OBJECT_IMAGE1D, CL_MEM_OBJECT_IMAGE1D_BUFFER, CL_MEM_OBJECT_IMAGE2D, CL_MEM_OBJECT_IMAGE3D, CL_MEM_OBJECT_IMAGE1D_ARRAY, CL_MEM_OBJECT_IMAGE2D_ARRAY, }; for (size_t i = 0; i < ARRAY_SIZE(image_types); i++) { cl_uint numImageFormats = 0; error = clGetSupportedImageFormats( context, CL_MEM_KERNEL_READ_AND_WRITE, image_types[i], 0, NULL, &numImageFormats); test_error(error, "Unable to query number of CL_MEM_KERNEL_READ_AND_WRITE " "image formats"); totalReadWriteImageFormats += numImageFormats; } if (maxReadWriteImageArgs == 0) { test_assert_error( totalReadWriteImageFormats == 0, "CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS returned 0 " "but clGetSupportedImageFormats(CL_MEM_KERNEL_READ_AND_WRITE) " "returned a non-empty set"); } else { test_assert_error( totalReadWriteImageFormats != 0, "CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS is non-zero " "but clGetSupportedImageFormats(CL_MEM_KERNEL_READ_AND_WRITE) " "returned an empty set"); } return TEST_PASS; } int test_consistency_2d_image_from_buffer(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { // clGetDeviceInfo, passing CL_DEVICE_IMAGE_PITCH_ALIGNMENT or // CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT // May return 0, indicating that device does not support Creating a 2D Image // from a Buffer. cl_int error; const cl_image_format imageFormat = { CL_RGBA, CL_UNORM_INT8 }; const size_t imageDim = 2; const size_t elementSize = 4; const size_t bufferSize = imageDim * imageDim * elementSize; clMemWrapper buffer; clMemWrapper image; cl_uint imagePitchAlignment = 0; error = clGetDeviceInfo(deviceID, CL_DEVICE_IMAGE_PITCH_ALIGNMENT, sizeof(imagePitchAlignment), &imagePitchAlignment, NULL); test_error(error, "Unable to query " "CL_DEVICE_IMAGE_PITCH_ALIGNMENT"); cl_uint imageBaseAddressAlignment = 0; error = clGetDeviceInfo(deviceID, CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT, sizeof(imageBaseAddressAlignment), &imageBaseAddressAlignment, NULL); test_error(error, "Unable to query " "CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT"); bool supports_cl_khr_image2d_from_buffer = is_extension_available(deviceID, "cl_khr_image2d_from_buffer"); if (imagePitchAlignment == 0 || imageBaseAddressAlignment == 0) { // This probably means that Creating a 2D Image from a Buffer is not // supported. // Test setup: buffer = clCreateBuffer(context, CL_MEM_READ_ONLY, bufferSize, NULL, &error); test_error(error, "Unable to create test buffer"); // Check that both queries return zero: test_assert_error( imagePitchAlignment == 0, "CL_DEVICE_IMAGE_PITCH_ALIGNMENT returned a non-zero value but " "CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT returned 0"); test_assert_error( imageBaseAddressAlignment == 0, "CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT returned a non-zero value " "but CL_DEVICE_IMAGE_PITCH_ALIGNMENT returned 0"); // clGetDeviceInfo, passing CL_DEVICE_EXTENSIONS // Will not describe support for the cl_khr_image2d_from_buffer // extension if device does not support Creating a 2D Image from a // Buffer. test_assert_error(supports_cl_khr_image2d_from_buffer == false, "Device does not support Creating a 2D Image from a " "Buffer but does support cl_khr_image2d_from_buffer"); // clCreateImage or clCreateImageWithProperties, passing image_type // equal to CL_MEM_OBJECT_IMAGE2D and mem_object not equal to // NULL // Returns CL_INVALID_OPERATION if no devices in context support // Creating a 2D Image from a Buffer. cl_image_desc imageDesc = { 0 }; imageDesc.image_type = CL_MEM_OBJECT_IMAGE2D; imageDesc.image_width = imageDim; imageDesc.image_height = imageDim; imageDesc.mem_object = buffer; image = clCreateImage(context, CL_MEM_READ_ONLY, &imageFormat, &imageDesc, NULL, &error); test_failure_error( error, CL_INVALID_OPERATION, "Device does not support Creating a 2D Image from a " "Buffer but clCreateImage did not return CL_INVALID_OPERATION"); image = clCreateImageWithProperties(context, NULL, CL_MEM_READ_ONLY, &imageFormat, &imageDesc, NULL, &error); test_failure_error(error, CL_INVALID_OPERATION, "Device does not support Creating a 2D Image from a " "Buffer but clCreateImageWithProperties did not " "return CL_INVALID_OPERATION"); } else { test_assert_error(supports_cl_khr_image2d_from_buffer, "Device supports Creating a 2D Image from a Buffer " "but does not support cl_khr_image2d_from_buffer"); } return TEST_PASS; } // Nothing needed for sRGB Images: // All of the sRGB Image Channel Orders (such as CL_​sRGBA) are optional for // devices supporting OpenCL 3.0. int test_consistency_depth_images(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { // The CL_DEPTH Image Channel Order is optional for devices supporting // OpenCL 3.0. cl_int error; cl_uint totalDepthImageFormats = 0; const cl_mem_flags mem_flags[] = { CL_MEM_WRITE_ONLY, CL_MEM_READ_WRITE, CL_MEM_KERNEL_READ_AND_WRITE, }; for (size_t i = 0; i < ARRAY_SIZE(mem_flags); i++) { cl_uint numImageFormats = 0; error = clGetSupportedImageFormats(context, mem_flags[i], CL_MEM_OBJECT_IMAGE2D, 0, NULL, &numImageFormats); test_error( error, "Unable to query number of CL_MEM_OBJECT_IMAGE2D image formats"); std::vector imageFormats(numImageFormats); error = clGetSupportedImageFormats( context, mem_flags[i], CL_MEM_OBJECT_IMAGE2D, imageFormats.size(), imageFormats.data(), NULL); test_error(error, "Unable to query CL_MEM_OBJECT_IMAGE2D image formats"); for (auto& imageFormat : imageFormats) { if (imageFormat.image_channel_order == CL_DEPTH) { totalDepthImageFormats++; } } } bool supports_cl_khr_depth_images = is_extension_available(deviceID, "cl_khr_depth_images"); if (totalDepthImageFormats == 0) { test_assert_error(supports_cl_khr_depth_images == false, "Device does not support Depth Images but does " "support cl_khr_depth_images"); } else { test_assert_error(supports_cl_khr_depth_images, "Device supports Depth Images but does not support " "cl_khr_depth_images"); } return TEST_PASS; } int test_consistency_device_and_host_timer(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { // clGetPlatformInfo, passing CL_PLATFORM_HOST_TIMER_RESOLUTION // May return 0, indicating that platform does not support Device and Host // Timer Synchronization. cl_int error; cl_platform_id platform = NULL; error = clGetDeviceInfo(deviceID, CL_DEVICE_PLATFORM, sizeof(platform), &platform, NULL); test_error(error, "Unable to query CL_DEVICE_PLATFORM"); cl_ulong hostTimerResolution = 0; error = clGetPlatformInfo(platform, CL_PLATFORM_HOST_TIMER_RESOLUTION, sizeof(hostTimerResolution), &hostTimerResolution, NULL); test_error(error, "Unable to query CL_PLATFORM_HOST_TIMER_RESOLUTION"); if (hostTimerResolution == 0) { // clGetDeviceAndHostTimer, clGetHostTimer // Returns CL_INVALID_OPERATION if the platform associated with device // does not support Device and Host Timer Synchronization. cl_ulong dt = 0; cl_ulong ht = 0; error = clGetDeviceAndHostTimer(deviceID, &dt, &ht); test_failure_error( error, CL_INVALID_OPERATION, "CL_PLATFORM_HOST_TIMER_RESOLUTION returned 0 but " "clGetDeviceAndHostTimer did not return CL_INVALID_OPERATION"); error = clGetHostTimer(deviceID, &ht); test_failure_error( error, CL_INVALID_OPERATION, "CL_PLATFORM_HOST_TIMER_RESOLUTION returned 0 but " "clGetHostTimer did not return CL_INVALID_OPERATION"); } return TEST_PASS; } int test_consistency_il_programs(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { // clGetDeviceInfo, passing CL_DEVICE_IL_VERSION or // CL_DEVICE_ILS_WITH_VERSION // May return an empty string and empty array, indicating that device does // not support Intermediate Language Programs. cl_int error; clProgramWrapper program; clKernelWrapper kernel; // Even if the device does not support Intermediate Language Programs the // size of the string query should not be zero. size_t sz = SIZE_MAX; error = clGetDeviceInfo(deviceID, CL_DEVICE_IL_VERSION, 0, NULL, &sz); test_error(error, "Unable to query CL_DEVICE_IL_VERSION"); test_assert_error(sz != 0, "CL_DEVICE_IL_VERSION should return a non-zero size"); std::string ilVersion = get_device_il_version_string(deviceID); error = clGetDeviceInfo(deviceID, CL_DEVICE_ILS_WITH_VERSION, 0, NULL, &sz); test_error(error, "Unable to query CL_DEVICE_ILS_WITH_VERSION"); if (ilVersion == "" || sz == 0) { // This probably means that Intermediate Language Programs are not // supported. // Check that both queries are consistent: test_assert_error( ilVersion == "", "CL_DEVICE_IL_VERSION returned a non-empty string but " "CL_DEVICE_ILS_WITH_VERSION returned no supported ILs"); test_assert_error(sz == 0, "CL_DEVICE_ILS_WITH_VERSION returned supported ILs " "but CL_DEVICE_IL_VERSION returned an empty string"); bool supports_cl_khr_il_program = is_extension_available(deviceID, "cl_khr_il_program"); test_assert_error(supports_cl_khr_il_program == false, "Device does not support IL Programs but does " "support cl_khr_il_program"); // Test setup: error = create_single_kernel_helper(context, &program, &kernel, 1, &test_kernel, "test"); test_error(error, "Unable to create test kernel"); // clGetProgramInfo, passing CL_PROGRAM_IL // Returns an empty buffer (such as param_value_size_ret equal to 0) if // no devices in the context associated with program support // Intermediate Language Programs. error = clGetProgramInfo(program, CL_PROGRAM_IL, 0, NULL, &sz); test_error(error, "Unable to query CL_PROGRAM_IL"); test_assert_error(sz == 0, "Device does not support IL Programs but " "CL_PROGRAM_IL returned a non-zero size"); // clCreateProgramWithIL // Returns CL_INVALID_OPERATION if no devices in context support // Intermediate Language Programs. cl_uint bogus = 0xDEADBEEF; clProgramWrapper ilProgram = clCreateProgramWithIL(context, &bogus, sizeof(bogus), &error); test_failure_error( error, CL_INVALID_OPERATION, "Device does not support IL Programs but clCreateProgramWithIL did " "not return CL_INVALID_OPERATION"); // clSetProgramSpecializationConstant // Returns CL_INVALID_OPERATION if no devices associated with program // support Intermediate Language Programs. cl_uint specConst = 42; error = clSetProgramSpecializationConstant( program, 0, sizeof(specConst), &specConst); test_failure_error(error, CL_INVALID_OPERATION, "Device does not support IL Programs but " "clSetProgramSpecializationConstant did not return " "CL_INVALID_OPERATION"); } return TEST_PASS; } int test_consistency_subgroups(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { // clGetDeviceInfo, passing CL_DEVICE_MAX_NUM_SUB_GROUPS // May return 0, indicating that device does not support Subgroups. cl_int error; clProgramWrapper program; clKernelWrapper kernel; cl_uint maxNumSubGroups = 0; error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_NUM_SUB_GROUPS, sizeof(maxNumSubGroups), &maxNumSubGroups, NULL); test_error(error, "Unable to query CL_DEVICE_MAX_NUM_SUB_GROUPS"); if (maxNumSubGroups == 0) { // Test setup: error = create_single_kernel_helper(context, &program, &kernel, 1, &test_kernel, "test"); test_error(error, "Unable to create test kernel"); // clGetDeviceInfo, passing // CL_DEVICE_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS // Returns CL_FALSE if device does not support Subgroups. cl_bool ifp = CL_FALSE; error = clGetDeviceInfo( deviceID, CL_DEVICE_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS, sizeof(ifp), &ifp, NULL); test_error( error, "Unable to query CL_DEVICE_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS"); test_assert_error(ifp == CL_FALSE, "Device does not support Subgroups but " "CL_DEVICE_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS " "did not return CL_FALSE"); // clGetDeviceInfo, passing CL_DEVICE_EXTENSIONS // Will not describe support for the cl_khr_subgroups extension if // device does not support Subgroups. bool supports_cl_khr_subgroups = is_extension_available(deviceID, "cl_khr_subgroups"); test_assert_error(supports_cl_khr_subgroups == false, "Device does not support Subgroups but does " "support cl_khr_subgroups"); // clGetKernelSubGroupInfo // Returns CL_INVALID_OPERATION if device does not support Subgroups. size_t sz = SIZE_MAX; error = clGetKernelSubGroupInfo(kernel, deviceID, CL_KERNEL_MAX_NUM_SUB_GROUPS, 0, NULL, sizeof(sz), &sz, NULL); test_failure_error( error, CL_INVALID_OPERATION, "Device does not support Subgroups but clGetKernelSubGroupInfo did " "not return CL_INVALID_OPERATION"); } return TEST_PASS; } static void CL_CALLBACK program_callback(cl_program, void*) {} int test_consistency_prog_ctor_dtor(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { cl_int error; clProgramWrapper program; clKernelWrapper kernel; // Test setup: error = create_single_kernel_helper(context, &program, &kernel, 1, &test_kernel, "test"); test_error(error, "Unable to create test kernel"); // clGetProgramInfo, passing CL_PROGRAM_SCOPE_GLOBAL_CTORS_PRESENT or // CL_PROGRAM_SCOPE_GLOBAL_DTORS_PRESENT // Returns CL_FALSE if no devices in the context associated with program // support Program Initialization and Clean-Up Kernels. cl_bool b = CL_FALSE; error = clGetProgramInfo(program, CL_PROGRAM_SCOPE_GLOBAL_CTORS_PRESENT, sizeof(b), &b, NULL); test_error(error, "Unable to query CL_PROGRAM_SCOPE_GLOBAL_CTORS_PRESENT"); test_assert_error( b == CL_FALSE, "CL_PROGRAM_SCOPE_GLOBAL_CTORS_PRESENT did not return CL_FALSE"); error = clGetProgramInfo(program, CL_PROGRAM_SCOPE_GLOBAL_DTORS_PRESENT, sizeof(b), &b, NULL); test_error(error, "Unable to query CL_PROGRAM_SCOPE_GLOBAL_DTORS_PRESENT"); test_assert_error( b == CL_FALSE, "CL_PROGRAM_SCOPE_GLOBAL_DTORS_PRESENT did not return CL_FALSE"); // clSetProgramReleaseCallback // Returns CL_INVALID_OPERATION if no devices in the context associated with // program support Program Initialization and Clean-Up Kernels. error = clSetProgramReleaseCallback(program, program_callback, NULL); test_failure_error( error, CL_INVALID_OPERATION, "clSetProgramReleaseCallback did not return CL_INVALID_OPERATION"); return TEST_PASS; } int test_consistency_3d_image_writes(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { // clGetSupportedImageFormats, passing CL_MEM_OBJECT_IMAGE3D and one of // CL_MEM_WRITE_ONLY, CL_MEM_READ_WRITE, or CL_MEM_KERNEL_READ_AND_WRITE // Returns an empty set (such as num_image_formats equal to 0), // indicating that no image formats are supported for writing to 3D // image objects, if no devices in context support Writing to 3D Image // Objects. cl_int error; cl_uint total3DImageWriteFormats = 0; const cl_mem_flags mem_flags[] = { CL_MEM_WRITE_ONLY, CL_MEM_READ_WRITE, CL_MEM_KERNEL_READ_AND_WRITE, }; for (size_t i = 0; i < ARRAY_SIZE(mem_flags); i++) { cl_uint numImageFormats = 0; error = clGetSupportedImageFormats(context, mem_flags[i], CL_MEM_OBJECT_IMAGE3D, 0, NULL, &numImageFormats); test_error( error, "Unable to query number of CL_MEM_OBJECT_IMAGE3D image formats"); total3DImageWriteFormats += numImageFormats; } bool supports_cl_khr_3d_image_writes = is_extension_available(deviceID, "cl_khr_3d_image_writes"); if (total3DImageWriteFormats == 0) { // clGetDeviceInfo, passing CL_DEVICE_EXTENSIONS // Will not describe support for the cl_khr_3d_image_writes extension if // device does not support Writing to 3D Image Objects. test_assert_error(supports_cl_khr_3d_image_writes == false, "Device does not support Writing to 3D Image Objects " "but does support cl_khr_3d_image_writes"); } else { test_assert_error(supports_cl_khr_3d_image_writes, "Device supports Writing to 3D Image Objects but " "does not support cl_khr_3d_image_writes"); } return TEST_PASS; }