Discussion :

[Letinono]

Bonjour,

J'ai un problème avec un code CUDA. J'utilise des textures pour récupérer plus rapidement les données d'une image copiée en mémoire globale. Je fais des actions dessus, puis je détruit le cudaArray ainsi que le bind de la texture sur cette array et je lance a nouveau la fonction (code ci-après).

Mon soucis vient du fait que les données copiées dans le GPU sont ok, mais au 2eme lancement, il semble que pour la récupération de la 1ère ligne de la texture, les données sont encore celles de l'ancienne image. Je ne suis pas très très clair je pense, je vous poste le code:

Code :

Sélectionner tout - Visualiser dans une fenêtre à part

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
// Executed several times
void gpuRunPartialZNCC( tile *tile1, tile *tile2, measure *listMeasure, int nbParam, int numTile,
                       double *h_sumIm1, double *h_sumIm2, double *h_squareSumIm1, double *h_squareSumIm2, double *h_crossSum, interpolateMode interMode)
{
    // Tuiles et imagette extraite GPU
    float *d_extract;
    cudaArray *d_tile2;
 
    // Taille des zones mémoires
    size_t sizeTemp;
 
    // Descritption de Canal pour la texture
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
 
    // Dimension du GPU
    dim3    blockSizeRED( NBTREDUC );
    dim3    gridSizeRED( ((tile1->size.x * tile1->size.y) / NBTREDUC) / 2 );
    int        smemSizeRED = NBTREDUC*sizeof(float);
    dim3    blockSizeEXT(NBTINTER,NBTINTER);
    dim3    gridSizeEXT(tile1->size.x/NBTINTER, tile1->size.y/NBTINTER);
 
    // Allocation mémoire et copie des images sur le GPU
    sizeTemp = tile1->size.x * tile1->size.y * sizeof(float);
    cudaMemcpy( d_tile1, tile1->data, sizeTemp, cudaMemcpyHostToDevice );
 
    cudaMalloc( (void **)&d_extract, sizeTemp );
 
    sizeTemp = tile2->size.x * tile2->size.y * sizeof(float);
    cudaMallocArray( &d_tile2, &channelDesc, tile2->size.x, tile2->size.y );
    cudaMemcpyToArray( d_tile2, 0, 0, (void*)(tile2->data), sizeTemp, cudaMemcpyHostToDevice);
 
    // Bind de la texture
    texture<float, 2, cudaReadModeElementType> &myTexture = getTexture();
    if (interMode == BILINEAR)
        myTexture.filterMode = cudaFilterModeLinear;    // Interpolation intégrée
    else
        myTexture.filterMode = cudaFilterModePoint;
    myTexture.normalized = false;
    cudaBindTextureToArray( myTexture, d_tile2, channelDesc);
 
    for ( int measureCurent = 0; measureCurent < nbParam; measureCurent++) 
    {
        // Extraction de l'imagette
        extractImg(d_extract, tile1->size, tile1->offset, tile2->offset, listMeasure[measureCurent].param, gridSizeEXT, blockSizeEXT);
 
    }
 
    cudaUnbindTexture( myTexture );
 
    // Libération mémoire
    cudaFree( d_extract );
    cudaFreeArray( d_tile2 );
}

Et l'appel au kernel :

Code :

Sélectionner tout - Visualiser dans une fenêtre à part

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
__global__ void extraction(float *d_out, int dimOut, int im1offsetX, int im1offsetY, int im2offsetX, int im2offsetY, float ax, float ay, float bx, float by, float cx, float cy) {
 
    int ix        = IMUL(blockDim.x, blockIdx.x) + threadIdx.x + im1offsetX;
    int iy        = IMUL(blockDim.y, blockIdx.y) + threadIdx.y + im1offsetY;
 
    float xp    = ax*ix + bx*iy + cx  + .5 - im2offsetX;    // +.5 => Correction de l'emplacement du pixel sur GPU
    float yp    = ay*ix + by*iy + cy  + .5 - im2offsetY;
 
    d_out[iy*dimOut + ix] = tex2D(getDeviceTexture(), xp, yp);
}
 
void extractImg(float *d_extract, coord2D extractSize, coord2D tile1Offset, coord2D tile2Offset, polynomialParam parameters, dim3 gridSize, dim3 blockSize)
{
    float ax = parameters.ax;
    float ay = parameters.ay;
    float bx = parameters.bx;
    float by = parameters.by;
    float cx = parameters.cx;
    float cy = parameters.cy;
 
    printf("\nOffset image1 : %d, %d\t Offset image2 : %d, %d\n\n", tile1Offset.x, tile1Offset.y, tile2Offset.x, tile2Offset.y);
 
    extraction<<< gridSize, blockSize >>>(d_extract, extractSize.x, tile1Offset.x, tile1Offset.y, tile2Offset.x, tile2Offset.y, ax, ay, bx, by, cx, cy);
    cudaThreadSynchronize();
}

Edit:

Voila une image pour etre un peu plus parlant :



On voit bien sur l'image de droite que la 1ere ligne de pixel est la meme qui sur l'image de gauche, ce qui ne devrait pas être le cas...

J'attends vos questions pour tenter d'être plus clair.

Merci pour votre aide.

[Letinono]

Erreur stupide comme d'habitude, je n'ai pas enlevé les offset avant de stocker les données...

d_out[(iy-im1offsetY)*dimOut + (ix-im1offset2)] = tex2D(getDeviceTexture(), xp, yp);