diff --git a/test_common/harness/testHarness.cpp b/test_common/harness/testHarness.cpp index c81acc2d..856ad75d 100644 --- a/test_common/harness/testHarness.cpp +++ b/test_common/harness/testHarness.cpp @@ -759,6 +759,12 @@ test_status callSingleTestFunction( test_definition test, cl_device_id deviceToU log_info("%s test currently not implemented\n", test.name); status = TEST_SKIP; } + else if (ret == TEST_SKIPPED_ITSELF) + { + /* Tests can also let us know they're not supported by the implementation */ + log_info("%s test not supported\n", test.name); + status = TEST_SKIP; + } else { /* Print result */ diff --git a/test_common/harness/threadTesting.h b/test_common/harness/threadTesting.h index 81a5757b..71d57973 100644 --- a/test_common/harness/threadTesting.h +++ b/test_common/harness/threadTesting.h @@ -23,6 +23,7 @@ #endif #define TEST_NOT_IMPLEMENTED -99 +#define TEST_SKIPPED_ITSELF -100 typedef int (*basefn)(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); extern int test_threaded_function( basefn fnToTest, cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); diff --git a/test_conformance/generic_address_space/advanced_tests.cpp b/test_conformance/generic_address_space/advanced_tests.cpp index f4025fef..8f796097 100644 --- a/test_conformance/generic_address_space/advanced_tests.cpp +++ b/test_conformance/generic_address_space/advanced_tests.cpp @@ -947,85 +947,51 @@ int test_generic_advanced_casting(cl_device_id deviceID, cl_context context, cl_ return test.Execute(deviceID, context, queue, num_elements); } -int test_generic_ptr_to_host_mem(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { +int test_generic_ptr_to_host_mem_svm(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { cl_int result = CL_SUCCESS; - const std::string GLOBAL_KERNEL_FUNCTION = common::CONFORMANCE_VERIFY_FENCE + - NL - NL "bool helperFunction(uint *ptr, uint tid) {" - NL " if (!isFenceValid(get_fence(ptr)))" - NL " return false;" - NL - NL " if (*ptr != tid)" - NL " return false;" - NL - NL " return true;" - NL "}" - NL - NL "__kernel void testKernel(__global uint *results, __global uint *buf) {" - NL " uint tid = get_global_id(0);" - NL - NL " results[tid] = helperFunction(&buf[tid], tid);" - NL "}" - NL; - - const std::string LOCAL_KERNEL_FUNCTION = common::CONFORMANCE_VERIFY_FENCE + - NL - NL "bool helperFunction(uint *ptr, uint tid) {" - NL " if (!isFenceValid(get_fence(ptr)))" - NL " return false;" - NL - NL " if (*ptr != tid)" - NL " return false;" - NL - NL " return true;" - NL "}" - NL - NL "__kernel void testKernel(__global uint *results, __local uint *buf) {" - NL " uint tid = get_global_id(0);" - NL " if (get_local_id(0) == 0) {" - NL " for (uint i = 0; i < get_local_size(0); ++i) {" - NL " uint idx = get_local_size(0) * get_group_id(0) + i;" - NL " buf[idx] = idx;" - NL " }" - NL " }" - NL - NL " work_group_barrier(CLK_LOCAL_MEM_FENCE);" - NL " results[tid] = helperFunction(&buf[tid], tid);" - NL "}" - NL; - - CAdvancedTest test_global_ptr(GLOBAL_KERNEL_FUNCTION, ARG_TYPE_HOST_PTR); - result |= test_global_ptr.Execute(deviceID, context, queue, num_elements); - - CAdvancedTest test_local_ptr(LOCAL_KERNEL_FUNCTION, ARG_TYPE_HOST_LOCAL); - result |= test_local_ptr.Execute(deviceID, context, queue, num_elements / 64); - /* Test SVM capabilities and select matching tests */ cl_device_svm_capabilities caps; + auto version = get_device_cl_version(deviceID); + auto expected_min_version = Version(2, 0); cl_int error = clGetDeviceInfo(deviceID, CL_DEVICE_SVM_CAPABILITIES, sizeof(caps), &caps, NULL); test_error(error, "clGetDeviceInfo(CL_DEVICE_SVM_CAPABILITIES) failed"); + if ((version < expected_min_version) || (version > Version(2,2) && caps == 0)) + return TEST_SKIPPED_ITSELF; + if (caps & CL_DEVICE_SVM_COARSE_GRAIN_BUFFER) { - CAdvancedTest test_global_svm_ptr(GLOBAL_KERNEL_FUNCTION, ARG_TYPE_COARSE_GRAINED_SVM); + CAdvancedTest test_global_svm_ptr(common::GLOBAL_KERNEL_FUNCTION, ARG_TYPE_COARSE_GRAINED_SVM); result |= test_global_svm_ptr.Execute(deviceID, context, queue, num_elements); } if (caps & CL_DEVICE_SVM_FINE_GRAIN_BUFFER) { - CAdvancedTest test_global_svm_ptr(GLOBAL_KERNEL_FUNCTION, ARG_TYPE_FINE_GRAINED_BUFFER_SVM); + CAdvancedTest test_global_svm_ptr(common::GLOBAL_KERNEL_FUNCTION, ARG_TYPE_FINE_GRAINED_BUFFER_SVM); result |= test_global_svm_ptr.Execute(deviceID, context, queue, num_elements); } if (caps & CL_DEVICE_SVM_FINE_GRAIN_SYSTEM) { - CAdvancedTest test_global_svm_ptr(GLOBAL_KERNEL_FUNCTION, ARG_TYPE_FINE_GRAINED_SYSTEM_SVM); + CAdvancedTest test_global_svm_ptr(common::GLOBAL_KERNEL_FUNCTION, ARG_TYPE_FINE_GRAINED_SYSTEM_SVM); result |= test_global_svm_ptr.Execute(deviceID, context, queue, num_elements); } if (caps & CL_DEVICE_SVM_ATOMICS) { - CAdvancedTest test_global_svm_ptr(GLOBAL_KERNEL_FUNCTION, ARG_TYPE_ATOMICS_SVM); + CAdvancedTest test_global_svm_ptr(common::GLOBAL_KERNEL_FUNCTION, ARG_TYPE_ATOMICS_SVM); result |= test_global_svm_ptr.Execute(deviceID, context, queue, num_elements); } return result; } + +int test_generic_ptr_to_host_mem(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { + cl_int result = CL_SUCCESS; + + CAdvancedTest test_global_ptr(common::GLOBAL_KERNEL_FUNCTION, ARG_TYPE_HOST_PTR); + result |= test_global_ptr.Execute(deviceID, context, queue, num_elements); + + CAdvancedTest test_local_ptr(common::LOCAL_KERNEL_FUNCTION, ARG_TYPE_HOST_LOCAL); + result |= test_local_ptr.Execute(deviceID, context, queue, num_elements / 64); + + return result; +} diff --git a/test_conformance/generic_address_space/base.h b/test_conformance/generic_address_space/base.h index 44d61fef..a8fbabdd 100644 --- a/test_conformance/generic_address_space/base.h +++ b/test_conformance/generic_address_space/base.h @@ -36,4 +36,49 @@ namespace common { NL " return false;" NL "}" NL; + + static std::string GLOBAL_KERNEL_FUNCTION = CONFORMANCE_VERIFY_FENCE + + NL + NL "bool helperFunction(uint *ptr, uint tid) {" + NL " if (!isFenceValid(get_fence(ptr)))" + NL " return false;" + NL + NL " if (*ptr != tid)" + NL " return false;" + NL + NL " return true;" + NL "}" + NL + NL "__kernel void testKernel(__global uint *results, __global uint *buf) {" + NL " uint tid = get_global_id(0);" + NL + NL " results[tid] = helperFunction(&buf[tid], tid);" + NL "}" + NL; + + static std::string LOCAL_KERNEL_FUNCTION = CONFORMANCE_VERIFY_FENCE + + NL + NL "bool helperFunction(uint *ptr, uint tid) {" + NL " if (!isFenceValid(get_fence(ptr)))" + NL " return false;" + NL + NL " if (*ptr != tid)" + NL " return false;" + NL + NL " return true;" + NL "}" + NL + NL "__kernel void testKernel(__global uint *results, __local uint *buf) {" + NL " uint tid = get_global_id(0);" + NL " if (get_local_id(0) == 0) {" + NL " for (uint i = 0; i < get_local_size(0); ++i) {" + NL " uint idx = get_local_size(0) * get_group_id(0) + i;" + NL " buf[idx] = idx;" + NL " }" + NL " }" + NL + NL " work_group_barrier(CLK_LOCAL_MEM_FENCE);" + NL " results[tid] = helperFunction(&buf[tid], tid);" + NL "}" + NL; } diff --git a/test_conformance/generic_address_space/main.cpp b/test_conformance/generic_address_space/main.cpp index 298f4d4f..5d648983 100644 --- a/test_conformance/generic_address_space/main.cpp +++ b/test_conformance/generic_address_space/main.cpp @@ -38,6 +38,7 @@ extern int test_generic_variable_gentype(cl_device_id deviceID, cl_context conte extern int test_builtin_functions(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); extern int test_generic_advanced_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); extern int test_generic_ptr_to_host_mem(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); +extern int test_generic_ptr_to_host_mem_svm(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); extern int test_max_number_of_params(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); test_definition test_list[] = { @@ -62,6 +63,7 @@ test_definition test_list[] = { ADD_TEST( builtin_functions ), ADD_TEST( generic_advanced_casting ), ADD_TEST( generic_ptr_to_host_mem ), + ADD_TEST( generic_ptr_to_host_mem_svm ), ADD_TEST( max_number_of_params ), };