- 304 名前：デフォルトの名無しさん mailto:sage [2012/01/10(火) 18:12:32.43 ]
- >>303
コンパイル後のバイナリを見たわけじゃないから予測だが、
s_v[threadIdx.x] = g_v[i+i%3];
x = s_v[(threadIdx.x+3)%32];
g_x[i] = x;
は
1.1:tmp=g_v[i+i%3]; //グローバルアクセスなので遅い
1.2:s_v[threadIdx.x]=tmp; //
2.1:x = s_v[(threadIdx.x+3)%32];
3.1:g_x[i] = x;
みたいに解釈される。
で、単純な式だからコンパイラは
s_v[threadIdx.x]とs_v[(threadIdx.x+3)%32];
が必ず別のアドレスになる事を検知して
1.1 2.1 3.1 1.2のように命令を入れ替えてしまう。
だから__threadfence_block()がいるはず。
