aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorGil Pitney <gil.pitney@linaro.org>2014-12-15 15:35:36 -0800
committerGil Pitney <gil.pitney@linaro.org>2014-12-15 15:35:36 -0800
commit4c383810169fe14b87a56d022b61f5155154f2e6 (patch)
tree8c65f1a2f3a31647f84cad65bd996e42432dbe16
parent92573c6df908d1845f28f0679ebc1ecefb9a7124 (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.txt2
-rw-r--r--tests/basic_parameter_types.cl4
-rw-r--r--tests/basic_parameter_types.txt158
-rw-r--r--tests/test_basic_parameter_types.cpp154
-rw-r--r--tests/test_basic_parameter_types.h43
-rw-r--r--tests/tests.c2
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");