OpenCL scalar vs vector

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

__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 GPU, AMD-APP-SDK-v2.6.

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

EDIT:
global_work_size = 1024*1024 (и больше)
локальный_размер_работы = 256
Время измеряется с помощью CL_PROFILING_COMMAND_START и CL_PROFILING_COMMAND_END.

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