Struttura CUDA rilegata zero zero

Question

Provo a leggere i valori da una trama e li scrivo nella memoria globale. Sono sicuro che la parte di scrittura funziona, beause posso mettere valori costanti nel kernel e posso vederli in uscita:

__global__ void 
bartureKernel(float* g_odata, int width, int height) 
{ 
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x; 
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y; 

    if(x < width && y < height) { 
      unsigned int idx = (y*width + x); 
      g_odata[idx] = tex2D(texGrad, (float)x, (float)y).x; 

    } 
} 

La texture che voglio usare è una texture galleggiante 2D con due canali, in modo da ho definito come:

texture<float2, 2, cudaReadModeElementType> texGrad; 

E il codice che chiama il kernel inizializza la trama con alcuni valori diversi da zero costanti:

float* d_data_grad = NULL; 

cudaMalloc((void**) &d_data_grad, gradientSize * sizeof(float)); 
CHECK_CUDA_ERROR; 

texGrad.addressMode[0] = cudaAddressModeClamp; 
texGrad.addressMode[1] = cudaAddressModeClamp; 
texGrad.filterMode = cudaFilterModeLinear; 
texGrad.normalized = false; 

cudaMemset(d_data_grad, 50, gradientSize * sizeof(float)); 
CHECK_CUDA_ERROR; 

cudaBindTexture(NULL, texGrad, d_data_grad, cudaCreateChannelDesc<float2>(), gradientSize * sizeof(float)); 

float* d_data_barture = NULL; 
cudaMalloc((void**) &d_data_barture, outputSize * sizeof(float)); 
CHECK_CUDA_ERROR; 

dim3 dimBlock(8, 8, 1); 
dim3 dimGrid(((width-1)/dimBlock.x)+1, ((height-1)/dimBlock.y)+1, 1); 

bartureKernel<<< dimGrid, dimBlock, 0 >>>(d_data_barture, width, height); 

lo so, l'impostazione del t esporre byte a tutti "50" non ha molto senso nel contesto dei float, ma dovrebbe almeno darmi dei valori diversi da zero da leggere.

riesco a leggere solo zeri anche se ...

Answer 1

Si utilizza cudaBindTexture di legare la vostra struttura per la memoria allocata da cudaMalloc. Nel kernel si utilizza la funzione per leggere i valori dalla trama. Questo è il motivo per cui sta leggendo degli zeri.

Se si associa la trama alla memoria lineare utilizzando cudaBindTexture, viene letto utilizzando tex1Dfetch all'interno del kernel.

tex2D viene utilizzato per leggere solo dai tessuti che sono tenuti a passo memoria lineare (assegnata dal cudaMallocPitch) utilizzando la funzione cudaBindTexture2D, o quelle trame che sono destinati a cudaArray utilizzando la funzione cudaBindTextureToArray

Ecco la tabella di base, il resto si può leggere dalla guida di programmazione:

Tipo di memoria ----------------- al trova utilizzando ----------------- associato utilizzando ----------------------- Lettura The Kernel By

Memoria lineare ................... cudaMalloc .................... .... cudaBindTexture ............................. tex1Dfetch

Passo memoria lineare ........ cudaMallocPitch ............. cudaBindTexture2D ........................ tex2D

cudaArray .... ........................ cudaMallocArray ............. ............. cudaBindTextureToArraytex1D o tex2D

3D cudaArray ............. ......... cudaMalloc3DArray ........ cudaBindTextureToArray ............. tex3D

Answer 2

di aggiungere, mediante accesso tex1Dfetch si riferiscono al indicizzazione intero. Tuttavia, il resto sono indicizzati in base a virgola mobile e devi aggiungere +0,5 per ottenere il valore esatto desiderato.

Sono curioso perché crei float e ti leghi a una trama float2? Può dare risultati ambigui. float2 non è una trama float 2D. Può essere effettivamente utilizzato per la rappresentazione di un numero complesso.

typedef struct {float x; float y;} float2; 

Penso che questo tutorial ti aiuterà a capire come utilizzare la memoria delle texture in cuda. http://www.drdobbs.com/parallel/cuda-supercomputing-for-the-masses-part/218100902

Il kernel che hai mostrato non beneficia molto dell'uso della trama. tuttavia, se utilizzato correttamente, sfruttando la località, la memoria delle texture può migliorare le prestazioni di un bel po '. Inoltre, è utile per l'interpolazione.