我正在尝试创建2D纹理对象,4x4 uint8_t。下面是代码:
__global__ void kernel(cudaTextureObject_t tex)
{
int x = threadIdx.x;
int y = threadIdx.y;
uint8_t val = tex2D<uint8_t>(tex, x, y);
printf("%d, ", val);
return;
}
int main(int argc, char **argv)
{
cudaTextureObject_t tex;
uint8_t dataIn[16] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15};
uint8_t* dataDev = 0;
cudaMalloc((void**)&dataDev, 16);
struct cudaResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = cudaResourceTypePitch2D;
resDesc.res.pitch2D.devPtr = dataDev;
resDesc.res.pitch2D.desc.x = 8;
resDesc.res.pitch2D.desc.y = 8;
resDesc.res.pitch2D.desc.f = cudaChannelFormatKindUnsigned;
resDesc.res.pitch2D.width = 4;
resDesc.res.pitch2D.height = 4;
resDesc.res.pitch2D.pitchInBytes = 4;
struct cudaTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
cudaCreateTextureObject(&tex, &resDesc, &texDesc, NULL);
cudaMemcpy(dataDev, &dataIn[0], 16, cudaMemcpyHostToDevice);
dim3 threads(4, 4);
kernel<<<1, threads>>>(tex);
cudaDeviceSynchronize();
return 0;
}
字符串
我预计结果会是这样的:
0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15,
型
即纹理对象的所有值(顺序无关紧要)。
但实际结果是:
0, 2, 4, 6, 0, 2, 4, 6, 0, 2, 4, 6, 0, 2, 4, 6,
型
我做错了什么?
1条答案
按热度按时间iih3973s1#
当你使用
pitch2D
变量进行纹理操作时,底层的分配应该是一个适当的 pitched 分配。我认为通常人们会使用cudaMallocPitch
来创建它。然而,要求是:cudaResourceDesc::res::pitch 2D::pitchInline以字节为单位指定两行之间的间距,并且必须与cudaDeviceProp::texturePitchAlignment对齐。
在我的GPU上,最后一个属性是32。我不知道你的GPU,但我打赌你的GPU的属性不是4。但是你在这里指定了4:
字符串
同样,我认为人们通常会使用使用
cudaMallocPitch
创建的间距分配来实现这一点。然而,如果行到行的维度(以字节为单位)可以被texturePitchAlignment
整除(在我的例子中是32),那么我确实可以传递普通的线性分配。我做的另一个改变是使用
cudaCreateChannelDesc<>()
而不是像你一样手动设置参数。这会创建一组不同的desc
参数,似乎也会影响结果。研究差异应该不难。当我调整你的代码来解决这些问题时,我得到的结果对我来说似乎是合理的:
型
正如评论中所问的,如果我要做类似的事情,但使用
cudaMallocPitch
和cudaMemcpy2D
,它可能看起来像这样:型
虽然我们这里有纹理对象,但很容易证明纹理引用也会发生类似的问题。你不能创建任意小的2D纹理引用,就像你不能创建任意小的2D纹理对象一样。我也不打算提供演示,因为它会在很大程度上重复上面的内容,人们不应该在新的开发工作中使用纹理引用-纹理对象是更好的方法。