風辰的cuda入門系列教程
1. cuda簡介
2. linux下cuda開發環境構建
3. cuda架構
4.cuda c語言
前一節，我已經說了怎樣在ubuntu linux上構建cuda開發環境了，對乙個語言來說，最簡單的，也是用得最多的當然是它的語法了，下面我簡單的介紹一下cuda的語法。
cuda c不是c語言，而是對c語言進行擴充套件。cuda對c的擴充套件主要包括以下四個方面：
函式型別限定符，用來確定函式是在cpu還是在gpu上執行，以及這個函式是從cpu呼叫還是從gpu呼叫。
變數型別限定符，用來規定變數儲存什麼位置上。在傳統的cpu程式上，這個任務由編譯器承擔。在cuda中，不僅要使用主機端的記憶體，還要使用裝置端的視訊記憶體和gpu片上的暫存器、共享儲存器和快取。在cuda儲存器模型中，一共抽象出來了8種不同的儲存器。複雜的儲存器模型使得必須要使用限定符要說明變數的儲存位置。
texture，texture表明其繫結的資料可以被紋理快取加速訪問，其實資料本身的存放位置並沒有改變，紋理是**於圖形學的一介概念，cuda使用它的原因一部分在於支援圖形處理，另一方面也可以利用它的一些特殊功能。
如果在gpu上執行的函式內部的變數沒有限定符，那表示它存放在暫存器或者本地儲存器中，在暫存器中的資料只歸執行緒所有，其它執行緒不可見。
如果sm的暫存器用完，那麼編譯器就會將本應放到暫存器中的變數放到本地儲存器中。
執行配置運算子<<< >>>，用來傳遞核心函式的執行引數。執行配置有四個引數，第乙個引數宣告網格的大小，第二個引數宣告塊的大小，第三個引數宣告動態分配的共享儲存器大小，預設為0，最後乙個引數宣告執行的流，預設為0。
五個內建變數，用於在執行時獲得網格和塊的尺寸及執行緒索引等資訊
其它的還有數學函式，原子函式，紋理讀取、繫結函式，內建柵欄，記憶體fence函式等。一般而言，知道這些就應該能夠寫出cuda程式了，當然要寫好的話，必須知道很多其它的細節。
6. 乙個例子
上一節，已經簡單的說了一下cuda c的基本語法；因而在本節，兄弟決定以乙個例子為基礎說明cuda程式的基本組成部分，不過說實話兄弟選擇的例子並不太好，這個例子就是採用積分法計算圓周率/π的值。其計算原理是：在[0,1]範圍內積分1/(1+x*x)
首先，讓我們看一下在cpu上的計算流程，其計算流程如下
/*序列計算pi的程式，基本思想為：將積分區間均分為num小塊，將每小塊的面積加起來。
*/ float cpupi(int num)
return sum/num;}
複製**
很明顯，我們可以將for迴圈分解，使用cuda處理。
有乙個問題就是：for內部對sum變數的更新是互斥的，而cuda中並沒有浮點原子函式，對於這個問題的解決方案是：將for迴圈內部的兩個語句拆開，分成兩個核心函式來做運算。核心函式英文名為kernel，就是乙個能夠在gpu上運算的模組。第乙個核心計算計算每個小積分塊面積，並將每個block內所有執行緒對應的積分塊面積加起來，存入全域性儲存器；第二個核心將前乙個核心存入的資料加起來。下面是kernel**：
/*在gpu上計算pi的程式，要求塊數和塊內線程數都是2的冪
前一部分為計算block內歸約，最後大小為塊數
後一部分為單個block歸約，最後儲存到*pi中。
*//*
在gpu上計算pi的程式，要求塊數和塊內線程數都是2的冪
前一部分為計算block內歸約，最後大小為塊數
後一部分為單個block歸約，最後儲存到*pi中。
*/__global__ void reducepi1(float *d_sum,int num)
for(int i=(blockdim.x>>1);i>0;i>>=1)
__syncthreads();
}if(threadidx.x==0)
d_sum[blockidx.x]=s_pi[0];
}__global__ void reducepi2(float *d_sum,int num,float *d_pi)
// printf("%d,%f\n",id,s_sum[id]);
if(id==0)}
複製**
其中__syncthreads()是cuda的內建命令，其作用是保證block內的所有執行緒都已經執行到呼叫__syncthreads()的位置。
由上面的**可以看出，使用使用cuda的主要阻礙在於資料相關性。
一般而言，cuda程式的基本模式是：
這個程式在我的機器(cpu 2.0ghz,gpu gtx295)上的加速比超過100，不知道在你們的機器上能夠加速多少？
7. cuda程式設計模式
cuda支援大量的執行緒級並行（thread level parallel），並在硬體中動態地建立、排程和執行這些執行緒，在cpu中，這些操作是重量級的，但是在cuda中，這些操作是輕量級的。cuda程式設計模型將cpu作為主機（host），而將gpu做為協處理器（coprocessor），或者裝置（device），以cpu來控制程式整體的序列邏輯和任務排程，而讓gpu來執行一些能夠被高度執行緒化的資料並行部分。即讓gpu與cpu協同工作，更確切的說是cpu控制gpu工作。gpu只有在計算高度資料並行任務時才發揮作用。
一般而言，ｃｕｄａ並行程式包括序列計算部分和平行計算部分，平行計算部分稱之為核心（kernel），核心只是乙個在ｇｐｕ上執行的資料並行**段。理想情況下，序列**的作用應該只是清理上個核心函式，並啟動下乙個核心函式，但由於目前的gpu的功能仍然十分有限，序列部分的工作量仍然十分可觀
8. cuda執行緒層次
gpu執行緒以網格（grid）的方式組織，而每個網格中又包含若干個執行緒塊，在g80/gt200系列中，每乙個執行緒塊最多可包含512/768/1024個執行緒，fermi架構中每個執行緒塊支援高達1536個執行緒。同一執行緒塊中的眾多執行緒擁有相同的指令位址，不僅能夠並行執行，而且能夠通過共享儲存器（shared memory）和柵欄（barrier）實現塊內通訊。這樣，同一網格內的不同塊之間存在不需要通訊的粗粒度並行，而乙個塊內的執行緒之間又形成了允許通訊的細粒度並行。這些就是cuda的關鍵特性：執行緒按照粗粒度的執行緒塊和細粒度的執行緒兩個層次進行組織、在細粒度並行的層次通過共享儲存器和柵欄同步實現通訊，這就是cuda的雙層執行緒模型。
在執行時，gpu的任務分配單元（global block scheduler）將網格分配到gpu晶元上。啟動cuda 核心時，需要將網格資訊從cpu傳輸到gpu。任務分配單元根據這些資訊將塊分配到sm上。任務分配單元使用的是輪詢策略：輪詢檢視sm是否還有足夠的資源來執行新的塊，如果有則給sm分配乙個新的塊，如果沒有則檢視下乙個sm。決定能否分配的因素有：每個塊使用的共享儲存器數量，每個塊使用的暫存器數量，以及其它的一些限制條件。任務分配單元在sm的任務分配中保持平衡，但是程式設計師可以通過更改塊內線程數，每個執行緒使用的暫存器數和共享儲存器數來隱式的控制，從而保證sm之間的任務均衡。任務以這種方式劃分能夠使程式獲得了可擴充套件性：由於每個子問題都能在任意乙個sm上執行，cuda程式在核心數量不同的處理器上都能正常執行，這樣就隱藏了硬體差異。
對於程式設計師來說，他們需要將任務劃分為互不相干的粗粒度子問題(最好是易平行計算)，再將每個子問題劃分為能夠使用執行緒處理的問題。
同一執行緒塊中的執行緒開始於相同的指令位址，理論上能夠以不同的分支執行。但實際上，在塊內的分支因為sm構架的原因被大大限制了。核心函式實質上是以塊為單位執行的。同一執行緒塊中的執行緒需要sm中的共享儲存器共享資料，因此它們必須在同乙個sm中發射。執行緒塊中的每乙個執行緒被發射到乙個sp上。
任務分配單元可以為每個sm分配最多8個塊。而sm中的執行緒排程單元又將分配到的塊進行細分，將其中的執行緒組織成更小的結構，稱為執行緒束（warp）。在cuda中，warp對程式設計師來說是透明的，它的大小可能會隨著硬體的發展發生變化，在當前版本的cuda中，每個warp是由32個執行緒組成的。sm中一條指令的延遲最小為4個指令週期。8個sp採用了發射一次指令，執行4次的流水線結構。所以由32個執行緒組成的warp是cuda程式執行的最小單位，並且同乙個warp是嚴格序列的，因此在warp內是無須同步的。在乙個sm中可能同時有來自不同塊的warp。當乙個塊中的warp在進行訪存或者同步等高延遲操作時，另乙個塊可以占用sm中的計算資源。這樣，在sm內就實現了簡單的亂序執行。不同塊之間的執行沒有順序，完全並行。無論是在一次只能處理乙個執行緒塊的gpu上,還是在一次能處理數十乃至上百個執行緒塊的gpu上，這一模型都能很好的適用。
目前，某一時刻只能有乙個核心函式正在執行，但是在fermi架構中，這一限制已被解除。如果在乙個核心訪問資料時，另乙個核心能夠進行計算，則可以有效的提高裝置的利用率。
9. cuda儲存器組織
10.cuda執行模式
CUDA學習系列教程（四）
gpu程式一般步驟 過程中，一般要盡量降低資料通訊的消耗，所以如果程式需要複製大量的資料到gpu，顯然不是很合適使用gpu運算，最理想的情況是，每次複製的資料很小，然後運算量很大，輸出的結果還是很小，複製回cpu。核函式是gpu每個thread上執行的程式。必須通過 gobal 函式型別限定符定義。...
kotlin入門潛修系列教程
kotlin入門潛修之開門篇 kotlin是什麼？kotlin入門潛修之基礎篇 基礎語法 kotlin入門潛修之基礎篇 packages 包 kotlin入門潛修之基礎篇 控制表示式 kotlin入門潛修之基礎篇 返回和跳轉 kotlin入門潛修之類和物件篇 類和物件 kotlin入門潛修之類和物件...
angular入門系列教程4
本篇主要目的就是繼續完善home頁下的index子頁面的內容，處理乙個列表，進行增刪改查過濾等操作。主要的更改有兩個，乙個是修改模板index.html，還有就是增加控制器indexctl的控制邏輯。先看模板 很簡單，就是在上邊增加了操作區，主要是乙個輸入框，乙個增加按鈕，還有就是過濾列表的輸入框 ...