Questo codice OpenCL può essere ottimizzato?

Question

Sto lavorando su un pezzo di codice OpenCL per una funzione di matrice specializzato: un Dx1 vettore v, due DxD matrici A e B e una costante c, ritorno 1xD vettore r dove r[i] = c * sum_over_j (v[j] * A[i][j] * B[i][j])

seguito è quello che ho così lontano, ma scorre incredibilmente lento. Una versione senza somma che restituisce una matrice DxD è circa dieci volte più veloce. Viene chiamato da PyOpenCL se questo fa alcuna differenza.

C'è qualcosa che non va? Potrebbe essere ottimizzato?

#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(1); 
     float sum = 0; 
     for(int k = 0; k < D; k++) 
     { 
      sum += vector[k] * matrix[(y*D) + k] 
      * matrix2[(y*D) + k ]; 
     } 
     result[y] = sum * factor; 
     } 

Cheers!

Answer 1

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