我自己无法弄明白,确保内核中使用的内存不变的最佳方法是什么.
http://stackoverflow…r-pleasant-way也有类似的问题.
我正在使用GTX580并仅编译2.0功能.我的内核看起来像
我正在使用GTX580并仅编译2.0功能.我的内核看起来像
__global__ Foo(const int *src,float *result) {...}
我在主机上执行以下代码：
cudamalloc(src,size); cudamemcpy(src,hostSrc,size,cudamemcpyHostToDevice); Foo<<<...>>>(src,result);
另一种方法是添加
__constant__ src[size];
到.cu文件,从内核中删除src指针并执行
cudamemcpyToSymbol("src",cudamemcpyHostToDevice);
Foo<<<...>>>(result);
这两种方式是等效的还是第一种不保证使用常量内存而不是全局内存？大小动态变化,所以第二种方式在我的情况下不方便.
解决方法
第二种方法是确保将数组编译为CUDA常量内存并通过常量内存缓存正确访问的唯一方法.但是你应该问自己如何在一个线程块中访问该数组的内容.如果每个线程都将统一访问数组,那么使用常量内存将具有性能优势,因为常量内存缓存有一个广播机制(它还节省了全局内存带宽,因为常量内存存储在offchip DRAM和缓存中减少DRAM事务计数).但是如果访问是随机的,那么可以对本地内存的访问进行序列化,这将对性能产生负面影响.
可能适合__constant__内存的典型事物是模型系数,权重和需要在运行时设置的其他常量值.例如,在Fermi GPU上,内核参数列表存储在常量存储器中.但是如果内容是非均匀访问的,并且成员的类型或大小在调用之间不是恒定的,则优选正常的全局内存.
另外请记住,每个GPU上下文的常量内存限制为64kb,因此将大量数据存储在常量内存中是不切实际的.如果您需要大量具有缓存的只读存储,则可能值得尝试将数据绑定到纹理并查看性能如何.在费米前卡上,它通常会产生方便的性能增益,在Fermi上,与全局内存相比,结果可预测性更低,因为该架构中的缓存布局得到改善.