一般CUDA中使用const memory的時機是該變數只要讀取不做寫入且資料量不大時使用
但是在計算版本3.5上，手冊提出說可以在函式中加入const (float*) restrict Var
要求編譯器把變數從不同的管線放入唯讀記憶體
據我的了解這兩個唯讀記憶體應該是由不同的硬體來支持
我比較想知道在計算版本3.5上是將資料放到哪裡去了？是在L1上嗎？
如果是的話會跟shared memory搶記憶體嗎？
另外，計算版本3.5的唯讀記憶體做的事情看起來比用__const__ 宣告還來的好用
因為不用同步性的複製資料也不用被限制大小！
楼主您好，
(1)使用__constant__将放入constant cache的特定bank
(2)程序中的常数(例如1234.5f), 如果无法放入指令中（作为立即数），也将放入constant cache, 但将是另外一个bank. 这个是自动的。
(3)在3.5上使用const float * restrcit p; 然后通过p访问，将使用ldg指令，该指令将会通过read-only cache来进行缓冲的。
您的第二个问题，数据在哪里，刚才第三点说了，将会被放置到read-only cache, 这个是以前的texture cache改善而来的，容量变的更大，使用也变得自动了(但是您依然可以用__ldg()的, 如果您不需要自动的话)。read-only cache (texture cache)不是L1 Data-Cache, 也不是L1 Instruction Cache, 它就是它。
因为它就是它，它不会跟shared memory抢空间的，它是独立的48KB/SMX。
您的第三个问题，
对__constant__(constant cache)和read-only cache，他们是不同的，也不能说谁更好。
对于warp一致的常数/只读访问，显然用constant cache较好，因为它具有超低的延迟(当命中的时候).
对于p[id]这种，显然用read only cache较好，在访问不同的数据的时候，它比constant cache有更大的带宽，以及，有大的多的延迟。
您的第四个问题，
不用同步性的资料复制不用被限制大小，这个无法理解。
因为无法理解，所以无法回答第四个问题，请其他会员、版主、NVIDIA原厂支持、总版主为您解答。
感謝版主的回應
我想在確認一下const restrcit 是用textrue cache的話
那這樣會不會跟其他的tex_Obj發生資源搶占的問題(例如頻寬)？
或者如果都是在read only的情況下，有沒有說用tex或const __restrict__誰會比較好？
據我了解tex沒辦法放下double的型態，所以我有在思考要把tex換回對齊的double
至於第四個問題當我沒問好了！
我只是想說複製一筆資料到const memory是sync的需要等待傳輸完後才能執行事情
在kernel執行時應該也是只是沒看到罷了！
楼主您好，
（1）我无法确定直接的__ldg或者const __restrict__的读取和普通的tex*()读取，究竟是哪一方会在对read only cache的使用上占据上风。换句话说，我目前没有关于read only cache的替换策略(cache replacement policy)的资料。因为此问题无法直接回答您。
以及，如果需要一个可能的答案的话，那可能是LRU策略的，您的最近的tex*()访问可能将驱逐cache中的__ldg得来的数据，或者反过来也亦然。
（2）使用哪个好，或者说哪个不好，您可以这样看，使用texture能带来更多的功能，例如对坐标的normalize, 以及插值的处理，以及越界的处理(wrap或者mirror或者border等)等等，以及对数据的临近性提前存储等等，这些高级的功能只有通过texture访问才能实现。如果您不需要这些功能，您可以简单的使用__ldg或者const restrict，后者相当与一个后备存储是整个global memory构成的线性内存的普通surface读取。
所以简单的说，如果不需要这些texture功能，建议直接使用__ldg。
甚至建议使用直接的普通读取(直接L2), 有的时候，直接读取也可能展现更好的性能。这个不一定。
（3）纹元的类型的确不可以是double, 但是如果您可以使用int2或者float2替代，这样可以完成直接读取的，能够利用cache. （当然，您这样就不能使用插值等功能了）。
（4）复制到__constant__的过程，您也可以和host异步的api的，不一定非要同步。
您可以考虑cudaGetSymbolAddress() + cudaMemcpyAsync()或者cudaMemcpyToSymbolAsync(), 这样复制过称可以和其他的kernel执行overlap, 或者在tesla卡上，存在2个copy engines的时候可以和另外的一个方向的复制同时进行。
现在已经过子夜12点了，祝您端午节愉快！
關於版主所提的第三點用float2實現double的部分可以再具體說明一下嗎？有點不懂，或者有沒有一些參考可以提供，謝謝！
另外我想再釐清一下在kepler架構下的L1就是僅能作為緩衝之用，使用者無法自行配置資料進去，除非是用shared memory，若此敘述為真的話，那nVidia提供的cudaFuncSetCacheConfig (PreferShared / L1)在kepler架構下是沒有意義的？原則上就是shared memory不夠用就加大這樣？加大L1並不會有任何幫助。
至於版主所言tex跟const __restrict__都是走tex的管線，但tex多了一些特異功能那這樣跟用surface memory有什麼區別？
我會想用surface是因為這個變數在某些kernel是唯獨變數，有些是要寫入的變數，只是我想要加速變數的讀取所以讓他走tex管線。可是const __restrict__也實現了讓變數走tex管線當作唯獨變數，不加上關鍵字的話又可以寫入資料到變數裡面。如果是這樣的話，那nVidia為什麼要在CC3.0上加入surface的型態，cuArray的型態有什麼有優勢而我卻沒有善用嗎？
還是說僅如版主總結的取決於有沒有用到tex的特異功能以及硬體採用不同的策略讀取變數刪除變數？
LZ您好：
1：因為float2和double都是連續的8B的空間，所以將一個double的變量在讀取的時候當做是float2讀取，在使用的時候，就地轉換為double的類型使用。
2：L1 cache，L2 cache在哪個架構下都是僅作緩衝之用的，使用者無法手工配置。NV的GPU提供了可以手工配置的cache，名為shared memory。
以及“cudaFuncSetCacheConfig (PreferShared / L1)”這個是有意義的，因為在fermi和kepler架構下，L1 cache和shared memory使用了相同的硬件資源，這一項配置是對該資源的劃分。
以及，不甚理解您“原則上就是shared memory不夠用就加大這樣”的含義，shared memory+L1 cache一共64KB/SM，按照前述配置劃分好以後，是不變的。
而L1 cache有自己的緩衝功能，從這一點說，加大L1 cache的容量可能會提升程序總體的效能。以及 L1 cache加大是以shared memory容量減小為代價的，一些使用shared memory比較多的程式可能無法執行。但是只要shared memory夠用，也並無負面影響。
其他部份，留待橫掃斑竹回來為您回答。
祝您好運~
好的。我继续补充ICE的说法。
（1）-话说楼主“特异功能”这词用的真是妙。
简单的说，用前者可以随意访问任何一个线性地址开始的任何一个元素。而用surface reference/object读取这个是做不到了。前者极大的简化了代码的写法。
（2）如果一个kernel里，对一个缓冲区又要读取，又要写入，当且仅当您读写的位置是能完全分离开的情况下，您才可以使用surface访问的，其他任何访问（例如本kernel内先写入，又读取），将会导致完全无法预测的结果。请注意不要这样使用surface.
（3）同理，const __restrict__指向的指针，您可以用另外一个指针（无此修饰）进行写入。但请注意的是，它一是和surface读写有同样的问题（不要同时在一个kernel进行), 二是在同一个kernel内进行本身就违反了__restrict__的本意。
（4）surface读写在1.x和2.x的老卡上也是支持的。这不是3.0的新加入功能。新加入的只是surface object, 原来还有surface reference的一套使用模型的。
（5）关于您的为何要用surface的问题，一是surface本身可以利用texture/read only cache, 二是它可以直接写入到cuArray。cuArray的优势本身是存在的，例如它可能是将数据提前按线性、平面、空间循序进行数据重拍的，以便读取方便（此点手册没直接说。但可以大致看出）。
关于您的最后一句话，表示未能看懂您的中文表达。
對於L1或有疑問是我曾經有把一些沒有用到shared的kernel去做L1/sheard的比例調整計時，發現幾乎對於程式的效率沒有什麼影響，且在手冊上似乎只有說預設shared不夠時可以調整，也沒有強調對於改動比例後造成的副作用，所以才會有疑問。
話說回來我還是不太會用surface就是了，用起來速度都變慢…
也感謝兩位版主熱情的回應，我會再仔細思考看看當中的差異。
(題外話，CUDA5.5的Profiler介面、功能改版改好多變得不太會用)