From: Andreas Beckmann <anbe@debian.org>
Date: Thu, 10 Dec 2020 18:47:30 +0000 (+0100)
Subject: [PATCH] add testcase for Debian bug #975931
X-Git-Tag: archive/raspbian/1.8-1+rpi1^2~5
X-Git-Url: https://dgit.raspbian.org/?a=commitdiff_plain;h=e90a78b3d8e12d42a4d03261b29a97a2bedb2ec6;p=pocl.git

[PATCH] add testcase for Debian bug #975931


Gbp-Pq: Name 2001-add-testcase-for-Debian-bug-975931.patch
---

diff --git a/tests/regression/CMakeLists.txt b/tests/regression/CMakeLists.txt
index 9220ae1..e53ad99 100644
--- a/tests/regression/CMakeLists.txt
+++ b/tests/regression/CMakeLists.txt
@@ -40,6 +40,7 @@ set(PROGRAMS_TO_BUILD test_barrier_between_for_loops test_early_return
   test_for_with_var_iteration_count test_id_dependent_computation
   test_locals test_loop_phi_replication test_multi_level_loops_with_barriers
   test_simple_for_with_a_barrier test_structs_as_args test_vectors_as_args
+  test_llvm_segfault_debian_bug_975931
   test_barrier_before_return test_infinite_loop test_constant_array
   test_undominated_variable test_setargs test_null_arg
   test_fors_with_var_iteration_counts test_issue_231 test_issue_445
@@ -80,6 +81,8 @@ add_test_pocl(NAME "regression/test_issue_893" COMMAND "test_issue_893")
 
 add_test_pocl(NAME "regression/test_flatten_barrier_subs" COMMAND "test_flatten_barrier_subs" EXPECTED_OUTPUT "test_flatten_barrier_subs.output")
 
+add_test_pocl(NAME "regression/test_llvm_segfault_debian_bug_975931" COMMAND "test_llvm_segfault_debian_bug_975931")
+
 # repl
 
 add_test_pocl(NAME "regression/phi_nodes_not_replicated_REPL" COMMAND "test_loop_phi_replication")
diff --git a/tests/regression/test_llvm_segfault_debian_bug_975931.cpp b/tests/regression/test_llvm_segfault_debian_bug_975931.cpp
new file mode 100644
index 0000000..220fe62
--- /dev/null
+++ b/tests/regression/test_llvm_segfault_debian_bug_975931.cpp
@@ -0,0 +1,173 @@
+/*
+
+   Copyright (c) 2013 Pekka Jääskeläinen and
+                      Kalray
+
+   Permission is hereby granted, free of charge, to any person obtaining a copy
+   of this software and associated documentation files (the "Software"), to deal
+   in the Software without restriction, including without limitation the rights
+   to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+   copies of the Software, and to permit persons to whom the Software is
+   furnished to do so, subject to the following conditions:
+
+   The above copyright notice and this permission notice shall be included in
+   all copies or substantial portions of the Software.
+
+   THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+   IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+   FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+   AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+   LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+   OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+   THE SOFTWARE.
+*/
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <CL/cl.h>
+#include <poclu.h>
+#include "config.h"
+
+const char source[] =
+"#ifdef DOUBLE_PRECISION\n"
+"    #ifdef cl_khr_fp64\n"
+"    #pragma OPENCL EXTENSION cl_khr_fp64 : enable\n"
+"    #else\n"
+"    #pragma OPENCL EXTENSION cl_amd_fp64 : enable\n"
+"    #endif\n"
+"#endif\n"
+"\n"
+"__kernel void Sdot_kernel( __global float *_X, __global float *_Y, __global float *scratchBuff,\n"
+"                                        uint N, uint offx, int incx, uint offy, int incy, int doConj )\n"
+"{\n"
+"    __global float *X = _X + offx;\n"
+"    __global float *Y = _Y + offy;\n"
+"    float dotP = (float) 0.0;\n"
+"\n"
+"    if ( incx < 0 ) {\n"
+"        X = X + (N - 1) * abs(incx);\n"
+"    }\n"
+"    if ( incy < 0 ) {\n"
+"        Y = Y + (N - 1) * abs(incy);\n"
+"    }\n"
+"\n"
+"    int gOffset;\n"
+"    for( gOffset=(get_global_id(0) * 4); (gOffset + 4 - 1)<N; gOffset+=( get_global_size(0) * 4 ) )\n"
+"    {\n"
+"        float4 vReg1, vReg2, res;\n"
+"\n"
+"        #ifdef INCX_NONUNITY\n"
+"             vReg1 = (float4)(  (X + (gOffset*incx))[0 + ( incx * 0)],  (X + (gOffset*incx))[0 + ( incx * 1)],  (X + (gOffset*incx))[0 + ( incx * 2)],  (X + (gOffset*incx))[0 + ( incx * 3)]);\n"
+"        #else\n"
+"            vReg1 = vload4(  0, (__global float *) (X + gOffset) );\n"
+"        #endif\n"
+"\n"
+"        #ifdef INCY_NONUNITY\n"
+"             vReg2 = (float4)(  (Y + (gOffset*incy))[0 + ( incy * 0)],  (Y + (gOffset*incy))[0 + ( incy * 1)],  (Y + (gOffset*incy))[0 + ( incy * 2)],  (Y + (gOffset*incy))[0 + ( incy * 3)]);\n"
+"        #else\n"
+"            vReg2 = vload4(  0, (__global float *) (Y + gOffset) );\n"
+"        #endif\n"
+"\n"
+"        ;\n"
+"         res =  vReg1 *  vReg2 ;\n"
+"        dotP +=  res .S0 +  res .S1 +  res .S2 +  res .S3;\n"
+";          // Add-up elements in the vector to give a scalar\n"
+"    }\n"
+"\n"
+"    // Loop for the last thread to handle the tail part of the vector\n"
+"    // Using the same gOffset used above\n"
+"    for( ; gOffset<N; gOffset++ )\n"
+"    {\n"
+"        float sReg1, sReg2, res;\n"
+"        sReg1 = X[gOffset * incx];\n"
+"        sReg2 = Y[gOffset * incy];\n"
+"\n"
+"        ;\n"
+"             res =  sReg1 *  sReg2 ;\n"
+"             dotP =  dotP +  res ;\n"
+"        }\n"
+"\n"
+"    // Note: this has to be called outside any if-conditions- because REDUCTION uses barrier\n"
+"    // dotP of work-item 0 will have the final reduced item of the work-group\n"
+"    __local float p1753 [ 64 ];\n"
+"	uint QKiD0 = get_local_id(0);\n"
+"	 p1753 [ QKiD0 ] =  dotP ;\n"
+"	barrier(CLK_LOCAL_MEM_FENCE);\n"
+"\n"
+"	if( QKiD0 < 32 ) {\n"
+"		 p1753 [ QKiD0 ] = p1753 [ QKiD0 ] + p1753 [ QKiD0 + 32 ];\n"
+"	}\n"
+"	barrier(CLK_LOCAL_MEM_FENCE);\n"
+"\n"
+"	if( QKiD0 < 16 ) {\n"
+"		 p1753 [ QKiD0 ] = p1753 [ QKiD0 ] + p1753 [ QKiD0 + 16 ];\n"
+"	}\n"
+"	barrier(CLK_LOCAL_MEM_FENCE);\n"
+"\n"
+"	if( QKiD0 < 8 ) {\n"
+"		 p1753 [ QKiD0 ] = p1753 [ QKiD0 ] + p1753 [ QKiD0 + 8 ];\n"
+"	}\n"
+"	barrier(CLK_LOCAL_MEM_FENCE);\n"
+"\n"
+"	if( QKiD0 < 4 ) {\n"
+"		 p1753 [ QKiD0 ] = p1753 [ QKiD0 ] + p1753 [ QKiD0 + 4 ];\n"
+"	}\n"
+"	barrier(CLK_LOCAL_MEM_FENCE);\n"
+"\n"
+"	if( QKiD0 < 2 ) {\n"
+"		 p1753 [ QKiD0 ] = p1753 [ QKiD0 ] + p1753 [ QKiD0 + 2 ];\n"
+"	}\n"
+"	barrier(CLK_LOCAL_MEM_FENCE);\n"
+"\n"
+"	if( QKiD0 == 0 ) {\n"
+"	 dotP  = p1753 [0] + p1753 [1];\n"
+"	}\n"
+"\n"
+"    if( (get_local_id(0)) == 0 ) {\n"
+"        scratchBuff[ get_group_id(0) ] = dotP;\n"
+"    }\n"
+"}\n"
+"\n"
+;
+
+#define MAX_PLATFORMS 32
+#define MAX_DEVICES   32
+#define MAX_BINARIES  32
+
+int main(){
+  cl_int err;
+  cl_platform_id platforms[MAX_PLATFORMS];
+  cl_uint nplatforms;
+  cl_device_id devices[MAX_DEVICES];
+  cl_uint ndevices;
+  cl_program program = NULL;
+  size_t binsizes[MAX_BINARIES];
+  size_t nbinaries;
+
+  CHECK_CL_ERROR(clGetPlatformIDs(MAX_PLATFORMS, platforms, &nplatforms));
+  TEST_ASSERT(nplatforms > 0);
+
+  CHECK_CL_ERROR(clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_ALL, MAX_DEVICES, devices, &ndevices));
+  TEST_ASSERT(ndevices > 0);
+
+  cl_context context = clCreateContext(NULL, 1, devices, NULL, NULL, &err);
+  CHECK_OPENCL_ERROR_IN("clCreateContext");
+
+  const char * src[] = {source};
+  program = clCreateProgramWithSource(context, 1, src, NULL, &err);
+  CHECK_OPENCL_ERROR_IN("clCreateProgramWithSource");
+
+  CHECK_CL_ERROR(clBuildProgram(program, 1, devices, "-g -DINCX_NONUNITY -DINCY_NONUNITY", NULL, NULL));
+
+  CHECK_CL_ERROR(clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES, sizeof(binsizes), binsizes, &nbinaries));
+  printf("binary size: %zd\n", binsizes[0]);
+
+  CHECK_CL_ERROR(clReleaseProgram(program));
+
+  CHECK_CL_ERROR (clReleaseContext (context));
+
+  printf ("OK\n");
+
+  return EXIT_SUCCESS;
+}