Khronos Basic Test Failures
===========================
Khronos Test Version: OpenCL 1.1: April 4, 2010.

Usage:  
% cd opencl_conformance/test_conformance/basic
% test_basic <test-name>


<test-name>: 
Failure Mode:
Analysis:

hiloeo
astype
======
Failure Mode:   
------------
Runs out of system memory, and crashes the test.  
However, the test is passing all of the subtests before it crashes.

Analysis:
--------
valgrind analysis on shamrock showed huge memory leaks around creating and
deleting programs, which were due to LLVM objects not getting freed.  This
could either be a usage problem, or leaks in LLVM MCJIT execution engine.

Note also in llvm-src/tools/clang/include/clang/Frontend/CompilerInstance this
comment:
   // FIXME: Eliminate the llvm_shutdown requirement, that should either be part
   // of the context or else not CompilerInstance specific.
   bool ExecuteAction(FrontendAction &Act);

The Khronos tests do not call llvm_shutdown (nor should they), but also often
do not call clReleaseProgram() after calling clCreateProgram() many times in a
loop .

async_copy_global_to_local.txt
async_copy_local_to_global.txt
async_strided_copy_global_to_local.txt
async_strided_copy_local_to_global.txt
======================================
Failure Mode:   
------------
All of the above 4 tests fail in the same way:  Due to the Khronos generated
CL file not being able to compile.  These also fail the same way on 
Keystone II EVM (which doesn't use MCJIT).

async_copy_global_to_local...
Testing char
program.cl:9:153: error: used type 'event_t' where arithmetic or pointer type is required

ERROR: clBuildProgram failed! (CL_BUILD_PROGRAM_FAILURE from /home/gpitney/opencl_conformance/test_common/harness/kernelHelpers.c:35)
Original source is: ------------

__kernel void test_fn( const __global char *src, __global char *dst, __local char *localBuffer, int copiesPerWorkgroup, int copiesPerWorkItem )
{
 int i;
 for(i=0; i<copiesPerWorkItem; i++)
	 localBuffer[ get_local_id( 0 )*copiesPerWorkItem+i ] = (char)(char)0;
	barrier( CLK_LOCAL_MEM_FENCE );
	event_t event;
	event = async_work_group_copy( (__local char*)localBuffer, (__global const char*)(src+copiesPerWorkgroup*get_group_id(0)), (size_t)copiesPerWorkgroup, (event_t)0 );
	wait_group_events( 1, &event );
 for(i=0; i<copiesPerWorkItem; i++)
  dst[ get_global_id( 0 )*copiesPerWorkItem+i ] = localBuffer[ get_local_id( 0 )*copiesPerWorkItem+i ];
}
Build not successful for device "Intel(R) Core(TM) i7-2620M CPU @ 2.70GHz", status: CL_BUILD_ERROR
Build log for device "Intel(R) Core(TM) i7-2620M CPU @ 2.70GHz" is: ------------
program.cl:9:153: error: used type 'event_t' where arithmetic or pointer type is required

Analysis:
--------
Note the cast of (event_t)0 in the kernel above.
Per the discussion here: http://comments.gmane.org/gmane.comp.compilers.clang.scm/93008 , it appears the spec is vague on this point, but the Khronos 
test nevertheless expects the cast to compile.   

It seems the Khronos test and clang are in conflict.

kernel_memory_alignment_constant.txt
====================================
Failure Mode:   
-------------
This fails due to inability to compile a Khronos test generated CL program.


kernel_memory_alignment_constant...
Device version string: "OpenCL 1.1 "
Testing char...
        Testing parameter kernel...
        Testing constant kernel...
program.cl:2:17: error: variable in constant address space must be initialized
program.cl:3:18: error: variable in constant address space must be initialized
program.cl:4:18: error: variable in constant address space must be initialized
program.cl:5:18: error: variable in constant address space must be initialized
program.cl:6:18: error: variable in constant address space must be initialized
program.cl:7:19: error: variable in constant address space must be initialized

ERROR: clBuildProgram failed! (CL_BUILD_PROGRAM_FAILURE from /home/gpitney/open
Original source is: ------------

  constant char mem0[3];
  constant char2 mem2[3];
  constant char3 mem3[3];
  constant char4 mem4[3];
  constant char8 mem8[3];
  constant char16 mem16[3];

kernel void test(global ulong *results)
{
   results[0] = (ulong)&mem0;
   results[1] = (ulong)&mem2;
   results[2] = (ulong)&mem3;
   results[3] = (ulong)&mem4;
   results[4] = (ulong)&mem8;
   results[5] = (ulong)&mem16;
}


Analysis:
--------
Interestingly, this generated CL code compiles on Keystone 
(not using MCJIT, using LLVM 3.3), and the test passes.

Some digging shows this clang error was added after LLVM 3.3 
(LLVM version used by TI Keystone, which explains why it passes there):
http://lists.cs.uiuc.edu/pipermail/cfe-commits/Week-of-Mon-20131230/096405.html

In this case, the LLVM clang compiler and the Khronos tests seem to be
in conflict.


local_kernel_scope
==================
Failure Mode:   
------------
The max of a set of unsigned integers computed from an OCL kernel differs
from the max of the same set computed by the host.

Error Message:
local_kernel_scope...
	Testing with 6 groups, 184 elements per group...
ERROR: Local max validation failed! (expected 4274779084, got 4290015211 for i=0)

Analysis:
--------
This test fails for numCPUs > 1, because a local variable defined in a kernel
is being allocated by clang into global memory, rather than thread local
storage.

For this OpenCL code:

     __kernel void test( __global unsigned int * input, __global unsigned int *outMaxes ) {
          __local unsigned int localStorage[256*4];
          [...]
           }

      The LLVM IR produced is:

             @test.localStorage = internal unnamed_addr addrspace(2) global [256 x i32] zeroinitializer, align 4


The expectation is that clang would have generated a thread_local attribute
on OpenCL __local variables, allowing the MCJIT/ARM backend to allocate
the variable localStorage into TLS at runtime.