绑定的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”并没有太多意义，但至少它应该给我一些非零值来读取。

然而，我只能读到零……

- Philipp F

---

你在哪里以及如何显示那些输出为零的值？ - talonmies

Answer 1

您正在使用cudaBindTexture将纹理绑定到由cudaMalloc分配的内存中。在内核中，您正在使用tex2D函数从纹理中读取值。这就是为什么它会读取零值。

如果使用cudaBindTexture将纹理绑定到线性内存，则在内核中使用tex1Dfetch读取。 tex2D只用于读取那些绑定到pitch linear memory(通过cudaMallocPitch分配）或绑定到cudaArray(通过cudaBindTextureToArray)的纹理。

以下是基本表格，如需更多信息请参考编程指南： Memory Type----------------- Allocated Using-----------------Bound Using-----------------------Read In The Kernel By Linear Memory...................cudaMalloc........................cudaBindTexture.............................tex1Dfetch Pitch Linear Memory.........cudaMallocPitch.............cudaBindTexture2D........................tex2D cudaArray............................cudaMallocArray.............cudaBindTextureToArray.............tex1D or tex2D 3D cudaArray......................cudaMalloc3DArray........cudaBindTextureToArray.............tex3D

- sgarizvi

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

- Hong Zhou