Generic Address Space: Skip tests utlilising SVM on devices that do n… (#712)

* Harness: Added an option for a test to skip itself.

New condtion in callSingleTestFunction. This allows a test to return TEST_SKIP if the device does not support the requested test.

* Split generic_ptr_to_host_mem into two tests.

The old generic_ptr_to_host_mem test tests two different device capabilities (SVM and not SVM). This is a prerequisite for the following commit.

* Generic Address Space: Skip tests utlilising SVM on devices that do not support SVM.

Where a device does not support SVM, do not run the generic address space test(s) that rely on SVM.

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 */

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 );

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;
+}

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;
 }

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 ),
 };