diff options
author | Gil Pitney <gil.pitney@linaro.org> | 2014-12-15 15:35:36 -0800 |
---|---|---|
committer | Gil Pitney <gil.pitney@linaro.org> | 2014-12-15 15:35:36 -0800 |
commit | 4c383810169fe14b87a56d022b61f5155154f2e6 (patch) | |
tree | 8c65f1a2f3a31647f84cad65bd996e42432dbe16 | |
parent | 92573c6df908d1845f28f0679ebc1ecefb9a7124 (diff) |
Remove test test_basic_parameter_types from sanity tests
This test was added as an aid to debug a particular Khronos test, which
is now fixed, so is being removed.
Signed-off-by: Gil Pitney <gil.pitney@linaro.org>
-rw-r--r-- | tests/CMakeLists.txt | 2 | ||||
-rw-r--r-- | tests/basic_parameter_types.cl | 4 | ||||
-rw-r--r-- | tests/basic_parameter_types.txt | 158 | ||||
-rw-r--r-- | tests/test_basic_parameter_types.cpp | 154 | ||||
-rw-r--r-- | tests/test_basic_parameter_types.h | 43 | ||||
-rw-r--r-- | tests/tests.c | 2 |
6 files changed, 0 insertions, 363 deletions
diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index a5d8d7a..3b4175f 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -8,7 +8,6 @@ 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 @@ -26,7 +25,6 @@ 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 deleted file mode 100644 index 66fb628..0000000 --- a/tests/basic_parameter_types.cl +++ /dev/null @@ -1,4 +0,0 @@ -__kernel void test_kernel(int4 i, __global float4 *result) -{ - result[0] = convert_float4(i); -} diff --git a/tests/basic_parameter_types.txt b/tests/basic_parameter_types.txt deleted file mode 100644 index 5bbb483..0000000 --- a/tests/basic_parameter_types.txt +++ /dev/null @@ -1,158 +0,0 @@ - -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 deleted file mode 100644 index b703b16..0000000 --- a/tests/test_basic_parameter_types.cpp +++ /dev/null @@ -1,154 +0,0 @@ -/* - * 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_int i[4]={0,1,2,3}; // input: one int4 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_int)*4, &i); - 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)i[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 deleted file mode 100644 index 8dd346d..0000000 --- a/tests/test_basic_parameter_types.h +++ /dev/null @@ -1,43 +0,0 @@ -/* - * 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 80be728..4db057b 100644 --- a/tests/tests.c +++ b/tests/tests.c @@ -30,7 +30,6 @@ #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" @@ -59,7 +58,6 @@ 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"); |