Attachment 133653 Details for Bug 102345 – Debug log

Debug log

Debug log.txt (text/plain), 5.49 KB, created by Chris Cummins on 2017-08-21 19:57:29 UTC

(hide)

Description:

Filename:

MIME Type:

Creator: Chris Cummins

Created: 2017-08-21 19:57:29 UTC

Size: 5.49 KB

patch

obsolete

>[cldrive] Platform: Intel Gen OCL Driver
>Build options:
>
>CL kernel source:
>__kernel void self_test(__global int *buf) {  __local int tmp[3];  tmp[get_local_id(0)] = buf[get_local_id(0)];  barrier(CLK_LOCAL_MEM_FENCE);  buf[get_global_id(0)] = tmp[2 - get_local_id(0)] + buf[get_global_id(0)];}
>.decl_function self_test
>## 71 registers ##
>.decl.dword %0 local_id_0
>.decl.dword %1 local_id_1
>.decl.dword %2 local_id_2
>.decl.dword %3 group_id_0
>.decl.dword %4 group_id_1
>.decl.dword %5 group_id_2
>.decl.dword %6 num_groups_0
>.decl.dword %7 num_groups_1
>.decl.dword %8 num_groups_2
>.decl.dword %9 local_size_0
>.decl.dword %10 local_size_1
>.decl.dword %11 local_size_2
>.decl.dword %12 enqueued_local_size_0
>.decl.dword %13 enqueued_local_size_1
>.decl.dword %14 enqueued_local_size_2
>.decl.dword %15 global_size_0
>.decl.dword %16 global_size_1
>.decl.dword %17 global_size_2
>.decl.dword %18 global_offset_0
>.decl.dword %19 global_offset_1
>.decl.dword %20 global_offset_2
>.decl.dword %21 stack_pointer
>.decl.qword %22 stack_buffer
>.decl.word %23 block_ip
>.decl.dword %24 barrier_id
>.decl.dword %25 thread_number
>.decl.dword %26 work_dimension
>.decl.dword %27 zero
>.decl.dword %28 one
>.decl.word %29 retVal
>.decl.dword %30 dwblockip
>.decl.qword %31 profiling_buffer_pointer
>.decl.dword %32 profiling_timestamps0
>.decl.dword %33 profiling_timestamps1
>.decl.dword %34 profiling_timestamps2
>.decl.dword %35 profiling_timestamps3
>.decl.dword %36 profiling_timestamps4
>.decl.dword %37 threadid
>.decl.qword %38 constant_addrspace_start
>.decl.qword %39 stack_size
>.decl.qword %40 enqueue_buffer_pointer
>.decl.dword %41
>.decl.dword %42
>.decl.dword %43
>.decl.dword %44
>.decl.dword %45
>.decl.dword %46
>.decl.dword %47
>.decl.dword %48
>.decl.dword %49
>.decl.dword %50
>.decl.dword %51
>.decl.dword %52
>.decl.dword %53
>.decl.dword %54
>.decl.dword %55
>.decl.dword %56
>.decl.dword %57
>.decl.dword %58
>.decl.dword %59
>.decl.dword %60
>.decl.dword %61
>.decl.dword %62
>.decl.dword %63
>.decl.dword %64
>.decl.dword %65
>.decl.dword %66
>.decl.dword %67
>.decl.dword %68
>.decl.dword %69
>.decl.dword %70
>## 1 input registers ##
>decl_input.global %41 buf
>## 0 output register ##
>## 0 pushed register
>## 3 blocks ##
>LABEL $0
>LOADI.uint32 %42 4
>
>LABEL $1
>LOADI.uint32 %59 2
>SHL.int32 %43 %0 %59
>ADD.int32 %44 %41 %43
>SUB.uint32 %60 %44 %41
>LOAD.int32.global.aligned {%45} %60 bti:2
>SHL.int32 %46 %0 %59
>ADD.int32 %47 %42 %46
>LOADI.uint32 %62 0
>STORE.int32.local.aligned %47 {%45} bti:254
>SYNC.workgroup.local_read.local_write
>SUB.int32 %48 %59 %0
>SHL.int32 %49 %48 %59
>ADD.int32 %50 %42 %49
>LOAD.int32.local.aligned {%51} %50 bti:254
>MUL.int32 %52 %3 %12
>ADD.int32 %53 %52 %0
>ADD.int32 %54 %53 %18
>SHL.int32 %55 %54 %59
>ADD.int32 %56 %41 %55
>SUB.uint32 %69 %56 %41
>LOAD.int32.global.aligned {%57} %69 bti:2
>ADD.int32 %58 %57 %51
>SUB.uint32 %70 %56 %41
>STORE.int32.global.aligned %70 {%58} bti:2
>
>LABEL $2
>RET
>
>.end_function
>
>self_test's SELECTION IR begin:
>WARNING: not completed yet, welcome for the FIX!
>[0]       L0:
>[2]     CMP.le(16)          arf         :   %23<8,8,1>:UW   0x0:UW  
>[4](f0.1)   IF(16)              :   0:D 
>[6]     MOV(1)              %42<0>:UD   :   0x4:UD  
>[8]       L4:
>[10]        ENDIF(16)           :   0:D 
>
>[12]          L1:
>[14]        CMP.le(16)          arf         :   %23<8,8,1>:UW   0x1:UW  
>[16](f0.1)  MOV(16)             %23<1>:UW   :   0xffff:UW   
>[18]        CMP.eq(16)          arf         :   %23<8,8,1>:UW   0xffff:UW   
>[20](f0.1)  JMPI(1)             :   0:D 
>[22](f0.1)  MOV(16)             %23<1>:UW   :   0x1:UW  
>[24]        SHL(16)             %43<1>:D    :   %0<8,8,1>:D 2:D 
>[26]        ADD(16)             %44<1>:D    :   %41<0,1,0>:D    %43<8,8,1>:D    
>[28]        ADD(16)             %60<1>:UD   :   %44<8,8,1>:UD   -%41<0,1,0>:UD  
>[30]        UNTYPED_READ(16)    %45<1>:UD   :   %60<8,8,1>:UD   0x2:UD  
>[32]        SHL(16)             %46<1>:D    :   %0<8,8,1>:D 2:D 
>[34]        ADD(16)             %47<1>:D    :   %42<0,1,0>:D    %46<8,8,1>:D    
>[36]        UNTYPED_WRITE(16)   :   %47<8,8,1>:UD   %45<8,8,1>:UD   0xfe:UD 
>[38]        BARRIER(16)         %72<1>:F    :   %71<8,8,1>:UD   
>[40]        ADD(16)             %48<1>:D    :   -%0<8,8,1>:D    2:D 
>[42]        SHL(16)             %49<1>:D    :   %48<8,8,1>:D    2:D 
>[44]        ADD(16)             %50<1>:D    :   %42<0,1,0>:D    %49<8,8,1>:D    
>[46]        UNTYPED_READ(16)    %51<1>:UD   :   %50<8,8,1>:UD   0xfe:UD 
>[48]        MUL(1)              %52<0>:D    :   %3<0,1,0>:D %12<0,1,0>:UD   
>[50]        ADD(16)             %53<1>:D    :   %52<0,1,0>:D    %0<8,8,1>:D 
>[52]        ADD(16)             %54<1>:D    :   %53<8,8,1>:D    %18<0,1,0>:D    
>[54]        SHL(16)             %55<1>:D    :   %54<8,8,1>:D    2:D 
>[56]        ADD(16)             %56<1>:D    :   %41<0,1,0>:D    %55<8,8,1>:D    
>[58]        ADD(16)             %69<1>:UD   :   %56<8,8,1>:UD   -%41<0,1,0>:UD  
>[60]        UNTYPED_READ(16)    %57<1>:UD   :   %69<8,8,1>:UD   0x2:UD  
>[62]        ADD(16)             %58<1>:D    :   %57<8,8,1>:D    %51<8,8,1>:D    
>[64]        ADD(16)             %70<1>:UD   :   %56<8,8,1>:UD   -%41<0,1,0>:UD  
>[66]        UNTYPED_WRITE(16)   :   %70<8,8,1>:UD   %58<8,8,1>:UD   0x2:UD  
>
>[68](f0.1)    L2:
>[70](f0.1)  EOT(16)             :   
>
>self_test's SELECTION IR end.
>
>[cldrive] Device: Intel(R) HD Graphics Haswell GT2 Desktop
>[cldrive] OpenCL optimizations: off
>CL kernel source:
>kernel void A(global float4* a) {
>  a[0][get_global_id(0)] = 0;
>}
>
>ASSERTION FAILED: Not implemented
>  at file ../backend/src/llvm/llvm_scalarize.cpp, function bool gbe::Scalarize::scalarizeInsert(llvm::InsertElementInst*), line 838

Actions: View

Attachments on bug 102345: 133653

Use of freedesktop.org services, including Bugzilla, is subject to our Code of Conduct. How we collect and use information is described in our Privacy Policy.