OpenCL escalar del vector vs

Question

tengo núcleo simple:

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

Por qué cuando cambio flotante a float4, kernel se ejecuta más de un 30% más lento?

Todos los tutoriales dice, que el uso de tipos de vectores acelera la computación ...

Por el lado de host, la memoria alocated de argumentos float4 es de 16 bytes alineados y global_work_size para clEnqueueNDRangeKernel es 4 veces más pequeños.

Kernel se ejecuta en AMD HD5770 GPU, AMD-APP-SDK-v2.6.

info Dispositivo para CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT devuelve 4.

EDIT:
global_work_size = 1024 * 1024 (y mayor)
local_work_size = 256
tiempo medido usando CL_PROFILING_COMMAND_START y CL_PROFILING_COMMAND_END.

Para global_work_size más pequeño (8196 para float/2048 para float4), la versión vectorizada es más rápida, pero me gustaría saber, ¿por qué?

Answer 1

No sé a qué tutoriales hace referencia, pero deben ser antiguos. Tanto ATI como NVIDIA usan arquitecturas escalares de gpu durante al menos media década. Hoy en día, el uso de vectores en su código es solo por conveniencia sintáctica, no tiene ningún beneficio de rendimiento sobre el código escalar simple. Resulta que la arquitectura escalar es mejor para las GPU que para las de tipo vectorial: es mejor utilizar los recursos de hardware.

Answer 2

No estoy seguro de por qué los vectores serían mucho más lentos para usted, sin saber más sobre el grupo de trabajo y el tamaño global. Yo esperaría que al menos el mismo rendimiento.

Si es adecuado para su kernel, ¿puede comenzar con C teniendo los valores en A? Esto reduciría el acceso a la memoria en un 33%. Tal vez esto se aplica a su situación?

__kernel vecadd(__global const float4 *B, 
       __global float4 *C) 
{ 
    int idx = get_global_id(0); 
    C[idx] += B[idx]; 
} 

Además, ¿has cansado de leer los valores en un vector privado, y luego agregar? O tal vez ambas estrategias.

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