Как векторизовать трехмерную свертку 3x3?

Question

Я пытаюсь написать оптимизированную свертку 2D-изображений 3x3 для изображения 1280x720.

Для простоты к краевому условию приближается входной сигнал 1284 * 724.

Вот мой код ядра:

__kernel 
    __attribute__((vec_type_hint(float4))) 
    void conv2d3x3(              
     __global const float* restrict input,       
     __global float* restrict output, 
     __constant const float4* restrict hL, 
     /* 3x3 kernel, padded with 3 zeros on the right, used to calculate 
     "left" output samples*/ 
     __constant const float4* restrict hR 
     /*same 3x3 kernel, padded with 3 samples on the left*/)   
    {           
     int j = get_global_id(0)*2; //[0,639] 
     int i = get_global_id(1)*2; //[0,359]     

     /* load a 4x4 block, note stride is 1284 because input is padded by 4*/ 
     float4 data0=vload4(0,input+1284*(i+0)+j); 
     float4 data1=vload4(0,input+1284*(i+1)+j); 
     float4 data2=vload4(0,input+1284*(i+2)+j); 
     float4 data3=vload4(0,input+1284*(i+3)+j); 

     /* sum(data[0:2,0:2].* h)*/ 
     float prodTL=dot(data0,hL[0])+dot(data1,hL[1])+dot(data2,hL[2]); 
     /* sum(data[0:2,1:3].* h)*/ 
     float prodTR=dot(data0,hR[0])+dot(data1,hR[1])+dot(data2,hR[2]); 
     /* sum(data[1:3,0:2].* h)*/ 
     float prodBL=dot(data1,hL[0])+dot(data2,hL[1])+dot(data3,hL[2]); 
     /* sum(data[1:3,1:3].* h)*/ 
     float prodBR=dot(data1,hR[0])+dot(data2,hR[1])+dot(data3,hR[2]); 

     output[1280*(i+0)+j]=prodTL; 
     output[1280*(i+0)+j+1]=prodTR; 
     output[1280*(i+1)+j]=prodBL; 
     output[1280*(i+1)+j+1]=prodBR; 
    } 

Рациональный этой конструкции, чтобы загрузить 4x4 блок данных, сделать четыре 3x4 извилины и генерировать 4 выходных отсчетов.

Этот код имеет несколько очевидных проблем:

1) вектор нагрузки не совпадает с вектором границы.

2) складирование продукции не векторизованы

3) производительность низка: 3ms на Intel XEON 1245v3 с P4600 (с Beignet OpenCL implenentation) и 27ms на Freescale IMX6Q с GC2000 (с Freescale OpenCL libOpenCL).

Вопрос:

1) то, что я сделал неправильно, и почему это так медленно?

2) Какую производительность следует ожидать в процентах от сырых FLOPS? (p4600 способен к 20EU * 2PFU/EU * SIMD8 = 320FLOPS/цикл между 350 МГц и 1,2 ГГц, в то время как GC2000 способен не менее 14GFLOPS)

3) в целом, как векторизовать фиксированный размер несепарабельного размера 2D-свертка без создания избыточного трафика памяти и конфликта кэша?

Answer 1

Во-первых, мой неоптимизированная результат:

AMD FX 8150 @ 3,3 ГГц (32 FP элементы => 1 добавить + 1 мул = 64 FLOPS за один цикл):

3.71ms включая копирование времени между отдельные буферы opencl и массивы C#.

2.05ms не включая массивные копии.

Использование одномерного выполнения ядра ndrange вместо 2D. [0,640x360]

__kernel 
    __attribute__((vec_type_hint(float4))) 
    void bench(              
     __global const float* restrict input,       
     __global float* restrict output, 
     __constant const float4* restrict hL, 

     __constant const float4* restrict hR 
    )   
    { 
       int gli=get_global_id(0); 
       int j = (gli%640) * 2 ; 
       int i = (gli/640) * 2; 

       /* load a 4x4 block*/ 
       float4 data0 = vload4(0, input + 1280 * (i + 0) + j); 
       float4 data1 = vload4(0, input + 1280 * (i + 1) + j); 
       float4 data2 = vload4(0, input + 1280 * (i + 2) + j); 
       float4 data3 = vload4(0, input + 1280 * (i + 3) + j); 


       float prodTL = dot(data0, hL[0]) + dot(data1, hL[1]) + dot(data2, hL[2]); 

       float prodTR = dot(data0, hR[0]) + dot(data1, hR[1]) + dot(data2, hR[2]); 

       float prodBL = dot(data1, hL[0]) + dot(data2, hL[1]) + dot(data3, hL[2]); 

       float prodBR = dot(data1, hR[0]) + dot(data2, hR[1]) + dot(data3, hR[2]); 

       output[1280 * (i + 0) + j] = prodTL; 
       output[1280 * (i + 0) + j + 1] = prodTR; 
       output[1280 * (i + 1) + j] = prodBL; 
       output[1280 * (i + 1) + j + 1] = prodBR; 
      } 

Принимающая сторона (C# массивы):

 float[] inp = new float[1280*720*2]; 
     float[] outp = new float[1280*720*2]; 
     float[] hL = new float[1024]; 
     float[] hR = new float[1024]; 

С упреждающей в частные регистры (я могу только надеяться, драйверы, используя регистры центрального процессора):

2мс

Оптимизированная часть:

 float4 hl2=hL[2]; 
     float4 hl1=hL[1]; 
     float4 hl0=hL[0]; 

     float4 hr2=hR[2]; 
     float4 hr1=hR[1]; 
     float4 hr0=hR[0]; 

     float prodTL = dot(data0, hl0) + dot(data1, hl1) + dot(data2, hl2); 

     float prodTR = dot(data0, hr0) + dot(data1, hr1) + dot(data2, hr2); 

     float prodBL = dot(data1, hl0) + dot(data2, hl1) + dot(data3, hl2); 

     float prodBR = dot(data1, hr0) + dot(data2, hr1) + dot(data3, hr2); 

Теперь с увеличенным параллелизмом на точечных изделиях:

Сумма трех точек будет равна одной большой точке.

float16 prodhl =(float16)(hl0, hl1, hl2, (float4)(0.0f,0.0f,0.0f,0.0f));      
       float16 prodhl =(float16)(hr0, hr1, hr2, (float4)(0.0f,0.0f,0.0f,0.0f)); 
       float16 prodTdata =(float16)(data0,data1,data2,(float4)(0.0f,0.0f,0.0f,0.0f));  

       float16 prodBdata=(float16)(data1,data2,data3,(float4)(0.0f,0.0f,0.0f,0.0f));  

       float prodTL = dot(prodTdata, prodhl); 
       float prodTR = dot(prodTdata, prodhr); 

       float prodBL = dot(prodBdata, prodhl); 
       float prodBR = dot(prodBdata, prodhr); 

Исполнение без каких-либо копий массива:

0,5412 мс

meybe его только способность AVX процессора. Если нет, то должен быть некоторый параллелизм уровня инструкций.

В этой части (последняя часть float4 float16) имеется 1/4 потраченной впустую мощность, поэтому должен достигнуть 0,4 мс.

Примечание: размер группы нитей был 256. Я не пытался увеличиваться до 1024, поскольку он не подходит для всех устройств, таких как amd gpu.

Вы можете попробовать параллелизм уровня процесса процесса, чтобы увеличить пропускную способность и избить один контекст opencl (если вы уже это сделали).