Flux.2 klein で vae 変換時に ggml 内でエラー

今までもろもろの条件により、vaeもcpuで実行していたのですが、--diffusion-fa指定しなきゃ出力できないなら、vaeもgpuでいいんじゃね？的なことを感じたのでちょっと試したところ、256x256では出力できました。

Flux.2のvaeはsafetensorsでも小さめなので、1024x1024でもいけるのではないかと実行してみました。

実行はできたのですが…

Provided range and/or offset does not fit in int. Pass `-fno-sycl-id-queries-fit-in-int' to remove this limit.Exception caught at file:/home/siriuth/stable-diffusion.cpp/ggml/src/ggml-sycl/ggml-sycl.cpp, line:4760
Error OP IM2COL

ggml内でエラーが起きてしまったようです…。

前に試したときはメモリーが足りなくてダメだった記憶がありますが今度は致命的な部分て死んでいるようです…って、行番号に記憶があります…

見てみると例外処理の中の行で

} catch (sycl::exception & e) {
    std::cerr << e.what() << "Exception caught at file:" << __FILE__ << ", line:" << __LINE__ << std::endl;
    std::cerr << "Error OP "<<ggml_op_name(dst->op)<< std::endl;
    std::exit(1);
}

の中でした。場所はggml_sycl_compute_forward関数内の例外なので、演算時に何らかの例外が発生してしまったようです。CPUだと問題ないので、おそらくこれはSYCL固有のバグの可能性が高いですね。

 

(追記 19:00)

geminiに相談してみてソースの確認を行って、手を加えてみたけど例外は収まらず。

ソース自体はCUDAのコードをoneAPIの変換ツールで変換されたコードの様でツール自体のコメントを見ると

DPCT1049:73: The work-group size passed to the SYCL kernel may exceed the limit. To get the device limit, query info::device::max_work_group_size. Adjust the work-group size if needed. 

DPCT1049:73: SYCLカーネルに渡されたワークグループサイズが制限を超えている可能性があります。デバイスの制限値を取得するには、info::device::max_work_group_sizeを照会してください。必要に応じてワークグループサイズを調整してください。 

（グーグル翻訳）

だ、そうで、結局この「デバイスの制限値」を超えてしまっているのが原因の様です。

実際の値はまだ確認していませんが、ggml側の処理は上限の制限(65535)は行っています。SYCLのコードとしてはこの値の乗算を行っているので、デバイスの上限値を超えた場合処理を分割する必要があるという事で…メンドクサイｗ

と言うか、CUDAのコードとoneAPIのSYCLのコード比べるとCUDAのコードがすごいスッキリしてそうなんですがｗ

と、実際に手を入れ始めると、SYCL_IM2COL_BLOCK_SIZE(256)で一応見てるみたいです。内蔵GPUがもっと小さい可能性がありますが、それよりも現実はさらに乗算しているので実際のブロックサイズを超えてしまうのでロジック自体破綻してるんで…ダメですねこれｗ