OpenCL: использование vload против векторного указателя

Question

Есть ли какое-либо преимущество использования vload против прямого назначения указателя вектора? Что будет быстрее в мобильных графических процессорах с меньшей вычислительной мощностью и пропускной способностью?

Ex: ** образец нагрузки **

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