AMD OpenCL コンパイラにはバグがある(Version 2348.3, 2442.8)

AMD OpenCL でとあるカーネルをコンパイルしようとしたところ

LLVM ERROR: cannot lower memory intrinsic in address space 1

とだけ出力してコンパイラがエラーを出すという現象に会いました.
(Linux(amdgpu-pro) ver 2348.3 および Windows ドライバ 2442.8 で確認しました. 2017/09/15 での概ね最新 AMD GPU ドライバ)

いろいろカーネルコードを取捨選択してデバッグしてみたところ, global memory から大きめの構造体をローカル変数領域にコピーすることで発生することがわかりました.

以下のような感じです.

typedef struct {
  float large_data[4096];
} MyStruct;

__kernel void main(__global MyStruct *s)
{
  MyStruct local_s = *s;
}

で, たぶん原因はこれっぽそうです.

Partially fix memcpy / memset / memmove lowering in SelectionDAG construction if address space != 0.
https://reviews.llvm.org/D7241

x86 でも大きめの構造体をローカル変数領域(スタック領域)に持つとダメというのは昔ありましたが, なにかしらもうちょっと原因がわかるようにレポートしてほしいところ.

AMD OpenCL コンパイラは, フロントエンド周りはしばらくアップデートされていないのかな?

回避策

• 構造体の大きさを減らすように設計しなおす.
• 自前で struct をコピーする関数を用意する.