Test: Added unit test to help debug Khronos basic_parameter_types vector test

See tests/basic_parameter_types.txt for usage.

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

6 files changed, 363 insertions, 0 deletions

diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt
index 3b4175f..a5d8d7a 100644
--- a/tests/CMakeLists.txt
+++ b/tests/CMakeLists.txt
@@ -8,6 +8,7 @@ set(OPENCL_TESTS_SOURCE
     test_context.cpp
     test_commandqueue.cpp
     test_mem.cpp
+    test_basic_parameter_types.cpp
     test_kernel.cpp
     test_program.cpp
     test_builtins.cpp
@@ -25,6 +26,7 @@ OPENCL_TEST(tests device)
 OPENCL_TEST(tests context)
 OPENCL_TEST(tests commandqueue)
 OPENCL_TEST(tests mem)
+OPENCL_TEST(tests basic_parameter_types)
 OPENCL_TEST(tests kernel)
 OPENCL_TEST(tests program)
 OPENCL_TEST(tests builtins)
diff --git a/tests/basic_parameter_types.cl b/tests/basic_parameter_types.cl
new file mode 100644
index 0000000..99ea4e7
--- /dev/null
+++ b/tests/basic_parameter_types.cl
@@ -0,0 +1,4 @@
+__kernel void test_kernel(char4 c, __global float4 *result)
+{
+  result[0] = convert_float4(c);
+}
diff --git a/tests/basic_parameter_types.txt b/tests/basic_parameter_types.txt
new file mode 100644
index 0000000..5bbb483
--- /dev/null
+++ b/tests/basic_parameter_types.txt
@@ -0,0 +1,158 @@
+
+Instructions for reproducing test_kernel vector store error in shamrock:
+
+1. Clone (new) public Linaro shamrock repo (gpgpu/shamrock.git) and checkout branch: basic_parameter_types.
+
+2. Build per shamrock readme.
+
+2. % cp <shamrock-src>/tests/basic_parameter_types.cl <shamrock-build>/tests
+
+3. Run the following gdb session (see comments lines with #### ):
+
+tests> gdb --args tests basic_parameter_types nofork
+GNU gdb (GDB) 7.5.91.20130417-cvs-ubuntu
+Copyright (C) 2013 Free Software Foundation, Inc.
+License GPLv3+: GNU GPL version 3 or later <http://gnu.org/licenses/gpl.html>
+This is free software: you are free to change and redistribute it.
+There is NO WARRANTY, to the extent permitted by law.  Type "show copying"
+and "show warranty" for details.
+This GDB was configured as "arm-linux-gnueabihf".
+For bug reporting instructions, please see:
+<http://www.gnu.org/software/gdb/bugs/>...
+Reading symbols from /home/user/shamrock_build/tests/tests...done.
+(gdb) b main
+Breakpoint 1 at 0xb45a: file /home/user/shamrock/tests/tests.c, line 44.
+(gdb) run
+Starting program: /home/user/shamrock_build/tests/tests basic_parameter_types nofork
+[Thread debugging using libthread_db enabled]
+Using host libthread_db library "/lib/arm-linux-gnueabihf/libthread_db.so.1".
+
+Breakpoint 1, main (argc=3, argv=0xbefff704)
+    at /home/user/shamrock/tests/tests.c:44
+44	    Suite *s = NULL;
+(gdb) b Coal::CPUKernelWorkGroup::run
+Breakpoint 2 at 0xb53e506a: file /home/user/shamrock/src/core/cpu/kernel.cpp, line 649.
+(gdb) c
+Continuing.
+Running suite(s): basic_parameter_types
+[New Thread 0xb4de4450 (LWP 9283)]
+[New Thread 0xb45e4450 (LWP 9284)]
+[Switching to Thread 0xb4de4450 (LWP 9283)]
+
+Breakpoint 2, Coal::CPUKernelWorkGroup::run (this=0xb3800468)
+    at /home/user/shamrock/src/core/cpu/kernel.cpp:649
+649	    std::vector<void *> locals_to_free;
+(gdb) list
+644	}
+645
+646	bool CPUKernelWorkGroup::run()
+647	{
+648	    // Get the kernel function to call
+649	    std::vector<void *> locals_to_free;
+650	    llvm::Function *kernel_func = p_kernel->callFunction();
+651
+652	    if (!kernel_func)
+653	        return false;
+(gdb)
+654
+655	    Program *p = (Program *)p_kernel->kernel()->parent();
+656	    CPUProgram *prog = (CPUProgram *)(p->deviceDependentProgram(p_kernel->device()));
+657
+658	    // Make object usable for execution:  (only applies to MCJIT):
+659	    prog->jit()->finalizeObject();
+660
+661		std::string kname = kernel_func->getName().str();
+662
+663		// original
+(gdb)
+664	    p_kernel_func_addr =
+665	        (void(*)(void *))prog->jit()->getPointerToFunction(kernel_func);
+666
+667		// TAG
+668		// llvm::Function *t_func = prog->jit()->FindFunctionNamed(p_kernel->p_kernel->p_name->str());
+669	//	llvm::Function *t_func = prog->jit()->FindFunctionNamed(p_kernel->kernel()->p_name.c_str());
+670	//	p_kernel_func_addr = (void(*)(void *))prog->jit()->getPointerToFunction(t_func);
+671		p_kernel_func_addr =(void(*)(void *)) prog->jit()->getFunctionAddress(kname);
+672
+673	    // Get the arguments
+(gdb) b 661   #### run to after finalizeObject(), test_kernel symbol now exists
+Breakpoint 3 at 0xb53e50ca: file /home/user/shamrock/src/core/cpu/kernel.cpp, line 661.
+(gdb) c
+Continuing.
+
+Breakpoint 3, Coal::CPUKernelWorkGroup::run (this=0xb3800468)
+    at /home/user/shamrock/src/core/cpu/kernel.cpp:661
+661		std::string kname = kernel_func->getName().str();
+(gdb) b test_kernel
+Breakpoint 4 at 0xb50cf008
+(gdb) c
+Continuing.
+
+Breakpoint 4, 0xb50cf008 in test_kernel ()
+(gdb) disass
+Dump of assembler code for function test_kernel:
+   0xb50cf000 <+0>:	sub	sp, sp, #4
+   0xb50cf004 <+4>:	str	r0, [sp]
+=> 0xb50cf008 <+8>:	mov	r0, sp
+   0xb50cf00c <+12>:	vld1.32	{d16[0]}, [r0 :32]
+   0xb50cf010 <+16>:	vmovl.u8	q8, d16
+   0xb50cf014 <+20>:	vmov.u16	r2, d16[2]
+   0xb50cf018 <+24>:	vmov.u16	r0, d16[0]
+   0xb50cf01c <+28>:	sxtb	r2, r2
+   0xb50cf020 <+32>:	sxtb	r0, r0
+   0xb50cf024 <+36>:	vmov.32	d20[0], r2
+   0xb50cf028 <+40>:	vmov.u16	r2, d16[1]
+   0xb50cf02c <+44>:	vmov.32	d18[0], r0
+   0xb50cf030 <+48>:	vmov.u16	r0, d16[3]
+   0xb50cf034 <+52>:	sxtb	r2, r2
+   0xb50cf038 <+56>:	sxtb	r0, r0
+   0xb50cf03c <+60>:	vmov.32	d18[1], r2
+   0xb50cf040 <+64>:	vmov.32	d20[1], r0
+   0xb50cf044 <+68>:	vcvt.f32.s32	q8, q9
+   0xb50cf048 <+72>:	vcvt.f32.s32	q9, q10
+   0xb50cf04c <+76>:	vext.8	q8, q8, q8, #8
+   0xb50cf050 <+80>:	vext.8	q8, q8, q9, #8
+   0xb50cf054 <+84>:	vst1.64	{d16-d17}, [r1 :128]
+   0xb50cf058 <+88>:	add	sp, sp, #4
+   0xb50cf05c <+92>:	bx	lr
+End of assembler dump.
+(gdb) stepi 18
+0xb50cf050 in test_kernel ()
+(gdb) stepi   ### step down to vst1.64 instruction.
+0xb50cf054 in test_kernel ()
+(gdb) display/x $r1   ### $r1 register should hold output float4 *result ptr
+2: /x $r1 = 0xf94bc8
+(gdb) x/4f $r1  #### hmmm, shouldn't these be all 0's?
+0xf94bc8:	-4.7186802e-07	2.39711496e-38	2.12743103e-07	7.56701171e-44
+(gdb) display/x $d16  ### show that convert_float4(char4) worked: {0,1,2,3}
+3: /x $d16 = {u8 = {0x0, 0x0, 0x0, 0x0, 0x0, 0x0, 0x80, 0x3f}, u16 = {0x0,
+    0x0, 0x0, 0x3f80}, u32 = {0x0, 0x3f800000}, u64 = 0x3f80000000000000,
+  f32 = {0x0, 0x1}, f64 = 0x0}
+(gdb) display/x $d17
+4: /x $d17 = {u8 = {0x0, 0x0, 0x0, 0x40, 0x0, 0x0, 0x40, 0x40}, u16 = {0x0,
+    0x4000, 0x0, 0x4040}, u32 = {0x40000000, 0x40400000},
+  u64 = 0x4040000040000000, f32 = {0x2, 0x3}, f64 = 0x20}
+(gdb) stepi ### now stepi vst1.64 instruction, should store {0,1,2,3} to [$r1]
+0xb50cf058 in test_kernel ()
+4: /x $d17 = {u8 = {0x0, 0x0, 0x0, 0x40, 0x0, 0x0, 0x40, 0x40}, u16 = {0x0,
+    0x4000, 0x0, 0x4040}, u32 = {0x40000000, 0x40400000},
+  u64 = 0x4040000040000000, f32 = {0x2, 0x3}, f64 = 0x20}
+3: /x $d16 = {u8 = {0x0, 0x0, 0x0, 0x0, 0x0, 0x0, 0x80, 0x3f}, u16 = {0x0,
+    0x0, 0x0, 0x3f80}, u32 = {0x0, 0x3f800000}, u64 = 0x3f80000000000000,
+  f32 = {0x0, 0x1}, f64 = 0x0}
+2: /x $r1 = 0xf940d9   #### OOPS!  $r1 register changed! NOT POSSIBLE?
+(gdb) x/4f $r1
+0xf940d9:	1.66694513e-24	-7.21760423e-29	1.77866814e-41	-3.85185989e-34
+(gdb) x/4f 0xf94bc8  #### and memory at old $r1 (result ptr) has garbage.
+0xf94bc8:	4.20389539e-45	2.39711496e-38	2.12743103e-07	7.56701171e-44
+(gdb) c
+Continuing.
+Conversion from char failed: got 4.2039e-45,expected 0
+Conversion from char failed: got 2.39711e-38,expected 1
+Conversion from char failed: got 2.12743e-07,expected 2
+Conversion from char failed: got 7.56701e-44,expected 3
+0%: Checks: 1, Failures: 1, Errors: 0
+/home/user/shamrock/tests/test_basic_parameter_types.cpp:139:F:basic_parameter_types:test_basic_parameter_types:0: the kernel hasn't done its job, the buffer is wrong
+[Thread 0xb4de4450 (LWP 9283) exited]
+[Thread 0xb45e4450 (LWP 9284) exited]
+[Inferior 1 (process 9280) exited with code 01]
diff --git a/tests/test_basic_parameter_types.cpp b/tests/test_basic_parameter_types.cpp
new file mode 100644
index 0000000..680ecaf
--- /dev/null
+++ b/tests/test_basic_parameter_types.cpp
@@ -0,0 +1,154 @@
+/*
+ * Copyright (c) 2011, Denis Steckelmacher <steckdenis@yahoo.fr>
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *     * Redistributions of source code must retain the above copyright
+ *       notice, this list of conditions and the following disclaimer.
+ *     * Redistributions in binary form must reproduce the above copyright
+ *       notice, this list of conditions and the following disclaimer in the
+ *       documentation and/or other materials provided with the distribution.
+ *     * Neither the name of the copyright holder nor the
+ *       names of its contributors may be used to endorse or promote products
+ *       derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL THE CONTRIBUTORS BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+#include <iostream>
+#include <stdio.h>
+#include <stdlib.h>
+
+#include "test_basic_parameter_types.h"
+#include "CL/cl.h"
+
+#define MAX_SOURCE_SIZE (0x51200)  // 50K chars max source file
+
+START_TEST (test_basic_parameter_types)
+{
+    cl_platform_id platform = 0;
+    cl_device_id device;
+    cl_context ctx;
+    cl_command_queue queue;
+    cl_program program;
+    cl_int error;
+    cl_kernel kernel;
+    cl_mem results;
+    cl_char c[4]={0,1,2,3};   // input: one char4 vector.
+    float results_back[1*4] = {0.0f, 0.0f, 0.0f, 0.0f};  // output: one float4 vector.
+    bool ok = true;
+    float expected;
+
+    const size_t global[3] = {1, 1, 1};
+    FILE *fp;
+    const char fileName[] = "basic_parameter_types.cl";  // Note: cp from ~/shamrock/tests/
+    size_t source_size;
+    char *source_str;
+
+    /* Load kernel source code */
+    fp = fopen(fileName, "r");
+    fail_if(
+        fp == NULL,
+        "unable to open CL file"
+    );
+    source_str = (char *)malloc(MAX_SOURCE_SIZE);
+    source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);
+    fclose(fp);
+
+    error = clGetDeviceIDs(platform, CL_DEVICE_TYPE_DEFAULT, 1, &device, 0);
+    fail_if(
+        error != CL_SUCCESS,
+        "unable to get the default device"
+    );
+
+    ctx = clCreateContext(0, 1, &device, 0, 0, &error);
+    fail_if(
+        error != CL_SUCCESS || ctx == 0,
+        "unable to create a valid context"
+    );
+
+    queue = clCreateCommandQueue(ctx, device, 0, &error);
+    fail_if(
+        error != CL_SUCCESS,
+        "unable to create a command queue"
+    );
+
+    program = clCreateProgramWithSource(ctx, 1, (const char **)&source_str,
+               (const size_t *)&source_size, &error);
+    fail_if(
+        error != CL_SUCCESS,
+        "cannot create a program from source with sane arguments"
+    );
+
+    error = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
+    fail_if(
+        error != CL_SUCCESS,
+        "cannot build a valid program"
+    );
+
+    kernel  = clCreateKernel(program, "test_kernel", &error);
+    fail_if(
+        error != CL_SUCCESS,
+        "unable to create a valid kernel"
+    );
+
+    // Create the results buffer
+    results = clCreateBuffer(ctx, CL_MEM_READ_WRITE, sizeof(cl_float)*1*4, NULL, &error);
+    fail_if(
+        error != CL_SUCCESS,
+        "cannot create a valid read-write buffer"
+    );
+
+    error = clSetKernelArg(kernel, 0, sizeof(cl_char)*4, &c);
+    fail_if(
+        error != CL_SUCCESS,
+        "cannot set kernel argument"
+    );
+    error = clSetKernelArg(kernel, 1, sizeof(cl_mem), &results);
+
+    error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global, 0, 0, 0, NULL);
+    fail_if(
+        error != CL_SUCCESS,
+        "unable to queue a NDRange kernel with local work size guessed"
+    );
+
+    // Read back the results
+    error = clEnqueueReadBuffer(queue, results, CL_TRUE, 0, sizeof(cl_float)*1*4, results_back, 0, NULL, NULL);
+
+    // Verify the results
+    for (int index = 0; index < 4; index++) {
+        expected = (float)c[index];
+        if (results_back[index] != expected) {
+            std::cout << "Conversion from char failed: got " << results_back[index] << ",expected " << expected << std::endl;
+	    ok = false;
+        }
+    }
+    fail_if(
+        ok == false,
+        "the kernel hasn't done its job, the buffer is wrong"
+    );
+
+    clReleaseMemObject(results);
+    clReleaseKernel(kernel);
+    clReleaseProgram(program);
+    clReleaseContext(ctx);
+}
+END_TEST
+
+TCase *cl_bpt_tcase_create(void)
+{
+    TCase *tc = NULL;
+    tc = tcase_create("basic_parameter_types");
+    tcase_add_test(tc, test_basic_parameter_types);
+    return tc;
+}
diff --git a/tests/test_basic_parameter_types.h b/tests/test_basic_parameter_types.h
new file mode 100644
index 0000000..8dd346d
--- /dev/null
+++ b/tests/test_basic_parameter_types.h
@@ -0,0 +1,43 @@
+/*
+ * Copyright (c) 2011, Denis Steckelmacher <steckdenis@yahoo.fr>
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *     * Redistributions of source code must retain the above copyright
+ *       notice, this list of conditions and the following disclaimer.
+ *     * Redistributions in binary form must reproduce the above copyright
+ *       notice, this list of conditions and the following disclaimer in the
+ *       documentation and/or other materials provided with the distribution.
+ *     * Neither the name of the copyright holder nor the
+ *       names of its contributors may be used to endorse or promote products
+ *       derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL THE CONTRIBUTORS BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+#ifndef __UTEST_BPT__
+#define __UTEST_BPT__
+
+#include <check.h>
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+TCase *cl_bpt_tcase_create(void);
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif
diff --git a/tests/tests.c b/tests/tests.c
index 4db057b..80be728 100644
--- a/tests/tests.c
+++ b/tests/tests.c
@@ -30,6 +30,7 @@
 #include "test_context.h"
 #include "test_commandqueue.h"
 #include "test_mem.h"
+#include "test_basic_parameter_types.h"
 #include "test_kernel.h"
 #include "test_program.h"
 #include "test_builtins.h"
@@ -58,6 +59,7 @@ int main(int argc, char **argv)
     TESTSUITE(context, "context");
     TESTSUITE(commandqueue, "commandqueue");
     TESTSUITE(mem, "mem");
+    TESTSUITE(bpt, "basic_parameter_types");
     TESTSUITE(kernel, "kernel");
     TESTSUITE(program, "program");
     TESTSUITE(builtins, "builtins");