v0.9: Shamrock README and TODO listsHEADmaster

and updated CREDITS

Signed-off-by: Gil Pitney <gil.pitney@linaro.org>

1 files changed, 38 insertions, 285 deletions

diff --git a/tests/basic_test_failures.lst b/tests/basic_test_failures.lst
index d5e4871..a0fff9c 100644
--- a/tests/basic_test_failures.lst
+++ b/tests/basic_test_failures.lst
@@ -12,6 +12,7 @@ Failure Mode:
 Analysis:
 
 hiloeo
+astype
 ======
 Failure Mode:   
 ------------
@@ -22,8 +23,17 @@ 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 a bug in LLVM MCJIT execution engine.
+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
@@ -34,7 +44,7 @@ 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 EVM (which doesn't use MCJIT). 
+Keystone II EVM (which doesn't use MCJIT).
 
 async_copy_global_to_local...
 Testing char
@@ -65,12 +75,12 @@ 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.   
 
-So, it seems a clang patch for OpenCL event_t casts of zero may be required.
+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.
 
 
@@ -86,7 +96,7 @@ 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/opencl_conformance/test_common/harness/kernelHelpers.c:35)
+ERROR: clBuildProgram failed! (CL_BUILD_PROGRAM_FAILURE from /home/gpitney/open
 Original source is: ------------
 
   constant char mem0[3];
@@ -116,297 +126,40 @@ 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 are in conflict.
-
-local_kernel_def.txt
-====================
-Failure Mode:   
-------------
-This fails due to inability to compile a Khronos test generated CL program.
-
-
-local_kernel_def...
-program.cl:3:23: error: 'tmp_sum' declared as an array with a negative size
-
-ERROR: clBuildProgram failed! (CL_BUILD_PROGRAM_FAILURE from /home/gpitney/opencl_conformance/test_common/harness/kernelHelpers.c:35)
-Original source is: ------------
-__kernel void compute_sum_with_localmem(__global int *a, int n, __global int *sum)
-{
-	 __local int tmp_sum[-2147483648];
-    int  tid = get_local_id(0);
-    int  lsize = get_local_size(0);
-    int  i;
+In this case, the LLVM clang compiler and the Khronos tests seem to be
+in conflict.
 
-[... snip ...]
 
-Analysis:
---------
-This test also fails on Keystone, but the negative number is (-4).
-
-The Khronos test is casting a size_t value for work group size to an int,
-and printing it into the kernel string using the %d printf() modifier.
-This does not appear to be the right printf() modifier for a size_t, so
-the test code appears to be in error.
- 
-parameter_types
-=================
+local_kernel_scope
+==================
 Failure Mode:   
 ------------
-Invalid results results returned from test-generated OCL kernel, which uses
-vector parameters of various sizes.
-
-[ ... snip ...]
-Testing vector size 4
-Kernel: __kernel void test_kernel(
-char4 c, uchar4 uc, short4 s, ushort4 us, int4 i, uint4 ui, float4 f,
-__global float4 *result)
-{
-  result[0] = convert_float4(c);
-  result[1] = convert_float4(uc);
-  result[2] = convert_float4(s);
-  result[3] = convert_float4(us);
-  result[4] = convert_float4(i);
-  result[5] = convert_float4(ui);
-  result[6] = f;
-}
-
-Conversion from char4 failed: index 0 got 4.28107e-38, expected 0.
-Conversion from char4 failed: index 2 got 16, expected 2.
-Conversion from char4 failed: index 3 got 1, expected -3.
-Conversion from uchar4 failed: index 0 got 4.28107e-38, expected 16.
-Conversion from uchar4 failed: index 1 got -1, expected 1.
-Conversion from uchar4 failed: index 2 got 18, expected 2.
-Conversion from uchar4 failed: index 3 got 1, expected 3.
-Conversion from short4 failed: index 0 got -19, expected -17.
-Conversion from short4 failed: index 2 got 20, expected 2.
-Conversion from short4 failed: index 3 got 1, expected -3.
-Conversion from ushort4 failed: index 0 got -23, expected 18.
-Conversion from ushort4 failed: index 1 got -1, expected 1.
-Conversion from ushort4 failed: index 2 got 0, expected 2.
-Conversion from ushort4 failed: index 3 got 0, expected 3.
-Conversion from int4 failed: index 0 got 0, expected -19.
-Conversion from int4 failed: index 1 got 0, expected -1.
-Conversion from int4 failed: index 2 got 0, expected 2.
-Conversion from int4 failed: index 3 got 0, expected -3.
-Conversion from uint4 failed: index 0 got 0, expected 20.
-Conversion from uint4 failed: index 1 got 0, expected 1.
-Conversion from uint4 failed: index 2 got 0, expected 2.
-Conversion from uint4 failed: index 3 got 0, expected 3.
-Conversion from float4 failed: index 0 got 0, expected -23.
-Conversion from float4 failed: index 1 got 0, expected -1.
-Conversion from float4 failed: index 2 got 0, expected 2.
-Conversion from float4 failed: index 3 got 0, expected -3.
-Testing vector size 8
-Kernel: __kernel void test_kernel(
-char8 c, uchar8 uc, short8 s, ushort8 us, int8 i, uint8 ui, float8 f,
-__global float8 *result)
-{
-  result[0] = convert_float8(c);
-  result[1] = convert_float8(uc);
-  result[2] = convert_float8(s);
-  result[3] = convert_float8(us);
-  result[4] = convert_float8(i);
-  result[5] = convert_float8(ui);
-  result[6] = f;
-}
-
-Conversion from char8 failed: index 0 got -5.99946e-08, expected 0.
-Conversion from char8 failed: index 2 got 16, expected 2.
-Conversion from char8 failed: index 3 got 1, expected -3.
-Conversion from char8 failed: index 4 got 4.28106e-38, expected 4.
-Conversion from char8 failed: index 5 got -1, expected -5.
-Conversion from char8 failed: index 6 got 18, expected 6.
-Conversion from char8 failed: index 7 got 1, expected -7.
-Conversion from uchar8 failed: index 0 got -19, expected 16.
-Conversion from uchar8 failed: index 1 got -1, expected 1.
-Conversion from uchar8 failed: index 2 got 20, expected 2.
-Conversion from uchar8 failed: index 3 got 1, expected 3.
-Conversion from uchar8 failed: index 4 got -5.99946e-08, expected 4.
-Conversion from uchar8 failed: index 5 got -1, expected 5.
-Conversion from uchar8 failed: index 6 got 0, expected 6.
-Conversion from uchar8 failed: index 7 got 0, expected 7.
-Conversion from short8 failed: index 0 got 0, expected -17.
-Conversion from short8 failed: index 1 got 0, expected -1.
-Conversion from short8 failed: index 2 got 0, expected 2.
-Conversion from short8 failed: index 3 got 0, expected -3.
-Conversion from short8 failed: index 4 got 0, expected 4.
-Conversion from short8 failed: index 5 got 0, expected -5.
-Conversion from short8 failed: index 6 got 0, expected 6.
-Conversion from short8 failed: index 7 got 0, expected -7.
-Conversion from ushort8 failed: index 0 got 0, expected 18.
-Conversion from ushort8 failed: index 1 got 0, expected 1.
-Conversion from ushort8 failed: index 2 got 0, expected 2.
-Conversion from ushort8 failed: index 3 got 0, expected 3.
-Conversion from ushort8 failed: index 4 got 0, expected 4.
-Conversion from ushort8 failed: index 5 got 0, expected 5.
-Conversion from ushort8 failed: index 6 got 0, expected 6.
-Conversion from ushort8 failed: index 7 got 0, expected 7.
-Conversion from int8 failed: index 0 got 0, expected -19.
-Conversion from int8 failed: index 1 got 0, expected -1.
-Conversion from int8 failed: index 2 got 0, expected 2.
-Conversion from int8 failed: index 3 got 0, expected -3.
-Conversion from int8 failed: index 4 got 0, expected 4.
-Conversion from int8 failed: index 5 got 0, expected -5.
-Conversion from int8 failed: index 6 got 0, expected 6.
-Conversion from int8 failed: index 7 got 0, expected -7.
-Conversion from uint8 failed: index 0 got 0, expected 20.
-Conversion from uint8 failed: index 1 got 0, expected 1.
-Conversion from uint8 failed: index 2 got 0, expected 2.
-Conversion from uint8 failed: index 3 got 0, expected 3.
-Conversion from uint8 failed: index 4 got 0, expected 4.
-Conversion from uint8 failed: index 5 got 0, expected 5.
-Conversion from uint8 failed: index 6 got 0, expected 6.
-Conversion from uint8 failed: index 7 got 0, expected 7.
-Conversion from float8 failed: index 0 got 0, expected -23.
-Conversion from float8 failed: index 1 got 0, expected -1.
-Conversion from float8 failed: index 2 got 0, expected 2.
-Conversion from float8 failed: index 3 got 0, expected -3.
-Conversion from float8 failed: index 4 got 0, expected 4.
-Conversion from float8 failed: index 5 got 0, expected -5.
-Conversion from float8 failed: index 6 got 0, expected 6.
-Conversion from float8 failed: index 7 got 0, expected -7.
-Testing vector size 16
-Kernel: __kernel void test_kernel(
-char16 c, uchar16 uc, short16 s, ushort16 us, int16 i, uint16 ui, float16 f,
-__global float16 *result)
-{
-  result[0] = convert_float16(c);
-  result[1] = convert_float16(uc);
-  result[2] = convert_float16(s);
-  result[3] = convert_float16(us);
-  result[4] = convert_float16(i);
-  result[5] = convert_float16(ui);
-  result[6] = f;
-}
+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.
 
-Conversion from char16 failed: index 0 got -7.22404e-06, expected 0.
-Conversion from char16 failed: index 2 got 16, expected 2.
-Conversion from char16 failed: index 3 got 1, expected -3.
-Conversion from char16 failed: index 4 got -3.96717e-07, expected 4.
-Conversion from char16 failed: index 5 got -1, expected -5.
-Conversion from char16 failed: index 6 got 18, expected 6.
-Conversion from char16 failed: index 7 got 1, expected -7.
-Conversion from char16 failed: index 8 got 0, expected 8.
-Conversion from char16 failed: index 9 got -1, expected -9.
-Conversion from char16 failed: index 10 got 20, expected 10.
-Conversion from char16 failed: index 11 got 1, expected -11.
-Conversion from char16 failed: index 12 got 4.28106e-38, expected 12.
-Conversion from char16 failed: index 13 got -1, expected -13.
-Conversion from char16 failed: index 14 got 0, expected 14.
-Conversion from char16 failed: index 15 got 0, expected -15.
-Conversion from uchar16 failed: index 0 got 0, expected 16.
-Conversion from uchar16 failed: index 1 got 0, expected 1.
-Conversion from uchar16 failed: index 2 got 0, expected 2.
-Conversion from uchar16 failed: index 3 got 0, expected 3.
-Conversion from uchar16 failed: index 4 got 0, expected 4.
-Conversion from uchar16 failed: index 5 got 0, expected 5.
-Conversion from uchar16 failed: index 6 got 0, expected 6.
-Conversion from uchar16 failed: index 7 got 0, expected 7.
-Conversion from uchar16 failed: index 8 got 0, expected 8.
-Conversion from uchar16 failed: index 9 got 0, expected 9.
-Conversion from uchar16 failed: index 10 got 0, expected 10.
-Conversion from uchar16 failed: index 11 got 0, expected 11.
-Conversion from uchar16 failed: index 12 got -3.96712e-07, expected 12.
-Conversion from uchar16 failed: index 13 got 0, expected 13.
-Conversion from uchar16 failed: index 14 got 0, expected 14.
-Conversion from uchar16 failed: index 15 got 0, expected 15.
-Conversion from short16 failed: index 0 got 0, expected -17.
-Conversion from short16 failed: index 1 got 0, expected -1.
-Conversion from short16 failed: index 2 got 0, expected 2.
-Conversion from short16 failed: index 3 got 0, expected -3.
-Conversion from short16 failed: index 4 got 0, expected 4.
-Conversion from short16 failed: index 5 got 0, expected -5.
-Conversion from short16 failed: index 6 got 0, expected 6.
-Conversion from short16 failed: index 7 got 0, expected -7.
-Conversion from short16 failed: index 8 got 0, expected 8.
-Conversion from short16 failed: index 9 got 0, expected -9.
-Conversion from short16 failed: index 10 got 0, expected 10.
-Conversion from short16 failed: index 11 got 0, expected -11.
-Conversion from short16 failed: index 12 got 0, expected 12.
-Conversion from short16 failed: index 13 got 0, expected -13.
-Conversion from short16 failed: index 14 got 0, expected 14.
-Conversion from short16 failed: index 15 got 0, expected -15.
-Conversion from ushort16 failed: index 0 got 0, expected 18.
-Conversion from ushort16 failed: index 1 got 0, expected 1.
-Conversion from ushort16 failed: index 2 got 0, expected 2.
-Conversion from ushort16 failed: index 3 got 0, expected 3.
-Conversion from ushort16 failed: index 4 got 0, expected 4.
-Conversion from ushort16 failed: index 5 got 0, expected 5.
-Conversion from ushort16 failed: index 6 got 0, expected 6.
-Conversion from ushort16 failed: index 7 got 0, expected 7.
-Conversion from ushort16 failed: index 8 got 0, expected 8.
-Conversion from ushort16 failed: index 9 got 0, expected 9.
-Conversion from ushort16 failed: index 10 got 0, expected 10.
-Conversion from ushort16 failed: index 11 got 0, expected 11.
-Conversion from ushort16 failed: index 12 got 0, expected 12.
-Conversion from ushort16 failed: index 13 got 0, expected 13.
-Conversion from ushort16 failed: index 14 got 0, expected 14.
-Conversion from ushort16 failed: index 15 got 0, expected 15.
-Conversion from int16 failed: index 0 got 0, expected -19.
-Conversion from int16 failed: index 1 got 0, expected -1.
-Conversion from int16 failed: index 2 got 0, expected 2.
-Conversion from int16 failed: index 3 got 0, expected -3.
-Conversion from int16 failed: index 4 got 0, expected 4.
-Conversion from int16 failed: index 5 got 0, expected -5.
-Conversion from int16 failed: index 6 got 0, expected 6.
-Conversion from int16 failed: index 7 got 0, expected -7.
-Conversion from int16 failed: index 8 got 0, expected 8.
-Conversion from int16 failed: index 9 got 0, expected -9.
-Conversion from int16 failed: index 10 got 0, expected 10.
-Conversion from int16 failed: index 11 got 0, expected -11.
-Conversion from int16 failed: index 12 got 0, expected 12.
-Conversion from int16 failed: index 13 got 0, expected -13.
-Conversion from int16 failed: index 14 got 0, expected 14.
-Conversion from int16 failed: index 15 got 0, expected -15.
-Conversion from uint16 failed: index 0 got 0, expected 20.
-Conversion from uint16 failed: index 1 got 0, expected 1.
-Conversion from uint16 failed: index 2 got 0, expected 2.
-Conversion from uint16 failed: index 3 got 0, expected 3.
-Conversion from uint16 failed: index 4 got 0, expected 4.
-Conversion from uint16 failed: index 5 got 0, expected 5.
-Conversion from uint16 failed: index 6 got 0, expected 6.
-Conversion from uint16 failed: index 7 got 0, expected 7.
-Conversion from uint16 failed: index 8 got 0, expected 8.
-Conversion from uint16 failed: index 9 got 0, expected 9.
-Conversion from uint16 failed: index 10 got 0, expected 10.
-Conversion from uint16 failed: index 11 got 0, expected 11.
-Conversion from uint16 failed: index 12 got 0, expected 12.
-Conversion from uint16 failed: index 13 got 0, expected 13.
-Conversion from uint16 failed: index 14 got 0, expected 14.
-Conversion from uint16 failed: index 15 got 0, expected 15.
-
-[ then crashes on next sub-tests]
+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:
 --------
-Some rather intense debugging found the culprit being the float<n> *
-results output kernel vector argument was being *modified* by the MCJIT
-generated ARM assembly kernel code!  This was determined by gdb debugging
-via assembly into the JIT'd kernel, and also inserting callbacks to 
-builtin funcitons to inspect the results pointer argument at entry and exit
-to and from the kernel function.
+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.
 
-After creating a simplified test case using lli, was able to reproduce the error 
-and fix the issue by modifying the intermediate IR of the test case.
+For this OpenCL code:
 
-However, the same modifications translated into shamrock did not resolve the issue there.
+     __kernel void test( __global unsigned int * input, __global unsigned int *outMaxes ) {
+          __local unsigned int localStorage[256*4];
+          [...]
+           }
 
-This issue may be the cause of many of the other basic test failures which
-involve vector parameters used in JIT'd ARM kernels.
+      The LLVM IR produced is:
 
-These other tests fail due to unexpected results being returned from the JIT'ed
-kernels on ARM:
+             @test.localStorage = internal unnamed_addr addrspace(2) global [256 x i32] zeroinitializer, align 4
 
-local_kernel_scope
-explicit_s2v_<type>
-fpmath_float4
-intmath_int4
-intmath_long2
-
-TODO:
-=====
-kernel_memory_alignment_local - clSetKenrelArg failed.
-vload_local - clSetKernelArg failed
-vstore_local - clSetKernelArg failed
-local_arg_def - clCreateBuffer failed.
 
+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.