OpenCL : Using vload Vs vector pointer

Question

使用 vload 与直接分配向量指针相比有什么优势吗？在计算能力和带宽较低的移动 GPU 中，哪个速度更快？

例如:vload 示例

__kernel vec_add(__global const float* a, __global const float* b, __global float* c){
    float4 a_sub;
    float4 b_sub;
    float4 c_sub;

    a_sub = vload4(0, &a[get_global_id(0)]);
    b_sub = vload4(0, &b[get_global_id(0)]);

    c_sub = a_sub + b_sub;
   
    vstore(c_sub, 0, &c[get_global_id(0)]);
}

矢量指针示例

__kernel vec_add(__global const float* a, __global const float* b, __global float* c){
    float4 a_sub;
    float4 b_sub;
    float4 c_sub;

    a_sub = ((global const float4*)a)[get_global_id(0)];
    b_sub = ((global const float4*)b)[get_global_id(0)];;

    c_sub = a_sub + b_sub;
   
    vstore(c_sub, 0, &c[get_global_id(0)]);
}

Answer 1

如评论中所述，它取决于目标硬件，这是加载数据的最快方式。换句话说，您应该对哪个是最好的进行基准测试，或者有什么不同。但是，我不记得通过更改语法来实现任何加速。

如果您需要使用 float* 缓冲区，那么尝试的第三个选项是像这样编写相同的负载:

    a_sub = (float4){
        a[get_global_id(0) * 4 + 0],
        a[get_global_id(0) * 4 + 1],
        a[get_global_id(0) * 4 + 2],
        a[get_global_id(0) * 4 + 3]
    };

但是，很多时候没有理由使用float* 缓冲区，您可以使用float4* 缓冲区。在那种情况下，编译器肯定可以知道负载将被对齐。此外，当我更改缓冲区类型时，我在移动平台上看到了显着的速度提升。因此，在您的情况下，内核签名如下所示:

__kernel vec_add(__global const float4* a, __global const float4* b, __global float4* c){

负载将是:

   a_sub = a[get_global_id(0)];