OpenCL скаляр против вектора

Question

У меня простое ядро:

__kernel vecadd(__global const float *A,
                __global const float *B,
                __global float *C)
{
    int idx = get_global_id(0);
    C[idx] = A[idx] + B[idx];
}

Почему, когда я изменяю float на float4 , ядро ​​работает более чем на 30% медленнее?

Во всех уроках говорится, что использование векторных типов ускоряет вычисления ...

На стороне хоста память, выделенная для аргументов float4, выровнена на 16 байт, а global_work_size для clEnqueueNDRangeKernel в 4 раза меньше.

Ядро работает на графическом процессоре AMD HD5770, AMD-APP-SDK-v2.6.

Информация об устройстве для CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT возвращает 4.

EDIT:
global_work_size = 1024 * 1024 (и выше)
local_work_size = 256
Время измеряется с помощью CL_PROFILING_COMMAND_START и CL_PROFILING_COMMAND_END.

Для меньшего global_work_size (8196 для float / 2048 для float4) векторизованная версия быстрее, но я хотел бы знать, почему?

Answer 1

Я не знаю, на какие уроки вы ссылаетесь, но они, должно быть, старые.И ATI, и NVIDIA используют архитектуры скалярных графических процессоров уже не менее полувека.В настоящее время использование векторов в вашем коде только для синтаксического удобства, оно не дает никакого преимущества в производительности по сравнению с простым скалярным кодом.Оказывается, скалярная архитектура лучше для графических процессоров, чем для векторных - она ​​лучше использует аппаратные ресурсы.

Answer 2

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

Если это подходит для вашего ядра, можете ли вы начать с того, что C имеет значения в A? Это сократит доступ к памяти на 33%. Может быть, это относится к вашей ситуации?

__kernel vecadd(__global const float4 *B,
                __global float4 *C)
{
    int idx = get_global_id(0);
    C[idx] += B[idx];
}

Кроме того, вы устали читать значения в приватном векторе, а потом добавлять? Или, может быть, обе стратегии.

__kernel vecadd(__global const float4 *A,
                __global const float4 *B,
                __global float4 *C)
{
    int idx = get_global_id(0);
    float4 tmp = A[idx] + B[idx];
    C[idx] = tmp;
}