OpenCL 标量与向量答案

【问题标题】：OpenCL scalar vs vectorOpenCL 标量与向量
【发布时间】：2012-02-14 13:12:21
【问题描述】：

我有简单的内核：

__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），矢量化版本更快，但我想知道，为什么？

【问题讨论】：

全局工作大小和工作组大小的值是多少？您测量的时间是什么？如何测量？
全局工作大小 = 1024*1024 本地工作大小 = 256，我使用 CL_PROFILING_COMMAND_START 和 CL_PROFILING_COMMAND_END 测量 clEnquueNDRangeKernel 的时间。对于较小的 global_work_size（float 为 8196 / float4 为 2048），矢量化版本更快，但我想知道，为什么？
较小和较大的工作大小之间的差异可能是由于您的恒定缓存。所以 2 个问题：1）如果你删除了 const，它是否仍然更快小而慢？ 2）如果你介于两者之间，比如浮点数为 65536，浮点数为 16384，那么会发生什么？

标签： opencl gpu gpgpu

【解决方案1】：

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

【讨论】：

【解决方案2】：

如果不了解工作组和全局大小的更多信息，我不确定为什么向量会慢得多。我希望它至少具有相同的性能。

如果它适合您的内核，您可以从具有 A 中的值的 C 开始吗？这将减少 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;
}

【讨论】：