Ottimizzazione n. 1: rendere il vettore __local.
Il mio primo passaggio in questo ha ottenuto un miglioramento decente nelle prestazioni. Ho notato che ogni vettore [k] viene letto per un totale di volte D, quindi l'ho copiato su un __local. Questo è possibile solo perché D è abbastanza piccolo da permetterlo. Il kernel come lo avete sopra soffre di un terribile ALU: rapporto di recupero di 0,08 su entrambi i 5870 e 6970 gpus. Anche i gpus più lenti stanno ancora aspettando l'accesso alla memoria.
#define D 1000
__kernel void element_mult(
__global float *result,
__global const float *vector,
__global const float *matrix,
__global const float *matrix2,
const float factor)
{
int y = get_global_id(0);
float sum = 0;
__local float vectCopy[D];
int ls = get_local_size(0);
int lid = get_local_id(0);
for(int i=0;i<D;i+=ls){
vectCopy[i+lid] = vector[i+lid];
}
mem_fence(CLK_LOCAL_MEM_FENCE);
for(int k = 0; k < D; k++)
{
sum += vectCopy[k] * matrix[(y*D) + k] * matrix2[(y*D) + k ];
}
result[y] = sum * factor;
}
Con questa modifica, APP profiler sta mostrando un nuovo ALU: prendere il rapporto di 0,20 per la GPU 5870 e 6970. I tempi medi sono cambiati da 1513 -> 1034 e 1261 -> 861 sulle stesse carte. I gpus di fascia bassa sono ora vincolati da ALU invece di recupero. (rapporto maggiore di 4: 1)
Opimization # 2: calcola ogni risultato [y] utilizzando un intero gruppo di lavoro.
Si dovrebbe fare questo ID D erano molto più grandi (100k +). L'idea è di ottenere il miglior pattern di accesso alla memoria usando il gruppo di lavoro per calcolare un singolo elemento del risultato alla volta. Ho definito ls (dimensione locale) per essere 64 qui, perché funziona sul mio hardware, così come la maggior parte dei venditori. La dimensione del gruppo di lavoro utilizzata dal lato host dovrà essere 64 a meno che non si modifichi tale definizione. Ha bisogno di essere definito per creare la somma [ls] di memoria come __local, e non mi piace passare __local vars di dimensioni variabili nei miei kernel.
risultati: 5870 ALU: lettura = 0,59: 1, media = 708. 6970 ALU: lettura = 0,72, media = 590. Secondo APP profiler, questo è circa il doppio della tua lista originale.
#define D 1000
#define ls 64
__kernel void element_mult(
__global float *result,
__global const float *vector,
__global const float *matrix,
__global const float *matrix2,
const float factor)
{
__local float vectCopy[D];
int lid = get_local_id(0);
for(int i=0;i<D;i+=ls){
vectCopy[i+lid] = vector[i+lid];
}
mem_fence(CLK_LOCAL_MEM_FENCE);
int ng = get_num_groups(0);
int gid = get_group_id(0);
int y, k;
__local float sum[ls];
for(y = gid; y < D; y+=ng){
for(k = lid; k < D; k+=ls)
{
sum[lid] += vectCopy[k] * matrix[(y*D) + k] * matrix2[(y*D) + k ];
}
if(lid==0){
result[y] = sum[0];
for(k=1;k<ls;k++){
result[y] += sum[k];
}
result[y] *= factor;
}
mem_fence(CLK_LOCAL_MEM_FENCE);
}
}
EDIT: APP profiler = AMD APP KernelAnalyzer
Siete sicuri, non è vero, che il calcolo di y * D viene sollevato fuori dal giro k dal compilatore? E che la sotto-espressione comune (y * D) + k è calcolata una sola volta in ogni iterazione? –
Lo stai facendo su una GPU NVIDIA, per caso? – talonmies
@talonmies, non posso esserne sicuro. Il calcolo non viene eseguito localmente sul mio computer; fondamentalmente deve essere OpenCL. – trolle3000