CUDAコンパイラの最適化バグ?
狙ってやっているわけではないのですが、四半期ぶりの日記です。
ここ半年ばかりGPGPUを使った研究に取り組んでいて、なかなかすばらしい性能を出してくれています。
ところが、ある日コードを書いていると突然コンパイラが落ちるようになりました。
2>Compiling with CUDA Build Rule... 2>"C:\CUDA\bin\nvcc.exe" ... ... 2>nvopencc ERROR: C:\CUDA\bin/../open64/lib//be.exe returned non-zero status -1073741819
その時は適当にコードを変えていたら出なくなったのですが、その後再発し開発が進まなくなってしまいました。
というわけで簡単に調べたメモ。
まず計算自体はShared Memory上のデータとConstant Memory上のフィルタの畳み込み積分を計算するもので、コードにするとこんな感じになります(実際はもう少し複雑です)。
float out[2][2];
out[0][0] = out[0][1] = ... = 0;
#pragma unroll
for(uint x = 0; x < SIZE; x++){
#pragma unroll
for(uint p = 0; p < 2; p++){
out[v][p] += smem[threadIdx.x+x]*cmem[v][x][p];
}
}
}問題が起きたのは、ここでoutの要素数やSIZEを増やした時。
CUDAではローカル変数はLocal Memoryというところに保存されるのですが、レジスタが空いている時はコンパイラが変数をレジスタに割り当ててくれます。
ただし配列の場合は、ループアンロールを行うなどしてアクセスする要素をコンパイル時に決定できるようにしないとレジスタ割当が行われません(ただの経験則です、念のため)。
そして問題のコンパイルエラーは配列の要素数を増やしたときもしくはSIZEを大きくしたときに発生しました。詳しくは解析していませんが、どうやら、配列への命令がある一定数(数千〜?)に達するとコンパイルできなくなるようです。
試行錯誤した結果、配列を二つに分割することで解決しました。先ほどの例でいえば
float out0[2], out1[2];
out0[0] = out1[0] = ... = 0;
#pragma unroll
for(uint x = 0; x < SIZE; x++){
out0[v] += smem[threadIdx.x+x]*cmem[v][x];
out1[v] += smem[threadIdx.x+x]*cmem[v][x];
}
}と、一部の配列を手動でアンロールして別々の配列に割り当てることで解決。
