[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