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(float 为 8196/float4 为 2048)，矢量化版本更快，但我想知道，为什么？

Answer 1

我不知道你指的是什么教程，但它们一定是旧的。
ATI 和 NVIDIA 都使用标量 GPU 架构至少已有五年了。
现在在你的代码中使用向量只是为了语法上的方便，与普通的标量代码相比，它没有任何性能优势。
事实证明，标量架构比向量架构更适合 GPU——它更擅长利用硬件资源。