Page MenuHomePhabricator

SchedDataCalculator crashes for hello_world on Kepler+
ArchivedPublic

Authored by pmoreau on Jan 28 2016, 9:39 PM.
clCreateCommandQueue() succeeded.
clCreateProgramWithSource() suceeded.
clBuildProgram() suceeded.
clCreateKernel() suceeded.
clCreateBuffer() succeeded.
clSetKernelArg() succeeded.
translating program of type 5
Parsing SPIR-V generated by 6 (version 10)
Version 1.0
ID bound: 13
pi:6 ()
BB:0 (2 instructions) - df = { }
-> BB:2 (tree)
0: ld u64 %r0d c0[0x0] (0)
1: bra BB:2 (0)
BB:2 (7 instructions) - df = { }
2: mov u64 %r1d 0x0000000000000000 (0)
3: mov u64 %r2d 0x0000000000000000 (0)
4: mov u64 %r3d 0x0000000000000008 (0)
5: mad u64 %r2d %r3d %r1d %r2d (0)
6: add u64 %r4d %r0d %r2d (0)
7: mov u32 %r5 0x40490fd0 (0)
8: st u32 # g[%r4d+0x0] %r5 exit (0)
MAIN:-1 ()
BB:0 (0 instructions) - df = { }
pi:6 ()
BB:0 (2 instructions) - df = { }
-> BB:2 (tree)
0: ld u64 %r6d c0[0x0] (0)
1: bra BB:2 (0)
BB:2 (7 instructions) - idom = BB:0, df = { }
2: mov u64 %r7d 0x0000000000000000 (0)
3: mov u64 %r8d 0x0000000000000000 (0)
4: mov u64 %r9d 0x0000000000000008 (0)
5: mad u64 %r10d %r9d %r7d %r8d (0)
6: add u64 %r11d %r6d %r10d (0)
7: mov u32 %r12 0x40490fd0 (0)
8: st u32 # g[%r11d+0x0] %r12 exit (0)
MAIN:-1 ()
BB:0 (0 instructions) - df = { }
PEEPHOLE: DeadCodeElim
PEEPHOLE: CopyPropagation
PEEPHOLE: MergeSplits
PEEPHOLE: GlobalCSE
PEEPHOLE: LocalCSE
PEEPHOLE: AlgebraicOpt
PEEPHOLE: ModifierFolding
PEEPHOLE: ConstantFolding
PEEPHOLE: LoadPropagation
PEEPHOLE: IndirectPropagation
PEEPHOLE: MemoryOpt
PEEPHOLE: LocalCSE
PEEPHOLE: DeadCodeElim
pi:6 ()
BB:0 (1 instructions) - df = { }
-> BB:2 (tree)
0: bra BB:2 (0)
BB:2 (5 instructions) - idom = BB:0, df = { }
1: mov u64 %r9d 0x0000000000000008 (0)
2: mad u64 %r10d %r9d 0x0000000000000000 0x0000000000000000 (0)
3: add u64 %r11d %r10d c0[0x0] (0)
4: mov u32 %r12 0x40490fd0 (0)
5: st u32 # g[%r11d+0x0] %r12 exit (0)
MAIN:-1 ()
BB:0 (0 instructions) - df = { }
pi:6 ()
BB:0 (1 instructions) - df = { }
-> BB:2 (tree)
0: bra BB:2 (0)
BB:2 (5 instructions) - idom = BB:0, df = { }
1: mov u64 %r9d 0x0000000000000008 (0)
2: mad u64 %r10d %r9d 0x0000000000000000 0x0000000000000000 (0)
3: add u64 %r11d %r10d c0[0x0] (0)
4: mov u32 %r12 0x40490fd0 (0)
5: st u32 # g[%r11d+0x0] %r12 exit (0)
buildLiveSets(BB:0)
buildLiveSets(BB:2)
BB:2 live set of out blocks:
BitSet of size 13:
BB:2 live set after propagation:
BitSet of size 13:
BB:0 live set of out blocks:
BitSet of size 13:
BB:0 live set after propagation:
BitSet of size 13:
BuildIntervals(BB:0)
BuildIntervals(BB:2)
%12 <- live range [4(4), 5)
%11 <- live range [3(3), 5)
%10 <- live range [2(2), 3)
%9 <- live range [1(1), 2)
allocateRegisters to 6 instructions
printing live intervals ...
livei(%9): [1 2)
livei(%10): [2 3)
livei(%11): [3 5)
livei(%12): [4 5)
RIG_Node[%0]($[1]-1): 2 colors, weight inf, deg 0/62
X
RIG_Node[%1]($[1]-1): 2 colors, weight inf, deg 0/62
X
RIG_Node[%2]($[1]-1): 2 colors, weight inf, deg 0/62
X
RIG_Node[%3]($[1]-1): 2 colors, weight inf, deg 0/62
X
RIG_Node[%4]($[1]-1): 2 colors, weight inf, deg 0/62
X
RIG_Node[%5]($[1]-1): 1 colors, weight inf, deg 0/63
X
RIG_Node[%6]($[1]-1): 2 colors, weight inf, deg 0/62
X
RIG_Node[%7]($[1]-1): 2 colors, weight inf, deg 0/62
X
RIG_Node[%8]($[1]-1): 2 colors, weight inf, deg 0/62
X
RIG_Node[%9]($[1]-1): 2 colors, weight 1.000000, deg 0/62
X
RIG_Node[%10]($[1]-1): 2 colors, weight 1.000000, deg 0/62
X
RIG_Node[%11]($[1]-1): 2 colors, weight 0.500000, deg 2/62
X %12
RIG_Node[%12]($[1]-1): 1 colors, weight 1.000000, deg 2/63
X %11
edge: (%12, deg 2/63) >-< (%11, deg 2/62)
SIMPLIFY: pushed %12
edge: (%11, deg 0/62) >-< (%12, deg 2/63)
SIMPLIFY: pushed %11
SIMPLIFY: pushed %10
SIMPLIFY: pushed %9
SELECT phase
NODE[%9, 2 colors]
GPR:BitSet of size 63:
assigned reg 0
NODE[%10, 2 colors]
GPR:BitSet of size 63:
assigned reg 0
NODE[%11, 2 colors]
GPR:BitSet of size 63:
assigned reg 0
NODE[%12, 1 colors]
(%12) X (%11): $r0 + 2
GPR:BitSet of size 63:
0 1
assigned reg 2
RegAlloc done: 1
MAIN:-1 ()
BB:0 (0 instructions) - df = { }
buildLiveSets(BB:0)
BB:0 live set of out blocks:
BitSet of size 0:
BB:0 live set after propagation:
BitSet of size 0:
BuildIntervals(BB:0)
allocateRegisters to 0 instructions
printing live intervals ...
SELECT phase
RegAlloc done: 1
PEEPHOLE: FlatteningPass
terminate called after throwing an instance of 'std::out_of_range'
what(): vector::_M_range_check: __n (which is 2) >= this->size() (which is 2)
Program received signal SIGABRT, Aborted.
0x00007ffff76bb5f8 in raise () from /usr/lib/libc.so.6
(gdb) bt
#0 0x00007ffff76bb5f8 in raise () from /usr/lib/libc.so.6
#1 0x00007ffff76bca7a in abort () from /usr/lib/libc.so.6
#2 0x00007fffed40aadd in __gnu_cxx::__verbose_terminate_handler ()
at /build/gcc/src/gcc-5.3.0/libstdc++-v3/libsupc++/vterminate.cc:95
#3 0x00007fffed408936 in __cxxabiv1::__terminate (handler=<optimized out>)
at /build/gcc/src/gcc-5.3.0/libstdc++-v3/libsupc++/eh_terminate.cc:47
#4 0x00007fffed408981 in std::terminate () at /build/gcc/src/gcc-5.3.0/libstdc++-v3/libsupc++/eh_terminate.cc:57
#5 0x00007fffed408b98 in __cxxabiv1::__cxa_throw (obj=obj@entry=0xc93960,
tinfo=0x7fffed6efad8 <typeinfo for std::out_of_range>, dest=0x7fffed41dd40 <std::out_of_range::~out_of_range()>)
at /build/gcc/src/gcc-5.3.0/libstdc++-v3/libsupc++/eh_throw.cc:87
#6 0x00007fffed4317c7 in std::__throw_out_of_range_fmt (__fmt=<optimized out>)
at /build/gcc/src/gcc-5.3.0/libstdc++-v3/src/c++11/functexcept.cc:104
#7 0x00007fffea9addde in std::vector<nv50_ir::SchedDataCalculator::RegScores, std::allocator<nv50_ir::SchedDataCalculator::RegScores> >::_M_range_check (this=0x7fffffffb6f8, __n=2) at /usr/include/c++/5.3.0/bits/stl_vector.h:803
#8 0x00007fffea9ada39 in std::vector<nv50_ir::SchedDataCalculator::RegScores, std::allocator<nv50_ir::SchedDataCalculator::RegScores> >::at (this=0x7fffffffb6f8, __n=2) at /usr/include/c++/5.3.0/bits/stl_vector.h:824
#9 0x00007fffea9abc51 in nv50_ir::SchedDataCalculator::visit (this=0x7fffffffb6d0, bb=0xc7a260)
at ../../../../../mesa/src/gallium/drivers/nouveau/codegen/nv50_ir_emit_nvc0.cpp:2873
#10 0x00007fffea8dcc11 in nv50_ir::Pass::doRun (this=0x7fffffffb6d0, func=0xc8e390, ordered=true, skipPhi=true)
at ../../../../../mesa/src/gallium/drivers/nouveau/codegen/nv50_ir_bb.cpp:495
#11 0x00007fffea8dcad7 in nv50_ir::Pass::run (this=0x7fffffffb6d0, func=0xc8e390, ordered=true, skipPhi=true)
at ../../../../../mesa/src/gallium/drivers/nouveau/codegen/nv50_ir_bb.cpp:477
#12 0x00007fffea9ac94f in nv50_ir::calculateSchedDataNVC0 (targ=0xc7f660, func=0xc8e390)
at ../../../../../mesa/src/gallium/drivers/nouveau/codegen/nv50_ir_emit_nvc0.cpp:3133
#13 0x00007fffea9ac9cc in nv50_ir::CodeEmitterNVC0::prepareEmission (this=0xc80d60, func=0xc8e390)
at ../../../../../mesa/src/gallium/drivers/nouveau/codegen/nv50_ir_emit_nvc0.cpp:3142
#14 0x00007fffea96340f in nv50_ir::CodeEmitter::prepareEmission (this=0xc80d60, prog=0xc8d7b0)
at ../../../../../mesa/src/gallium/drivers/nouveau/codegen/nv50_ir_target.cpp:206
#15 0x00007fffea963ea8 in nv50_ir::Program::emitBinary (this=0xc8d7b0, info=0xc7e8f0)
at ../../../../../mesa/src/gallium/drivers/nouveau/codegen/nv50_ir_target.cpp:363
#16 0x00007fffea8d37c7 in nv50_ir_generate_code (info=0xc7e8f0)
at ../../../../../mesa/src/gallium/drivers/nouveau/codegen/nv50_ir.cpp:1212
#17 0x00007fffea9c29c5 in nvc0_program_translate (prog=prog@entry=0xc88930, chipset=<optimized out>,
debug=debug@entry=0x65d438) at ../../../../../mesa/src/gallium/drivers/nouveau/nvc0/nvc0_program.c:577
#18 0x00007fffea9bf931 in nvc0_compute_validate_program (nvc0=nvc0@entry=0x65d110)
at ../../../../../mesa/src/gallium/drivers/nouveau/nvc0/nvc0_compute.c:122
#19 0x00007fffea9e3f9d in nve4_compute_state_validate (nvc0=0x65d110)
at ../../../../../mesa/src/gallium/drivers/nouveau/nvc0/nve4_compute.c:305
#20 nve4_launch_grid (pipe=0x65d110, block_layout=0xc8bec0, grid_layout=0xc7a990, label=0, input=0xc70be0)
at ../../../../../mesa/src/gallium/drivers/nouveau/nvc0/nve4_compute.c:452
#21 0x00007ffff7aabda5 in clover::kernel::launch (this=0x6a49b0, q=...,
grid_offset=std::vector of length 1, capacity 1 = {...}, grid_size=std::vector of length 1, capacity 1 = {...},
block_size=std::vector of length 1, capacity 1 = {...})
at ../../../../../mesa/src/gallium/state_trackers/clover/core/kernel.cpp:81
#22 0x00007ffff7a63234 in <lambda(clover::event&)>::operator()(clover::event &) const (__closure=0xc93b30)
at ../../../../../mesa/src/gallium/state_trackers/clover/api/kernel.cpp:292
#23 0x00007ffff7a6484f in std::_Function_handler<void(clover::event&), clEnqueueNDRangeKernel(cl_command_queue, cl_kernel, cl_uint, const size_t*, const size_t*, const size_t*, cl_uint, _cl_event* const*, _cl_event**)::<lambda(clover::event&)> >::_M_invoke(const std::_Any_data &, clover::event &) (__functor=..., __args#0=...)
at /usr/include/c++/5.3.0/functional:1871
#24 0x00007ffff7aa537b in std::function<void (clover::event&)>::operator()(clover::event&) const (this=0xc70920,
__args#0=...) at /usr/include/c++/5.3.0/functional:2271
#25 0x00007ffff7aa33cd in clover::event::trigger (this=0xc708e0)
at ../../../../../mesa/src/gallium/state_trackers/clover/core/event.cpp:55
#26 0x00007ffff7aa3aca in clover::hard_event::hard_event(clover::command_queue&, unsigned int, clover::ref_vector<clover::event> const&, std::function<void (clover::event&)>) (this=0xc708e0, q=..., command=4592, deps=..., action=...)
at ../../../../../mesa/src/gallium/state_trackers/clover/core/event.cpp:126
#27 0x00007ffff7a63c40 in clover::create<clover::hard_event, clover::command_queue&, int, clover::ref_vector<clover::event>&, clEnqueueNDRangeKernel(cl_command_queue, cl_kernel, cl_uint, const size_t*, const size_t*, const size_t*, cl_uint, _cl_event* const*, _cl_event**)::<lambda(clover::event&)> >(void) ()
at ../../../../../mesa/src/gallium/state_trackers/clover/util/pointer.hpp:230
#28 0x00007ffff7a63414 in clEnqueueNDRangeKernel (d_q=0x65ce08, d_kern=0x6a49b8, dims=1, d_grid_offset=0x0,
d_grid_size=0x7fffffffe638, d_block_size=0x7fffffffe638, num_deps=0, d_deps=0x0, rd_ev=0x0)
at ../../../../../mesa/src/gallium/state_trackers/clover/api/kernel.cpp:293
#29 0x000000000040118d in main (argc=1, argv=0x7fffffffe778) at hello_world.c:171

Event Timeline

pmoreau created this paste.Jan 28 2016, 9:39 PM
pmoreau created this object with visibility "Public (No Login Required)".
pmoreau created this object with edit policy "Nouveau (Project)".
pmoreau changed the title of this paste from SchedDataCalculator crashes for hello_world on Fermi+ to SchedDataCalculator crashes for hello_world on Kepler+.
pmoreau archived this paste.Feb 6 2016, 3:28 PM