如何了解__syncthread()的使用時機：試誤法

在使用CUDA shared memory常會使用__syncthread()來避免資料不同步處理造成的破壞。問題是一般人無法判定何時會產生這個問題，所以會先加以做為保險。

可是加了之後確實會影響執行速度，Bee這邊測到的就差了四倍多。
所以要適當加入__syncthread()才是合理。

剛好之前bee使用CPU及GPU同時開發程式的方法，因為一邊為single-thread另一個為multi-thread，所以可以保證single-thread是正確的。剛好可以比對結果。

若是不加__syncthread()會使結果比對產生錯誤，那就要加。沒有錯誤就照CPU版本的直接移入GPU版本。

這種情況玩個幾次，基本上就可以知道何時要加，何時不要加。

第二個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，不過就是不好寫。

3D Webcam

想當時比CUDA比賽時，為了做webcam解3D。花了時間及錢去買二個Webcam。

結果二個獨立webcam有很多問題。包括顏色不相同，雜訊狀況也不一樣。又要解決不平行問題，結果每次都要進行調整。

不久前看到3D webcam才覺得有點做白工，3D Webcam不見得二個webcam狀況完全相同，但差異性小了許多，且機構的關係是固定的。

一個USB線同時可以進行二個webcam的影像傳輸。不會一次用掉二個USB。

另外還可以用紅藍眼鏡玩玩3D影像。
有了這個，以後要自己試用CUDA解3D要方便多了。