技术开发 频道

CUDA Thread Block:transpose

  【IT168 文档】在 Heresy 寫的前兩篇 sample 程式(VectorAdd、DeviceInfo)裡,都是很簡單的程式;像 VectorAdd 裡,也是刻意把 vector size 設小,避掉 thread 數目超過 block 限制的問題,以避免要用到複數個 block。但是實際上,應該都是會超過 thread block 的大小限制的(畢竟 G80 的 block 大小只有到 512…)~

  這一篇主要打算用 CUDA SDK 的範例程式 transpose(一般會在 C:\Program Files\NVIDIA Corporation\NVIDIA CUDA SDK\projects\transpose)來針對 CUDA 中,將 thread 切割成數個 thread block 的處理做個簡單的整哩,並對 device 上的 shared memory 來做一些簡單的研究。

  專案簡介

  在此專案的 Header 檔裡,對此專案的說明如下:

  Matrix transpose with Cuda

  This example transposes arbitrary-size matrices. It compares a naive transpose kernel that suffers from non-coalesced writes, to an optimized transpose with fully coalesced memory access and no bank conflicts. On a G80 GPU, the optimized transpose can be more than 10x faster for large matrices.

  簡單講,這個範例是在進行矩陣轉置的計算。而他的比較,是透過 shared memory 來對轉置時記憶體的存取動作非常好的化,讓存取由 non-coalesced 寫入變成 fully coalesced 存取;在 G80 上,這樣可以獲得十倍的效率增長~(Heresy 自己測試沒有增加那麼多就是了)

  首先,這個專案有三個檔案:

  • transpose.cu

  是最主要的檔案,main 和 runTest 這兩個主要的函式都在這個檔案。

  • transpose_kernel.cu

  轉置矩陣的 device kernel 程式。transpose_naive 是最簡單的寫法,transpose 則是透過 shared memory 來做非常好的化的方法。

  • transpose_kernel.cu

  轉置矩陣的 CPU 計算函式 computeGold。主要是用來比對 device 計算出來的結果的正確性。

  在 runTest 中,他會先宣告一個 size_x * size_y (實際設定是 256 * 4096)的一維 float 陣列 h_idata 來代表一個 size_x * size_y 的二維矩陣;然後再將一些值填進去,並把資料複製到 device memory d_idata。而他的 grid 和 block 的大小,則是使用

dim3 grid(size_x / BLOCK_DIM, size_y / BLOCK_DIM, 1);dim3 threads(BLOCK_DIM, BLOCK_DIM, 1);

   的設定。而由於 BLOCK_DIM 是定義成 16,所以再搭配 256 * 4096 來看,實際上執行的時候,是會把總共需要執行的 1048576(256*4096) 個 thread,在 grid 中分成 16 * 256 個 thread block,而每個 block 裡有16*16 個 thread。而執行的呼叫方法,就是

transpose<<< grid, threads >>>(d_odata, d_idata, size_x, size_y);
0
相关文章