The integer tests use genrand_int32 to fill the input buffers used by the tests. This is generally not a problem, except that for the larger scalar types, int and long, the probability that the relatively few samples that are tested contain zero is low. This is particularly unfortunate for functions such as popcount(), clz(), and ctz() where getting the correct result for a zero input is especially important. We should have a way to ensure that 0 and other interesting values are guaranteed to be tested.

Looks like fix in spir tests is missed in opensource and present in internal gitlab

cl20_trunk have this fix.

Hi,
Is it possible that enable support " cl_khr_create_command_queue” extension in bellow code path ?
test_conformance/compiler/test_compiler_defines_for_extensions.cpp

At least four of tests' kernels run into data race:

• enqueue_flags_wait_work_group_simple
• enqueue_flags_wait_work_group_event
• enqueue_flags_wait_work_group_local
• enqueue_flags_wait_work_group_event_local

res[gid] is being written by get_num_groups()/get_local_size() threads. And if one of these thread put 2 to the array cell, it still may be overwritten with 1 by another thread. Thus, the test doesn't really test anything.

I may be missing something, but I don't find memory_scope_sub_group mentioned anywhere. I don't believe there is coverage of this scope in any of the tests.

for branches cl21+ because requires spirv support

An implementation that treats them as signed can pass the tests. Discovered on cl12_trunk.

They use online compilation and completely bypass the SPIR-V offline path. A few examples:

• basic/int*
• basic/createkernelsinprogram
• api/null_buffer_arg
• api/kernels
• api/create_kernels
• generic_address_space/advanced_tests
• ...

It appears that many main functions return the result from runTestHarness directly. https://github.com/KhronosGroup/OpenCL-CTS/search?q=runTestHarness&unscoped_q=runTestHarness

