diff --git a/CMakeLists.txt b/CMakeLists.txt index ce71604c..d8206637 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -17,6 +17,7 @@ else(CMAKE_BUILD_TYPE STREQUAL "release") endif(CMAKE_BUILD_TYPE STREQUAL "release") add_definitions(-DCL_TARGET_OPENCL_VERSION=300) +add_definitions(-DCL_USE_DEPRECATED_OPENCL_2_2_APIS=1) add_definitions(-DCL_USE_DEPRECATED_OPENCL_2_1_APIS=1) add_definitions(-DCL_USE_DEPRECATED_OPENCL_2_0_APIS=1) add_definitions(-DCL_USE_DEPRECATED_OPENCL_1_2_APIS=1) diff --git a/test_common/harness/deviceInfo.cpp b/test_common/harness/deviceInfo.cpp index c9816fa2..96a04f9a 100644 --- a/test_common/harness/deviceInfo.cpp +++ b/test_common/harness/deviceInfo.cpp @@ -45,7 +45,8 @@ static std::string get_device_info_string(cl_device_id device, throw std::runtime_error("clGetDeviceInfo failed\n"); } - return std::string(info.begin(), info.end()); + /* The returned string does not include the null terminator. */ + return std::string(info.data(), size - 1); } /* Determines if an extension is supported by a device. */ diff --git a/test_common/harness/errorHelpers.cpp b/test_common/harness/errorHelpers.cpp index 5cd87171..e17888b7 100644 --- a/test_common/harness/errorHelpers.cpp +++ b/test_common/harness/errorHelpers.cpp @@ -85,7 +85,13 @@ const char *IGetErrorString( int clErrorCode ) case CL_INVALID_IMAGE_DESCRIPTOR: return "CL_INVALID_IMAGE_DESCRIPTOR"; case CL_INVALID_COMPILER_OPTIONS: return "CL_INVALID_COMPILER_OPTIONS"; case CL_INVALID_LINKER_OPTIONS: return "CL_INVALID_LINKER_OPTIONS"; - case CL_INVALID_DEVICE_PARTITION_COUNT: return "CL_INVALID_DEVICE_PARTITION_COUNT"; + case CL_INVALID_DEVICE_PARTITION_COUNT: + return "CL_INVALID_DEVICE_PARTITION_COUNT"; + case CL_INVALID_PIPE_SIZE: return "CL_INVALID_PIPE_SIZE"; + case CL_INVALID_DEVICE_QUEUE: return "CL_INVALID_DEVICE_QUEUE"; + case CL_INVALID_SPEC_ID: return "CL_INVALID_SPEC_ID"; + case CL_MAX_SIZE_RESTRICTION_EXCEEDED: + return "CL_MAX_SIZE_RESTRICTION_EXCEEDED"; default: return "(unknown)"; } } diff --git a/test_common/harness/errorHelpers.h b/test_common/harness/errorHelpers.h index 53b74bf8..3238a956 100644 --- a/test_common/harness/errorHelpers.h +++ b/test_common/harness/errorHelpers.h @@ -82,6 +82,21 @@ #define test_failure_warning_ret(errCode, expectedErrCode, msg, retValue) { if( errCode != expectedErrCode ) { print_failure_warning( errCode, expectedErrCode, msg ); warnings++ ; } } #define print_failure_warning(errCode, expectedErrCode, msg) log_error( "WARNING: %s! (Got %s, expected %s from %s:%d)\n", msg, IGetErrorString( errCode ), IGetErrorString( expectedErrCode ), __FILE__, __LINE__ ); +// generate an error when an assertion is false (not error code related) +#define test_assert_error(condition, msg) \ + test_assert_error_ret(condition, msg, TEST_FAIL) +#define test_assert_error_ret(condition, msg, retValue) \ + { \ + if (!(condition)) \ + { \ + print_assertion_error(condition, msg); \ + return retValue; \ + } \ + } +#define print_assertion_error(condition, msg) \ + log_error("ERROR: %s! (!(%s) from %s:%d)\n", msg, #condition, __FILE__, \ + __LINE__); + #define ASSERT_SUCCESS(expr, msg) \ do \ { \ diff --git a/test_conformance/api/main.cpp b/test_conformance/api/main.cpp index 9308fa8d..8d8d20ad 100644 --- a/test_conformance/api/main.cpp +++ b/test_conformance/api/main.cpp @@ -132,6 +132,16 @@ test_definition test_list[] = { ADD_TEST_VERSION(consistency_memory_model, Version(3, 0)), ADD_TEST_VERSION(consistency_device_enqueue, Version(3, 0)), ADD_TEST_VERSION(consistency_pipes, Version(3, 0)), + ADD_TEST_VERSION(consistency_progvar, Version(3, 0)), + ADD_TEST_VERSION(consistency_non_uniform_work_group, Version(3, 0)), + ADD_TEST_VERSION(consistency_read_write_images, Version(3, 0)), + ADD_TEST_VERSION(consistency_2d_image_from_buffer, Version(3, 0)), + ADD_TEST_VERSION(consistency_depth_images, Version(3, 0)), + ADD_TEST_VERSION(consistency_device_and_host_timer, Version(3, 0)), + ADD_TEST_VERSION(consistency_il_programs, Version(3, 0)), + ADD_TEST_VERSION(consistency_subgroups, Version(3, 0)), + ADD_TEST_VERSION(consistency_prog_ctor_dtor, Version(3, 0)), + ADD_TEST_VERSION(consistency_3d_image_writes, Version(3, 0)), }; const int test_num = ARRAY_SIZE(test_list); diff --git a/test_conformance/api/procs.h b/test_conformance/api/procs.h index b55f5a3b..21dca3f2 100644 --- a/test_conformance/api/procs.h +++ b/test_conformance/api/procs.h @@ -150,3 +150,39 @@ extern int test_consistency_device_enqueue(cl_device_id deviceID, int num_elements); extern int test_consistency_pipes(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); +extern int test_consistency_progvar(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_consistency_non_uniform_work_group(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_consistency_read_write_images(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_consistency_2d_image_from_buffer(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_consistency_depth_images(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_consistency_device_and_host_timer(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_consistency_il_programs(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_consistency_subgroups(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_consistency_prog_ctor_dtor(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_consistency_3d_image_writes(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); diff --git a/test_conformance/api/test_api_consistency.cpp b/test_conformance/api/test_api_consistency.cpp index 6b9f4771..99d17e05 100644 --- a/test_conformance/api/test_api_consistency.cpp +++ b/test_conformance/api/test_api_consistency.cpp @@ -15,8 +15,9 @@ // #include "testBase.h" #include "harness/testHarness.h" +#include "harness/deviceInfo.h" -const char* test_kernel = R"CLC( +static const char* test_kernel = R"CLC( __kernel void test(__global int* dst) { dst[0] = 0; } @@ -28,7 +29,7 @@ int test_consistency_svm(cl_device_id deviceID, cl_context context, // clGetDeviceInfo, passing CL_DEVICE_SVM_CAPABILITIES: // May return 0, indicating that device does not support Shared Virtual // Memory. - int error; + cl_int error; const size_t allocSize = 16; clMemWrapper mem; @@ -60,24 +61,18 @@ int test_consistency_svm(cl_device_id deviceID, cl_context context, clGetMemObjectInfo(mem, CL_MEM_USES_SVM_POINTER, sizeof(usesSVMPointer), &usesSVMPointer, NULL); test_error(error, "Unable to query CL_MEM_USES_SVM_POINTER"); - if (usesSVMPointer != CL_FALSE) - { - log_error("CL_DEVICE_SVM_CAPABILITIES returned 0 but " - "CL_MEM_USES_SVM_POINTER did not return CL_FALSE\n"); - return TEST_FAIL; - } + 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); - if (ptr0 != NULL || ptr1 != NULL) - { - log_error("CL_DEVICE_SVM_CAPABILITIES returned 0 but clSVMAlloc " - "returned a non-NULL value\n"); - return TEST_FAIL; - } + 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 @@ -182,7 +177,7 @@ static int check_atomic_capabilities(cl_device_atomic_capabilities atomicCaps, int test_consistency_memory_model(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { - int error; + cl_int error; cl_device_atomic_capabilities atomicCaps = 0; error = clGetDeviceInfo(deviceID, CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES, @@ -221,7 +216,7 @@ int test_consistency_device_enqueue(cl_device_id deviceID, cl_context context, // clGetDeviceInfo, passing CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES // May return 0, indicating that device does not support Device-Side Enqueue // and On-Device Queues. - int error; + cl_int error; cl_device_device_enqueue_capabilities dseCaps = 0; error = clGetDeviceInfo(deviceID, CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES, @@ -239,13 +234,10 @@ int test_consistency_device_enqueue(cl_device_id deviceID, cl_context context, sizeof(devQueueProps), &devQueueProps, NULL); test_error(error, "Unable to query CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES"); - if (devQueueProps != 0) - { - log_error("CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 but " - "CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES returned a " - "non-zero value\n"); - return TEST_FAIL; - } + 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, @@ -262,74 +254,54 @@ int test_consistency_device_enqueue(cl_device_id deviceID, cl_context context, sizeof(u), &u, NULL); test_error(error, "Unable to query CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE"); - if (u != 0) - { - log_error("CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 but " - "CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE returned a " - "non-zero value\n"); - return TEST_FAIL; - } + 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"); - if (u != 0) - { - log_error("CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 but " - "CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE returned a " - "non-zero value\n"); - return TEST_FAIL; - } + 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"); - if (u != 0) - { - log_error("CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 but " - "CL_DEVICE_MAX_ON_DEVICE_QUEUES returned a " - "non-zero value\n"); - return TEST_FAIL; - } + 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"); - if (u != 0) - { - log_error("CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 but " - "CL_DEVICE_MAX_ON_DEVICE_EVENTS returned a " - "non-zero value\n"); - return TEST_FAIL; - } + 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 or - // CL_QUEUE_DEVICE_DEFAULT - // Returns 0 or NULL if the device associated with command_queue does - // not support On-Device Queues. + // 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); - // TODO: is this a valid query? See: - // https://github.com/KhronosGroup/OpenCL-Docs/issues/402 - // test_error(error, "Unable to query CL_QUEUE_SIZE"); - if (error == CL_SUCCESS && u != 0) - { - log_error("CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 but " - "CL_QUEUE_SIZE returned a non-zero value\n"); - return TEST_FAIL; - } + 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"); - if (q != NULL) - { - log_error("CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 but " - "CL_QUEUE_DEVICE_DEFAULT returned a non-NULL value\n"); - return TEST_FAIL; - } + 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 @@ -377,13 +349,10 @@ int test_consistency_device_enqueue(cl_device_id deviceID, cl_context context, test_error( error, "Unable to query CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT"); - if (b != CL_TRUE) - { - log_error("DEVICE_QUEUE_SUPPORTED is set but " - "CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT returned " - "CL_FALSE\n"); - return TEST_FAIL; - } + test_assert_error( + b == CL_TRUE, + "DEVICE_QUEUE_SUPPORTED is set but " + "CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT returned CL_FALSE"); } } @@ -395,7 +364,7 @@ int test_consistency_pipes(cl_device_id deviceID, cl_context context, { // clGetDeviceInfo, passing CL_DEVICE_PIPE_SUPPORT // May return CL_FALSE, indicating that device does not support Pipes. - int error; + cl_int error; cl_bool pipeSupport = CL_FALSE; error = clGetDeviceInfo(deviceID, CL_DEVICE_PIPE_SUPPORT, @@ -415,36 +384,27 @@ int test_consistency_pipes(cl_device_id deviceID, cl_context context, error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_PIPE_ARGS, sizeof(u), &u, NULL); test_error(error, "Unable to query CL_DEVICE_MAX_PIPE_ARGS"); - if (u != 0) - { - log_error("CL_DEVICE_PIPE_SUPPORT returned CL_FALSE, but " - "CL_DEVICE_MAX_PIPE_ARGS returned a non-zero value\n"); - return TEST_FAIL; - } + 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"); - if (u != 0) - { - log_error("CL_DEVICE_PIPE_SUPPORT returned CL_FALSE, but " - "CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS returned a " - "non-zero value\n"); - return TEST_FAIL; - } + 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"); - if (u != 0) - { - log_error( - "CL_DEVICE_PIPE_SUPPORT returned CL_FALSE, but " - "CL_DEVICE_PIPE_MAX_PACKET_SIZE returned a non-zero value\n"); - return TEST_FAIL; - } + 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. @@ -455,13 +415,17 @@ int test_consistency_pipes(cl_device_id deviceID, cl_context context, "clCreatePipe did not return CL_INVALID_OPERATION"); // clGetPipeInfo - // Returns CL_INVALID_MEM_OBJECT since pipe cannot be a valid pipe - // object. - error = clGetPipeInfo(mem, 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"); + // Returns CL_INVALID_OPERATION if no devices in the context associated + // with pipe support Pipes. + 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_assert_error(error == CL_INVALID_OPERATION, + "CL_DEVICE_PIPE_SUPPORT returned CL_FALSE but " + "clGetPipeInfo did not return CL_INVALID_OPERATION"); } else { @@ -473,13 +437,687 @@ int test_consistency_pipes(cl_device_id deviceID, cl_context context, sizeof(b), &b, NULL); test_error(error, "Unable to query CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT"); - if (b != CL_TRUE) - { - log_error("CL_DEVICE_PIPE_SUPPORT returned CL_TRUE but " - "CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT returned " - "CL_FALSE\n"); - return TEST_FAIL; - } + 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 (int 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 (int 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 (int 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;