OpenCL scalare vs vettore

Question

ho semplice kernel:

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

Perché quando cambio galleggiante float4, kernel viene eseguito più lentamente oltre il 30%?

Tutti i tutorial dice, che utilizzando i tipi di vettore accelera calcolo ...

Sul lato host, memoria alocated per argomenti float4 è 16 byte allineati e global_work_size per clEnqueueNDRangeKernel è 4 volte più piccolo.

Il kernel gira su GPU AMD HD5770, AMD-APP-SDK-v2.6.

Info dispositivo per CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT torna 4.

EDIT:
global_work_size = 1024 * 1024 (e versioni successive)
local_work_size = 256
Tempo misurata utilizzando CL_PROFILING_COMMAND_START e CL_PROFILING_COMMAND_END.

Per le dimensioni globali global_work_size (8196 per float/2048 per float4), la versione vettoriale è più veloce, ma mi piacerebbe sapere perché?

Answer 1

Non so quali sono i tutorial a cui si fa riferimento, ma devono essere vecchi. Sia ATI che NVIDIA utilizzano architetture di gpu scalari per almeno metà decade. Oggigiorno usare i vettori nel codice è solo per comodità sintattica, non ha alcun vantaggio sulle prestazioni rispetto al semplice codice scalare. Si scopre che l'architettura scalare è migliore per le GPU rispetto a quella vettoriale: è meglio utilizzare le risorse hardware.

Answer 2

Non sono sicuro del motivo per cui i vettori sarebbero molto più lenti per voi, senza saperne di più sul gruppo di lavoro e le dimensioni globali. Mi aspetterei che almeno la stessa prestazione.

Se è adatto per il proprio kernel, si può iniziare con C con i valori in A? Ciò ridurrebbe l'accesso alla memoria del 33%. Forse questo si applica alla tua situazione?

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

Inoltre, hai stanco di leggere i valori in un vettore privato, quindi aggiungere? O forse entrambe le strategie.

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