纹理声明

system · 2013 年1 月 23 日 07:29

纹理不能再主机端声明吗？为啥会出错呢，在设备端声明好像没有问题

system · 2013 年1 月 23 日 08:27

是的，纹理引用需要在设备端。

5.0新引入了纹理对象这个概念，可以在host端创立，但需要3.x的卡。

1.x和3.x推荐使用纹理。以及，如果是3.5, 可以使用__ldg()这个来通过纹理缓存读取数据。

system · 2013 年1 月 23 日 08:37

我纹理绑定不是要在主机端就绑定吗？主机端用纹理出现未声明的标识符的错误

system · 2013 年1 月 23 日 09:15

是这样的，

你需要在1个.cu里面进行纹理引用的声明的，虽然你只是简单的写texture<…> your_tex_ref;
并且和你的其他host代码在一个文件里。但实际上，这是在device上的。

试图甩开.cu, 只在host端(例如.cpp)上声明一个texture是不可以的。这个文件必须被cuda编译器给过滤一遍。

你的第二个问题，你的绑定的确是在host代码里进行的，但不能在纯的host代码里，你需要将你的host code放入.cu才行。

system · 2013 年1 月 23 日 09:19

就是说，所有关于纹理声明纹理绑定的都必须在CU文件里实现吧？

system · 2013 年1 月 23 日 09:31

也不是“必须”，因为非要在cpp里也是可以的。

但是你不能使用常规的直接使用texture reference的做法了，此时使用cudaGetTextureReference()来取得。（实际上，这会从当前的module里尝试寻找）。

这样可以规避纹理引用到声明和绑定代码必须在一起，但多了一句取得texture reference的过程。

system · 2013 年1 月 23 日 09:40

好的，明白了，顺便问一下，常量存储器是怎么用的？在主机端声明还是在设备端声明？是不是用cudaMemcpyToSymbol 复制后就直接在kernel函数里能使用了？

system · 2013 年1 月 23 日 09:43

建议您另发一贴询问此问题，这样比较清晰。

system · 2013 年1 月 23 日 09:47

好的，谢谢，你们很热心，小白很感谢:lol

system · 2013 年1 月 23 日 09:51

是的。楼主你说的对，__constant__是在device端声明的。您的用法也没错误。

其次，还有另外一种用法，可以通过ldu指令，在确保了warp访问地址一致的情况下。可以使用constatn cache来访问显存中的内容。

最后，在2.x和3.x上，还有第三种用法：如果一个kernel需要一个在执行时候才知道的值，或者是一些值，可以作为形参。形参的值也是保存在constant cache的。（例如你就需要一个4B值，直接作为形参即可，无需反复cudaMemcpyToSymbol折腾，这样简单点）。

system · 2013 年1 月 23 日 09:56

那个，上文忘记说明，第二种用法也需要2.x和3.x。

换句话说，如果你用的是1.x, 那么你只能老老实实的用__constant__。

后2种用法，需要2.x和3.x的硬件。在这些硬件上，constant cache表现的很强大，如果能确保访问行为一致，推荐使用。（举个例子，你有1个4B的形参，这个形参(你的constant cache里的固定位置的值)可以直接作为某些指令的操作数）。

system · 2013 年1 月 23 日 10:06

__constant__在device端就声明了，__constant__常量如何初始化呢？它的值不是要由主机端传递进去的吗？

system · 2013 年1 月 23 日 10:06

[

嗯，我的是2.1的设备。就是想把一个数组常量放在__constant__里使用

system · 2013 年1 月 23 日 10:13

那么建议您使用__constant__, 这个用起来非常方便。是不错的选择。

system · 2013 年1 月 23 日 10:16

楼主您真健忘。

您可以采取您在7#的cudaGetSymbolAddress和cudaMemcpyToSymbol方式。

除此之外，您还可以使用__constant__ type your_arrray_name[length] = {…values…};的方式进行初始化，如果它们的值终生不变的话。

system · 2013 年1 月 23 日 10:19

对了，需要说明一下的是，一般常用的cudaMemcpyToSymbol()是上文前者cudaGetSymbolAddress的扩展，即它实际上是上文取得地址+然后复制的组合操作。

system · 2013 年1 月 24 日 08:33

[

我的纹理声明和绑定都是在一个CU文件里进行的：
主要代码如下：
texture<float, 2, cudaReadModeElementType> texRef;

extern “C” void cuda_rainbrow(float table,float src,uchar dest)
{
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0 ,cudaChannelFormatKindUnsigned);
cudaArray cuArray;
cudaMallocArray(&cuArray, &channelDesc, 256, 3); //为cuArray开辟空间
cudaMemcpyToArray(cuArray, 0, 0, table,25634, cudaMemcpyHostToDevice);//
texRef.addressMode[0] = cudaAddressModeClamp;//注意此tex为设备端声明
texRef.addressMode[1] = cudaAddressModeClamp;
texRef.filterMode = cudaFilterModeLinear;
texRef.normalized = false;
cudaBindTextureToArray(texRef, cuArray, channelDesc);

dim3 blocks(div(IMG_W,BlockX),div(IMG_H,BlockY));
dim3 threads(BlockX,BlockY);
rainbow<<<blocks,threads>>>(table,src,dest);
}
但rainbow 函数在另一个cu里文件面，我要在rainbow使用这个纹理怎么使用？用extern吗？

system · 2013 年1 月 24 日 08:52

上文已经给过建议了(通过在cudaGetTextureReference()在另外一个.cpp里使用）。

（1）你可以依然使用这个建议(只不过是在另外一个.cu（而不是.cpp)中使用cudaGetTextureReference()。）。

（2）如果你不愿意接受这个建议，你可以在一个.cu里写上你的texture<float, 2, cudaReadModeElementType> texRef;。
在你另外的rainbow.cu里面：
extern texture<float, 2, cudaReadModeElementType> texRef;
然后void cuda_rainbow(…){…不变…};

请酌情采纳上述(1)(2)建议里面的另外一种。
请注意的是：第二个建议要求另外一个文件也是.cu。而第一个建议无此要求。

EDIT: 如果采取第一个建议，请注意一些函数，例如你用到的cudaBindTextureToArray()具有多个重载版本，有些版本需要作为.cu(即被cuda编译器编译)时才有效. 所以如果你采用第一个建议，并且使用的不是.cu, 而是.cpp, 请注意一些函数版本不可用。不过你可以通过简单的改变参数类型以规避。