CUDA中文教程03之心得體會

定義在GPU上的變數：

1、使用關鍵字__device__ __local__  int X，則意味著該變數是定義在thread中的，它的生存週期跟它所在的thread一致。實際上定義為__local__ 的變數會存在global memory 中，所以速度也會很慢，一般不採用__local__關鍵字定義變數。關鍵字預設情況下的變數是存在register中的，速度比存在global memory中快，只有當register存滿了變數之後，系統才會自動把變數定義為__local__的，所以不要隨意的採用__local__關鍵字，這其實是下下策。

2、使用關鍵字__device__ __shared__  int X，則意味著該變數是定義在block中的，它的生存週期跟它所在的block一致，並且為該block裡的512個thread共享，都可以訪問到這個變數。

3、如果預設了第二個關鍵字，即__device__              int X，則是定義在grid中的，不僅在GPU中可見，CPU也可見。

4、使用關鍵字__device__ __constant__  int X，則意味著該變數是定義在grid中的，是一常量，在run的過程中不會改變其值，且GPU及CPU均可見，在CPU中可見，是因為CPU要把該值傳入GPU中。

*使用這些關鍵字時，如果是用了__local__，__shared__，__constant__，則不必要在前面加__device__，系統就會知道定義的是GPU上的變數，如果是定義了global memory 的變數，則需寫__device__關鍵字即可。

*自動變數，即沒有任何限定詞的，會自動的放到register中，除了陣列，陣列會存在local memory中，所以當宣告陣列時，必須存到thread中去run。

怎樣選擇關鍵字呢？

第一步，考慮該關鍵字是否被CPU可見：“是”，進入第二步；“不是”，進入第三步。

第二步，如果要被CPU可見，則選擇關鍵字__global__或者__constant__。而且在宣告變數時，必須寫在所有函式體外，保證全域性性。

第三步，如果不被CPU可見，則選擇關鍵字__shared__或__local__或者預設不寫關鍵字（存在register中），這時變數宣告必須在kernel 函式中。

shared memory 是一個很重要的概念，因為如果我們每次都去access global memory的話，就要很久的時間，所以我們要提高速度，就得想怎樣換到shared memory中去。這裡我們採用的是tile data ，即把資料切片的方法，變成一個個subset，使其剛好滿足shared memory的大小，處理完資料之後再從shared memory傳到global memory去。

*指標只能指向定義在global memory中的函式或變數。

源自：xieyan blog

來自 “ ITPUB部落格 ” ，連結：http://blog.itpub.net/22785983/viewspace-619745/，如需轉載，請註明出處，否則將追究法律責任。