Wie lese ich eine CUDA Textur zum Testen zurück?

Question

Ok, so weit, ich kann ein Array auf dem Host-Computer (vom Typ float) erstellen, und kopieren Sie es auf die GPU, dann bringt es zurück zum Host als ein anderes Array (um zu testen, ob die Kopie erfolgreich war das Original).

Ich erstelle dann ein CUDA-Array aus dem Array auf der GPU. Dann binde ich dieses Array an eine CUDA-Textur.

Ich möchte jetzt diese Textur zurück und vergleichen Sie mit dem ursprünglichen Array (erneut zu testen, dass es richtig kopiert). Ich habe einen Beispielcode gesehen, der die unten gezeigte Funktion readTexel() verwendet. Es scheint nicht für mich zu funktionieren ... (im Grunde funktioniert alles außer für den Abschnitt in der Funktion bindToTexture (float * deviceArray) ab der Zeile readTexels (SIZE, testArrayDevice)).

Irgendwelche Vorschläge für eine andere Möglichkeit, dies zu tun? Oder gibt es einige offensichtliche Probleme, die ich in meinem Code verpasst habe?

Danke für die Hilfe Jungs!

#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

kurz am Ende:

------------- in Ihrem main.cu -------- -------------------------------------------------- -----------------------------

-1. Definiert die Textur als globlal variable

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

in der Hauptfunktion:

-2. Verwenden Sie kombinierten Arrays mit Textur

 
    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 
 

Jetzt Prozess, den Sie lesen und die Ergebnisse auf andere Zone der Vorrichtung globalen Speicher schreiben, nicht auf die Textur selbst!

Answer 2

readTexels() ist ein Kernel (__global__) Funktion, das heißt, es läuft auf der GPU. Daher müssen Sie zum Starten eines Kernels die korrekte Syntax verwenden.

Werfen Sie einen Blick in den CUDA Programmierleitfaden und einige der SDK-Beispiele, die beide über die NVIDIA CUDA site verfügbar sind, um zu sehen, wie ein Kernel gestartet wird.

Hinweis: Es wird so etwas wie readTexels<<<grid,block>>>(...)