OpenCL的标VS矢量

我有简单的内核：OpenCL的标VS矢量

__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]; 
}

为什么当我改变浮到个float4，内核运行慢30％以上？

所有教程说，使用矢量类型的加速计算...

在主机端，内存为alocated参数float4变量为16个字节对齐，并global_work_size为clEnqueueNDRangeKernel小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/float4为2048），向量化版本更快，但我想知道为什么？

来源

2012-01-19 ldanko

全球工作规模和工作组规模的价值是什么？你什么时候测量，以及如何？ –

全局工作大小= 1024 * 1024本地工作大小= 256，我使用CL_PROFILING_COMMAND_START和CL_PROFILING_COMMAND_END测量clEnquueNDRangeKernel的时间。对于较小的global_work_size（8196代表float/2048代表float4），向量化版本更快，但我想知道，为什么？ – ldanko

较小和较大的工作量之间的差异可能是由于您的缓存不断。所以2个问题： 1）如果你删除了const，它对于小而言还是比较快，对于较大的则是较慢？ 2）如果你去中间的某个地方，比如说浮动的65536和浮动的16384，那么会发生什么？ – user1111929

我不知道你提到的教程是什么，但它们必须是旧的。 ATI和NVIDIA现在都使用标量GPU架构至少五年。现在，在您的代码中使用矢量仅用于语法方便，与普通标量代码相比，它没有性能优势。事实证明，对于GPU，标量架构比矢量更好 - 它更好地利用硬件资源。

来源

2012-04-19 11:03:00 lucho

我不确定为什么矢量对于你来说要慢得多，而不了解工作组和全局大小。我希望它至少有相同的表现。

如果它适合你的内核，你可以从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; 
}

来源

2012-01-20 16:00:26 mfa

OpenCL的标VS矢量

回答

相关问题