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

Fix generic address space OpenCL 2.0 assumption (#1575)

Browse Source 
Conformance tests should not assume support for OpenCL 2.0, or specific
features. This patch removes uses of
`__opencl_c_program_scope_global_variables` for devices which do not
support it in the generic address space conformance tests. Additionally,
the `clCompileProgram` options have been updated to select between
`CL2.0` and `CL3.0` depending on which the device under test supports.

Co-authored-by: Amy Worthington <amy@codeplay.com>
This commit is contained in:
Kenneth Benzie
2022-11-23 14:08:30 +00:00
committed by GitHub
parent 640f6e66c7
commit def23dd6e6
1 changed files with 61 additions and 28 deletions
Download Patch File Download Diff File Expand all files Collapse all files

89
test_conformance/generic_address_space/advanced_tests.cpp
View File

@@ -107,7 +107,32 @@ public:
program = clCreateProgramWithSource(context, 1, &srcPtr, NULL, &error);
test_error(error, "clCreateProgramWithSource failed");
error = clCompileProgram(program, 1, &deviceID, "-cl-std=CL2.0", 0, NULL, NULL, NULL, NULL);
// Use the latest OpenCL-C version supported by the device. This
// allows calling code to force a particular CL C version if it is
// required, but also means that callers need not specify a version
// if they want to assume the most recent CL C.
auto version = get_max_OpenCL_C_for_context(context);
const char* cl_std = nullptr;
if (version >= Version(3, 0))
{
cl_std = "-cl-std=CL3.0";
}
else if (version >= Version(2, 0) && version < Version(3, 0))
{
cl_std = "-cl-std=CL2.0";
}
else
{
// If the -cl-std build option is not specified, the highest
// OpenCL C 1.x language version supported by each device is
// used when compiling the program for each device.
cl_std = "";
}
error = clCompileProgram(program, 1, &deviceID, cl_std, 0, NULL,
NULL, NULL, NULL);
if (error != CL_SUCCESS)
PrintCompilationLog(program, deviceID);
@@ -118,7 +143,8 @@ public:
preCompiledLibrary = clCreateProgramWithSource(context, 1, &srcPtrLibrary, NULL, &error);
test_error(error, "clCreateProgramWithSource failed");
error = clCompileProgram(preCompiledLibrary, 1, &deviceID, "-cl-std=CL2.0", 0, NULL, NULL, NULL, NULL);
error = clCompileProgram(preCompiledLibrary, 1, &deviceID, cl_std,
0, NULL, NULL, NULL, NULL);
if (error != CL_SUCCESS)
PrintCompilationLog(preCompiledLibrary, deviceID);
@@ -136,7 +162,9 @@ public:
}
else {
if (create_single_kernel_helper_with_build_options(context, &program, &kernel, 1, &srcPtr, "testKernel", "-cl-std=CL2.0")) {
if (create_single_kernel_helper(context, &program, &kernel, 1,
&srcPtr, "testKernel"))
{
log_error("create_single_kernel_helper failed\n");
return -1;
}
@@ -276,31 +304,36 @@ int test_library_function(cl_device_id deviceID, cl_context context, cl_command_
NL "}"
NL;
const std::string KERNEL_FUNCTION =
NL
NL "extern bool helperFunction(float *floatp, float val);"
NL
NL "__global float gfloat = 1.0f;"
NL
NL "__kernel void testKernel(__global uint *results) {"
NL " uint tid = get_global_id(0);"
NL
NL " __global float *gfloatp = &gfloat;"
NL " __local float lfloat;"
NL " lfloat = 2.0f;"
NL " __local float *lfloatp = &lfloat;"
NL " float pfloat = 3.0f;"
NL " __private float *pfloatp = &pfloat;"
NL
NL " uint failures = 0;"
NL
NL " failures += helperFunction(gfloatp, gfloat) ? 0 : 1;"
NL " failures += helperFunction(lfloatp, lfloat) ? 0 : 1;"
NL " failures += helperFunction(pfloatp, pfloat) ? 0 : 1;"
NL
NL " results[tid] = failures == 0;"
NL "}"
NL;
const std::string KERNEL_FUNCTION = R"OpenCLC(
extern bool helperFunction(float *floatp, float val);
#ifdef __opencl_c_program_scope_global_variables
__global float gfloat = 1.0f;
#endif
__kernel void testKernel(__global uint *results) {
uint tid = get_global_id(0);
#ifdef __opencl_c_program_scope_global_variables
__global float *gfloatp = &gfloat;
#endif
__local float lfloat;
lfloat = 2.0f;
__local float *lfloatp = &lfloat;
float pfloat = 3.0f;
__private float *pfloatp = &pfloat;
uint failures = 0;
#ifdef __opencl_c_program_scope_global_variables
failures += helperFunction(gfloatp, gfloat) ? 0 : 1;
#endif
failures += helperFunction(lfloatp, lfloat) ? 0 : 1;
failures += helperFunction(pfloatp, pfloat) ? 0 : 1;
results[tid] = failures == 0;
};
)OpenCLC";
CAdvancedTest test(LIBRARY_FUNCTION, KERNEL_FUNCTION);