<html>
<head>
<base href="https://bugs.llvm.org/">
</head>
<body><table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Bug ID</th>
<td><a class="bz_bug_link
bz_status_NEW "
title="NEW - llvm-spir crashed"
href="https://bugs.llvm.org/show_bug.cgi?id=51996">51996</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>llvm-spir crashed
</td>
</tr>
<tr>
<th>Product</th>
<td>tools
</td>
</tr>
<tr>
<th>Version</th>
<td>12.0
</td>
</tr>
<tr>
<th>Hardware</th>
<td>PC
</td>
</tr>
<tr>
<th>OS</th>
<td>Linux
</td>
</tr>
<tr>
<th>Status</th>
<td>NEW
</td>
</tr>
<tr>
<th>Severity</th>
<td>enhancement
</td>
</tr>
<tr>
<th>Priority</th>
<td>P
</td>
</tr>
<tr>
<th>Component</th>
<td>llvm-as
</td>
</tr>
<tr>
<th>Assignee</th>
<td>unassignedbugs@nondot.org
</td>
</tr>
<tr>
<th>Reporter</th>
<td>david.livshin@dalsoft.com
</td>
</tr>
<tr>
<th>CC</th>
<td>llvm-bugs@lists.llvm.org
</td>
</tr></table>
<p>
<div>
<pre>couldn't find entry for llvm-spirv that crashed; used the following:
clang -cc1 -O3 -emit-llvm -cl-std=CL1.2 -triple spir64-unknown-unknown
-finclude-default-header -cl-ext=cl_khr_fp64 -o x.bc x.cl
llvm-as-12 x.bc -o x.obj
llvm-spirv x.obj -o x.spr
!.../david/wd/dpl/openCL>make x.spr
clang -cc1 -O3 -emit-llvm -cl-std=CL1.2 -triple spir64-unknown-unknown
-finclude-default-header -cl-ext=cl_khr_fp64 -o x.bc x.cl
llvm-as-12 x.bc -o x.obj
llvm-spirv x.obj -o x.spr
PLEASE submit a bug report to <a href="https://bugs.llvm.org/">https://bugs.llvm.org/</a> and include the crash
backtrace.
Stack dump:
0. Program arguments: llvm-spirv x.obj -o x.spr
1. Running pass 'LLVMToSPIRV' on module 'x.obj'.
/lib/x86_64-linux-gnu/libLLVM-11.so.1(_ZN4llvm3sys15PrintStackTraceERNS_11raw_ostreamE+0x1f)[0x7f31b488802f]
/lib/x86_64-linux-gnu/libLLVM-11.so.1(_ZN4llvm3sys17RunSignalHandlersEv+0x50)[0x7f31b4886380]
/lib/x86_64-linux-gnu/libLLVM-11.so.1(+0xaed500)[0x7f31b4888500]
/lib/x86_64-linux-gnu/libc.so.6(+0x41040)[0x7f31b39d7040]
/lib/libLLVMSPIRVLib.so.11(_ZN5SPIRV11LLVMToSPIRV9transTypeEPN4llvm4TypeE+0xde)[0x7f31b8fefafe]
/lib/libLLVMSPIRVLib.so.11(_ZN5SPIRV11LLVMToSPIRV13transAsmINTELEPN4llvm9InlineAsmE+0xb0)[0x7f31b8ff2eb0]
/lib/libLLVMSPIRVLib.so.11(_ZN5SPIRV11LLVMToSPIRV27transValueWithoutDecorationEPN4llvm5ValueEPNS_15SPIRVBasicBlockEbNS0_13FuncTransModeE+0xc1e)[0x7f31b8feda2e]
/lib/libLLVMSPIRVLib.so.11(_ZN5SPIRV11LLVMToSPIRV10transValueEPN4llvm5ValueEPNS_15SPIRVBasicBlockEbNS0_13FuncTransModeE+0x13a)[0x7f31b8fef83a]
/lib/libLLVMSPIRVLib.so.11(_ZN5SPIRV11LLVMToSPIRV13transFunctionEPN4llvm8FunctionE+0xd9)[0x7f31b8ff51f9]
/lib/libLLVMSPIRVLib.so.11(_ZN5SPIRV11LLVMToSPIRV9translateEv+0x2d0)[0x7f31b8ff7b90]
/lib/libLLVMSPIRVLib.so.11(_ZN5SPIRV11LLVMToSPIRV11runOnModuleERN4llvm6ModuleE+0x70)[0x7f31b8ff7d50]
/lib/x86_64-linux-gnu/libLLVM-11.so.1(_ZN4llvm6legacy15PassManagerImpl3runERNS_6ModuleE+0x3e0)[0x7f31b4997700]
/lib/libLLVMSPIRVLib.so.11(_ZN4llvm10writeSpirvEPNS_6ModuleERKN5SPIRV14TranslatorOptsERSoRNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE+0xd2)[0x7f31b8ff7e32]
llvm-spirv(main+0x1bcb)[0x55b33a4d1f9b]
/lib/x86_64-linux-gnu/libc.so.6(__libc_start_main+0xd5)[0x7f31b39be565]
llvm-spirv(_start+0x2e)[0x55b33a4d229e]
make: *** [makefile:45: x.spr] Segmentation fault (core dumped)
make: *** Deleting file 'x.spr'
clang -v
Ubuntu clang version 12.0.0-3ubuntu1~21.04.2
Target: x86_64-pc-linux-gnu
Thread model: posix
InstalledDir: /usr/bin
Found candidate GCC installation: /usr/bin/../lib/gcc/x86_64-linux-gnu/10
Found candidate GCC installation: /usr/lib/gcc/x86_64-linux-gnu/10
Selected GCC installation: /usr/bin/../lib/gcc/x86_64-linux-gnu/10
Candidate multilib: .;@m64
Selected multilib: .;@m64
llvm-as-12 --version
LLVM (<a href="http://llvm.org/">http://llvm.org/</a>):
LLVM version 12.0.0
Optimized build.
Default target: x86_64-pc-linux-gnu
Host CPU: tigerlake
llvm-spirv --version
LLVM (<a href="http://llvm.org/">http://llvm.org/</a>):
LLVM version 11.0.1
Optimized build.
Default target: x86_64-pc-linux-gnu
Host CPU: tigerlake
x.cl - the watered down version of the real kernel; removing some lines make it
pass
__kernel void dpl_stencil_1p( __const unsigned long n, __global double *a,
__global double *b, __global double *x,
__local double *Save, __local double *Intrmd, __local double *Init
)
{
size_t globalID, localSz, globalSz, numGroups;
uint workDim;
globalID = get_global_id( 0 );
// The number of work-items per work-group = WI/WG in dimension 0
localSz = get_local_size ( 0 );
// The total number of work-items (WI) in dimension 0
globalSz = get_global_size( 0 );
// The total number of work-groups (WG) in dimension 0
numGroups = get_num_groups( 0 );
// The number of dimensions
workDim = get_work_dim();
size_t sz;
sz = localSz;
if ( globalID == 0 ) Save[sz-10] = 1.23;
barrier( CLK_LOCAL_MEM_FENCE ); // CLK_GLOBAL_MEM_FENCE
if ( globalID == 0 ) x[globalID] = localSz;
else if ( globalID == 1 ) x[globalID] = globalSz;
else if ( globalID == 2 ) x[globalID] = workDim;
else if ( globalID == 3 ) x[globalID] = numGroups;
else if ( globalID == 4 ) x[globalID] = Save[0];
else if ( globalID == 5 ) x[globalID] = Save[sz-10];
else x[globalID] = a[globalID] + b[globalID];
int ThreadsNum;
int nOfThreads, nOfThread;
nOfThreads = localSz;
nOfThread = globalID;
double Save_, Intrmd_;
// Stage1
unsigned long nn, iCnt, iRem, loopCnt, loopIters;
unsigned long i, j;
double xx, xx1;
nn = n - 1;
iCnt = nn / nOfThreads;
iRem = nn % nOfThreads;
loopCnt = iCnt;
loopIters = nOfThread*iCnt;
if ( nOfThread < iRem )
{
loopIters++;
}
// if ( nOfThread != ( nOfThreads - 1 ) )
{
/*
** we dont need the data in the last thread,
** but it seems tha on GPU it is easier just to execute the code
*/
Save_ = 0.;
Intrmd_ = 1.;
for ( j = loopIters + 1, i = j + loopCnt; j < i; j++ )
{
Save_ = xx*Save_ + xx1;
}
Save[nOfThread] = Save_;
Intrmd[nOfThread] = Intrmd_;
} // if ( nOfThread != ( nOfThreads - 1 ) )
}</pre>
</div>
</p>
<hr>
<span>You are receiving this mail because:</span>
<ul>
<li>You are on the CC list for the bug.</li>
</ul>
</body>
</html>