2012-09-08 32 views
5

我對CUDA的工作原理有些困惑,線程是否執行同一條指令(SIMT),但是使用不同索引訪問單個數據?或者它被認爲是「不同的數據」(所以它也是SIMD)?CUDA線程,SMX,SP和塊,它們是如何工作的?

SMX是整個GPU芯片嗎?一個SMX應該由多個SP組成,每個SP一次執行一個線程,是一個只分配給一個SP的線程塊?

我現在

回答

12

網格發射是1-3維推出的線程塊的有點糊塗。線程塊是1-3維的線程組。 CUDA工作分配器將線程塊分配給SMX單元。低端設備可能有1個SMX單元。高端設備可能具有> 10個SMX單元。

SMX單元以32個稱爲warps線程的組的形式分割線程塊。 SMX裝置最多可以分配64個經紗或16個區塊。由於資源限制(塊,扭曲,每個線程的寄存器,每塊的共享內存或障礙),數量可能會更少。

每個SMX設備都有4個變形調度器,每個調度器負責一部分變形。在每個循環中,warp調度程序將選擇合格的warp併發出1或2條指令。爲了雙重問題,兩條指令必須是獨立的,並使用不同的執行單元。例如,一條指令可以派發到浮點單元,另一條指令可以派發到加載存儲單元。

除了雙重發布之外,一個warp調度器可以對一個warp發回獨立的指令。當檢測到依賴關係或下一條指令的執行單元繁忙時,或者warp沒有指令(等待提取)時,warp調度程序將選擇不同的warp(如果有的話)。

每個線程都有自己的一組通用寄存器,條件代碼,謂詞代碼和本地存儲器。每個線程都是線程塊的成員。所有線程都可以訪問包含共享內存和屏障的線程塊資源。網格啓動中的所有線程都可以訪問包含常量內存,紋理綁定和表面綁定的網格資源。所有線程都可以訪問全局內存。

+0

非常好的總結。所以障礙是有限的資源?這是爲什麼? –

+0

PTX條指令支持每個線程塊(CTA)16個屏障。爲了高效,這些是有限的硬件資源。大多數計算程序使用1個屏障(__syncthreads())。 MAX_BLOCKS_PER_SM支持16的好處非常小,不可能發生在使用CUDA C/C++的情況下。 –

相關問題