6

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

__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”在浮点数的上下文中没有多大意义，但它至少应该给我一些非零值来读取。

我只能读零...

4

2 回答 2

8

您正在使用cudaBindTexture将纹理绑定到由cudaMalloc. 在内核中，您使用tex2D函数从纹理中读取值。这就是它读取零的原因。

如果使用 将纹理绑定到线性内存cudaBindTexture，则使用tex1Dfetch内核内部读取它。

tex2D用于仅从使用该函数绑定到间距线性内存（由 分配cudaMallocPitch）的那些纹理中读取，或者使用该函数从cudaBindTexture2D那些绑定到cudaArray的纹理中读取cudaBindTextureToArray

这是基本表，其余的可以从编程指南中阅读：

内存类型-----------------分配使用-----绑定使用---------- -------------读入内核

线性存储器cudaMalloc........................ cudaBindTexture………………………………………………………………………………tex1Dfetch

音高线性记忆cudaMallocPitch………………………………………………………………………………………………cudaBindTexture2Dtex2D

cudaArray cudaMallocArray....................... cudaBindTextureToArray........ .....tex1D 或tex2D

3D cudaArray cudaMalloc3DArray……………………………………………………………………………………………………………………………………………………cudaBindTextureToArraytex3D

于 2012-10-29T11:27:18.667 回答
3

此外，使用 tex1Dfetch 的访问基于整数索引。但是，其余的都是基于浮点的索引，您必须添加 +0.5 才能获得您想要的确切值。

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

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

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

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

于 2012-10-29T17:38:02.197 回答