comparaison Matlab vs CUDA corrélation et réduction sur un tableau 2D

Question

J'essaie de comparer la corrélation croisée en utilisant FFT vs en utilisant la méthode de fenêtrage.

Mon code Matlab est:

isize = 20; 
n = 7; 
for i = 1:n %%7x7 xcorr 
    for j = 1:n 
    xcout(i,j) = sum(sum(ffcorr1 .* ref(i:i+isize-1,j:j+isize-1))); %%ref is 676 element array and ffcorr1 is a 400 element array 
    end 
end 

similaire noyau CUDA:

__global__ void xc_corr(double* in_im, double* ref_im, int pix3, int isize, int n, double* out1, double* temp1, double* sum_temp1) 
{ 

    int p = blockIdx.x * blockDim.x + threadIdx.x; 
    int q = 0; 
    int i = 0; 
    int j = 0; 
    int summ = 0; 

    for(i = 0; i < n; ++i) 
    { 
     for(j = 0; j < n; ++j) 
     { 
      summ = 0; //force update 
      for(p = 0; p < pix1; ++p) 
      { 
       for(q = 0; q < pix1; ++q) 
       { 
        temp1[((i*n+j)*pix1*pix1)+p*pix1+q] = in_im[p*pix1+q] * ref_im[(p+i)*pix1+(q+j)];    
        sum_temp1[((i*n+j)*pix1*pix1)+p*pix1+q] += temp1[((i*n+j)*pix1*pix1)+p*pix1+q]; 
        out1[i*n+j] = sum_temp1[((i*n+j)*pix1*pix1)+p*pix1+q]; 
       } 
      }  
     } 
    } 

Je l'ai appelé dans mon noyau comme

int blocksize = 64; //multiple of 32 
int nblocks = (pix3+blocksize-1)/blocksize; //round to max pix3 = 400 
xc_corr <<< nblocks,blocksize >>> (ffcorr1, ref_d, pix3, isize, npix, xcout, xc_partial); 
cudaThreadSynchronize(); 

D'une certaine façon, quand je fais un diff sur le fichier de sortie, je vois que le noyau CUDA ne calcule que les 400 premiers éléments.

Quelle est la bonne façon d'écrire ce noyau ??

Aussi, quelle est la différence en déclarant i, j comme indiqué ci-dessous dans mon noyau ??

int i = blockIdx.x * blockDim.y + threadIdx.x * threadIdx.y; 
int j = blockIdx.y * blockDim.x + threadIdx.x * threadIdx.y; 

Answer 1

int blocksize = 64; //multiple of 32 
int nblocks = (pix3+blocksize-1)/blocksize; //round to max pix3 = 400 
xc_corr <<< nblocks,blocksize >>> (ffcorr1, ref_d, pix3, isize, npix, xcout, xc_partial); 

signifie que vous lancent 64 fils par bloc, et le nombre de threadblocks égal à 1 plus que nécessaire pour traiter des éléments de PIX3. Si pix3 est en effet 400, alors vous traitez 400 éléments car vous lancerez 7 threadblocks, chacun d'entre eux faisant 64 points, et 48 d'entre eux ne font rien.

Je ne sais pas trop quel est le problème ici.

En outre,

int i = blockIdx.x * blockDim.y + threadIdx.x * threadIdx.y; 
int j = blockIdx.y * blockDim.x + threadIdx.x * threadIdx.y; 

blocksize et nblocks sont effectivement convertie en DIM3 vecteurs, de sorte qu'ils aient un (x, y, z) de valeur. Si vous appelez un noyau avec < < 64,7 >>, qui va se traduire à

dim3 blocksize(64,1,1); 
dim3 nblocks(7,1,1); 
kernel<<blocksize,nblocks>>(); 

donc pour chaque appel du noyau, le blockIdx a 3 composants, l'ID de fil x, y et z, correspondant à la grille 3d des threads dans lesquels vous êtes. Dans votre cas, puisque vous avez seulement un composant x, blockIdx.y et threadIdx.y vont tous être 1 n'importe quoi. Donc essentiellement, ils sont inutiles. Honnêtement, vous semblez devoir relire les bases de CUDA à partir du mode d'emploi, car il vous manque beaucoup de bases. L'expliquer ici ne serait pas économique since it's all written down in a nice documentation you can get here. Et si vous voulez juste avoir une FFT plus rapide avec cuda, il y a un certain nombre de bibliothèques que vous pouvez simplement télécharger et installer sur la zone CUDA de Nvidia qui le fera pour vous si vous ne vous souciez pas de l'apprentissage de CUDA.

Bonne chance mate.

PS. vous n'avez pas besoin d'appeler cudaThreadSynchronize après chaque noyau;)