La texture CUDA lue lit zéro

Question

J'essaie de lire les valeurs d'une texture et de les réécrire dans la mémoire globale. Je suis sûr que la partie d'écriture fonctionne, beause je peux mettre des valeurs constantes dans le noyau et je peux les voir dans la sortie:

__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 Je veux utiliser est une texture flottante 2D avec deux canaux, donc je définissais comme:

texture<float2, 2, cudaReadModeElementType> texGrad; 

Et le code qui appelle le noyau initialise la texture avec des valeurs non nulles constantes:

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); 

Je sais, la mise en t Exposer des octets à tout "50" n'a pas beaucoup de sens dans le contexte des flottants, mais il devrait au moins me donner des valeurs non nulles à lire.

Je ne peux lire que des zéros ...

Answer 1

Vous utilisez cudaBindTexture pour lier la texture à la mémoire allouée par cudaMalloc. Dans le noyau, vous utilisez la fonction tex2D pour lire les valeurs de la texture. C'est pourquoi il lit des zéros.

Si vous liez la texture à la mémoire linéaire en utilisant cudaBindTexture, il est lu en utilisant tex1Dfetch dans le noyau.

tex2D est utilisé pour lire uniquement à partir de ces textures qui sont liés à la mémoire linéaire de pas (qui sont affectés par cudaMallocPitch) en utilisant la fonction cudaBindTexture2D, ou les textures qui sont liés à cudaArray en utilisant la fonction cudaBindTextureToArray

Voici la table de base, le reste, vous pouvez lire le guide de programmation:

type de mémoire ----------------- Al situé à l'aide de ----------------- Lié à l'aide de ----------------------- Read In Le noyau Par

Mémoire linéaire ................... cudaMalloc .................... .... cudaBindTexture ............................. tex1Dfetch

Mémoire linéaire de hauteur ........ cudaMallocPitch ............. cudaBindTexture2D ........................ tex2D

cudaArray .... ........................ cudaMallocArray ............. cudaBindTextureToArray ............. tex1D ou tex2D

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

Answer 2

Pour ajouter, l'accès à l'aide de tex1Dfetch est basé sur l'indexation d'entier. Cependant, le reste est indexé en fonction du nombre à virgule flottante, et vous devez ajouter +0.5 pour obtenir la valeur exacte que vous voulez.

Je suis curieux de savoir pourquoi créez-vous float et lier à une texture float2? Cela peut donner des résultats ambigus. float2 n'est pas une texture flottante 2D. Il peut effectivement être utilisé pour la représentation d'un nombre complexe. Je pense que ce tutoriel va vous aider à comprendre comment utiliser la mémoire de texture dans cuda. http://www.drdobbs.com/parallel/cuda-supercomputing-for-the-masses-part/218100902

Le noyau que vous avez montré ne profite pas beaucoup de l'utilisation de la texture. Cependant, si elle est utilisée correctement, en exploitant la localité, la mémoire de texture peut améliorer considérablement les performances. En outre, il est utile pour l'interpolation.