在使用CUDA shared memory常會使用__syncthread()來避免資料不同步處理造成的破壞。問題是一般人無法判定何時會產生這個問題,所以會先加以做為保險。
可是加了之後確實會影響執行速度,Bee這邊測到的就差了四倍多。
所以要適當加入__syncthread()才是合理。
剛好之前bee使用CPU及GPU同時開發程式的方法,因為一邊為single-thread另一個為multi-thread,所以可以保證single-thread是正確的。剛好可以比對結果。
若是不加__syncthread()會使結果比對產生錯誤,那就要加。沒有錯誤就照CPU版本的直接移入GPU版本。
這種情況玩個幾次,基本上就可以知道何時要加,何時不要加。
2009年11月30日 星期一
2009年11月19日 星期四
第二個CUDA程式,使用Shared memory
因為有共同讀取的來源資料,所以先行載入shared memory。然後經過運算,之後回存。
因為Windows下的CUDA沒有除錯器,所以先發展CPU下可以執行的GPU架構程式。先將變數之間的關係弄清楚再移進GPU。
至於shared memory也是一種區域性變數,所以用local variable來模擬。
CPU版本程式:
void CPU_mask_rotate(unsigned int *src, unsigned int *dest, int gridDim_x, int gridDim_y)
{
unsigned int idata[HEAD_LENGTH];
for(int blockIdx_y=0; blockIdx_y < gridDim_y ; blockIdx_y++){
for(int blockIdx_x=0; blockIdx_x < gridDim_x; blockIdx_x++){
// Read to local var.
for(int i=0; i<HEAD_LENGTH; i++){
int t = (blockIdx_y*HEAD_LENGTH+i)*gridDim_x + blockIdx_x;
idata[i] = src[t];
}
// convert format
for(int threadIdx_x=0; threadIdx_x<TRANSPOSE_SIZE; threadIdx_x++){
for(int j=0; j<HEAD_BLOCK; j++){
unsigned int r;
r=0;
for(int k=0; k<TRANSPOSE_SIZE; k++){
unsigned int u,v;
u = idata[j*TRANSPOSE_SIZE + k];
v = (u>>threadIdx_x & 0x1)<<k;
r |= v;
}
int t = (blockIdx_y * gridDim_x + blockIdx_x)*HEAD_LENGTH + threadIdx_x*HEAD_BLOCK + j;
dest[t] = r;
}
}
}
}
}
GPU版本程式
__global__ void Kernelmask_rotate(unsigned int *src, unsigned int *dest)
{
__shared__ unsigned int idata[HEAD_LENGTH];
for(int i=0; i<HEAD_LENGTH; i++){
int t = (blockIdx.y * HEAD_LENGTH+i)*gridDim.x + blockIdx.x;
idata[i] = src[t];
}
__syncthreads();
for(int j=0; j<HEAD_BLOCK; j++){
unsigned int r;
r=0;
for(int k=0; k<TRANSPOSE_SIZE; k++){
unsigned int u,v;
u = idata[j*TRANSPOSE_SIZE + k];
v = (u>>threadIdx.x & 0x1)<<k;
r |= v;
}
int t = (blockIdx.y * gridDim.x + blockIdx.x)*HEAD_LENGTH + threadIdx.x*HEAD_BLOCK + j;
dest[t] = r;
}
}
CPU寫好時,就改一下變數名( _改為. )就可以copy到GPU的版本。
只是在發展中發現一個奇怪現象。
因為讀入shared memory可以使各thread少去對global memory的讀取,而減少執行時間。
但在寫程式的時候,去算各位址不好判定,所以先用一塊shared memory存結果,再從shared memory存回global memory,因為這樣比較好寫。
後來發現CUDA SDK內的example都是直接回存。這樣可以省去shared memory,不過就是不好寫。
因為Windows下的CUDA沒有除錯器,所以先發展CPU下可以執行的GPU架構程式。先將變數之間的關係弄清楚再移進GPU。
至於shared memory也是一種區域性變數,所以用local variable來模擬。
CPU版本程式:
void CPU_mask_rotate(unsigned int *src, unsigned int *dest, int gridDim_x, int gridDim_y)
{
unsigned int idata[HEAD_LENGTH];
for(int blockIdx_y=0; blockIdx_y < gridDim_y ; blockIdx_y++){
for(int blockIdx_x=0; blockIdx_x < gridDim_x; blockIdx_x++){
// Read to local var.
for(int i=0; i<HEAD_LENGTH; i++){
int t = (blockIdx_y*HEAD_LENGTH+i)*gridDim_x + blockIdx_x;
idata[i] = src[t];
}
// convert format
for(int threadIdx_x=0; threadIdx_x<TRANSPOSE_SIZE; threadIdx_x++){
for(int j=0; j<HEAD_BLOCK; j++){
unsigned int r;
r=0;
for(int k=0; k<TRANSPOSE_SIZE; k++){
unsigned int u,v;
u = idata[j*TRANSPOSE_SIZE + k];
v = (u>>threadIdx_x & 0x1)<<k;
r |= v;
}
int t = (blockIdx_y * gridDim_x + blockIdx_x)*HEAD_LENGTH + threadIdx_x*HEAD_BLOCK + j;
dest[t] = r;
}
}
}
}
}
GPU版本程式
__global__ void Kernelmask_rotate(unsigned int *src, unsigned int *dest)
{
__shared__ unsigned int idata[HEAD_LENGTH];
for(int i=0; i<HEAD_LENGTH; i++){
int t = (blockIdx.y * HEAD_LENGTH+i)*gridDim.x + blockIdx.x;
idata[i] = src[t];
}
__syncthreads();
for(int j=0; j<HEAD_BLOCK; j++){
unsigned int r;
r=0;
for(int k=0; k<TRANSPOSE_SIZE; k++){
unsigned int u,v;
u = idata[j*TRANSPOSE_SIZE + k];
v = (u>>threadIdx.x & 0x1)<<k;
r |= v;
}
int t = (blockIdx.y * gridDim.x + blockIdx.x)*HEAD_LENGTH + threadIdx.x*HEAD_BLOCK + j;
dest[t] = r;
}
}
CPU寫好時,就改一下變數名( _改為. )就可以copy到GPU的版本。
只是在發展中發現一個奇怪現象。
因為讀入shared memory可以使各thread少去對global memory的讀取,而減少執行時間。
但在寫程式的時候,去算各位址不好判定,所以先用一塊shared memory存結果,再從shared memory存回global memory,因為這樣比較好寫。
後來發現CUDA SDK內的example都是直接回存。這樣可以省去shared memory,不過就是不好寫。
2009年11月6日 星期五
3D Webcam
想當時比CUDA比賽時,為了做webcam解3D。花了時間及錢去買二個Webcam。
結果二個獨立webcam有很多問題。包括顏色不相同,雜訊狀況也不一樣。又要解決不平行問題,結果每次都要進行調整。
不久前看到3D webcam才覺得有點做白工,3D Webcam不見得二個webcam狀況完全相同,但差異性小了許多,且機構的關係是固定的。
一個USB線同時可以進行二個webcam的影像傳輸。不會一次用掉二個USB。
另外還可以用紅藍眼鏡玩玩3D影像。
有了這個,以後要自己試用CUDA解3D要方便多了。
訂閱:
文章 (Atom)