y-twist cont.

y-twist-test

上次只用一個 thread 算 y-twist 的時候成效不彰,這次便把它拆成一堆 thread blocks 來試試看!!

先前提到要讓 SM 中平行處理的 wraps 數目越多越好,使其盡可能地增進平行化的效能。但在這次實驗的發現,y-twist 的 kernel 在 block size 設定為 512 或是 64 個 threads 去 deform 99 萬個 vertices 時,運算時間的差距其實只在 1 ms 左右。所以接下來便以每個 block  含 512 個 threads 為前提,將 memory access 的部分作為變因,再繼續作一些小測試。

將上次 kernel 中的 for loop 展開成 1D 的 grid ,每個 grid cell 含有一個 2D block 會變成下面這個新的 kernel:

__global__ void gpu_twist(float4 pts[], float magnitude, float env)
{
	int k = blockIdx.x * BLOCK_WIDTH * BLOCK_WIDTH + threadIdx.y * BLOCK_WIDTH + threadIdx.x;

	float ff = magnitude * pts[k].y * env;
	if( ff != 0.0f )
	{
		double cct = __cosf(ff);
		double cst = __sinf(ff);
		double tt = pts[k].x * cct - pts[k].z * cst;

		pts[k].z = pts[k].x * cst + pts[k].z * cct;
		pts[k].x = tt;
	}
}

拆成多個 threads 之後,運算時間便從 1620 ms 大幅降為 19 ms,而對照組 CPU 版本的計算時間則是大約在 120 ms 左右(可見在 GPU 只用一個 thread 作運算有多麼搞笑)。現在 kernel 裡的變數除了幾個是 automatic variable 之外,其餘都是來自 global memory。因為 y-twist 的運算單純到沒變數需要 share,所以接著要作的事便是把一些多次存取 global memory 的部分轉至 automatic variable 中。不過我在這過程中犯了一個錯誤:我原先認為 pts[k].y 只被讀取一次,若將其放進 automatic variable 中,反倒多浪費一個 register,似乎不太划算,因此我最初只把重複存取的 pts[k].x 和 .z  放進 automatic variable 中:

float ptX = pts[k].x;
float ptZ = pts[k].z;

雖然這樣也可獲得一些成效,使計算時間降至 15.7 ms 左右,但後來在翻閱 programming guide 時發現了下面這一段話:

the device is capable of reading 32-bit, 64-bit, or 128-bit words from global memory into registers in a single instruction.

這代表直接把 pts 讀進 float4 其實比拆成兩個要有效率:

float4 pt = pts[k];

另一個重點是 kernel automatic variable 的個數其實並不直接等於 register 的個數!! 雖然它一般都會被放進 register 裡頭,但若遇到比較複雜的 structure 或是 array 時,這些變數是會放至 local memory 裡的 (其存取的 cost 和存取 global memory 是差不多的)。但這些資訊很難直接從程式碼中推測出來,有時稍微變更一下 expression 的順序,便會影響 register 的使用數目。若需精確查詢 kernel 裡面 register 的使用數目或是 shared/local/const memory 的使用量,有以下兩種方式:

  • PtxAsOptionV 設為 Yes (–ptxas-options=-v)
    之後在 build 的時候,便可以在 output panel 看到每個 kernel 的資源使用狀況。
  • 開啟 keep preprocessed files
    相關資訊都寫在 .cubin 檔案中每個 code 區塊的標頭裡。

此外, SDK 也提供了一個很實用的 excel 檔案 — “CUDA_Occupancy_calculator”。只要先選擇現在使用的 compute capability,並將剛剛那些查詢得來的數值一一輸入相對應的欄位,便可掌握目前 SIMT 配置 threads 的概況。

結果把 pts 的讀取方式作修改之後,每個 kernel 中使用的 register 數目雖由 6 個增為 8 個,但此數值仍是在安全範圍內。最後計算的時間則是再快個 1 ms,時間約為 13.7 ms。哈! Hello World 差不多告一段落了,該是找實戰練習題的時候了!! 🙂

Advertisements

Leave a Reply

Fill in your details below or click an icon to log in:

WordPress.com Logo

You are commenting using your WordPress.com account. Log Out / Change )

Twitter picture

You are commenting using your Twitter account. Log Out / Change )

Facebook photo

You are commenting using your Facebook account. Log Out / Change )

Google+ photo

You are commenting using your Google+ account. Log Out / Change )

Connecting to %s