绑定 CUDA 纹理读取为零

Question

我尝试从纹理读取值并将它们写回全局内存。 我确信编写部分是有效的，因为我可以将常量值放入内核中，并且可以在输出中看到它们：

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

    }
}

我要使用的纹理是具有两个通道的2D浮动纹理，所以我将其定义为：

texture<float2, 2, cudaReadModeElementType> texGrad;

调用内核的代码使用一些常量非零值初始化纹理：

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

我知道，将纹理字节设置为全部“50”在浮点数的上下文中没有多大意义，但它至少应该给我一些非零值来读取。

虽然我只能读零...

Answer 1

您正在使用

cudaBindTexture

将纹理绑定到

cudaMalloc

分配的内存。在内核中，您使用

tex2D

函数从纹理中读取值。这就是为什么它读数为零。

如果使用

cudaBindTexture

将纹理绑定到线性内存，则会在内核内部使用

tex1Dfetch

读取纹理。

tex2D

用于仅读取使用函数 cudaMallocPitch 绑定到

pitch 线性内存

（由

cudaBindTexture2D

分配）的纹理，或使用函数绑定到 cudaArray 的纹理

cudaBindTextureToArray

这是基本的表格，剩下的你可以阅读编程指南：

读入内核

内存类型	分配使用	绑定使用	通过
线性记忆	cudaMalloc	cudaBindTexture	tex1Dfetch
音调线性记忆	cudaMallocPitch	cudaBindTexture2D	tex2D
cudaArray	cudaMallocArray	cudaBindTextureToArray	tex1D 或 tex2D
3D cudaArray	cudaMalloc3DArray	cudaBindTextureToArray	tex3D

Answer 2

补充一下，使用 tex1Dfetch 的访问是基于整数索引的。 然而，其余的都是基于浮点数索引的，你必须加上+0.5才能得到你想要的精确值。

我很好奇为什么要创建 float 并绑定到 float2 纹理？它可能会给出不明确的结果。 float2 不是 2D 浮动纹理。它实际上可以用于表示复数。

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

我认为本教程将帮助您了解如何在 cuda 中使用纹理内存。 http://www.drdobbs.com/parallel/cuda-supercomputing-for-the-masses-part/218100902

您显示的内核并没有从使用纹理中受益匪浅。然而，如果利用得当，通过利用局部性，纹理内存可以大大提高性能。此外，它对于插值很有用。