CUDAのバンクコンフリクトを再現・検出してみる

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <stdlib.h>



///////////////////////////////////////////////
// GPU側 ////////////////////////////////////// 

// 32個のスレッドで一回ずつ呼び出される。
__global__ void test_kernel(unsigned int* c) {


    // 32 × 32　のシェアードメモリ確保
    // この時、 mem[0] , mem[32] , mem[64] , ... を同時にアクセスするとバンクコンフリクトする
    __shared__ float mem[32*32];


    //このスレッドの番号を取得
    size_t xpos = blockIdx.x * blockDim.x + threadIdx.x;
    
    // シェアードメモリに書き込み xposが0,1,2,3...なので、0,32,64,...（＝同じバンク）に書き込む
    mem[xpos*32] = xpos;

    // これを入れないと最適化？されるのか検出できない
    c[xpos] = mem[xpos];

}

////////////////////////////////////////////////////////////
////////////////////////////////////////////////////////////
////////////////////////////////////////////////////////////

int main()
{

    dim3 grid(1, 1);   //グリッド数

    dim3 block(32, 1); //ブロック数

    unsigned int* p_gpu;

    cudaMalloc( (void**)&p_gpu,32*4);

    test_kernel << <grid, block >> > (p_gpu); // 32個のスレッドを実行

    unsigned int p_cpu[32];
    cudaMemcpy(&p_cpu, p_gpu, 32*4, cudaMemcpyDeviceToHost);//GPU側から実行結果を取得

    cudaFree(p_gpu);//GPU側のメモリを解放


    for (size_t i = 0; i < 32; i++)
        printf("%d\n", p_cpu[i]);

    return 0;
}

// 32個のスレッドで一回ずつ呼び出される。
__global__ void test_kernel(unsigned int* c) {


    // 32 × 32　のメモリ
    // この時、 mem[0] , mem[32] , mem[64] , ... を同時にアクセスするとバンクコンフリクトする
    __shared__ float mem[32*32];

    //アクセス法
    //このスレッドが担当する画素の位置を二次元座標で求める
    size_t xpos = blockIdx.x * blockDim.x + threadIdx.x;
    
    // シェアードメモリに書き込み xposが0,1,2,3,... なので、違うバンクに書き込む
    mem[xpos] = xpos;

    // これを入れないと最適化？されるのか検出できない
    c[xpos] = mem[xpos];

}