c - 从 cudaBindTexture2D 读取

Question

我一直在尝试将 2D 数组存储在纹理内存中并通过 cudaBindTexture2D 从中读取 但返回的值是0，但我不确定这是否是cudaBindTexture2D和tex2D()的正确用法；

我编写了一个非常简单的代码来尝试一下:

#include <cuda.h>
#include <stdio.h>
#include <stdlib.h>
texture<uint, cudaTextureType2D, cudaReadModeElementType> tex;
__global__ 
void texture2DTest(int *x){
*x = tex2D(tex,0,0);

}

void initTable(int textureTable[][9]){
int i=0;
int j=0;

for(i=0; i<10; i++){
    for(j=0; j<9; j++){
        textureTable[i][j]=0;
    }
}

textureTable[0][0] = 12;

}

int main (int argc, char ** argv){

int textureTable[10][9];

int *d_x;
int x=2;

size_t pitch;

initTable(textureTable);

cudaMalloc(&d_x, sizeof(int)); 
cudaMemcpy(d_x, &x, sizeof(int), cudaMemcpyHostToDevice);

cudaMallocPitch( (void**)textureTable,&pitch, 9, 10);
cudaChannelFormatDesc desc = cudaCreateChannelDesc<uint>(); 
cudaBindTexture2D(NULL, tex, textureTable, desc, 9, 10, pitch) ;

texture2DTest<<<1,1>>>(d_x);

cudaThreadSynchronize();

cudaMemcpy(&x,d_x, sizeof(int), cudaMemcpyDeviceToHost);

printf(" \n %d \n",x);

cudaUnbindTexture(tex);

return 0;
}

谢谢。

Answer 1

所提供的代码存在不少问题。

• 使用 cudaMallocPitch 的设备内存分配完全损坏。您正在尝试将设备内存分配给已在主机上分配的二维数组。 尝试这样做将导致内存损坏和未定义的行为。设备内存分配需要一个单独的指针变量，分配后应将内存从主机复制到设备。

• cudaMallocPitch 的第三个参数期望内存宽度(以字节为单位)；不是元素。

• 纹理只能绑定(bind)到设备内存，因此 cudaBindTexture 期望设备内存指针作为输入。

解决上述所有问题后，您的最终 main 将如下所示:

int main (int argc, char ** argv)
{

int textureTable[10][9];

int *d_x;
int x = 2;

size_t pitch;

initTable(textureTable);

cudaMalloc(&d_x, sizeof(int)); 
cudaMemcpy(d_x, &x, sizeof(int), cudaMemcpyHostToDevice);

int* d_textureTable; //Device texture table

//Allocate pitch linear memory to device texture table
cudaMallocPitch((void**)&d_textureTable,&pitch, 9 * sizeof(int), 10);

//Use Memcpy2D as the pitch of host and device memory may be different
cudaMemcpy2D(d_textureTable, pitch, textureTable, 9 * sizeof(int), 9 *sizeof(int), 10, cudaMemcpyHostToDevice); 

cudaChannelFormatDesc desc = cudaCreateChannelDesc<uint>(); 
cudaBindTexture2D(NULL, tex, d_textureTable, desc, 9, 10, pitch) ;

texture2DTest<<<1,1>>>(d_x);

cudaThreadSynchronize();

cudaMemcpy(&x,d_x, sizeof(int), cudaMemcpyDeviceToHost);

printf(" \n %d \n",x);

cudaUnbindTexture(tex);
//Don't forget to free the allocated memory
cudaFree(d_textureTable);
cudaFree(d_x);
return 0;
}