ahmed/OpenCL-CTS
1
0
Fork 0
mirror of https://github.com/KhronosGroup/OpenCL-CTS.git synced 2026-03-19 06:09:01 +00:00
Code Issues Packages Projects Releases Wiki Activity

Migrate generic_address_space suite to the new test registration framework (#2333)

Browse Source 
Contributes to https://github.com/KhronosGroup/OpenCL-CTS/issues/2181.

Signed-off-by: Michael Rizkalla <michael.rizkalla@arm.com>
This commit is contained in:
Michael Rizkalla
2025-03-18 17:11:43 +00:00
committed by GitHub
parent 2131aff0e9
commit 7cd13e1e93
5 changed files with 97 additions and 123 deletions
Download Patch File Download Diff File Expand all files Collapse all files

67
test_conformance/generic_address_space/advanced_tests.cpp
View File

@@ -293,7 +293,8 @@ private:
const ExtraKernelArgMemType _extraKernelArgMemType;
};
int test_library_function(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(library_function)
{
const std::string LIBRARY_FUNCTION = common::CONFORMANCE_VERIFY_FENCE +
NL
NL "bool helperFunction(float *floatp, float val) {"
@@ -340,10 +341,11 @@ __kernel void testKernel(__global uint *results) {
CAdvancedTest test(LIBRARY_FUNCTION, KERNEL_FUNCTION);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_generic_variable_volatile(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(generic_variable_volatile)
{
std::vector<std::string> KERNEL_FUNCTIONS;
KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE +
@@ -420,10 +422,11 @@ int test_generic_variable_volatile(cl_device_id deviceID, cl_context context, cl
CAdvancedTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_generic_variable_const(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(generic_variable_const)
{
std::vector<std::string> KERNEL_FUNCTIONS;
KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE +
@@ -474,10 +477,11 @@ int test_generic_variable_const(cl_device_id deviceID, cl_context context, cl_co
CAdvancedTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_generic_variable_gentype(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(generic_variable_gentype)
{
const std::string KERNEL_FUNCTION_TEMPLATE = common::CONFORMANCE_VERIFY_FENCE +
NL
NL "%s"
@@ -531,7 +535,8 @@ int test_generic_variable_gentype(cl_device_id deviceID, cl_context context, cl_
const std::string cl_khr_fp64_pragma = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable";
// Add double floating types if they are supported
if (is_extension_available(deviceID, "cl_khr_fp64")) {
if (is_extension_available(device, "cl_khr_fp64"))
{
for (size_t j = 0; j < sizeof(vector_sizes) / sizeof(vector_sizes[0]); j++) {
for (size_t k = 0; k < sizeof(address_spaces) / sizeof(address_spaces[0]); k++) {
char temp_kernel[1024];
@@ -551,7 +556,8 @@ int test_generic_variable_gentype(cl_device_id deviceID, cl_context context, cl_
const std::string cl_khr_fp16_pragma = "#pragma OPENCL EXTENSION cl_khr_fp16 : enable";
// Add half floating types if they are supported
if (is_extension_available(deviceID, "cl_khr_fp16")) {
if (is_extension_available(device, "cl_khr_fp16"))
{
for (size_t j = 0; j < sizeof(vector_sizes) / sizeof(vector_sizes[0]); j++) {
for (size_t k = 0; k < sizeof(address_spaces) / sizeof(address_spaces[0]); k++) {
char temp_kernel[1024];
@@ -586,7 +592,7 @@ int test_generic_variable_gentype(cl_device_id deviceID, cl_context context, cl_
CAdvancedTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
void create_math_kernels(std::vector<std::string>& KERNEL_FUNCTIONS) {
@@ -919,19 +925,21 @@ void create_vstore_kernels(std::vector<std::string>& KERNEL_FUNCTIONS, cl_device
}
}
int test_builtin_functions(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(builtin_functions)
{
std::vector<std::string> KERNEL_FUNCTIONS;
create_math_kernels(KERNEL_FUNCTIONS);
create_vload_kernels(KERNEL_FUNCTIONS, deviceID);
create_vstore_kernels(KERNEL_FUNCTIONS, deviceID);
create_vload_kernels(KERNEL_FUNCTIONS, device);
create_vstore_kernels(KERNEL_FUNCTIONS, device);
CAdvancedTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_generic_advanced_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(generic_advanced_casting)
{
std::vector<std::string> KERNEL_FUNCTIONS;
KERNEL_FUNCTIONS.push_back(
@@ -980,18 +988,20 @@ int test_generic_advanced_casting(cl_device_id deviceID, cl_context context, cl_
CAdvancedTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_generic_ptr_to_host_mem_svm(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(generic_ptr_to_host_mem_svm)
{
cl_int result = CL_SUCCESS;
/* Test SVM capabilities and select matching tests */
cl_device_svm_capabilities caps;
auto version = get_device_cl_version(deviceID);
auto version = get_device_cl_version(device);
auto expected_min_version = Version(2, 0);
cl_int error = clGetDeviceInfo(deviceID, CL_DEVICE_SVM_CAPABILITIES, sizeof(caps), &caps, NULL);
cl_int error = clGetDeviceInfo(device, CL_DEVICE_SVM_CAPABILITIES,
sizeof(caps), &caps, NULL);
test_error(error, "clGetDeviceInfo(CL_DEVICE_SVM_CAPABILITIES) failed");
if ((version < expected_min_version)
@@ -1000,35 +1010,40 @@ int test_generic_ptr_to_host_mem_svm(cl_device_id deviceID, cl_context context,
if (caps & CL_DEVICE_SVM_COARSE_GRAIN_BUFFER) {
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);
result |=
test_global_svm_ptr.Execute(device, context, queue, num_elements);
}
if (caps & CL_DEVICE_SVM_FINE_GRAIN_BUFFER) {
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);
result |=
test_global_svm_ptr.Execute(device, context, queue, num_elements);
}
if (caps & CL_DEVICE_SVM_FINE_GRAIN_SYSTEM) {
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);
result |=
test_global_svm_ptr.Execute(device, context, queue, num_elements);
}
if (caps & CL_DEVICE_SVM_ATOMICS) {
CAdvancedTest test_global_svm_ptr(common::GLOBAL_KERNEL_FUNCTION, ARG_TYPE_ATOMICS_SVM);
result |= test_global_svm_ptr.Execute(deviceID, context, queue, num_elements);
result |=
test_global_svm_ptr.Execute(device, 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) {
REGISTER_TEST(generic_ptr_to_host_mem)
{
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);
result |= test_global_ptr.Execute(device, 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);
result |= test_local_ptr.Execute(device, context, queue, num_elements / 64);
return result;
}

14
test_conformance/generic_address_space/atomic_tests.cpp
View File

@@ -91,10 +91,9 @@ kernel void testKernel(global atomic_int* globalPtr, local atomic_int* localPtr)
)OpenCLC";
}
int test_generic_atomics_invariant(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int)
REGISTER_TEST(generic_atomics_invariant)
{
const auto version = get_device_cl_version(deviceID);
const auto version = get_device_cl_version(device);
if (version < Version(2, 0)) return TEST_SKIPPED_ITSELF;
@@ -108,7 +107,7 @@ int test_generic_atomics_invariant(cl_device_id deviceID, cl_context context,
size_t wgSize, retSize;
// Attempt to find the simd unit size for the device.
err = clGetKernelWorkGroupInfo(kernel, deviceID,
err = clGetKernelWorkGroupInfo(kernel, device,
CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE,
sizeof(wgSize), &wgSize, &retSize);
test_error(err, "clGetKernelWorkGroupInfo failed");
@@ -154,10 +153,9 @@ int test_generic_atomics_invariant(cl_device_id deviceID, cl_context context,
return CL_SUCCESS;
}
int test_generic_atomics_variant(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int)
REGISTER_TEST(generic_atomics_variant)
{
const auto version = get_device_cl_version(deviceID);
const auto version = get_device_cl_version(device);
if (version < Version(2, 0)) return TEST_SKIPPED_ITSELF;
@@ -171,7 +169,7 @@ int test_generic_atomics_variant(cl_device_id deviceID, cl_context context,
size_t wgSize, retSize;
// Attempt to find the simd unit size for the device.
err = clGetKernelWorkGroupInfo(kernel, deviceID,
err = clGetKernelWorkGroupInfo(kernel, device,
CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE,
sizeof(wgSize), &wgSize, &retSize);
test_error(err, "clGetKernelWorkGroupInfo failed");

60
test_conformance/generic_address_space/basic_tests.cpp
View File

@@ -97,7 +97,8 @@ private:
const std::vector<std::string> _kernels;
};
int test_function_get_fence(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(function_get_fence)
{
const std::string KERNEL_FUNCTION = common::CONFORMANCE_VERIFY_FENCE +
NL
NL "__global int gint = 1;"
@@ -142,10 +143,11 @@ int test_function_get_fence(cl_device_id deviceID, cl_context context, cl_comman
CBasicTest test(KERNEL_FUNCTION);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_function_to_address_space(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(function_to_address_space)
{
const std::string KERNEL_FUNCTION =
NL
NL "__global int gint = 1;"
@@ -190,10 +192,11 @@ int test_function_to_address_space(cl_device_id deviceID, cl_context context, cl
CBasicTest test(KERNEL_FUNCTION);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_variable_get_fence(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(variable_get_fence)
{
const std::string KERNEL_FUNCTION = common::CONFORMANCE_VERIFY_FENCE +
NL
NL "__global int gint = 1;"
@@ -223,10 +226,11 @@ int test_variable_get_fence(cl_device_id deviceID, cl_context context, cl_comman
CBasicTest test(KERNEL_FUNCTION);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_variable_to_address_space(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(variable_to_address_space)
{
const std::string KERNEL_FUNCTION =
NL
NL "__global int gint = 1;"
@@ -256,10 +260,11 @@ int test_variable_to_address_space(cl_device_id deviceID, cl_context context, cl
CBasicTest test(KERNEL_FUNCTION);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(casting)
{
std::vector<std::string> KERNEL_FUNCTIONS;
// pointers to global, local or private are implicitly convertible to generic
@@ -345,10 +350,11 @@ int test_casting(cl_device_id deviceID, cl_context context, cl_command_queue que
CBasicTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_conditional_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(conditional_casting)
{
const std::string KERNEL_FUNCTION = common::CONFORMANCE_VERIFY_FENCE +
NL
NL "__global int gint = 1;"
@@ -376,10 +382,11 @@ int test_conditional_casting(cl_device_id deviceID, cl_context context, cl_comma
CBasicTest test(KERNEL_FUNCTION);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_chain_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(chain_casting)
{
const std::string KERNEL_FUNCTION = common::CONFORMANCE_VERIFY_FENCE +
NL
NL "__global int gint = 1;"
@@ -407,10 +414,11 @@ int test_chain_casting(cl_device_id deviceID, cl_context context, cl_command_que
NL;
CBasicTest test(KERNEL_FUNCTION);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_ternary_operator_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(ternary_operator_casting)
{
const std::string KERNEL_FUNCTION = common::CONFORMANCE_VERIFY_FENCE +
NL
NL "__global int gint = 1;"
@@ -435,10 +443,11 @@ int test_ternary_operator_casting(cl_device_id deviceID, cl_context context, cl_
CBasicTest test(KERNEL_FUNCTION);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_language_struct(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(language_struct)
{
std::vector<std::string> KERNEL_FUNCTIONS;
// implicit private struct
@@ -588,10 +597,11 @@ int test_language_struct(cl_device_id deviceID, cl_context context, cl_command_q
CBasicTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_language_union(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(language_union)
{
std::vector<std::string> KERNEL_FUNCTIONS;
KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE +
@@ -726,10 +736,11 @@ int test_language_union(cl_device_id deviceID, cl_context context, cl_command_qu
CBasicTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_multiple_calls_same_function(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(multiple_calls_same_function)
{
const std::string KERNEL_FUNCTION =
NL
NL "int shift2(const int *ptr, int arg) {"
@@ -759,10 +770,11 @@ int test_multiple_calls_same_function(cl_device_id deviceID, cl_context context,
CBasicTest test(KERNEL_FUNCTION);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_compare_pointers(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(compare_pointers)
{
std::vector<std::string> KERNEL_FUNCTIONS;
KERNEL_FUNCTIONS.push_back(
@@ -884,5 +896,5 @@ int test_compare_pointers(cl_device_id deviceID, cl_context context, cl_command_
CBasicTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}

66
test_conformance/generic_address_space/main.cpp
View File

@@ -18,66 +18,8 @@
#include <iostream>
// basic tests
extern int test_function_get_fence(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_function_to_address_space(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_variable_get_fence(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_variable_to_address_space(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_conditional_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_chain_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_ternary_operator_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_language_struct(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_language_union(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_multiple_calls_same_function(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_compare_pointers(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
// advanced tests
extern int test_library_function(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_generic_variable_volatile(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_generic_variable_const(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_generic_variable_gentype(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
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);
// atomic tests
int test_generic_atomics_invariant(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
int test_generic_atomics_variant(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
test_definition test_list[] = {
// basic tests
ADD_TEST(function_get_fence),
ADD_TEST(function_to_address_space),
ADD_TEST(variable_get_fence),
ADD_TEST(variable_to_address_space),
ADD_TEST(casting),
ADD_TEST(conditional_casting),
ADD_TEST(chain_casting),
ADD_TEST(ternary_operator_casting),
ADD_TEST(language_struct),
ADD_TEST(language_union),
ADD_TEST(multiple_calls_same_function),
ADD_TEST(compare_pointers),
// advanced tests
ADD_TEST(library_function),
ADD_TEST(generic_variable_volatile),
ADD_TEST(generic_variable_const),
ADD_TEST(generic_variable_gentype),
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),
// atomic tests
ADD_TEST(generic_atomics_invariant),
ADD_TEST(generic_atomics_variant),
};
const int test_num = ARRAY_SIZE( test_list );
test_status InitCL(cl_device_id device) {
test_status InitCL(cl_device_id device)
{
auto version = get_device_cl_version(device);
auto expected_min_version = Version(2, 0);
@@ -134,5 +76,7 @@ test_status InitCL(cl_device_id device) {
int main(int argc, const char *argv[])
{
return runTestHarnessWithCheck(argc, argv, test_num, test_list, false, false, InitCL);
return runTestHarnessWithCheck(
argc, argv, test_registry::getInstance().num_tests(),
test_registry::getInstance().definitions(), false, false, InitCL);
}

13
test_conformance/generic_address_space/stress_tests.cpp
View File

@@ -99,15 +99,20 @@ private:
const std::vector<std::string> _kernels;
};
int test_max_number_of_params(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(max_number_of_params)
{
cl_int error;
size_t deviceMaxParameterSize;
error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_PARAMETER_SIZE, sizeof(deviceMaxParameterSize), &deviceMaxParameterSize, NULL);
error = clGetDeviceInfo(device, CL_DEVICE_MAX_PARAMETER_SIZE,
sizeof(deviceMaxParameterSize),
&deviceMaxParameterSize, NULL);
test_error(error, "clGetDeviceInfo failed");
size_t deviceAddressBits;
error = clGetDeviceInfo(deviceID, CL_DEVICE_ADDRESS_BITS, sizeof(deviceAddressBits), &deviceAddressBits, NULL);
error =
clGetDeviceInfo(device, CL_DEVICE_ADDRESS_BITS,
sizeof(deviceAddressBits), &deviceAddressBits, NULL);
test_error(error, "clGetDeviceInfo failed");
size_t maxParams = deviceMaxParameterSize / (deviceAddressBits / 8);
@@ -174,5 +179,5 @@ int test_max_number_of_params(cl_device_id deviceID, cl_context context, cl_comm
CStressTest test(KERNEL_FUNCTION);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}