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字节对齐,而clEnqueueNDRangeKernel的global_work_size是4倍.

内核运行在AMD HD5770 GPU,AMD-APP-SDK-v2.6上.

CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT的设备信息返回4.

编辑:
global_work_size = 1024*1024(及更高版本)
local_work_size = 256
使用CL_PROFILING_COMMAND_START和CL_PROFILING_COMMAND_END测量的时间.

对于较小的global_work_size(8196为浮点数/ 2048为浮动4),矢量化版本更快,但我想知道,为什么？

Answer 1

我不知道你提到的教程是什么,但它们必须陈旧.ATI和NVIDIA现在都使用标量gpu架构至少五年.现在在代码中使用向量只是为了方便语法,它没有普通标量代码的性能优势.事实证明,标量架构对于GPU而言比向量更好 - 它更好地利用了硬件资源.