http://journal.mycom.co.jp/articles/2007/09/13/nvidia/index.html
“GeForceの父” David Kirk博士、東大で並列コンピューティングについて講演
這篇的重點有：
1. 每個Shader Multi-processor(SM)有768個thread、8個SP、每個TPC有兩個SM、所以總共有12288個data-thread可以on-the-fly執行。
2. 每32個data-thread組成一個warp，只要以這個單位來切換，就可以達到無延遲。
3. 每個SM有自己的share memory。(G80是16KB)
其次是CUDA programming的技巧
4. programmer可以設定grid，每個grid則由1~512個block來組成。
5. 每個block內部的資料可以以1~3次元的方式排列，這個block稱為Cooperative Thread Array(CTA)。
到這邊已經很清楚了：這樣一段一段的把資料結構切出來，就是希望programmer替GPU做好data的localize。
這樣的話每個SM對每個CTA進行處理，SM的share memory也只以CTA內部的單位來處理，理論上只要CTA有設好 = localize有做好，就不會有太多cache miss。
然後等到要進行記憶體存取的時候，就會發生latency，這時則把整個CTA(data set)給卸載換掉，來確保GPU維持滿載。
所以，實際在寫CUDA程式的人，除了上述的要點之外，不需要再深入顧忌GPU實際上的結構；
或者說，上面的要點實際上就是G8x這個GPU最終的可擴充範圍。
在ISA固定(保留擴充性前提)的狀況下，CTA的資料結構固定住之後，接下來NVIDIA只要擴充SM & TPC的規模即可。
可以增加shader processor數量與增設指令、增加share memory容量、增加TMU數量與功能(G9x已經做了)….在這些範疇內都還算是G8x的結構延伸。
於是在維持CUDA相容性的狀況下，其實還是有很多東西可以增補，調整。
為了培育CUDA的軟體資源，NVIDIA短期內的GPU結構變化，大概是離不開「TPC內部的擴充」、和「GPU總TPC數量的擴充」這兩項了。
當然這樣講的話，好像GPU就只剩shader變化一樣，只是shader之外的變化其實真的就不容易看出來，因為大家也不知道shader以外的繪圖技術進展會朝向什麼方向。
比方說NVIDIA資助發展的Logarithmic Shadow Maps，讓shadow map的效率進一步提升，而只需要Triangle Setup Engine支援對數值域即可，這不是研究的最尖端本來就很難知道。
http://www.xbitlabs.com/articles/video/display/radeon-hd3870-hd3850.html
Xbit針對RV670與G92/G80的benchmark。
可以看出G92相對於G80的底層還是有很多改善….比方說分支性能就有改善。
我想請問warp和shader是什麼呢？能解釋一下這兩個名詞嗎？我看官方CUDA guide好像沒提到shader，warp我則是不太清楚，網站上對8 series的規格，shader和core的頻率也不一樣，另外不知道block、thread，對應到硬體上面是什麼呢，我知道一個Multi processors似乎有8個stream processors，而在硬體上multi processor內的processor可共用一塊share memory，而從軟體來看則是一個Block內的thread可共用share memory，似乎可以對應？還是說是無法對應的？因為在附錄也有提到：The maximum number of active blocks per multiprocessor is 8;The maximum number of active threads per multiprocessor is 768;Each multiprocessor is composed of eight processors, so that a multiprocessor is able to process the 32 threads of a warp in four clock cycles.
是可以對應沒錯啊。
基本上block(Cooperative Thread Array，CTA，軟體) 是每個Threaded Processor Cluster(TPC，硬體)處理的單位，thread在軟體上是每個32bit的數字(1D~4D一組)，在硬體上則是一塊用來記憶這個數字的空間，每個TPC有辦法記憶768個thread，但是每個軟體block(CTA)要處理的東西勢必有超過768這個數字的數量。
你可以把TPC裡面能放的thread當成某種快取，實際上那是類似CPU的FGMT般的東西，只要multiprocessor遇到需要記憶體存取的地方，會引起管線停頓，為了避免停頓，會以warp為單位跳到晶片上下個可以繼續run的”一堆warp”，繼續run，直到剛剛停頓的warp有辦法繼續之後再跳回來。
所以只要你會用到互相存取….比方說，一個n*(n-1) /2的式子，那麼我就會需要我現在在run的n”旁邊”的數字。為了加速，你可以一口氣把128個數字讀到share memory上，根據CUDA guide的說法，所有的指令都有4cycle的延遲，在這狀況下其實share memory的速度是和register一樣的，所以算起來就像一次算128個數字了，這樣就可以掩蓋住GPU存取記憶體的時間。
block指的就是會這樣交互存取的範圍，只要是這個範圍內就稱為一個block，也叫做CTA、每個CTA交給一個TPC處理。
所以式子運算的速度要考量到硬體裡面可以一次存取的數量，所以才要告訴你硬體上有幾個thread，因為可以讓你算目前的程式是不是會卡在記憶體上。
謝謝，我整理一下我的理解，假設我的硬體有8個multi processors，軟體grid內有32個block，一個block有128個thread，當硬體在跑時，會把32個block分給8個multi processors(平均分？)，對於單個multi processor而言，它會以warps為單位執行，直到128個(4個warp)threads都跑完嗎？若以上都正確，那不知道當我只有一個block時，是否就只會使用到一個multi processor？
上面有一點筆誤，應該是4*4個warps都跑完..另外我還有一些問題，除了shader是什麼以外，我今天在看CUDA提供的example檔，注意到Bank conflicts的問題，在硬體圖來看就是不同threads同時存取同一個bank造成的，但是我不懂他上面提到，當有下面這式子時：
__shared__ float shared[32];
float data = shared[BaseIndex + s * tid];
為什麼只有s為odd才不會有bank conflicts？
> 若以上都正確，那不知道當我只有一個block時，是否就只會使用到一個multi processor？
是的，只會用到一個，所以你要盡量把block的執行時間均分，否則分配到mp的時候就不會均勻。
> shader是什麼
CUDA裡面沒有這個東西。_A_
硬要說的話，shader = CUDA 的 kernal program本身。
> __shared__ float shared[32];
> float data = shared[BaseIndex + s * tid];
> 為什麼只有s為odd才不會有bank conflicts？
其實你要故意錯開….
今天假設我們程式是先跑row、再跑column，所以是bank0在讀share memory的前幾個數字(比方說16個)，然後後幾個數字放在bank1。
但是一開始跑row沒事，接著跑column就完蛋了，剛好所有thread都讀到剛剛的bank0上頭，全部都炸掉了。
所以有一個方法是故意錯開，比方說16×16的矩陣，你就故意寫成17×16，這樣data[0][0]是在bank0、data[1][0]會變成bank1。
謝謝，不過我聽不太懂您舉的例子，所謂的”前幾個數字”是什麼意思呢？資料存在bank時不知道是依什麼規則排列的呢？