mirror of
https://github.com/KhronosGroup/OpenCL-CTS.git
synced 2026-03-19 14:09:03 +00:00
26cacf056e
* Fail feature macro compare if compiler has more features than runtime Because a C++11 `std::equal` only iterates over the first container, and matches with items in the second, if the second container contains more items the check can still pass even though they're not identical. Just use `==` instead. Fixes #979 * Move an expression to its point of use
762 lines
26 KiB
C++
762 lines
26 KiB
C++
//
|
|
// 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)
|
|
{
|
|
cl_bool compilerAvailable = CL_FALSE;
|
|
cl_int error =
|
|
clGetDeviceInfo(deviceID, CL_DEVICE_COMPILER_AVAILABLE,
|
|
sizeof(compilerAvailable), &compilerAvailable, NULL);
|
|
test_error(error, "Unable to query CL_DEVICE_COMPILER_AVAILABLE");
|
|
if (compilerAvailable == CL_FALSE)
|
|
{
|
|
// Note: Not checking that the feature array is empty 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".
|
|
log_info("Skipping test - no compiler is available.\n");
|
|
return TEST_SKIPPED_ITSELF;
|
|
}
|
|
|
|
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;
|
|
}
|