cuda - 绑定(bind)的 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;

    }
}

我要使用的纹理是具有两个 channel 的 2D float 纹理，因此我将其定义为:

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将纹理绑定(bind)到 cudaMalloc 分配的内存.在您使用的内核中 tex2D从纹理中读取值的函数。这就是它读取零的原因。

如果您使用 cudaBindTexture 将纹理绑定(bind)到线性内存，使用 tex1Dfetch 读取内核内部。
tex2D用于仅从绑定(bind)到 的纹理中读取音高线性内存 (由 cudaMallocPitch 分配)使用函数 cudaBindTexture2D ，或者那些绑定(bind)到 的纹理cudaArray 使用函数cudaBindTextureToArray
这是基本表，其余的可以从编程指南中阅读:

内存类型 ----------------- 使用 分配----------------- 使用 绑定(bind)----------------------- 通过 读入内核

线性内存.......cudaMalloc ...................... cudaBindTexture ...................... tex1Dfetch
音高线性内存...... cudaMallocPitch ............. cudaBindTexture2D ...................... tex2D
cudaArray....................... cudaMallocArray ............. cudaBindTextureToArray ............. tex1D或 tex2D
3D cudaArray...................... cudaMalloc3DArray ........ cudaBindTextureToArray ............. tex3D