While not being documented, runTestHarness returns the number of failed tests. This number is not constraint to a specific domain (e.g. [0, 255[).

Because this number is returned from main functions, this leads to undefined behaviour when this number is outside the domain specified by the C or C++ standard. In practice, this means that returning values bigger than 256 on some Linux machine could be truncated to 0 and therefore lead the test infrastructure to think everything went okay.

Instead, to avoid any undefined behaviour, the main function should probably return either EXIT_SUCCESS or EXIT_FAILURE to err on the side of caution.

I have a question about the clEnqueueFillImage tests. It looks like the test is always passing in floating point types for the fill_color, but it uses a variety of different cl_channel_type enums.

According to the spec:

"The fill color is a four component RGBA floating-point color value if
the image channel data type is not an unnormalized signed and unsigned integer type, is a four
component signed integer value if the image channel data type is an unnormalized signed integer
type and is a four component unsigned integer value if the image channel data type is an
unnormalized unsigned integer type. The fill color will be converted to the appropriate image
channel format and order associated with image"

So as far as I can tell, the driver isn't expected to perform a conversion from float to anything else, but that appears to be what the test expects. Can somebody please point me to an explanation for the behavior, or flag this as a bug in the test? I am looking at the CL 1.2 branch. Thanks in advance.

Neither the 2.1 nor 2.2 tests exercise this entrypoint.

In case cl_khr_fp64 not supported test spir printf should skip without errors now have:
test1.testCaseFloat...

error: double type is not supported on this platform
error: backend compiler failed build.

In cl20_trunk fix is present

In non_uniform_work_group test case, the required buffer size to create a global buffer may exceed available device memory size on some platform, for example IA32 platform. This will lead to failure.

All the test_write_image_* functions declare int totalErrors=0 as the first line, then redeclare it in the main for loop. The outer variable is returned at the end and is always 0.

The only way to fail the test is by early exit; e.g. finding numTries (5) failures in a single image.

Present only on cl20 branch. Lack in cl21 and cl22

lost fixes in in test_migrate.cpp

There are a number of places where we should make use of the SPIR-V validator:

• to check the output of the offline compilation flow
• to check that the checked-in SPIR-V binaries are correct as part of the continuous integration.

Internal MR (102) should be moved to opensource to fix subgroups testing to be aligned with specification OpenCL2.1
IFP - independent forward progress

• current version looks like only for platforms supporting OpenCL2.0
• condition for test execution - if implementation shows cl_khr_subgroups extension then use OpenCL 2.0 subgroups in device is OpenCL 2.X and higher then core subgroups should be used
• in OpenCL2.0 version we require extension cl_khr_subgrups and IFP is expected to be supported so test IFP should be executed
• in core subgroups we should not use function/enums with KHR suffixes (proper only for OpenCL2.0).
• in core subgroups version no extension, because this is a core feature and IFP is not obligatory . Only if capability CL_DEVICE_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS is set to 1 then start IFP testing

I see some lost fixes in files:
compatibility\test_common\harness\compat.h
compatibility\test_common\harness\msvc9.c
Lack of this fixes cause erros while compilation on VS2015, VS2017

Function verifyOutputBuffer has two 'return false' statements which indicate mismatches. But this function should return non-zero value in case of mismatch.

no fix only on cl22_trunk branch

The code in the test_min_max_parameter_size inconsistently assumes that gEmbeddedProfile implies !gHasLong. In the code snippet below line 999 will calculate the numberOfIntParametersToTry assuming the kernels will have int arguments. However, line 1054 will create kernel files with long arguments if longs are supported. Therefore in the case where there is an Embedded Profile device that supports the cles_khr_int64 extension the test will overrun CL_DEVICE_MAX_PARAMETER_SIZE.

The fix would be to change the if statements on the following lines to check gHasLong instead of gIsEmbedded

Here's a (possibly non-exhaustive) list of built-in functions that aren't tested:

• get_kernel_sub_group_count_for_ndrange
• get_kernel_max_sub_group_size_for_ndrange

A few other features aren't tested at all:

• memory_scope_sub_group

Fix is present only on cl20_trunk branch

Fix available only on branch cl20_trunk cl12_trunk Need to fix in other branches as well.

As I can see in the code half data type is tested, but only if "cl_khr_fp16" extension is supported.

If I understand correctly, a few test cases should be tested even if "cl_khr_fp16" extension is not supported by device.

From OpenCL C spec ver. 2.0 rev. 33:

The half data type can only be used to declare a pointer to a buffer that contains half values. A few valid examples are given below:

void
bar (__global half *p)
   ...
}
__kernel void
foo (__global half *pg, __local half *pl)
{
    __global half *ptr;
    int offset;
    ptr = pg + offset;
    bar(ptr);
}

So, half_scalar_p and half_scalar_p2 test cases listed here should be always tested. Is my understanding right?

If so, I can submit a PR: which branches should it affect?

This can generate test errors when extension string is bigger than test expects (then CL_INVALID_VALUE will be returned).
The following tests will be affected:

• all spir tests
  char extensions[1024] = {0}; size_t size; // Querying the device for its supported extensions cl_int errcode = clGetDeviceInfo(devId, CL_DEVICE_EXTENSIONS, sizeof(extensions), extensions, &size);
• computeinfo
  For example computeinfo tests defines config_data union where one of the filed is :
  char string[1024];

The goal here is to make tests more robust and trustworthy.

Currently, the conformance tests generate inputs randomly. For example, for ternary functions such as mad_sat the inputs are generated in test_three_param_integer_kernel as follows:

While, theoretically, running tests many, many times with different PRNG seeds can give a good confidence in the correctness of the implementation, in practice this is not feasible for large input spaces. E.g. mad_sat on long inputs has more than 6 octodecillion possible inputs to be tested, leaving plenty of room for bugs to hide.

In addition to random testing, the conformance test should probably provide mechanisms to check specific, representative inputs. Those inputs would then test different behaviours/corner cases.

Here are a few examples of inputs for mad_sat(x, y, z), which for reference purpose is clamp(x * y + z, domain_lower_bound, domain_upper_bound) if we assume integer capable of representing infinitely large values:

• x, y, z s.t. x * y + z isn't saturated;
• x, y, z s.t. x * y saturated but x * y + z doesn't;
• x, y, z s.t. x * y doesn't saturate but x * y + z does;
• x, y, z s.t. x * y is {positive, negative} but x * y + z is {negative, positive};
• ...

Since there is currently no OpenCL C construct that can target this, the test would probable have to use a hand-crafted SPIR-V module.

test_printf use file /tmp/tmpfile or tmpfile to capture OpenCL kernel outputs, which are compared with correct results. If we run multiple printf tests in parallel, they fail because they write to the same file and one test could read another one's outputs.

reading OpenCL-CTS/readme-spir-v-binaries.txt says:

To run the 2.2 conformance tests test suite for the C++ features you need need 
SPIR-V binaries.
can be picked from:
https://cvs.khronos.org/svn/repos/OpenCL/trunk/Khronos/spirv/spirv10_2015.11.25.zip

but that link can't be accessed for a non Khronos user so can upload this file?
I say that because txt file also says this:
Alternatively you can check out and build all of the below repositories.
problem is all this links also are to Khronos members and these projects altough exist on github are missing required branches so note I also asked
SPIRV-LLVM (KhronosGroup/SPIRV-LLVM#218)
and
libclcxx (KhronosGroup/libclcxx#15)
projects on github to be updated..

Needs a few tweaks to the Travis config (apparently additional libraries are needed on OSX).

It's currently checking for sizeof(half) == 8.

https://github.com/KhronosGroup/OpenCL-CTS/blob/cl12_trunk/test_conformance/basic/test_sizeof.c#L359

The root cause seems to be a violation in this test of OpenCL's type-based aliasing rules:

Consider the following kernel derived from Test_vLoadHalf.c:

__kernel void test( const __global half p, __global float2 f ) {
__local int data[8/2];
__local half hdata_p = (__local half) data;
__global int* i_p = (__global int*)p;
size_t i = get_global_id(0);
size_t lid = get_local_id(0);
int k;
for (k=0; k<2/2; k++)
data[lid2/2+k] = i_p[i2/2+k];
f[i] = vload_half2( lid, hdata_p );
}
The kernel stores to the local int array but loads half values from it, effectively accessing it using two distinct data types. This seems to go against C99’s aliasing rules which OpenCL complies with (section 6.1.8 in spec 1.2) and can lead compilers to generate incorrect code by relying on type-based alias analysis.

Such tests can be legalized by converting them to use (aligned) ‘char’ instead of ‘short’, ‘int’ etc. for the local buffers as converting to ‘char’ is allowed, e.g.:

__kernel void test( const __global half p, __global float2 f )
{
attribute ((aligned (4))) __local char data[48/2];
__local half hdata_p = (__local half*) data;
__global char* i_p = (__global char*)p;
size_t i = get_global_id(0);
size_t lid = get_local_id(0);
int k;
for (k=0; k<42/2; k++)
data[4lid2/2+k] = i_p[4i*2/2+k];
f[i] = vload_half2( lid, hdata_p );
}

The write tests use a normalized error fabsf( ( expected[ j ] - actual[ j ] ) / expected[ j ] ) for non-zero expected results. For denormals that get flushed to zero, this will compute a failing error value (e.g. 1 for CL_R images).

Additional issues in the error computation (limited here to CL_FLOAT tests in test_write_image_1D()):

1. CL_FLOAT images are tested against a relative error of 0.005f, which has no correlation to spec requirements.
2. The computed "err" is averaged over all channels, which is not provide for in the spec
3. NaN/Inf are not considered in the relative error calculation, and can lead to both false fails and false passes.

The overall test can still pass due to #27

The extension cl_khr_spirv_no_integer_wrap_decoration has been recently ratified. The change in test_compiler is required to extend known extension array of strings (otherwise the test will failed for platforms that supports this extension).

This triggers warnings with some compilers, requires CMake trickery and is bad practise. We ought to rename all the files to their rightful extension.

I'm anxious. so posting here to ask you.

A few issues to discuss:

• Do we really want to commit the binaries? Shouldn't we just generate them from the assembly as part of the build system?
• The assembly doesn't use the same syntax as the SPIR-V disassembler which makes it hard to check consistency between the binaries and assembly or check that PRs modify the binaries and assembly in the same manner.

Khronos Bug 15347 fixed code review issues + fix for a similar issue in basic constant test

The extension cl_khr_spirv_no_integer_wrap_decoration has its corresponding spir-v extension SPV_KHR_no_integer_wrap_decoration. New test is needed to verify if offline compiler supports this extensions. The test proposal appear soon after sync with internal Khronos repository. New test will extend set of tests in spirv_new directory

The basic async_strided_copy test contains the following:
200 for (int i=0; i<(int)globalBufferSize; i+=(int)elementSize*(int)stride)
201 {
202 if (memcmp( ((char )inBuffer)+i, ((char )outBuffer)+i, typeSize) != 0 )
203 {
204 unsigned char * inchar = (unsigned char)inBuffer + i;
205 unsigned char * outchar = (unsigned char)outBuffer + i;
206 char values[4096];
207 values[0] = 0;
208
209 log_error( "ERROR: Results of copy did not validate!\n" );
210 sprintf(values + strlen( values), "%d -> [", i);
211 for (int j=0; j<(int)elementSize; j++)
212 sprintf(values + strlen( values), "%2x ", inchar[ielementSize+j]);
213 sprintf(values + strlen(values), "] != [");
214 for (int j=0; j<(int)elementSize; j++)
215 sprintf(values + strlen( values), "%2x ", outchar[ielementSize+j]);
216 sprintf(values + strlen(values), "]");
217 log_error("%s\n", values);
218
219 return -1;
220 }
221 }

Lines 212 and 215 can access out of bounds since inchar and outchar have already been offset by "i". The simple fix is to simply remove the "i*elementSize" in the subscripting expressions. This problem appears in all branches I checked.

Changes for CRC calculation when offline compiler files (spirv) are used.