Come leggere una Texture CUDA per testare?

Question

Ok, finora, posso creare un array sul computer host (di tipo float), e copiarlo nella gpu, quindi riportarlo all'host come un altro array (per verificare se la copia ha avuto successo confrontando l'originale).

Quindi creo un array CUDA dall'array sulla GPU. Quindi lego quella matrice a una trama CUDA.

Ora voglio leggere quella trama e confrontarla con l'array originale (di nuovo per verificare che sia stata copiata correttamente). Ho visto alcuni esempi di codice che utilizzano la funzione readTexel() mostrata di seguito. Non sembra funzionare per me ... (fondamentalmente tutto funziona tranne la sezione nella funzione bindToTexture (float * deviceArray) che inizia alla riga readTexels (SIZE, testArrayDevice)).

Qualche suggerimento su un altro modo di farlo? O ci sono alcuni problemi evidenti che ho perso nel mio codice?

Grazie per l'aiuto ragazzi!

#include <stdio.h> 
#include <assert.h> 
#include <cuda.h> 

#define SIZE 20; 

//Create a channel description to use. 
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat); 

//Create a texture to use. 
texture<float, 2, cudaReadModeElementType> cudaTexture; 
//cudaTexture.filterMode = cudaFilterModeLinear; 
//cudaTexture.normalized = false; 

__global__ void readTexels(int amount, float *Array) 
{ 
    int index = blockIdx.x * blockDim.x + threadIdx.x; 

    if (index < amount) 
    { 
    float x = tex1D(cudaTexture, float(index)); 
    Array[index] = x; 
    } 
} 

float* copyToGPU(float* hostArray, int size) 
{ 
    //Create pointers, one for the array to be on the device, and one for bringing it back to the host for testing. 
    float* deviceArray; 
    float* testArray; 

    //Allocate some memory for the two arrays so they don't get overwritten. 
    testArray = (float *)malloc(sizeof(float)*size); 

    //Allocate some memory for the array to be put onto the GPU device. 
    cudaMalloc((void **)&deviceArray, sizeof(float)*size); 

    //Actually copy the array from hostArray to deviceArray. 
    cudaMemcpy(deviceArray, hostArray, sizeof(float)*size, cudaMemcpyHostToDevice); 

    //Copy the deviceArray back to testArray in host memory for testing. 
    cudaMemcpy(testArray, deviceArray, sizeof(float)*size, cudaMemcpyDeviceToHost); 

    //Make sure contents of testArray match the original contents in hostArray. 
    for (int i = 0; i < size; i++) 
    { 
    if (hostArray[i] != testArray[i]) 
    { 
     printf("Location [%d] does not match in hostArray and testArray.\n", i); 
    } 
    } 

    //Don't forget free these arrays after you're done! 
    free(testArray); 

    return deviceArray; //TODO: FREE THE DEVICE ARRAY VIA cudaFree(deviceArray); 
} 

cudaArray* bindToTexture(float* deviceArray) 
{ 
    //Create a CUDA array to translate deviceArray into. 
    cudaArray* cuArray; 

    //Allocate memory for the CUDA array. 
    cudaMallocArray(&cuArray, &cudaTexture.channelDesc, SIZE, 1); 

    //Copy the deviceArray into the CUDA array. 
    cudaMemcpyToArray(cuArray, 0, 0, deviceArray, sizeof(float)*SIZE, cudaMemcpyHostToDevice); 

    //Release the deviceArray 
    cudaFree(deviceArray); 

    //Bind the CUDA array to the texture. 
    cudaBindTextureToArray(cudaTexture, cuArray); 

    //Make a test array on the device and on the host to verify that the texture has been saved correctly. 
    float* testArrayDevice; 
    float* testArrayHost; 

    //Allocate memory for the two test arrays. 
    cudaMalloc((void **)&testArray, sizeof(float)*SIZE); 
    testArrayHost = (float *)malloc(sizeof(float)*SIZE); 

    //Read the texels of the texture to the test array in the device. 
    readTexels(SIZE, testArrayDevice); 

    //Copy the device test array to the host test array. 
    cudaMemcpy(testArrayHost, testArrayDevice, sizeof(float)*SIZE, cudaMemcpyDeviceToHost); 

    //Print contents of the array out. 
    for (int i = 0; i < SIZE; i++) 
    { 
    printf("%f\n", testArrayHost[i]); 
    } 

    //Free the memory for the test arrays. 
    free(testArrayHost); 
    cudaFree(testArrayDevice); 

    return cuArray; //TODO: UNBIND THE CUDA TEXTURE VIA cudaUnbindTexture(cudaTexture); 
    //TODO: FREE THE CUDA ARRAY VIA cudaFree(cuArray); 
} 


int main(void) 
{ 
    float* hostArray; 

    hostArray = (float *)malloc(sizeof(float)*SIZE); 

    for (int i = 0; i < SIZE; i++) 
    { 
    hostArray[i] = 10.f + i; 
    } 

    float* deviceAddy = copyToGPU(hostArray, SIZE); 

    free(hostArray); 

    return 0; 
} 

Answer 1

In breve:

------------- nella vostra main.cu -------- -------------------------------------------------- -----------------------------

-1. Definire la struttura come variabile globlal

     texture refTexture; // global variable ! 
     // meaning: address the texture with (x,y) (2D) and get an unsinged int 

Nella funzione principale:

-2. Utilizzare le matrici in combinazione con trama

 
    cudaArray* myArray; // declar. 
    // ask for memory 
    cudaMallocArray ( &myArray,
 
         &refTex.channelDesc, /* with this you don't need to fill a channel descriptor */ 
         width,
 
         height); 

-3. copy data from CPU to GPU (to the array)

 
cudaMemcpyToArray (arrayCudaEntrada, // destination: the array
 
         0, 0, // offsets 
         sourceData, // pointer uint* 
         widthheightsizeof(uint), // total amount of bytes to be copied 
         cudaMemcpyHostToDevice); 

-4. bind texture and array

 
    cudaBindTextureToArray(refTex,arrayCudaEntrada) 

-5. change some parameters in the texture

 
    refTextura_In.normalized = false; // don't automatically convert fetched data to [0,1[ 
     refTextura_In.addressMode[0] = cudaAddressModeClamp; // if my indexing is out of bounds: automatically use a valid indexing (0 if negative index, last if too great index) 
     refTextura_In.addressMode[1] = cudaAddressModeClamp;
 

---------- in the kernel --------------------------------------------------------

 
    // find out indexes (f,c) to process by this thread 
    uint f = (blockIdx.x * blockDim.x) + threadIdx.x; 
    uint c = (blockIdx.y * blockDim.y) + threadIdx.y;
 

// this is curious and necessary: indexes for reading from a texture 
    // are floats !. Even if you are certain to access (4,5) you have 
    // match the "center" this is (4.5, 5.5) 
    uint read = tex2D(refTex, c+0.5f, f+0.5f); // texRef is a global variable 
 

Now You processo di lettura e scrivere i risultati ad altre zone del dispositivo memoria globale, non la trama stessa!

Answer 2

readTexels() è un kernel (__global__) la funzione, vale a dire che gira sulla GPU. Pertanto è necessario utilizzare la sintassi corretta per avviare un kernel.

Dai un'occhiata alla Guida alla programmazione CUDA e ad alcuni campioni SDK, entrambi disponibili tramite NVIDIA CUDA site per vedere come avviare un kernel.

Suggerimento: Si finirà per qualcosa come readTexels<<<grid,block>>>(...)