diff --git a/test_conformance/compiler/main.cpp b/test_conformance/compiler/main.cpp index f0a9ef3a..e31300c8 100644 --- a/test_conformance/compiler/main.cpp +++ b/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), diff --git a/test_conformance/compiler/procs.h b/test_conformance/compiler/procs.h index 10ae142b..4a425ffe 100644 --- a/test_conformance/compiler/procs.h +++ b/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, diff --git a/test_conformance/compiler/test_feature_macro.cpp b/test_conformance/compiler/test_feature_macro.cpp index 7858c3c2..2f973eec 100644 --- a/test_conformance/compiler/test_feature_macro.cpp +++ b/test_conformance/compiler/test_feature_macro.cpp @@ -17,6 +17,7 @@ #include #include #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,56 @@ 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) +{ + 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; +}