ускорение при использовании float4, opencl

Question

У меня есть следующая функция ядра opencl, чтобы получить сумму столбца изображения.

__kernel void columnSum(__global float* src,__global float* dst,int srcCols, 
          int srcRows,int srcStep,int dstStep) 
{ 

    const int x = get_global_id(0); 
    srcStep >>= 2; 
    dstStep >>= 2; 

    if (x < srcCols) 
    { 
     int srcIdx = x ; 
     int dstIdx = x ; 

     float sum = 0; 

     for (int y = 0; y < srcRows; ++y) 
     { 
      sum += src[srcIdx]; 
      dst[dstIdx] = sum; 
      srcIdx += srcStep; 
      dstIdx += dstStep; 
     } 
    } 
} 

поручает, что каждый процесс нити столбца здесь, так что много нитей может получить column_sum каждой колонки параллельно.

Я также использую float4, чтобы переписать указанное выше ядро, чтобы каждый поток мог считывать 4 элемента в строке за один раз из исходного изображения, что показано ниже.

__kernel void columnSum(__global float* src,__global float* dst,int srcCols, 
          int srcRows,int srcStep,int dstStep) 
{ 

    const int x = get_global_id(0); 

    srcStep >>= 2; 
    dstStep >>= 2; 
    if (x < srcCols/4) 
    { 
     int srcIdx = x ; 
     int dstIdx = x ; 

     float4 sum = (float4)(0.0f, 0.0f, 0.0f, 0.0f); 

     for (int y = 0; y < srcRows; ++y) 
     { 
      float4 temp2; 
      temp2 = vload4(0, &src[4 * srcIdx]); 
      sum = sum + temp2; 

      vstore4(sum, 0, &dst[4 * dstIdx]); 

      srcIdx += (srcStep/4); 
      dstIdx += (dstStep/4); 
     } 
    } 
} 

В этом случае, теоретически, я считаю, что время, затрачиваемое на втором ядре, чтобы обработать изображение должно быть 1/4 от времени, потребляемого первой функцией ядра. Однако, независимо от того, насколько велико изображение, два ядра почти потребляют одно и то же время. Я не знаю почему. Можете ли вы, ребята, дать мне несколько идей? T

Answer 1

Я думаю, что на GPU float4 не является SIMD-операцией в OpenCL. Другими словами, если вы добавите два значения float4, сумма будет выполнена в четыре шага, а не сразу. Floatn действительно предназначен для процессора. На графическом процессоре floatn служит лишь удобным синтаксисом, по крайней мере, на картах Nvidia. Каждый поток на GPU действует так, как если бы он был скалярным процессором без SIMD. Но потоки в warp не являются независимыми, как в CPU. Правильный способ подумать о моделях GPGPU - это однопользовательские несколько потоков (SIMT). http://www.yosefk.com/blog/simd-simt-smt-parallelism-in-nvidia-gpus.html

Вы пытались запустить свой код на CPU? Я думаю, что код с float4 должен работать быстрее (потенциально в четыре раза быстрее), чем скалярный код на CPU. Также, если у вас есть процессор с AVX, тогда вы должны попробовать float8. Если код float4 работает быстрее на CPU, а float8 должен быть еще быстрее на процессоре с AVX.

Answer 2

какой архитектура вы используете?

Использование Float4 имеет более высокий параллелизм на уровне команд (и затем требуют 4 раза меньше потоков) так что теоретически должно быть быстрее (см http://www.cs.berkeley.edu/~volkov/volkov10-GTC.pdf)

Однако сделал я правильно вас понимаю ядро ​​вы делаете префиксов сумму (вам магазин частичная сумма на каждой итерации у)? Если это так, из-за магазинов узкое место находится в памяти.

Answer 3

Вы непосредственно читаете значения из массива «src» (глобальная память). Как правило, это в 400 раз медленнее частной памяти. Ваше узкое место определенно зависит от доступа к памяти, а не от операции «добавить».

Когда вы переходите из float в float4, операция вектора (добавление/умножение/...) более эффективна благодаря способности графического процессора работать с векторами. Однако чтение/запись в глобальную память остается неизменным. И поскольку это основное узкое место, вы не увидите никакого ускорения.

Если вы хотите скорректировать алгоритм, вы должны перейти в локальную память. Однако вам необходимо вручную разрешить управление памятью и правильный размер блока.

Answer 4

Форматы векторных данных OpenCL, такие как float4, лучше подходят для старых архитектур GPU, особенно для графических процессоров AMD. Современные графические процессоры не имеют регистров SIMD, доступных для отдельных рабочих элементов, в этом отношении они являются скалярными. CL_DEVICE_PREFERRED_VECTOR_WIDTH_* равен 1 для драйвера OpenCL на графических процессорах NVIDIA Kepler и интегрированной графике Intel HD. Поэтому добавление векторов float4 на современный графический процессор должно потребовать 4 операции. С другой стороны, драйвер OpenCL на процессоре Intel Core имеет CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT, равный 4, поэтому эти векторы могут быть добавлены за один шаг.

Answer 5

попытка определить атрибут __ __ ядра и видеть изменения в сроках выполнения , например попытаться определить:

__ ядро ​​аннулируются __ атрибут __ ((vec_type_hint (целое)))

или

__ __ ядро ​​недействительным атрибут __ ((vec_type_hint (int4)))

или некоторые floatN, как вы хотите

чтения мор e: https://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/functionQualifiers.html