Jak odczytywać teksturę CUDA do testowania?

Question

Ok, do tej pory mogę utworzyć tablicę na komputerze hosta (typu float) i skopiować ją do procesora GPU, a następnie przekazać ją hostowi jako kolejną tablicę (aby sprawdzić, czy kopia się powiodła, porównując do oryginalny).

Następnie tworzę tablicę CUDA z tablicy na GPU. Następnie wiążę tę tablicę z teksturą CUDA.

Chciałbym teraz odczytać tę teksturę i porównać ją z oryginalną tablicą (ponownie, aby przetestować, czy została poprawnie skopiowana). Widziałem przykładowy kod wykorzystujący funkcję readTexel() pokazaną poniżej. Wydaje się, że nie działa to dla mnie ... (zasadniczo wszystko działa z wyjątkiem sekcji w funkcji bindToTexture (float * deviceArray), począwszy od linii readTexels (SIZE, testArrayDevice)).

Wszelkie sugestie w inny sposób to zrobić? Czy są jakieś oczywiste problemy, które przeoczyłem w moim kodzie?

Dzięki za pomoc!

#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

skrócie:

------------- w main.cu -------- -------------------------------------------------- -----------------------------

-1. Określić strukturę jako globlal zmiennej

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

W głównej funkcji:

-2. Użyj tablic w połączeniu z fakturą

 
    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 
 

teraz przetwarzać odczytu i zapisu wyników do drugiej strefy globalnej urządzenia pamięci, a nie do samej tekstury!

Answer 2

readTexels() jest jądro (__global__) funkcji, to znaczy biegnie w graficznego. Dlatego aby uruchomić jądro, musisz użyć poprawnej składni.

Zapoznaj się z Przewodnikiem programistycznym CUDA i niektórymi próbkami SDK, dostępnymi za pośrednictwem NVIDIA CUDA site, aby zobaczyć, jak uruchomić jądro.

Podpowiedź: To będzie w końcu coś readTexels<<<grid,block>>>(...)