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