ggml SYCL 実際に留まっている場所

単純に出力ログを追加したり、エラーになるように実行したりログ収集から始めました。

/ggml/src/ggml-sycl/ggml-sycl.cpp
4535 static void ggml_backend_sycl_synchronize(ggml_backend_t backend) try {
4536     GGML_SYCL_DEBUG("[SYCL] call %s begin\n", __func__);
4537     ggml_backend_sycl_context * sycl_ctx = (ggml_backend_sycl_context *)backend->context;
4538     const queue_ptr stream = sycl_ctx->stream(sycl_ctx->device, 0);
4539     GGML_SYCL_DEBUG("[SYCL] call %s wait()\n", __func__);
4540     SYCL_CHECK(CHECK_TRY_ERROR((stream)->wait()));
4541
4542     GGML_SYCL_DEBUG("[SYCL] call %s GGM_UNUSED\n", __func__);
4543     GGML_UNUSED(backend);
4544     GGML_SYCL_DEBUG("[SYCL] call %s end\n", __func__);
4545 }
4546 catch (sycl::exception const &exc) {
4547   std::cerr << exc.what() << "Exception caught at file:" << __FILE__
4548             << ", line:" << __LINE__ << std::endl;
4549   std::exit(1);
4550 }
 

この結果、出力は

[SYCL][OP] call ggml_sycl_cpy done
[SYCL][OP] call ggml_sycl_dup done
[SYCL] call ggml_backend_sycl_synchronize begin
[SYCL] call ggml_backend_sycl_synchronize wait()

となった。具体的に止まったのは、

4540     SYCL_CHECK(CHECK_TRY_ERROR((stream)->wait()));

という事になりました。さて、これって単純な同期処理だと思うが…詳しく見てみましょうかね。

ログを取り直してみたところ上位のstable-diffusion.cppでは 

[INFO ] stable-diffusion.cpp:3401 - generating image: 1/1 - seed 42
[DEBUG] ggml_extend.hpp:1924 - flux compute buffer size: 2361.25 MB(VRAM)

と言うメッセージのあと息をしなくなるので、もう少しデバッグ出力を追加してどこで止まっているのかはっきりさせてみました。

[INFO ] stable-diffusion.cpp:3168 - get_learned_condition completed, taking 26.51s
[INFO ] stable-diffusion.cpp:3401 - generating image: 1/1 - seed 42
[INFO ] stable-diffusion.cpp:3404 - generating image: manual_seed()
[INFO ] stable-diffusion.cpp:3406 - generating image: noise
[INFO ] stable-diffusion.cpp:3409 - generating image: x_0
[DEBUG] ggml_extend.hpp:1924 - flux compute buffer size: 973.07 MB(VRAM)

デバッグ出力で行をズラしたので行数はずれていますが、3407行の

sd::Tensor<float> x_0 = sd_ctx->sd->sample(sd_ctx->sd->diffusion_model, ～

の呼び出しで止まっていました。ループの中ですが、画像出力はこの呼び出しで画像を取得している感じです。 結局、ggml内部でハングしている感じに見えますね。

気になるのでちょっと調べてみる。

ggml-sycl/common.hpp
typedef sycl::queue *queue_ptr;

気になるのはソース内で結構揺らぎがあってコンテキスト変数ctxのstreamメソッドの扱いが微妙に揺れてるんです。あと、stream->wait()だったり(stream)->wait()だったり。他にも記述揺れが…。

実際のところキューも幾つか存在し並列動作を前提に組まれているのですが、それも揺れてる感じを受けます。

１日かけてわからないながらも、ざっくりと見た感じどこかの処理で対象となる終了を待たずにスルーしちゃっている部分があるんじゃないかなぁというのが感触としてあるわけですが、いきなり見てもわからんわなｗｗ

まぁ、今までメモリ不足で死んでる可能性を気にしていましたが、現状としてはハッキリと同期処理で落ちているというのが分かっただけでもすっきりしたり。

デバッグログを出す為に、わざとエラーになるパターンで実行してたら、ログ出力で処理が遅くなったためにflux1-schnellで512x512ですがStep20通っちゃったというハプニングがあったり…ｗ

しっかし思っていた以上にソースが見れるような環境はやっぱり面白いですね。歯がゆいですがｗパッケージだと肝心な部分が結構ブラックボックスでどのように実装しているかは推定するしかないので…。