塊，線程，warpSize

Question

關於如何選擇#blocks & blockSize已經有很多討論，但我仍然缺少一些東西。我的很多顧慮解決這個問題：How CUDA Blocks/Warps/Threads map onto CUDA Cores?（爲了簡化討論，有足夠的perThread & perBlock內存內存限制是不是一個問題在這裏。）

kernelA<<<nBlocks, nThreads>>>(varA,constB, nThreadsTotal); 

1）爲了保持SM儘可能忙，我應該將nThreads設置爲warpSize的倍數。真正？

2）SM一次只能執行一個內核。那是所有SM的HWcores只執行kernelA。 （不是一些運行kernelA的HWcore，而有的運行kernelB。）所以如果我只有一個線程運行，我會「浪費」其他HWcores。真正？ 3）如果warp-scheduler問題的工作單位是warpSize（32個線程），並且每個SM有32個HWCore，那麼SM將被充分利用。當SM有48個HWcores時會發生什麼？當調度程序以32個塊的形式發佈工作時，如何保持所有48個內核的全部利用率？ （如果前面的段落是真的，如果調度器以HWcore大小爲單位發佈工作會不會更好？）

4）它看起來像warp-scheduler一次排隊2個任務。因此，當正在執行的內核停頓或阻塞時，第二個內核被交換。（不清楚，但我猜這裏的隊列深度超過2個內核。）這是正確的嗎？ 5）如果我的硬件上限爲512個線程（nThreadsMax），這並不意味着512個線程的內核在一個塊上運行速度最快。 （同樣，mem也不是問題。）如果我將512線程內核分佈在多個塊中，而不僅僅是一個，那麼我很有可能會獲得更好的性能。該塊在一個或多個SM上執行。真正？我認爲越小越好，但是我有多小我想要做什麼nBlocks？問題是，如何選擇nBlocks的價值體面呢？ （不一定是最佳的。）是否有數學方法來選擇nBlocks，還是僅僅是試用n-err。

Answer 1

讓我試着一一回答你的問題。

1. 這是正確的。
2. 「HWcores」完全代表什麼意思？你的發言的第一部分是正確的。

3. 根據NVIDIA Fermi Compute Architecture Whitepaper：「SM以32個並行線程組調度線程，稱爲warps。每個SM具有兩個warp調度器和兩個指令調度單元，允許兩個warp同時發佈和執行。兩個warp，並從每個warp發出一個指令到一組16個內核，16個加載/存儲單元或4個SFU。由於warps獨立執行，因此Fermi的調度程序不需要從指令流內檢查依賴性。

   此外，NVIDIA Keppler Architecture Whitepaper指出：「開普勒的四軸變形調度器選擇四個變形，並且每個變形週期可以分派兩個獨立的指令。」

   因此，「多餘」核心通過一次調度多個warp來使用。

4. warp調度程序調度相同內核的warps，而不是不同的kernel。

5. 不完全正確：每個塊被鎖定到一個SM，因爲這是它的共享內存所在。
6. 這是一個棘手的問題，取決於你的內核是如何實現的。你可能想看看瓦西里沃爾科夫的nVidia網絡研討會Better Performance at Lower Occupancy，它解釋了一些更重要的問題。不過，主要是，我建議你選擇你的線程數來提高入住率，使用CUDA Occupancy Calculator。

Answer 2

1）是的。

2）CC 2.0 - 3.0設備可以同時執行多達16個網格。每個SM被限制爲8個塊，因此爲了達到完全併發，設備必須至少有2個SM。

3）是，warp調度程序一次選擇併發出warp。忘記它們無關CUDA核心的概念。爲了隱藏延遲，您需要具有高指令級並行性或高佔用率。推薦CC 1.x> 25％，CC> = 2.0> 50％。一般而言，CC 3.0需要比2.0設備更高的佔用率，因爲調度程序翻倍，但每個SM僅增加33％。 Nsight VSE問題效率實驗是確定您是否有足夠的warps隱藏指令和內存延遲的最佳方法。不幸的是，Visual Profiler沒有這個指標。

4）沒有記錄warp調度算法;然而，它並沒有考慮線程塊產生的網格。對於CC 2.x和3.0設備，CUDA工作分配器將在分配來自下一個網格的塊之前分配來自網格的所有塊;然而，編程模型並不能保證這一點。

5）爲了保持SM忙，你必須有足夠的塊來填充設備。之後，你要確保你有足夠的經紗達到合理的佔用。使用大型線程塊有利有弊。大型線程塊通常使用較少的指令高速緩存並且在高速緩存上具有較小的佔用空間;然而，大型線程塊在syncthreads處停頓（SM可能變得效率較低，因爲選擇的warp較少），並傾向於使指令在類似的執行單元上執行。我建議每個線程塊嘗試128或256個線程來啓動。對於較大和較小的線程塊都有很好的理由。 5a）使用佔用率計算器。挑選太大的線程塊大小通常會導致您受限於寄存器。選擇太小的線程塊大小會發現受限於共享內存或每個SM限制的8個塊。