SharedメモリのAlignment調整

目次

問題と解決

WMMA APIなどではメモリのAlignment制約が結構厳しい.
cudaMallocしたGlobalメモリはAlignmentをよしなに調整してくれる (512bitにおいてくれている気がする).
一方Sharedメモリはというと、例えば
__global__ void kernel(/*float* a*/){
        int tid = threadIdx.x;
        // a[tid] = 1.0f;
        __shared__ float s0;
        __shared__ float s1;
        if(tid == 0){
                printf("%p,%p\n", &s0, &s1);
        }
}
というカーネルの実行結果は 0x7fc141000000,0x7fc141000004 となり、Alignmentの調整は行われない (Sharedメモリは小容量だし頷ける).
そこでSharedメモリでAlignmentを制約したい場合どうすればいいかというと、
__shared__ __align__(16) float s1;
とすればいい.
カテゴリー:CUDA
記事作成日:2018-09-03