¿Cómo leer una Textura CUDA para probar?

Question

Ok, hasta ahora, puedo crear una matriz en la computadora host (de tipo float), y copiarla a la gpu, luego traerla de regreso al host como otra matriz (para probar si la copia fue exitosa comparándola el original).

Luego creo una matriz CUDA de la matriz en la GPU. Luego uniré esa matriz a una textura CUDA.

Ahora quiero volver a leer esa textura y compararla con la matriz original (una vez más para comprobar que se copió correctamente). Vi un código de muestra que usa la función readTexel() que se muestra a continuación. No parece funcionar para mí ... (básicamente todo funciona, excepto la sección en la función bindToTexture (float * deviceArray) que comienza en la línea readTexels (SIZE, testArrayDevice)).

¿Alguna sugerencia de una manera diferente de hacer esto? ¿O hay algunos problemas obvios que perdí en mi código?

Gracias por la ayuda chicos!

#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

Brevemente:

------------- en su main.cu -------- -------------------------------------------------- -----------------------------

-1. Definir la textura como variable globlal

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

En la función principal:

-2. Utilizar matrices combinadas con la textura

 
    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 
 

Ahora procesar leer y escribir los resultados a otra zona del dispositivo de memoria global , no a la propia textura!

Answer 2

readTexels() es un kernel función (__global__), es decir que se ejecuta en la GPU. Por lo tanto, debe usar la sintaxis correcta para iniciar un kernel.

Eche un vistazo a la Guía de programación CUDA y algunas de las muestras de SDK, ambas disponibles a través de NVIDIA CUDA site para ver cómo iniciar un kernel.

Pista: va a terminar algo así como readTexels<<<grid,block>>>(...)