| // |
| // Copyright (c) 2020 The Khronos Group Inc. |
| // |
| // Licensed under the Apache License, Version 2.0 (the "License"); |
| // you may not use this file except in compliance with the License. |
| // You may obtain a copy of the License at |
| // |
| // http://www.apache.org/licenses/LICENSE-2.0 |
| // |
| // Unless required by applicable law or agreed to in writing, software |
| // distributed under the License is distributed on an "AS IS" BASIS, |
| // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. |
| // See the License for the specific language governing permissions and |
| // limitations under the License. |
| // |
| #include "testBase.h" |
| #include <vector> |
| #include <algorithm> |
| #include "errorHelpers.h" |
| |
| const char* macro_supported_source = R"(kernel void enabled(global int * buf) { |
| int n = get_global_id(0); |
| buf[n] = 0; |
| #ifndef %s |
| #error Feature macro was not defined |
| #endif |
| })"; |
| |
| const char* macro_not_supported_source = |
| R"(kernel void not_enabled(global int * buf) { |
| int n = get_global_id(0); |
| buf[n] = 0; |
| #ifdef %s |
| #error Feature macro was defined |
| #endif |
| })"; |
| |
| template <typename T> |
| cl_int check_api_feature_info_capabilities(cl_device_id deviceID, |
| cl_context context, cl_bool& status, |
| cl_device_info check_property, |
| cl_bitfield check_cap) |
| { |
| cl_int error = CL_SUCCESS; |
| T response; |
| error = clGetDeviceInfo(deviceID, check_property, sizeof(response), |
| &response, NULL); |
| test_error(error, "clGetDeviceInfo failed.\n"); |
| |
| if ((response & check_cap) == check_cap) |
| { |
| status = CL_TRUE; |
| } |
| else |
| { |
| status = CL_FALSE; |
| } |
| return error; |
| } |
| |
| cl_int check_api_feature_info_support(cl_device_id deviceID, cl_context context, |
| cl_bool& status, |
| cl_device_info check_property) |
| { |
| cl_int error = CL_SUCCESS; |
| cl_bool response; |
| error = clGetDeviceInfo(deviceID, check_property, sizeof(response), |
| &response, NULL); |
| test_error(error, "clGetDeviceInfo failed.\n"); |
| status = response; |
| return error; |
| } |
| |
| template <typename T> |
| cl_int check_api_feature_info_number(cl_device_id deviceID, cl_context context, |
| cl_bool& status, |
| cl_device_info check_property) |
| { |
| cl_int error = CL_SUCCESS; |
| T response; |
| error = clGetDeviceInfo(deviceID, check_property, sizeof(response), |
| &response, NULL); |
| test_error(error, "clGetDeviceInfo failed.\n"); |
| if (response > 0) |
| { |
| status = CL_TRUE; |
| } |
| else |
| { |
| status = CL_FALSE; |
| } |
| return error; |
| } |
| |
| cl_int check_api_feature_info_supported_image_formats(cl_device_id deviceID, |
| cl_context context, |
| cl_bool& status) |
| { |
| cl_int error = CL_SUCCESS; |
| cl_uint response = 0; |
| cl_uint image_format_count; |
| error = clGetSupportedImageFormats(context, CL_MEM_WRITE_ONLY, |
| CL_MEM_OBJECT_IMAGE3D, 0, NULL, |
| &image_format_count); |
| test_error(error, "clGetSupportedImageFormats failed"); |
| response += image_format_count; |
| error = clGetSupportedImageFormats(context, CL_MEM_READ_WRITE, |
| CL_MEM_OBJECT_IMAGE3D, 0, NULL, |
| &image_format_count); |
| test_error(error, "clGetSupportedImageFormats failed"); |
| response += image_format_count; |
| error = clGetSupportedImageFormats(context, CL_MEM_KERNEL_READ_AND_WRITE, |
| CL_MEM_OBJECT_IMAGE3D, 0, NULL, |
| &image_format_count); |
| test_error(error, "clGetSupportedImageFormats failed"); |
| response += image_format_count; |
| if (response > 0) |
| { |
| status = CL_TRUE; |
| } |
| else |
| { |
| status = CL_FALSE; |
| } |
| return error; |
| } |
| |
| cl_int check_compiler_feature_info(cl_device_id deviceID, cl_context context, |
| std::string feature_macro, cl_bool& status) |
| { |
| cl_int error = CL_SUCCESS; |
| clProgramWrapper program_supported; |
| clProgramWrapper program_not_supported; |
| char kernel_supported_src[1024]; |
| char kernel_not_supported_src[1024]; |
| sprintf(kernel_supported_src, macro_supported_source, |
| feature_macro.c_str()); |
| const char* ptr_supported = kernel_supported_src; |
| const char* build_options = "-cl-std=CL3.0"; |
| |
| error = create_single_kernel_helper_create_program( |
| context, &program_supported, 1, &ptr_supported, build_options); |
| test_error(error, "create_single_kernel_helper_create_program failed.\n"); |
| |
| sprintf(kernel_not_supported_src, macro_not_supported_source, |
| feature_macro.c_str()); |
| const char* ptr_not_supported = kernel_not_supported_src; |
| error = create_single_kernel_helper_create_program( |
| context, &program_not_supported, 1, &ptr_not_supported, |
| "-cl-std=CL3.0"); |
| test_error(error, "create_single_kernel_helper_create_program failed.\n"); |
| |
| cl_int status_supported = CL_SUCCESS; |
| cl_int status_not_supported = CL_SUCCESS; |
| status_supported = clBuildProgram(program_supported, 1, &deviceID, |
| build_options, NULL, NULL); |
| status_not_supported = clBuildProgram(program_not_supported, 1, &deviceID, |
| build_options, NULL, NULL); |
| if (status_supported != status_not_supported) |
| { |
| if (status_not_supported == CL_SUCCESS) |
| { |
| // kernel which verifies not supporting return passed |
| status = CL_FALSE; |
| } |
| else |
| { |
| // kernel which verifies supporting return passed |
| status = CL_TRUE; |
| } |
| } |
| else |
| { |
| log_error("Error: The macro feature is defined and undefined " |
| "in the same time\n"); |
| error = OutputBuildLogs(program_supported, 1, &deviceID); |
| test_error(error, "OutputBuildLogs failed.\n"); |
| error = OutputBuildLogs(program_not_supported, 1, &deviceID); |
| test_error(error, "OutputBuildLogs failed.\n"); |
| return TEST_FAIL; |
| } |
| return error; |
| } |
| |
| int feature_macro_verify_results(std::string test_macro_name, |
| cl_bool api_status, cl_bool compiler_status, |
| cl_bool& supported) |
| { |
| cl_int error = TEST_PASS; |
| log_info("Feature status: API - %s, compiler - %s\n", |
| api_status == CL_TRUE ? "supported" : "not supported", |
| compiler_status == CL_TRUE ? "supported" : "not supported"); |
| if (api_status != compiler_status) |
| { |
| log_info("%s - failed\n", test_macro_name.c_str()); |
| supported = CL_FALSE; |
| return TEST_FAIL; |
| } |
| else |
| { |
| log_info("%s - passed\n", test_macro_name.c_str()); |
| } |
| supported = api_status; |
| return error; |
| } |
| |
| int test_feature_macro_atomic_order_acq_rel(cl_device_id deviceID, |
| cl_context context, |
| std::string test_macro_name, |
| cl_bool& supported) |
| { |
| cl_int error = TEST_FAIL; |
| cl_bool api_status; |
| cl_bool compiler_status; |
| log_info("\n%s ...\n", test_macro_name.c_str()); |
| error = check_api_feature_info_capabilities<cl_device_atomic_capabilities>( |
| deviceID, context, api_status, CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES, |
| CL_DEVICE_ATOMIC_ORDER_ACQ_REL); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| error = check_compiler_feature_info(deviceID, context, test_macro_name, |
| compiler_status); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| return feature_macro_verify_results(test_macro_name, api_status, |
| compiler_status, supported); |
| } |
| |
| int test_feature_macro_atomic_order_seq_cst(cl_device_id deviceID, |
| cl_context context, |
| std::string test_macro_name, |
| cl_bool& supported) |
| { |
| cl_int error = TEST_FAIL; |
| cl_bool api_status; |
| cl_bool compiler_status; |
| log_info("\n%s ...\n", test_macro_name.c_str()); |
| |
| error = check_api_feature_info_capabilities<cl_device_atomic_capabilities>( |
| deviceID, context, api_status, CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES, |
| CL_DEVICE_ATOMIC_ORDER_SEQ_CST); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| error = check_compiler_feature_info(deviceID, context, test_macro_name, |
| compiler_status); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| return feature_macro_verify_results(test_macro_name, api_status, |
| compiler_status, supported); |
| } |
| |
| int test_feature_macro_atomic_scope_device(cl_device_id deviceID, |
| cl_context context, |
| std::string test_macro_name, |
| cl_bool& supported) |
| { |
| cl_int error = TEST_FAIL; |
| cl_bool api_status; |
| cl_bool compiler_status; |
| log_info("\n%s ...\n", test_macro_name.c_str()); |
| error = check_api_feature_info_capabilities<cl_device_atomic_capabilities>( |
| deviceID, context, api_status, CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES, |
| CL_DEVICE_ATOMIC_SCOPE_DEVICE); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| error = check_compiler_feature_info(deviceID, context, test_macro_name, |
| compiler_status); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| return feature_macro_verify_results(test_macro_name, api_status, |
| compiler_status, supported); |
| } |
| |
| int test_feature_macro_atomic_scope_all_devices(cl_device_id deviceID, |
| cl_context context, |
| std::string test_macro_name, |
| cl_bool& supported) |
| { |
| cl_int error = TEST_FAIL; |
| cl_bool api_status; |
| cl_bool compiler_status; |
| log_info("\n%s ...\n", test_macro_name.c_str()); |
| error = check_api_feature_info_capabilities<cl_device_atomic_capabilities>( |
| deviceID, context, api_status, CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES, |
| CL_DEVICE_ATOMIC_SCOPE_ALL_DEVICES); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| error = check_compiler_feature_info(deviceID, context, test_macro_name, |
| compiler_status); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| return feature_macro_verify_results(test_macro_name, api_status, |
| compiler_status, supported); |
| } |
| |
| int test_feature_macro_3d_image_writes(cl_device_id deviceID, |
| cl_context context, |
| std::string test_macro_name, |
| cl_bool& supported) |
| { |
| cl_int error = TEST_FAIL; |
| cl_bool api_status; |
| cl_bool compiler_status; |
| log_info("\n%s ...\n", test_macro_name.c_str()); |
| error = check_api_feature_info_supported_image_formats(deviceID, context, |
| api_status); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| error = check_compiler_feature_info(deviceID, context, test_macro_name, |
| compiler_status); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| return feature_macro_verify_results(test_macro_name, api_status, |
| compiler_status, supported); |
| } |
| |
| int test_feature_macro_device_enqueue(cl_device_id deviceID, cl_context context, |
| std::string test_macro_name, |
| cl_bool& supported) |
| { |
| cl_int error = TEST_FAIL; |
| cl_bool api_status; |
| cl_bool compiler_status; |
| log_info("\n%s ...\n", test_macro_name.c_str()); |
| error = check_api_feature_info_capabilities< |
| cl_device_device_enqueue_capabilities>( |
| deviceID, context, api_status, CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES, |
| CL_DEVICE_QUEUE_SUPPORTED); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| error = check_compiler_feature_info(deviceID, context, test_macro_name, |
| compiler_status); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| return feature_macro_verify_results(test_macro_name, api_status, |
| compiler_status, supported); |
| } |
| |
| int test_feature_macro_generic_address_space(cl_device_id deviceID, |
| cl_context context, |
| std::string test_macro_name, |
| cl_bool& supported) |
| { |
| cl_int error = TEST_FAIL; |
| cl_bool api_status; |
| cl_bool compiler_status; |
| log_info("\n%s ...\n", test_macro_name.c_str()); |
| error = check_api_feature_info_support( |
| deviceID, context, api_status, CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| error = check_compiler_feature_info(deviceID, context, test_macro_name, |
| compiler_status); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| return feature_macro_verify_results(test_macro_name, api_status, |
| compiler_status, supported); |
| } |
| |
| int test_feature_macro_pipes(cl_device_id deviceID, cl_context context, |
| std::string test_macro_name, cl_bool& supported) |
| { |
| cl_int error = TEST_FAIL; |
| cl_bool api_status; |
| cl_bool compiler_status; |
| log_info("\n%s ...\n", test_macro_name.c_str()); |
| error = check_api_feature_info_support(deviceID, context, api_status, |
| CL_DEVICE_PIPE_SUPPORT); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| error = check_compiler_feature_info(deviceID, context, test_macro_name, |
| compiler_status); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| return feature_macro_verify_results(test_macro_name, api_status, |
| compiler_status, supported); |
| } |
| |
| int test_feature_macro_program_scope_global_variables( |
| cl_device_id deviceID, cl_context context, std::string test_macro_name, |
| cl_bool& supported) |
| { |
| cl_int error = TEST_FAIL; |
| cl_bool api_status; |
| cl_bool compiler_status; |
| log_info("\n%s ...\n", test_macro_name.c_str()); |
| error = check_api_feature_info_number<size_t>( |
| deviceID, context, api_status, CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| error = check_compiler_feature_info(deviceID, context, test_macro_name, |
| compiler_status); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| return feature_macro_verify_results(test_macro_name, api_status, |
| compiler_status, supported); |
| } |
| |
| int test_feature_macro_read_write_images(cl_device_id deviceID, |
| cl_context context, |
| std::string test_macro_name, |
| cl_bool& supported) |
| { |
| cl_int error = TEST_FAIL; |
| cl_bool api_status; |
| cl_bool compiler_status; |
| log_info("\n%s ...\n", test_macro_name.c_str()); |
| error = check_api_feature_info_number<cl_uint>( |
| deviceID, context, api_status, CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| error = check_compiler_feature_info(deviceID, context, test_macro_name, |
| compiler_status); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| return feature_macro_verify_results(test_macro_name, api_status, |
| compiler_status, supported); |
| } |
| |
| int test_feature_macro_subgroups(cl_device_id deviceID, cl_context context, |
| std::string test_macro_name, |
| cl_bool& supported) |
| { |
| cl_int error = TEST_FAIL; |
| cl_bool api_status; |
| cl_bool compiler_status; |
| log_info("\n%s ...\n", test_macro_name.c_str()); |
| error = check_api_feature_info_number<cl_uint>( |
| deviceID, context, api_status, CL_DEVICE_MAX_NUM_SUB_GROUPS); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| error = check_compiler_feature_info(deviceID, context, test_macro_name, |
| compiler_status); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| return feature_macro_verify_results(test_macro_name, api_status, |
| compiler_status, supported); |
| } |
| |
| int test_feature_macro_work_group_collective_functions( |
| cl_device_id deviceID, cl_context context, std::string test_macro_name, |
| cl_bool& supported) |
| { |
| cl_int error = TEST_FAIL; |
| cl_bool api_status; |
| cl_bool compiler_status; |
| log_info("\n%s ...\n", test_macro_name.c_str()); |
| error = check_api_feature_info_support( |
| deviceID, context, api_status, |
| CL_DEVICE_WORK_GROUP_COLLECTIVE_FUNCTIONS_SUPPORT); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| error = check_compiler_feature_info(deviceID, context, test_macro_name, |
| compiler_status); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| return feature_macro_verify_results(test_macro_name, api_status, |
| compiler_status, supported); |
| } |
| |
| int test_feature_macro_images(cl_device_id deviceID, cl_context context, |
| std::string test_macro_name, cl_bool& supported) |
| { |
| cl_int error = TEST_FAIL; |
| cl_bool api_status; |
| cl_bool compiler_status; |
| log_info("\n%s ...\n", test_macro_name.c_str()); |
| error = check_api_feature_info_support(deviceID, context, api_status, |
| CL_DEVICE_IMAGE_SUPPORT); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| error = check_compiler_feature_info(deviceID, context, test_macro_name, |
| compiler_status); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| return feature_macro_verify_results(test_macro_name, api_status, |
| compiler_status, supported); |
| } |
| |
| int test_feature_macro_fp64(cl_device_id deviceID, cl_context context, |
| std::string test_macro_name, cl_bool& supported) |
| { |
| cl_int error = TEST_FAIL; |
| cl_bool api_status; |
| cl_bool compiler_status; |
| log_info("\n%s ...\n", test_macro_name.c_str()); |
| error = check_api_feature_info_capabilities<cl_device_fp_config>( |
| deviceID, context, api_status, CL_DEVICE_DOUBLE_FP_CONFIG, |
| CL_FP_FMA | CL_FP_ROUND_TO_NEAREST | CL_FP_INF_NAN | CL_FP_DENORM); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| error = check_compiler_feature_info(deviceID, context, test_macro_name, |
| compiler_status); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| return feature_macro_verify_results(test_macro_name, api_status, |
| compiler_status, supported); |
| } |
| |
| int test_feature_macro_int64(cl_device_id deviceID, cl_context context, |
| std::string test_macro_name, cl_bool& supported) |
| { |
| cl_int error = TEST_FAIL; |
| cl_bool api_status; |
| cl_bool compiler_status; |
| cl_int full_profile = 0; |
| log_info("\n%s ...\n", test_macro_name.c_str()); |
| size_t ret_len; |
| char profile[32] = { 0 }; |
| error = clGetDeviceInfo(deviceID, CL_DEVICE_PROFILE, sizeof(profile), |
| profile, &ret_len); |
| test_error(error, "clGetDeviceInfo(CL_DEVICE_PROFILE) failed"); |
| if (ret_len < sizeof(profile) && strcmp(profile, "FULL_PROFILE") == 0) |
| { |
| full_profile = 1; |
| } |
| else if (ret_len < sizeof(profile) |
| && strcmp(profile, "EMBEDDED_PROFILE") == 0) |
| { |
| full_profile = 0; |
| } |
| else |
| { |
| log_error("Unknown device profile: %s\n", profile); |
| return TEST_FAIL; |
| } |
| |
| if (full_profile) |
| { |
| api_status = CL_TRUE; |
| } |
| else |
| { |
| if (is_extension_available(deviceID, "cles_khr_int64")) |
| { |
| api_status = CL_TRUE; |
| } |
| else |
| { |
| cl_bool double_supported = CL_FALSE; |
| error = check_api_feature_info_capabilities<cl_device_fp_config>( |
| deviceID, context, double_supported, CL_DEVICE_DOUBLE_FP_CONFIG, |
| CL_FP_FMA | CL_FP_ROUND_TO_NEAREST | CL_FP_INF_NAN |
| | CL_FP_DENORM); |
| test_error(error, "checking CL_DEVICE_DOUBLE_FP_CONFIG failed"); |
| if (double_supported == CL_FALSE) |
| { |
| api_status = CL_FALSE; |
| } |
| else |
| { |
| log_error("FP double type is supported and cles_khr_int64 " |
| "extension not supported\n"); |
| return TEST_FAIL; |
| } |
| } |
| } |
| |
| error = check_compiler_feature_info(deviceID, context, test_macro_name, |
| compiler_status); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| return feature_macro_verify_results(test_macro_name, api_status, |
| compiler_status, supported); |
| } |
| |
| int test_consistency_c_features_list(cl_device_id deviceID, |
| std::vector<std::string> vec_to_cmp) |
| { |
| log_info("\nComparison list of features: CL_DEVICE_OPENCL_C_FEATURES vs " |
| "API/compiler queries.\n"); |
| cl_int error; |
| size_t config_size; |
| std::vector<cl_name_version> vec_device_feature; |
| std::vector<std::string> vec_device_feature_names; |
| error = clGetDeviceInfo(deviceID, CL_DEVICE_OPENCL_C_FEATURES, 0, NULL, |
| &config_size); |
| |
| test_error( |
| error, |
| "clGetDeviceInfo asking for CL_DEVICE_OPENCL_C_FEATURES failed.\n"); |
| if (config_size == 0) |
| { |
| log_info("Empty list of CL_DEVICE_OPENCL_C_FEATURES returned by " |
| "clGetDeviceInfo on this device.\n"); |
| } |
| else |
| { |
| int vec_elements = config_size / sizeof(cl_name_version); |
| vec_device_feature.resize(vec_elements); |
| error = clGetDeviceInfo(deviceID, CL_DEVICE_OPENCL_C_FEATURES, |
| config_size, vec_device_feature.data(), 0); |
| test_error( |
| error, |
| "clGetDeviceInfo asking for CL_DEVICE_OPENCL_C_FEATURES failed.\n"); |
| } |
| for (auto each_f : vec_device_feature) |
| { |
| vec_device_feature_names.push_back(each_f.name); |
| } |
| sort(vec_to_cmp.begin(), vec_to_cmp.end()); |
| sort(vec_device_feature_names.begin(), vec_device_feature_names.end()); |
| |
| if (vec_device_feature_names == vec_to_cmp) |
| { |
| log_info("Comparison list of features - passed\n"); |
| } |
| else |
| { |
| log_info("Comparison list of features - failed\n"); |
| error = TEST_FAIL; |
| } |
| log_info( |
| "Supported features based on CL_DEVICE_OPENCL_C_FEATURES API query:\n"); |
| for (auto each_f : vec_device_feature_names) |
| { |
| log_info("%s\n", each_f.c_str()); |
| } |
| |
| log_info("\nSupported features based on queries to API/compiler :\n"); |
| for (auto each_f : vec_to_cmp) |
| { |
| log_info("%s\n", each_f.c_str()); |
| } |
| |
| return error; |
| } |
| |
| #define NEW_FEATURE_MACRO_TEST(feat) \ |
| test_macro_name = "__opencl_c_" #feat; \ |
| error |= test_feature_macro_##feat(deviceID, context, test_macro_name, \ |
| supported); \ |
| if (supported) supported_features_vec.push_back(test_macro_name); |
| |
| |
| int test_features_macro(cl_device_id deviceID, cl_context context, |
| cl_command_queue queue, int num_elements) |
| { |
| |
| // Note: Not checking that the feature array is empty for the compiler not |
| // available case because the specification says "For devices that do not |
| // support compilation from OpenCL C source, this query may return an empty |
| // array." It "may" return an empty array implies that an implementation |
| // also "may not". |
| check_compiler_available(deviceID); |
| |
| int error = TEST_PASS; |
| cl_bool supported = CL_FALSE; |
| std::string test_macro_name = ""; |
| std::vector<std::string> supported_features_vec; |
| NEW_FEATURE_MACRO_TEST(program_scope_global_variables); |
| NEW_FEATURE_MACRO_TEST(3d_image_writes); |
| NEW_FEATURE_MACRO_TEST(atomic_order_acq_rel); |
| NEW_FEATURE_MACRO_TEST(atomic_order_seq_cst); |
| NEW_FEATURE_MACRO_TEST(atomic_scope_device); |
| NEW_FEATURE_MACRO_TEST(atomic_scope_all_devices); |
| NEW_FEATURE_MACRO_TEST(device_enqueue); |
| NEW_FEATURE_MACRO_TEST(generic_address_space); |
| NEW_FEATURE_MACRO_TEST(pipes); |
| NEW_FEATURE_MACRO_TEST(read_write_images); |
| NEW_FEATURE_MACRO_TEST(subgroups); |
| NEW_FEATURE_MACRO_TEST(work_group_collective_functions); |
| NEW_FEATURE_MACRO_TEST(images); |
| NEW_FEATURE_MACRO_TEST(fp64); |
| NEW_FEATURE_MACRO_TEST(int64); |
| |
| error |= test_consistency_c_features_list(deviceID, supported_features_vec); |
| |
| return error; |
| } |
