Page MenuHomePhabricator

SPIR-V vec-load fails on Tesla
ArchivedPublic

Authored by pmoreau on Jan 26 2016, 10:12 PM.
./vec-load: /home/pmoreau/userspaces/llvm_spirv/lib/libOpenCL.so.1: no version information available (required by ./vec-load)
There are 1 platforms.
There are 1 GPU devices.
clCreateProgramWithSource() suceeded.
clBuildProgram() suceeded.
clCreateKernel() suceeded.
clSetKernelArg() succeeded.
clSetKernelArg() succeeded.
translating program of type 5
Parsing SPIR-V generated by 6 (version 10)
Version 1.0
ID bound: 26
vec_load:8 ()
BB:0 (3 instructions) - df = { }
-> BB:2 (tree)
0: ld u32 %r0 s[0x10] (0)
1: ld u32 %r1 s[0x14] (0)
2: bra BB:2 (0)
BB:2 (36 instructions) - df = { }
3: mov u32 %r2 0x00000000 (0)
4: mov u32 %r3 0x00000000 (0)
5: mov u32 %r4 0x00000004 (0)
6: mad u32 %r3 %r4 %r2 %r3 (0)
7: add u32 %r5 %r1 %r3 (0)
8: ld u32 %r6 g[0x0] (0)
9: ld u32 %r7 g[0x4] (0)
10: ld u32 %r8 g[0x8] (0)
11: ld u32 %r9 g[0xc] (0)
12: mov u32 %r10 %r6 (0)
13: mov u32 %r11 0x00000000 (0)
14: mov u32 %r12 0x00000004 (0)
15: mad u32 %r11 %r12 %r2 %r11 (0)
16: add u32 %r13 %r0 %r11 (0)
17: st u32 # g[%r13+0x0] %r10 (0)
18: mov u32 %r14 %r7 (0)
19: mov u32 %r15 0x00000001 (0)
20: mov u32 %r16 0x00000000 (0)
21: mov u32 %r17 0x00000004 (0)
22: mad u32 %r16 %r17 %r15 %r16 (0)
23: add u32 %r18 %r0 %r16 (0)
24: st u32 # g[%r18+0x4] %r14 (0)
25: mov u32 %r19 %r8 (0)
26: mov u32 %r20 0x00000002 (0)
27: mov u32 %r21 0x00000000 (0)
28: mov u32 %r22 0x00000004 (0)
29: mad u32 %r21 %r22 %r20 %r21 (0)
30: add u32 %r23 %r0 %r21 (0)
31: st u32 # g[%r23+0x8] %r19 (0)
32: mov u32 %r24 %r9 (0)
33: mov u32 %r25 0x00000003 (0)
34: mov u32 %r26 0x00000000 (0)
35: mov u32 %r27 0x00000004 (0)
36: mad u32 %r26 %r27 %r25 %r26 (0)
37: add u32 %r28 %r0 %r26 (0)
38: st u32 # g[%r28+0xc] %r24 exit (0)
MAIN:-1 ()
BB:0 (0 instructions) - df = { }
vec_load:8 (in $r0)
BB:0 (4 instructions) - df = { }
-> BB:2 (tree)
0: mov u32 %r32 $r0 (0)
1: ld u32 %r33 s[0x10] (0)
2: ld u32 %r34 s[0x14] (0)
3: bra BB:2 (0)
BB:2 (36 instructions) - idom = BB:0, df = { }
4: mov u32 %r35 0x00000000 (0)
5: mov u32 %r36 0x00000000 (0)
6: mov u32 %r37 0x00000004 (0)
7: mad u32 %r38 %r37 %r35 %r36 (0)
8: add u32 %r39 %r34 %r38 (0)
9: ld u32 %r40 g[0x0] (0)
10: ld u32 %r41 g[0x4] (0)
11: ld u32 %r42 g[0x8] (0)
12: ld u32 %r43 g[0xc] (0)
13: mov u32 %r44 %r40 (0)
14: mov u32 %r45 0x00000000 (0)
15: mov u32 %r46 0x00000004 (0)
16: mad u32 %r47 %r46 %r35 %r45 (0)
17: add u32 %r48 %r33 %r47 (0)
18: st u32 # g[%r48+0x0] %r44 (0)
19: mov u32 %r49 %r41 (0)
20: mov u32 %r50 0x00000001 (0)
21: mov u32 %r51 0x00000000 (0)
22: mov u32 %r52 0x00000004 (0)
23: mad u32 %r53 %r52 %r50 %r51 (0)
24: add u32 %r54 %r33 %r53 (0)
25: st u32 # g[%r54+0x4] %r49 (0)
26: mov u32 %r55 %r42 (0)
27: mov u32 %r56 0x00000002 (0)
28: mov u32 %r57 0x00000000 (0)
29: mov u32 %r58 0x00000004 (0)
30: mad u32 %r59 %r58 %r56 %r57 (0)
31: add u32 %r60 %r33 %r59 (0)
32: st u32 # g[%r60+0x8] %r55 (0)
33: mov u32 %r61 %r43 (0)
34: mov u32 %r62 0x00000003 (0)
35: mov u32 %r63 0x00000000 (0)
36: mov u32 %r64 0x00000004 (0)
37: mad u32 %r65 %r64 %r62 %r63 (0)
38: add u32 %r66 %r33 %r65 (0)
39: st u32 # g[%r66+0xc] %r61 exit (0)
MAIN:-1 (in $r0)
BB:0 (1 instructions) - df = { }
40: mov u32 %r3 $r0 (0)
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
vec_load:8 (in $r0)
BB:0 (2 instructions) - df = { }
-> BB:2 (tree)
0: ld u32 %r33 s[0x10] (0)
1: bra BB:2 (0)
BB:2 (12 instructions) - idom = BB:0, df = { }
2: ld b128 { %r40 %r41 %r42 %r43 } g[0x0] (0)
3: mov u32 %r46 0x00000004 (0)
4: mov u32 %r48 %r33 (0)
5: st u32 # g[%r48+0x0] %r40 (0)
6: add u32 %r54 s[0x10] %r46 (0)
7: st u32 # g[%r54+0x4] %r41 (0)
8: mov u32 %r59 0x00000008 (0)
9: add u32 %r60 s[0x10] %r59 (0)
10: st u32 # g[%r60+0x8] %r42 (0)
11: mov u32 %r65 0x0000000c (0)
12: add u32 %r66 s[0x10] %r65 (0)
13: st u32 # g[%r66+0xc] %r43 exit (0)
MAIN:-1 (in $r0)
BB:0 (0 instructions) - df = { }
vec_load:8 (in $r0)
BB:0 (3 instructions) - df = { }
-> BB:2 (tree)
0: nop - # (0)
1: ld u32 %r33 s[0x10] (0)
2: bra BB:2 (0)
BB:2 (13 instructions) - idom = BB:0, df = { }
3: ld b128 %r67q g[0x0] (0)
4: split b128 { %r40 %r41 %r42 %r43 } %r67q (0)
5: mov u32 %r46 0x00000004 (0)
6: mov u32 %r48 %r33 (0)
7: st u32 # g[%r48+0x0] %r40 (0)
8: add u32 %r54 s[0x10] %r46 (0)
9: st u32 # g[%r54+0x4] %r41 (0)
10: mov u32 %r59 0x00000008 (0)
11: add u32 %r60 s[0x10] %r59 (0)
12: st u32 # g[%r60+0x8] %r42 (0)
13: mov u32 %r65 0x0000000c (0)
14: add u32 %r66 s[0x10] %r65 (0)
15: st u32 # g[%r66+0xc] %r43 exit (0)
buildLiveSets(BB:0)
buildLiveSets(BB:2)
BB:2 live set of out blocks:
BitSet of size 68:
BB:2 live set after propagation:
BitSet of size 68:
33
BB:0 live set of out blocks:
BitSet of size 68:
33
BB:0 live set after propagation:
BitSet of size 68:
BuildIntervals(BB:0)
%33 <- live range [1(1), 3)
BuildIntervals(BB:2)
%43 <- live range [4(4), 15)
%66 <- live range [14(14), 15)
%65 <- live range [13(13), 14)
%42 <- live range [4(4), 12)
%60 <- live range [11(11), 12)
%59 <- live range [10(10), 11)
%41 <- live range [4(4), 9)
%54 <- live range [8(8), 9)
%46 <- live range [5(5), 8)
%40 <- live range [4(4), 7)
%48 <- live range [6(6), 7)
%33 <- live range [3(1), 6)
%67 <- live range [3(3), 4)
allocateRegisters to 16 instructions
joining %67($-1) <- %40
joining %67($-1) <- %41
joining %67($-1) <- %42
joining %67($-1) <- %43
makeCompound(split = 1): split b128 { %r40 %r41 %r42 %r43 } %r67q (0)
compound: %67:ff <- %40:03
compound: %67:ff <- %41:0c
compound: %67:ff <- %42:30
compound: %67:ff <- %43:c0
joining %48($-1) <- %33
printing live intervals ...
livei(%31): [0 1)
livei(%33): [1 6)
livei(%40): [4 7)
livei(%41): [4 9)
livei(%42): [4 12)
livei(%43): [4 15)
livei(%46): [5 8)
livei(%48): [6 7)
livei(%54): [8 9)
livei(%59): [10 11)
livei(%60): [11 12)
livei(%65): [13 14)
livei(%66): [14 15)
livei(%67): [3 4)
RIG_Node[%0]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%1]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%2]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%3]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%4]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%5]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%6]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%7]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%8]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%9]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%10]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%11]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%12]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%13]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%14]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%15]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%16]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%17]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%18]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%19]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%20]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%21]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%22]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%23]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%24]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%25]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%26]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%27]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%28]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%29]($[1]0): 2 colors, weight inf, deg 0/255
X
RIG_Node[%30]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%31]($[1]0): 2 colors, weight inf, deg 0/255
X
RIG_Node[%32]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%33]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%34]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%35]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%36]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%37]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%38]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%39]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%40]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%41]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%42]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%43]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%44]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%45]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%46]($[1]-1): 2 colors, weight 0.333333, deg 10/255
X %67 %48
RIG_Node[%47]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%48]($[1]-1): 2 colors, weight 0.666667, deg 10/255
X %46 %67
RIG_Node[%49]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%50]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%51]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%52]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%53]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%54]($[1]-1): 2 colors, weight 1.000000, deg 8/255
X %67
RIG_Node[%55]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%56]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%57]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%58]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%59]($[1]-1): 2 colors, weight 1.000000, deg 8/255
X %67
RIG_Node[%60]($[1]-1): 2 colors, weight 1.000000, deg 8/255
X %67
RIG_Node[%61]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%62]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%63]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%64]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%65]($[1]-1): 2 colors, weight 1.000000, deg 8/255
X %67
RIG_Node[%66]($[1]-1): 2 colors, weight 1.000000, deg 8/255
X %67
RIG_Node[%67]($[1]-1): 8 colors, weight 2.083333, deg 56/249
X %48 %66 %65 %60 %59 %54 %46
edge: (%66, deg 8/255) >-< (%67, deg 56/249)
SIMPLIFY: pushed %66
edge: (%65, deg 8/255) >-< (%67, deg 48/249)
SIMPLIFY: pushed %65
edge: (%60, deg 8/255) >-< (%67, deg 40/249)
SIMPLIFY: pushed %60
edge: (%59, deg 8/255) >-< (%67, deg 32/249)
SIMPLIFY: pushed %59
edge: (%54, deg 8/255) >-< (%67, deg 24/249)
SIMPLIFY: pushed %54
edge: (%48, deg 10/255) >-< (%46, deg 10/255)
edge: (%48, deg 10/255) >-< (%67, deg 16/249)
SIMPLIFY: pushed %48
edge: (%46, deg 8/255) >-< (%67, deg 8/249)
edge: (%46, deg 8/255) >-< (%48, deg 10/255)
SIMPLIFY: pushed %46
edge: (%67, deg 0/249) >-< (%48, deg 8/255)
edge: (%67, deg 0/249) >-< (%66, deg 8/255)
edge: (%67, deg 0/249) >-< (%65, deg 8/255)
edge: (%67, deg 0/249) >-< (%60, deg 8/255)
edge: (%67, deg 0/249) >-< (%59, deg 8/255)
edge: (%67, deg 0/249) >-< (%54, deg 8/255)
edge: (%67, deg 0/249) >-< (%46, deg 8/255)
SIMPLIFY: pushed %67
SELECT phase
NODE[%67, 8 colors]
GPR:BitSet of size 256:
assigned reg 0
NODE[%46, 2 colors]
(%46) X (%67): no overlap
(%46)ff X (%40)03 & ff: $r0.03
(%46)ff X (%41)0c & ff: $r0.0c
(%46)ff X (%42)30 & ff: $r0.30
(%46)ff X (%43)c0 & ff: $r0.c0
GPR:BitSet of size 256:
0 1 2 3 4 5 6 7
assigned reg 8
NODE[%48, 2 colors]
(%48) X (%46): $r8 + 2
(%48) X (%67): no overlap
(%48)ff X (%40)03 & ff: $r0.03
(%48)ff X (%41)0c & ff: $r0.0c
(%48)ff X (%42)30 & ff: $r0.30
(%48)ff X (%43)c0 & ff: $r0.c0
(%33)ff X (%67)ff & ff: $r0.ff
(%33)ff X (%40)03 & ff: $r0.03
(%33)ff X (%41)0c & ff: $r0.0c
(%33)ff X (%42)30 & ff: $r0.30
(%33)ff X (%43)c0 & ff: $r0.c0
GPR:BitSet of size 256:
0 1 2 3 4 5 6 7 8 9
assigned reg 10
NODE[%54, 2 colors]
(%54) X (%67): no overlap
(%54) X (%40): no overlap
(%54)ff X (%41)0c & ff: $r0.0c
(%54)ff X (%42)30 & ff: $r0.30
(%54)ff X (%43)c0 & ff: $r0.c0
GPR:BitSet of size 256:
2 3 4 5 6 7
assigned reg 0
NODE[%59, 2 colors]
(%59) X (%67): no overlap
(%59) X (%40): no overlap
(%59) X (%41): no overlap
(%59)ff X (%42)30 & ff: $r0.30
(%59)ff X (%43)c0 & ff: $r0.c0
GPR:BitSet of size 256:
4 5 6 7
assigned reg 0
NODE[%60, 2 colors]
(%60) X (%67): no overlap
(%60) X (%40): no overlap
(%60) X (%41): no overlap
(%60)ff X (%42)30 & ff: $r0.30
(%60)ff X (%43)c0 & ff: $r0.c0
GPR:BitSet of size 256:
4 5 6 7
assigned reg 0
NODE[%65, 2 colors]
(%65) X (%67): no overlap
(%65) X (%40): no overlap
(%65) X (%41): no overlap
(%65) X (%42): no overlap
(%65)ff X (%43)c0 & ff: $r0.c0
GPR:BitSet of size 256:
6 7
assigned reg 0
NODE[%66, 2 colors]
(%66) X (%67): no overlap
(%66) X (%40): no overlap
(%66) X (%41): no overlap
(%66) X (%42): no overlap
(%66)ff X (%43)c0 & ff: $r0.c0
GPR:BitSet of size 256:
6 7
assigned reg 0
RegAlloc done: 1
MAIN:-1 (in $r0)
BB:0 (1 instructions) - df = { }
0: nop - # (0)
buildLiveSets(BB:0)
BB:0 live set of out blocks:
BitSet of size 4:
BB:0 live set after propagation:
BitSet of size 4:
BuildIntervals(BB:0)
allocateRegisters to 1 instructions
printing live intervals ...
livei(%2): [0 1)
RIG_Node[%0]($[1]0): 2 colors, weight inf, deg 0/255
X
RIG_Node[%1]($[1]-1): 2 colors, weight inf, deg 0/255
X
RIG_Node[%2]($[1]0): 2 colors, weight inf, deg 0/255
X
RIG_Node[%3]($[1]-1): 2 colors, weight inf, deg 0/255
X
SELECT phase
RegAlloc done: 1
PEEPHOLE: FlatteningPass
PEEPHOLE: NV50PostRaConstantFolding
vec_load:8 (in $r0)
BB:0 (1 instructions) - df = { }
-> BB:2 (tree)
0: ld u32 $r5 s[0x10] (8)
BB:2 (11 instructions) - idom = BB:0, df = { }
1: ld b128 $r0q g[0x0] (8)
2: mov u32 $r4 0x00000004 (8)
3: st u32 # g[$r5+0x0] $r0 (8)
4: add u32 $r0 s[0x10] $r4 (8)
5: st u32 # g[$r0+0x4] $r1 (8)
6: mov u32 $r0 0x00000008 (8)
7: add u32 $r0 s[0x10] $r0 (8)
8: st u32 # g[$r0+0x8] $r2 (8)
9: mov u32 $r0 0x0000000c (8)
10: add u32 $r0 s[0x10] $r0 (8)
11: st u32 # g[$r0+0xc] $r3 exit (8)
MAIN:-1 (in $r0)
BB:0 (0 instructions) - df = { }
EMIT: ld u32 $r5 s[0x10] (8)
EMIT: ld b128 $r0q g[0x0] (8)
zsh: segmentation fault (core dumped) NV50_PROG_DEBUG=255 ./vec-load
Program received signal SIGSEGV, Segmentation fault.
0x00007fffe84156e8 in nv50_ir::ValueRef::get (this=0x0)
at ../../../../../src/gallium/drivers/nouveau/codegen/nv50_ir.h:504
504 inline Value *get() const { return value; }
#0 0x00007fffe84156e8 in nv50_ir::ValueRef::get (this=0x0) at ../../../../../src/gallium/drivers/nouveau/codegen/nv50_ir.h:504
#1 0x00007fffe842de07 in nv50_ir::CodeEmitterNV50::srcId (this=0xc3e080, src=..., pos=9) at ../../../../../src/gallium/drivers/nouveau/codegen/nv50_ir_emit_nv50.cpp:139
#2 0x00007fffe8427f40 in nv50_ir::CodeEmitterNV50::emitLOAD (this=0xc3e080, i=0x66f1a0) at ../../../../../src/gallium/drivers/nouveau/codegen/nv50_ir_emit_nv50.cpp:682
#3 0x00007fffe842c7e1 in nv50_ir::CodeEmitterNV50::emitInstruction (this=0xc3e080, insn=0x66f1a0) at ../../../../../src/gallium/drivers/nouveau/codegen/nv50_ir_emit_nv50.cpp:1840
#4 0x00007fffe84a57fc in nv50_ir::Program::emitBinary (this=0xc3e410, info=0xc53340) at ../../../../../src/gallium/drivers/nouveau/codegen/nv50_ir_target.cpp:385
#5 0x00007fffe8414a63 in nv50_ir_generate_code (info=0xc53340) at ../../../../../src/gallium/drivers/nouveau/codegen/nv50_ir.cpp:1212
#6 0x00007fffe84aeb31 in nv50_program_translate (prog=0xc4ef40, chipset=172, debug=0x63ea88) at ../../../../../src/gallium/drivers/nouveau/nv50/nv50_program.c:374
#7 0x00007fffe84a9746 in nv50_compute_validate_program (nv50=0x63e760) at ../../../../../src/gallium/drivers/nouveau/nv50/nv50_compute.c:167
#8 0x00007fffe84a9898 in nv50_compute_state_validate (nv50=0x63e760) at ../../../../../src/gallium/drivers/nouveau/nv50/nv50_compute.c:204
#9 0x00007fffe84a9c49 in nv50_launch_grid (pipe=0x63e760, block_layout=0xc3e330, grid_layout=0xc4f320, label=0, input=0xc318c0) at ../../../../../src/gallium/drivers/nouveau/nv50/nv50_compute.c:283
#10 0x00007ffff7a9a251 in clover::kernel::launch (this=0xc3f0b0, 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 ../../../../../src/gallium/state_trackers/clover/core/kernel.cpp:81
#11 0x00007ffff7a516e0 in <lambda(clover::event&)>::operator()(clover::event &) const (__closure=0xc31400) at ../../../../../src/gallium/state_trackers/clover/api/kernel.cpp:292
#12 0x00007ffff7a52cfb 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
#13 0x00007ffff7a93827 in std::function<void (clover::event&)>::operator()(clover::event&) const (this=0xc30e50, __args#0=...) at /usr/include/c++/5.3.0/functional:2271
#14 0x00007ffff7a91879 in clover::event::trigger (this=0xc30e10) at ../../../../../src/gallium/state_trackers/clover/core/event.cpp:55
#15 0x00007ffff7a91f76 in clover::hard_event::hard_event(clover::command_queue&, unsigned int, clover::ref_vector<clover::event> const&, std::function<void (clover::event&)>) (this=0xc30e10, q=..., command=4592, deps=..., action=...)
at ../../../../../src/gallium/state_trackers/clover/core/event.cpp:126
#16 0x00007ffff7a520ec 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 ../../../../../src/gallium/state_trackers/clover/util/pointer.hpp:230
#17 0x00007ffff7a518c0 in clEnqueueNDRangeKernel (d_q=0x63e458, d_kern=0xc3f0b8, dims=1, d_grid_offset=0x0, d_grid_size=0x7fffffffdf20, d_block_size=0x7fffffffdf20, num_deps=0, d_deps=0x0, rd_ev=0x0)
at ../../../../../src/gallium/state_trackers/clover/api/kernel.cpp:293
#18 0x00000000004012b4 in clSimpleEnqueueNDRangeKernel (command_queue=0x63e458, kernel=0xc3f0b8, work_dim=1, global_work_size=0x7fffffffdf20, local_work_size=0x7fffffffdf20) at cl_simple.c:111
#19 0x0000000000400dac in main (argc=1, argv=0x7fffffffe088) at vec_load.c:45

Event Timeline

pmoreau created this paste.Jan 26 2016, 10:12 PM
pmoreau created this object with visibility "Public (No Login Required)".
pmoreau created this object with edit policy "Nouveau (Project)".
pmoreau archived this paste.Feb 23 2016, 1:27 PM