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;

とすればいい.