1. 執行緒束分化
執行緒束是sm中基本的執行單元。 當乙個執行緒塊的網格被啟動後， 網格中的執行緒塊分布在sm中。 一旦執行緒塊被排程到乙個sm上， 執行緒塊中的執行緒會被進一步劃分為執行緒束。乙個執行緒束由32個連續的執行緒組成， 在乙個執行緒束中， 所有的執行緒按照單指令多執行緒（simt） 方式執行； 也就是說， 所有執行緒都執行相同的指令， 每個執行緒在私有資料上進行操作。 下圖展示了執行緒塊的邏輯檢視和硬體檢視之間的關係。
gpu是相對簡單的裝置， 它沒有複雜的分支**機制。 乙個執行緒束中的所有執行緒在同一週期中必須執行相同的指令， 如果乙個執行緒執行一條指令， 那麼執行緒束中的所有執行緒都必須執行該指令。 如果在同一執行緒束中的執行緒使用不同的路徑通過同乙個應用程式， 這可能會產生問題。思考下面語句：
if(cond)
else
如果乙個執行緒束中的執行緒產生分化， 執行緒束將連續執行每乙個分支路徑， 而禁用不執行這一路徑的執行緒。 執行緒束分化會導致效能明顯地下降。 在前面的例子中可以看到， 執行緒束中並行執行緒的數量減少了一半： 只有16個執行緒同時活躍地執行， 而其他16個被禁用了。條件分支越多， 並行性削弱越嚴重。
注意， 執行緒束分化只發生在同乙個執行緒束中。 在不同的執行緒束中， 不同的條件值不會引起執行緒束分化。
2. 簡單的執行緒束分化**：
#include #include #include "device_launch_parameters.h"
#include #include // kernel1
__global__ void mathkernel1(float *c)
else // 奇數執行緒
c[tid] = ia + ib;
}// kernel2
__global__ void mathkernel2(float *c)
else // 奇數
c[tid] = ia + ib;
}// kernel3
__global__ void mathkernel3(float *c)
if (!ipred)
c[tid] = ia + ib;
}// kernel4
__global__ void mathkernel4(float *c)
else // 奇數
c[tid] = ia + ib;
}// kernelwarm
__global__ void warmingup(float *c)
else
c[tid] = ia + ib;
}// 主函式
int main(int argc, char **ar**)
執行緒束分化
在英偉達gpu中，執行緒束是最小的執行單位，執行緒束的大小代表了乙個sm同時併發的執行緒數目。一般在sm實際排程中，sm獲取到當前指令後，會將該指令廣播到sm內所有的硬體core中即sm中，這樣做可以減少從記憶體中讀取指令的次數，提高效率。而在cpu中，每個核是單獨執行乙個任務，在執行前每個核各自從...
CUDA 學習（七） 執行緒束
一 概述 執行緒束是gpu的基本執行單元。gpu是一組simd 向量處理器的集合。每一組執行緒或執行緒束中的執行緒同時執行。在理想狀態下，獲取當前指令只需要一次訪存，然後將指令廣播到這個執行緒所占用的所有sp中。當使用gpu進行程式設計時，必須使用向量型別指令，因為gpu採用的是向量體系結構，只有讓...
CUDA執行緒配置
1 每個sm上面失少要有192個啟用執行緒，暫存器寫後讀的資料依賴才能被掩蓋 2 將 暫存器 的bank衝突降到最低，應盡量使每個block含有的執行緒數是64的倍數 3 block的數量應設定得令可用的計算資源被充分的利用。由於每個block對映到乙個sm上面，所以至少應該讓block的數目跟sm...