openCL riduzione e passaggio array 2d

Question

Ecco il ciclo che voglio convertire in openCL.

for(n=0; n < LargeNumber; ++n) {  
    for (n2=0; n2< SmallNumber; ++n2) { 
     A[n]+=B[n2][n]; 
    }               
    Re+=A[n];  
} 

Ed ecco quello che ho finora, anche se, so che non è corretto e manca alcune cose.

__kernel void openCL_Kernel(__global int *A, 
         __global int **B, 
         __global int *C, 
         __global _int64 Re, 
            int D) 
{ 

int i=get_global_id(0); 
int ii=get_global_id(1); 

A[i]+=B[ii][i]; 

//barrier(..); ? 

Re+=A[i]; 

} 

Sono un principiante assoluto per questo tipo di cose. Prima di tutto so che non posso passare un doppio puntatore globale a un kernel openCL. Se puoi, aspetta un paio di giorni prima di pubblicare la soluzione, voglio capirlo da solo, ma se puoi aiutarmi a indicarmi la direzione giusta sarei grato.

Answer 1

Per quanto riguarda il problema con il passaggio dei doublepoint: questo tipo di problema viene in genere risolto copiando l'intera matrice (o qualsiasi altra cosa si sta lavorando) in un blocco continuo di memoria e, se i blocchi hanno lunghezze diverse che passano un'altra matrice, che contiene gli offset per le singole righe (quindi il tuo accesso sarà simile a B[index[ii]+i]).

Ora per la riduzione fino a Re: poiché non hai menzionato il tipo di dispositivo su cui stai lavorando, assumerò la sua GPU. In tal caso eviterei di fare la riduzione dello stesso kernel, dal momento che sarà lento come il modo in cui lo hai postato (dovresti serializzare l'accesso a Re su migliaia di thread (e anche l'accesso a A[i]). Invece vorrei scrivere il kernel necessario, che somma tutti in A[i] e inserire la riduzione da A in Re in un altro kernel e farlo in diversi passaggi, cioè si usa un kernel di riduzione che opera sull'elemento n e li riduce a qualcosa come n/16 (o qualsiasi altro numero). Quindi si chiama iterativamente quel kernel finché non si è su un elemento, che è il risultato (sto rendendo questa descrizione volutamente vaga, dal momento che hai detto che volevi pensare a te stesso.)

Come nota a margine: ci si rende conto che il codice originale non ha esattamente un modello di accesso alla memoria piacevole? Supponendo che B sia relativamente grande (e molto più grande di A a causa della seconda dimensione) avendo il ciclo interno iterato sull'indice esterno si creeranno molte cachemisses. Questo è ancora peggio quando il porting per la GPU, che è molto sensibile circa l'accesso di memoria coerente

Quindi riordino come questo può aumentare a dismisura le prestazioni:

for (n2=0; n2< SmallNumber; ++n2) 
    for(n=0; n < LargeNumber; ++n)  
    A[n]+=B[n2][n]; 
for(n=0; n < LargeNumber; ++n)             
    Re+=A[n];  

Ciò è particolarmente vero se si dispone di un compilatore che è bravo a autovectorization, dal momento che potrebbe essere in grado di vettorializzare tale costrutto, ma è molto improbabile che sia in grado di farlo per il codice originale (e se non può dimostrare che A e B[n2] non possono riferirsi alla stessa memoria che può 't trasformare il codice originale in questo).