了解和使用共享記憶體(1)
rob farber 是西北太平洋國家實驗室（pacific northwest national laboratory）的高階科研人員。他在多個國家級的實驗室進行大型並行運算的研究，並且是幾個新創企業的合夥人。大家可以發郵件到[email protected]與他溝通和交流。
cuda執行模型
為了提高效能，每個硬體多處理器都能同時積極處理多個塊。塊的數量取決於每個執行緒的暫存器的數量，每個塊的共享記憶體是由給定的核心提出要求的。在同一時間被同一多處理器處理的塊被稱為「活動的」。對資源要求最低的核心可以更好地利用（或占用）每個多處理器，因為多處理器的暫存器和共享記憶體在所有的活動的塊的執行緒之間被分開。使用cuda占有率計算器來尋求執行緒數量和活動的塊與暫存器數量和共享記憶體量之間的平衡。找到正確的結合點可以極為有效地加強您的核心的效能。如果每個多處理器上沒有足夠的暫存器或共享記憶體可以處理至少乙個塊，核心可能就不能啟動（參見第三節中關於cudagetlasterror() 的討論，以確定如何發現和處理這類故障）。
每個活動塊被分割成執行緒simd（單指令多資料）群，稱為warp：每個warp包含同樣數量的執行緒，稱為warp size，被多處理器以simd方式執行。這意味著warp中的每個執行緒傳遞的都是指令庫中的相同指令，指導執行緒執行一些操作或操縱本地和/或全域性記憶體。從硬體角度看，simd模型效率高而且經濟有效，但從軟體角度來看，很遺憾，它會使條件操作序列化（比如，使條件句的兩個分支都必須被求值）。請注意條件操作對核心執行有深刻的影響。謹慎使用的話，一般是可以控制的，但也有可能引起某些問題。
活動warp（比如所有活動塊的所有warp）是以時間片分配的：執行緒排程程式定期從乙個warp切換到另乙個warp以最大限度利用多處理器的計算資源。塊之間或塊中的warp之間的執行次序是不確定的，也就是說可以是任何次序。但是，執行緒可以通過__syncthreads()進行同步。請注意只有執行__syncthreads()後，才能保證寫入共享（和全域性）記憶體是可見的。除非變數被宣告為易失性（volatile）變數，否則編譯器可以被用於優化（例如重新排序或消除）記憶體的讀和寫，以提高效能。__syncthreads()允許在乙個條件句範圍內被呼叫，但只有當條件的求值方式在整個執行緒塊中都相同時才可以。否則的話，**執行過程可能會被掛起或產生意想不到的***。可喜的是，__syncthreads()的開銷很低，因為在沒有執行緒等待任何其他執行緒的情況下，它僅花費4個時鐘週期為乙個warp傳送指令。half-warp指的是warp的前一半或者後一半，這是記憶體訪問包括本期後面將要討論的合併記憶體訪問的乙個重要概念。
上面的討論有幾條內容需要繼續記住：
多處理器資源如共享記憶體是有限的，並且較有價值；
有效管理有限的多執行緒資源，如共享記憶體，因為諸多的cuda啟動裝置配置實際上對cuda開發人員來說都是至關重要的。
注意，條件操作（如if 語句）會對你的核心執行時產生重要影響。
cuda占有率計算器和nvcc編譯器是非常重要的共圖，我們要好好使用和學習－特別是當我們琢磨執行配置的時候。
cuda記憶體模式
表1 圖示了在裝置上執行的執行緒可讀取全域性記憶體和片上記憶體。
每個多處理器，如上面的塊(0, 0) 和塊(1, 0)所示，包含有以下四種記憶體型別：
每執行緒一套本地暫存器；
平行資料快取或共享記憶體由所有執行緒共享，執行共享記憶體空間；
唯讀常量快取由所有執行緒共享，從常量記憶體空間加速讀取。這是作為裝置記憶體的唯讀區執行的（常量記憶體將在下一專欄裡討論。屆時請參閱cuda programming guide 的5.1.2.2章節）。
所有處理器共享的唯讀紋理快取，加速從紋理儲存器空間的讀取。這是作為裝置記憶體的唯讀區執行的（在接下來的文章裡會**紋理儲存器。屆時請參閱cuda programming guide 的5.1.2.3章節）。
不要弄混淆的是，該圖示包括乙個在多處理器內的標識為「本地記憶體」的塊。本地記憶體意味著「在每個執行緒範圍之內」。它是記憶體抽象，而不是多處理器的乙個實際的硬體元件。事實上，本地記憶體通過編譯器在全域性記憶體內進行分配，和其它任何全域性記憶體區一樣提供完全相同的效能。本地記憶體基本上由編譯器使用，儲存任何程式設計師認為相對執行緒來說為「本地」的內容，這些內容因為一些原因不適合更快的記憶體。在正常情況下，在核心裡宣布的自動變數存於暫存器中，可快速讀取。在有些情況下，編譯器可能會選擇將這些變數放置在本地記憶體中，譬如當有太多暫存器變數，陣列包含有四個以上的元素，一些結構或陣列可能會小號太多暫存器空間，或當編譯器不能確定陣列是否根據常量索引定址。
注意，因為本地記憶體可能會導致效能降低。對ptx彙編**的檢查（通過使用-ptx 或 -keep option獲得）可判斷在第一編譯階段，變數是否已經存於本地記憶體中，因為變數是通過使用.local助憶符來宣告的，使用ld.local和st.local助憶符讀取。如果在最初階段沒有將它放置到本地記憶體，那麼在隨後的編譯階段如果發現它可能會消耗目標架構太多的暫存器空間也可能會決定將它放置在那裡。
一直到下乙個專欄之前，我都推薦使用占有率計算器，以扎扎實實地了解執行模型和核心啟動執行配置是如何影響暫存器的數量和共享記憶體的數量的。
CUDA, 用於大量資料的超級運算 第四節
了解和使用共享記憶體 1 rob farber 是西北太平洋國家實驗室 pacific northwest national laboratory 的高階科研人員。他在多個國家級的實驗室進行大型並行運算的研究，並且是幾個新創企業的合夥人。大家可以發郵件到rmfarber gmail.com與他溝通和...
安裝用於深度學習的CUDA環境
下面正式介紹在64位win8.1系統下安裝theano的過程 在anaconda下安裝mingw包即可，不需要安裝tdm gcc或者mingw steup.exe等，在cmd下輸入 conda install mingw libpython 即可在anaconda的安裝目錄下看到mingw包，然後在...
大量資料的問題
這些問題一般有著記憶體限制，使用hashmap和位 決不實際。1.只用2gb內存在20億個整數中找到出現次數最多的數？將20億個整數的大檔案用hash函式分為16小檔案 這個時候同乙個數一般分到了同乙個小檔案上，小檔案的數最好不要超過2億 這個時候每個小檔案用hash函式計算出現次數，這個時候得到1...