Added test to verify required features coupling

test_conformance/compiler/main.cpp

@@ -91,6 +91,8 @@ test_definition test_list[] = {
     ADD_TEST_VERSION(pragma_unroll, Version(2, 0)),
 
     ADD_TEST_VERSION(features_macro, Version(3, 0)),
+    ADD_TEST(features_macro_coupling),
+
     ADD_TEST(unload_valid),
     // ADD_TEST(unload_invalid), // disabling temporarily, see GitHub #977
     ADD_TEST(unload_repeated),

test_conformance/compiler/procs.h

@@ -235,6 +235,10 @@ extern int test_pragma_unroll(cl_device_id deviceID, cl_context context,
                               cl_command_queue queue, int num_elements);
 extern int test_features_macro(cl_device_id deviceID, cl_context context,
                                cl_command_queue queue, int num_elements);
+extern int test_features_macro_coupling(cl_device_id deviceID,
+                                        cl_context context,
+                                        cl_command_queue queue,
+                                        int num_elements);
 extern int test_unload_valid(cl_device_id deviceID, cl_context context,
                              cl_command_queue queue, int num_elements);
 extern int test_unload_invalid(cl_device_id deviceID, cl_context context,

test_conformance/compiler/test_feature_macro.cpp

@@ -17,6 +17,7 @@
 #include <vector>
 #include <algorithm>
 #include "errorHelpers.h"
+#include "harness/featureHelpers.h"
 
 const char* macro_supported_source = R"(kernel void enabled(global int * buf) {
         int n = get_global_id(0);
@@ -833,3 +834,58 @@ int test_features_macro(cl_device_id deviceID, cl_context context,
 
     return error;
 }
+
+// This test checks that a supported feature comes with other required features
+int test_features_macro_coupling(cl_device_id deviceID, cl_context context,
+                                 cl_command_queue queue, int num_elements)
+{
+    log_info("  testing for required features coupling\n");
+
+    OpenCLCFeatures features;
+    int error = get_device_cl_c_features(deviceID, features);
+    if (error)
+    {
+        log_error("Couldn't query OpenCL C features for the device!\n");
+        return TEST_FAIL;
+    }
+
+    if (features.supports__opencl_c_3d_image_writes
+        && !features.supports__opencl_c_images)
+    {
+        log_error("OpenCL C compilers that define the feature macro "
+                  "__opencl_c_3d_image_writes must also define the feature "
+                  "macro __opencl_c_images!\n");
+        return TEST_FAIL;
+    }
+
+    if (features.supports__opencl_c_device_enqueue
+        && !(features.supports__opencl_c_program_scope_global_variables
+             && features.supports__opencl_c_generic_address_space))
+    {
+        log_error("OpenCL C compilers that define the feature macro "
+                  "__opencl_c_device_enqueue must also define "
+                  "__opencl_c_generic_address_space and "
+                  "__opencl_c_program_scope_global_variables!\n");
+        return TEST_FAIL;
+    }
+
+    if (features.supports__opencl_c_pipes
+        && !features.supports__opencl_c_generic_address_space)
+    {
+        log_error("OpenCL C compilers that define the feature macro "
+                  "__opencl_c_pipes must also define the feature macro "
+                  "__opencl_c_generic_address_space!\n");
+        return TEST_FAIL;
+    }
+
+    if (features.supports__opencl_c_read_write_images
+        && !features.supports__opencl_c_images)
+    {
+        log_error("OpenCL C compilers that define the feature macro "
+                  "__opencl_c_read_write_images must also define the feature "
+                  "macro __opencl_c_images!\n");
+        return TEST_FAIL;
+    }
+
+    return TEST_PASS;
+}