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

代碼下載：

稍後提供