[Pkg-opencl-devel] Bug#913141: beignet: Segmentation fault while running opencv_perf_dnn
Rebecca N. Palmer
rebecca_palmer at zoho.com
Sun Nov 18 16:09:34 GMT 2018
Control: reassign -1 beignet-opencl-icd
A somewhat (but not very) cut-down version of the test case.
This produces a slightly different backtrace (the "No symbol table info
available" is LLVM bug #914021):
Thread 1 "bug913141" received signal SIGSEGV, Segmentation fault.
0x00007ffff19f68c0 in llvm::Value::getName() const ()
from /usr/lib/x86_64-linux-gnu/libLLVM-6.0.so.1
(gdb) bt full
#0 0x00007ffff19f68c0 in llvm::Value::getName() const ()
from /usr/lib/x86_64-linux-gnu/libLLVM-6.0.so.1
No symbol table info available.
#1 0x00007ffff26d1ae3 in
llvm::LPPassManager::runOnFunction(llvm::Function&) () from
/usr/lib/x86_64-linux-gnu/libLLVM-6.0.so.1
No symbol table info available.
#2 0x00007ffff19b4078 in
llvm::FPPassManager::runOnFunction(llvm::Function&) () from
/usr/lib/x86_64-linux-gnu/libLLVM-6.0.so.1
No symbol table info available.
#3 0x00007ffff2619833 in ?? ()
from /usr/lib/x86_64-linux-gnu/libLLVM-6.0.so.1
No symbol table info available.
#4 0x00007ffff19b39b6 in llvm::legacy::PassManagerImpl::run(llvm::Module&)
() from /usr/lib/x86_64-linux-gnu/libLLVM-6.0.so.1
No symbol table info available.
#5 0x00007ffff4fd6be4 in gbe::runModulePass(llvm::Module&,
llvm::TargetLibraryInfoImpl*, llvm::DataLayout const&, int, bool) ()
at ./backend/src/llvm/llvm_to_gen.cpp:222
MPM = <incomplete type>
#6 0x00007ffff4fd6fcd in gbe::llvmToGen(gbe::ir::Unit&, void const*,
int, bool, int, std::__cxx11::basic_string<char, std::char_traits<char>,
std::allocator<char> >&) () at ./backend/src/llvm/llvm_to_gen.cpp:341
errInfo = ""
o = <optimized out>
cl_mod = <optimized out>
passes__ = <incomplete type>
-------------- next part --------------
A non-text attachment was scrubbed...
Name: bug913141.cpp
Type: text/x-c++src
Size: 968 bytes
Desc: not available
URL: <http://alioth-lists.debian.net/pipermail/pkg-opencl-devel/attachments/20181118/3ff3edd9/attachment.cpp>
-------------- next part --------------
//This is originally taken from opencv_perf_dnn (part of https://github.com/opencv/opencv)
//but has been heavily cut down for use as a compiler-bug test case and is no longer expected to be functional
//Original's copyright notice:
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2017, Intel Corporation, all rights reserved.
// Copyright (c) 2016-2017 Fabian David Tschopp, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's 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.
//
// * The name of the copyright holders may not 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 Intel Corporation or 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.
//
//M*/
__kernel void U_GEMM_LIKE_CONV_k3x3_cn576_g1_s1x1_d1x1_b1_in64x48_p1x1_num1_M512_activ0_eltwise0_FP32_5_1_8_32_SIMD8( const __global float *src0, const __global float *src1, __global float * biases_base, __global float *dst_base, const int dst_offset, const ushort input_width, const ushort input_height, const ushort output_width, const ushort output_height, const int out_pitch_y, const int out_pitch_z, const int aligned_input_size, const int slice_pitch)
{
__global float *dst = dst_base + dst_offset;
const int group_x = get_group_id(0);
const int group_y = get_group_id(1);
const int global_x = get_global_id(0);
const int global_y = get_global_id(1);
const int global_z = get_global_id(2);
{
float8 blockC00 = 0.f;
float8 blockC10 = 0.f;
float8 blockC20 = 0.f;
float8 blockC30 = 0.f;
int curr_x = ( global_y % output_width );
int curr_y = ( global_y / output_width );
int saved_y = curr_y;
const __global float *src0_read = src0
+ aligned_input_size * global_z
+ (curr_y - 1) * input_width
+ (curr_x - 1);
const __global float *src1_read = src1 + ( global_x * 32 * 2);
int patch_depth = 0;
do//removing this (outer) while-loop makes it not crash
{
int patch_row = 0;
curr_y = saved_y;
do
{
float3 blockA00;//crashes whether this is an array or a vector
float* pblockA00 = (float*)(&blockA00);
/*doesn't crash if these 3 are made 'if 1', 'if 0' or a mix of those*/
if (curr_y >= 1 && curr_y < input_height + 1) pblockA00[0] = src0_read[0];
else pblockA00[0] = 0;
if (curr_y >= 1 && curr_y < input_height + 1) pblockA00[1] = src0_read[1];
else pblockA00[1] = 0;
if (curr_y >= 1 && curr_y < input_height + 1) pblockA00[2] = src0_read[2];
else pblockA00[2] = 0;
curr_y += 1;
src0_read += (input_width);
{ blockC00.s0 = mad( pblockA00[0], 6, blockC00.s0 );/*doesn't crash if either all the pblockA00 reads or all the blockC* reads are replaced with immediate numbers; doesn't crash if all the pblockA00 reads are replaced with blockA00.x/y/z, but _does_ crash if all of them are made pblockA00[0] not 0/1/2*/
blockC00.s1 = mad( pblockA00[0], 6, blockC00.s1 );
blockC00.s2 = mad( pblockA00[0], 6, blockC00.s2 );
blockC00.s3 = mad( pblockA00[0], 6, blockC00.s3 );
blockC00.s4 = mad( pblockA00[0], 6, blockC00.s4 );
blockC00.s5 = mad( pblockA00[0], 6, blockC00.s5 );
blockC00.s6 = mad( pblockA00[0], 6, blockC00.s6 );
blockC00.s7 = mad( pblockA00[0], 6, blockC00.s7 );
};
{ blockC00.s0 = mad( pblockA00[1], 6, blockC00.s0 );
blockC00.s1 = mad( pblockA00[1], 6, blockC00.s1 );
blockC00.s2 = mad( pblockA00[1], 6, blockC00.s2 );
blockC00.s3 = mad( pblockA00[1], 6, blockC00.s3 );
blockC00.s4 = mad( pblockA00[1], 6, blockC00.s4 );
blockC00.s5 = mad( pblockA00[1], 6, blockC00.s5 );
blockC00.s6 = mad( pblockA00[1], 6, blockC00.s6 );
blockC00.s7 = mad( pblockA00[1], 6, blockC00.s7 );
};
{ blockC10.s0 = mad( pblockA00[0], 6, blockC10.s0 );
blockC10.s1 = mad( pblockA00[0], 6, blockC10.s1 );
blockC10.s2 = mad( pblockA00[0], 6, blockC10.s2 );
blockC10.s3 = mad( pblockA00[0], 6, blockC10.s3 );
blockC10.s4 = mad( pblockA00[0], 6, blockC10.s4 );
blockC10.s5 = mad( pblockA00[0], 6, blockC10.s5 );
blockC10.s6 = mad( pblockA00[0], 6, blockC10.s6 );
blockC10.s7 = mad( pblockA00[0], 6, blockC10.s7 );
};
{ blockC10.s0 = mad( pblockA00[1], 6, blockC10.s0 );
blockC10.s1 = mad( pblockA00[1], 6, blockC10.s1 );
blockC10.s2 = mad( pblockA00[1], 6, blockC10.s2 );
blockC10.s3 = mad( pblockA00[1], 6, blockC10.s3 );
blockC10.s4 = mad( pblockA00[1], 6, blockC10.s4 );
blockC10.s5 = mad( pblockA00[1], 6, blockC10.s5 );
blockC10.s6 = mad( pblockA00[1], 6, blockC10.s6 );
blockC10.s7 = mad( pblockA00[1], 6, blockC10.s7 );
};/*after removing blockCxx... beyond here, doesn't crash at all*/
{ blockC20.s0 = mad( pblockA00[0], 6, blockC20.s0 );
blockC20.s1 = mad( pblockA00[0], 6, blockC20.s1 );
blockC20.s2 = mad( pblockA00[0], 6, blockC20.s2 );
blockC20.s3 = mad( pblockA00[0], 6, blockC20.s3 );
blockC20.s4 = mad( pblockA00[0], 6, blockC20.s4 );
blockC20.s5 = mad( pblockA00[0], 6, blockC20.s5 );
blockC20.s6 = mad( pblockA00[0], 6, blockC20.s6 );
blockC20.s7 = mad( pblockA00[0], 6, blockC20.s7 );
};
{ blockC20.s0 = mad( pblockA00[1], 6, blockC20.s0 );
blockC20.s1 = mad( pblockA00[1], 6, blockC20.s1 );
blockC20.s2 = mad( pblockA00[1], 6, blockC20.s2 );
blockC20.s3 = mad( pblockA00[1], 6, blockC20.s3 );
blockC20.s4 = mad( pblockA00[1], 6, blockC20.s4 );
blockC20.s5 = mad( pblockA00[1], 6, blockC20.s5 );
blockC20.s6 = mad( pblockA00[1], 6, blockC20.s6 );
blockC20.s7 = mad( pblockA00[1], 6, blockC20.s7 );
};//does crash after removing beyond here
{ blockC30.s0 = mad( pblockA00[0], 6, blockC30.s0 );
blockC30.s1 = mad( pblockA00[0], 6, blockC30.s1 );
blockC30.s2 = mad( pblockA00[0], 6, blockC30.s2 );
blockC30.s3 = mad( pblockA00[0], 6, blockC30.s3 );
blockC30.s4 = mad( pblockA00[0], 6, blockC30.s4 );
blockC30.s5 = mad( pblockA00[0], 6, blockC30.s5 );
blockC30.s6 = mad( pblockA00[0], 6, blockC30.s6 );
blockC30.s7 = mad( pblockA00[0], 6, blockC30.s7 );
};
{ blockC30.s0 = mad( pblockA00[1], 6, blockC30.s0 );
blockC30.s1 = mad( pblockA00[1], 6, blockC30.s1 );
blockC30.s2 = mad( pblockA00[1], 6, blockC30.s2 );
blockC30.s3 = mad( pblockA00[1], 6, blockC30.s3 );
blockC30.s4 = mad( pblockA00[1], 6, blockC30.s4 );
blockC30.s5 = mad( pblockA00[1], 6, blockC30.s5 );
blockC30.s6 = mad( pblockA00[1], 6, blockC30.s6 );
blockC30.s7 = mad( pblockA00[1], 6, blockC30.s7 );
};
{ blockC00.s0 = mad( pblockA00[2], 6, blockC00.s0 );
blockC00.s1 = mad( pblockA00[2], 6, blockC00.s1 );
blockC00.s2 = mad( pblockA00[2], 6, blockC00.s2 );
blockC00.s3 = mad( pblockA00[2], 6, blockC00.s3 );
blockC00.s4 = mad( pblockA00[2], 6, blockC00.s4 );
blockC00.s5 = mad( pblockA00[2], 6, blockC00.s5 );
blockC00.s6 = mad( pblockA00[2], 6, blockC00.s6 );
blockC00.s7 = mad( pblockA00[2], 6, blockC00.s7 );
};
{ blockC10.s0 = mad( pblockA00[2], 6, blockC10.s0 );
blockC10.s1 = mad( pblockA00[2], 6, blockC10.s1 );
blockC10.s2 = mad( pblockA00[2], 6, blockC10.s2 );
blockC10.s3 = mad( pblockA00[2], 6, blockC10.s3 );
blockC10.s4 = mad( pblockA00[2], 6, blockC10.s4 );
blockC10.s5 = mad( pblockA00[2], 6, blockC10.s5 );
blockC10.s6 = mad( pblockA00[2], 6, blockC10.s6 );
blockC10.s7 = mad( pblockA00[2], 6, blockC10.s7 );
};
{ blockC20.s0 = mad( pblockA00[2], 6, blockC20.s0 );
blockC20.s1 = mad( pblockA00[2], 6, blockC20.s1 );
blockC20.s2 = mad( pblockA00[2], 6, blockC20.s2 );
blockC20.s3 = mad( pblockA00[2], 6, blockC20.s3 );
blockC20.s4 = mad( pblockA00[2], 6, blockC20.s4 );
blockC20.s5 = mad( pblockA00[2], 6, blockC20.s5 );
blockC20.s6 = mad( pblockA00[2], 6, blockC20.s6 );
blockC20.s7 = mad( pblockA00[2], 6, blockC20.s7 );
};
{ blockC30.s0 = mad( pblockA00[2], 6, blockC30.s0 );
blockC30.s1 = mad( pblockA00[2], 6, blockC30.s1 );
blockC30.s2 = mad( pblockA00[2], 6, blockC30.s2 );
blockC30.s3 = mad( pblockA00[2], 6, blockC30.s3 );
blockC30.s4 = mad( pblockA00[2], 6, blockC30.s4 );
blockC30.s5 = mad( pblockA00[2], 6, blockC30.s5 );
blockC30.s6 = mad( pblockA00[2], 6, blockC30.s6 );
blockC30.s7 = mad( pblockA00[2], 6, blockC30.s7 );
};
}
while( ++patch_row < 3 );
}
while ( ++patch_depth < 576 );
}
}
More information about the Pkg-opencl-devel
mailing list