cudaAddressModeBorder 是否适用于非标准化坐标？答案

【问题标题】：Does cudaAddressModeBorder work with non-normalized coordinates?cudaAddressModeBorder 是否适用于非标准化坐标？
【发布时间】：2013-09-20 17:06:14
【问题描述】：

我正在使用纹理对象来访问 PGM 图像像素。我的愿望是让纹理获取给定坐标中的像素值，如果我超出边界，则为 0。

这是我的纹理描述：

unsigned char *device_input=NULL;
size_t input_pitch;
checkCudaErrors(cudaMallocPitch(&device_input, &input_pitch, sizeof(unsigned char)*IMAGE_WIDTH, IMAGE_HEIGHT));
checkCudaErrors(cudaMemcpy2D(device_input, input_pitch, image, sizeof(unsigned char)*IMAGE_WIDTH, sizeof(unsigned char)*IMAGE_WIDTH, IMAGE_HEIGHT, cudaMemcpyHostToDevice));

cudaResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = cudaResourceTypePitch2D;
resDesc.res.pitch2D.devPtr = device_input; // 
resDesc.res.pitch2D.pitchInBytes =  input_pitch;
resDesc.res.pitch2D.width = IMAGE_WIDTH;
resDesc.res.pitch2D.height = IMAGE_HEIGHT;
resDesc.res.pitch2D.desc = cudaCreateChannelDesc<unsigned char>();

cudaTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
texDesc.readMode = cudaReadModeElementType;
texDesc.normalizedCoords=false;
texDesc.addressMode[0]=cudaAddressModeBorder;
texDesc.addressMode[1]=cudaAddressModeBorder;

cudaTextureObject_t tex;
cudaCreateTextureObject(&tex, &resDesc, &texDesc, NULL);

但是，在我的内核中：

tex2D<unsigned char>(tex_inputImage,-100,-100)

这显然在图像边界之外返回图像[0,0]处的值而不是值0。

同样适用：

tex2D<unsigned char>(tex_inputImage,IMAGE_WIDTH+1,IMAGE_HEIGHT+1)

返回 image[IMAGE_WIDTH,IMAGE_HEIGHT] 处的值，而不是 0。

请注意，通过使用标准化坐标，cudaAddressModeBorder 可以按预期工作，但我不想使用标准化坐标。根据 nvidia 的编程指南 (Here)，非归一化坐标支持 cudaAddressModeBorder。

我是不是做错了什么？

【问题讨论】：

经过进一步调查，这似乎是一个已知的驱动程序错误，希望 nvidia 尽快修复它：devtalk.nvidia.com/default/topic/549214/…
您运行的是什么驱动程序版本？你试过最新的吗？

标签： cuda textures nvidia

【解决方案1】：

这是我自己问题的答案：

该程序在驱动程序版本为 319.32 的机器上运行，显然驱动程序存在一个错误，在使用普通坐标 (More on the problem here - check the last couple of replies) 时，该驱动程序将 cudaAddressModeBorder 视为 cudaAddressModeClamp。

该错误已在版本 319.49 中得到修复，cudaAddressModeBorder 在规范化和非规范化坐标下均按预期工作。

【讨论】：