OpenCL 學習step by step (10) 矩陣轉置

      本章學習一下在opencl中如何實作矩陣的轉置，主要的技巧還是利用好local memory，防止bank conflit以及使得全局memory的讀寫盡量是合并(coalensing)讀寫。

     我們的矩陣是一副二維灰階圖像256*256，矩陣的轉置也就是圖像的轉置。每個thread處理16(4*4)個pixel(uchar)，workgroup的size是(16,16)。

     下面直接看shader代碼：

uint wiWidth  = get_global_size(0);

uint gix_t = get_group_id(0);

uint giy_t = get_group_id(1);   

uint num_of_blocks_x = get_num_groups(0);

uint giy = gix_t;

uint gix = (gix_t+giy_t)%num_of_blocks_x;

uint lix = get_local_id(0);

uint liy = get_local_id(1);

uint blockSize = get_local_size(0);

uint ix = gix*blockSize + lix;

uint iy = giy*blockSize + liy;

int index_in = ix + (iy)*wiWidth*4;

// 通過合并讀寫把輸入資料裝入到lds中

int ind = liy*blockSize*4+lix;

block[ind]        = input[index_in];

block[ind+blockSize]    = input[index_in+wiWidth];

block[ind+blockSize*2] = input[index_in+wiWidth*2];

block[ind+blockSize*3] = input[index_in+wiWidth*3];

     因為workgroup size是（16,16），是以lix,liy的取值範圍都是0-15，下面我們通過圖檔看下，lix=0 liy=0,lix=1 liy=0時候，ind,以及index_in的值，進而得到輸入圖像資料如何映射到local memory中。

lix=0 liy=0

lix=1 liy=0

      下面是影射關系，(0,0) thread處理的16個pixel用血紅色表示，它們映射到lds的0bank和16bank，(1,0)thread處理的像素用綠色表示，它們映射到lds的bank1和bank17，有效的避免了bank conflit，而全局memory的通路不同thread對應連續的全局memory空間，可以實作合并讀寫，進而提高程式性能。

  把轉置的資料寫到全局memory中的代碼如下：

ix = giy*blockSize + lix;

iy = gix*blockSize + liy;

int index_out = ix + (iy)*wiWidth*4;

ind = lix*blockSize*4+liy;

uchar4 v0 = block[ind];

uchar4 v1 = block[ind+blockSize];

uchar4 v2 = block[ind+blockSize*2];

uchar4 v3 = block[ind+blockSize*3];

// 通過合并讀寫把lds中資料寫回到全局memory中

output[index_out]            = (uchar4)(v0.x, v1.x, v2.x, v3.x);

output[index_out+wiWidth]    = (uchar4)(v0.y, v1.y, v2.y, v3.y);

output[index_out+wiWidth*2]    = (uchar4)(v0.z, v1.z, v2.z, v3.z);

output[index_out+wiWidth*3]    = (uchar4)(v0.w, v1.w, v2.w, v3.w);

對應copy關系圖如下：

完整的代碼請參考：

工程檔案gclTutorial9

代碼下載下傳：

稍後提供