工作總結 2010 3 16

在cuda中，一般的資料複製到的顯示卡記憶體的部分，稱為global memory。這些記憶體時沒有cache 的，而且，訪問global memory所需要的時間（即latency）是非常長的，通常是數百個cycles. 如果我們的程式只有乙個thread，所以每次它讀取global memory的內容，就要等到實際讀取到資料、累加到sum之後，才能進行下一步，這就是為什麼如果採取乙個執行緒搞的話表現會很差。

由於global memory並沒有cache，所以要避免開巨大的latency的方法，就是要利用大量的threads。假設現在又大量的threads在同時執行，那麼當乙個thread讀取記憶體，開始等待結果的時候，gpu就可以立刻切換到下乙個thread，並讀取下乙個記憶體位置。因此，理想上當thread的數目夠多的時候，就可以完全把global memory的巨大latency隱藏起來了。

想一下如果把計算平方和的程式並行換額？其中的乙個比較好的方法是把數字分成若干組，把各組數字分別計算平方和後，最好再把每組的和加總起來就可以了，所以，一開始，我們可以把最後加總的動作，由cpu來進行。

現有一kernel函式__global__ static void sumofquery(int*, int* clock_t*)

n 第一種方法「偽多執行緒法」

即在主函式中呼叫<<<1, 1, 0>>>sumofquery(…)

即採用乙個block 乙個執行緒序列實現。

for (int i= 0, i < element_size; ++i)

sum += elem[i]

n 第二種採用多執行緒法

即在主函式中呼叫<<<1, thread_size, 0>>>sumofquery

sum += result[i] (i從0到thread_size) //host

for (i = tid * size; i < (tid + 1) * size; ++i) //kernel

存在缺點為記憶體延遲大，因為當第乙個執行緒等到記憶體資料時，第二個執行緒就開始執行，所以資料沒有成塊。

n 第三種採用改進型多執行緒法

size = data_size / thread_num

for (i = tid; i < data_size; i += thread_num)

result[tid] = sum;

此種方法的採用的很大一部分原因在於顯示卡上的記憶體是dram，因此基於它的操作（最有效的訪問方式）是以連續的方式訪問。上乙個程式中，雖然表面看起來是連續儲存記憶體位置，但是在實際上是與想法違背的。

實驗結果發現雖然latency有所減少，但沒有取得明顯效果，主要原因在於thread數目較小，即n個thread最多只能隱藏n個latency。所以，如果增加thread數目，就可以看到更好的效率，例如可以把thread的數量調的更高，但不能超過gpu所能執行的最大執行緒數目。

失之東隅收之桑榆，thread數目過多，同樣會導致cpu端的工作也可能家走。

n 第四種是採用塊的形式

即新增block塊，在cuda中，thread是可以分組的，乙個block中的thread，具有乙個共享的share memory，也可以進行同步工作。不同block之間的thread則不行，在計算平方和這個專案中，則並不太需要進行thread的同步動作，因此我們可以使用多個block來進一步增加thread的數目。

for (i = bid * thread_num + tid; i < data_size;

i += block_num * thread_num)

sum += num[i] * num[i];

result[bid * thread_num + tid] = sum;

n 第五種是採取塊內同步演算法與樹狀演算法，正在看中

工作總結 2010 3 16

找工作總結

2011工作總結

周工作總結

工作總結 2010 3 16

找工作總結

2011工作總結

周工作總結

相關推薦