問題と解決
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
