IGC 2.41.5: "Internal Compiler Error: Floating point exception" on a SPV_KHR_cooperative_matrix (SYCL joint_matrix, fp16 x fp16 -> fp32) kernel for mtl-u / arl-s / dg2 (a diagnostic expected)

オープン
#475 コメント 0 件 リアクション 0 件 担当者 0 名 GitHub で見る

まだ誰も着手していません。

評価

難易度
4/5
見積もり時間
3〜5日
初心者へのやさしさ
49/100
issue の種類
バグ
明瞭さ
おおむね明確
活発さ
活発
技術スタック
cpp
領域
compilers

調査の方向性

Start by reproducing the offline SPIR-V compilation with the supplied image and ocloc compile -spirv_input -file img.0 -device mtl-u; compare the result with bmg-g31 and tgllp. The issue does not name compiler source files or tests, so trace the failing cooperative matrix handling in IGC and locate the relevant diagnostics. Done when unsupported matrix input produces a clean diagnostic instead of an internal compiler error.

索引モデルが issue の本文から書いたものです。

説明

A small SYCL kernel that uses joint_matrix (fp16 a 8x16, fp16 b 16x16, fp32 accumulator 8x16, joint_matrix_mad) compiles and runs on Battlemage. Compiling the same SPIR-V offline for platforms that cannot run that combination crashes the compiler instead of reporting an error.

$ ocloc compile -spirv_input -file img.0 -device mtl-u
Compilation from IR - skipping loading of FCL
[mtl-u] IGC: Internal Compiler Error: Floating point exception
Build failed with error code: -11
-device Result
bmg-g31 Build succeeded
lnl-m Build succeeded
mtl-u Internal Compiler Error: Floating point exception
arl-s Internal Compiler Error: Floating point exception
dg2-g10 Internal Compiler Error: Floating point exception
tgllp clean error: input SPIR-V module uses extension 'SPV_KHR_cooperative_matrix' which were disabled by --spirv-ext option

Expected: a diagnostic like the Tiger Lake one (unsupported extension, or an unsupported matrix type combination for the platform), so a caller can catch it and fall back. On the JIT path the SYCL runtime refuses first ("no matrix hardware on the target device, joint_matrix is not supported"), but ahead-of-time compilation, and any tool that walks a binary's device images, hits the crash.

Reproduce
icpx -fsycl -O2 -o repro repro.cpp
clang-offload-extract --stem=img repro        # writes img.0 (2,080 bytes, also below as base64)
ocloc compile -spirv_input -file img.0 -device mtl-u
repro.cpp (38 lines)
// Minimal reproducer: a SYCL kernel that uses joint_matrix (fp16 x fp16 -> fp32, 8x16 . 16x16), compiled to SPIR-V by icpx,
// then compiled offline with ocloc for a GPU with no matrix hardware (Xe-LPG: -device mtl-u).
// Expected: a diagnostic ("matrix extensions are not supported on this device"). Observed: an internal compiler error.
//   icpx -fsycl -O2 -o repro repro.cpp
//   clang-offload-extract --stem=img repro
//   ocloc compile -spirv_input -file img.0 -device mtl-u     # crashes;  -device bmg-g31 compiles
#include <sycl/sycl.hpp>
#include <sycl/ext/oneapi/matrix/matrix.hpp>
namespace mat = sycl::ext::oneapi::experimental::matrix;
int main() {
    sycl::queue q;
    constexpr size_t TM = 8, TN = 16, TK = 16;
    sycl::half* A = sycl::malloc_shared<sycl::half>(TM * TK, q);
    sycl::half* B = sycl::malloc_shared<sycl::half>(TK * TN, q);
    float* C = sycl::malloc_shared<float>(TM * TN, q);
    for (size_t i = 0; i < TM * TK; ++i) A[i] = sycl::half(1.f);
    for (size_t i = 0; i < TK * TN; ++i) B[i] = sycl::half(2.f);
    q.submit([&](sycl::handler& h) {
        sycl::local_accessor<sycl::half, 1> As(TM * TK, h), Bs(TK * TN, h);
        h.parallel_for(sycl::nd_range<1>(16, 16), [=](sycl::nd_item<1> it) [[sycl::reqd_sub_group_size(16)]] {
            const size_t l = it.get_local_id(0);
            for (size_t i = l; i < TM * TK; i += 16) As[i] = A[i];
            for (size_t i = l; i < TK * TN; i += 16) Bs[i] = B[i];
            sycl::group_barrier(it.get_group());
            auto sg = it.get_sub_group();
            mat::joint_matrix<sycl::sub_group, float, mat::use::accumulator, TM, TN> acc;
            mat::joint_matrix_fill(sg, acc, 0.0f);
            mat::joint_matrix<sycl::sub_group, sycl::half, mat::use::a, TM, TK, mat::layout::row_major> a;
            mat::joint_matrix<sycl::sub_group, sycl::half, mat::use::b, TK, TN, mat::layout::row_major> b;
            mat::joint_matrix_load(sg, a, As.get_multi_ptr<sycl::access::decorated::no>(), TK);
            mat::joint_matrix_load(sg, b, Bs.get_multi_ptr<sycl::access::decorated::no>(), TN);
            mat::joint_matrix_mad(sg, acc, a, b, acc);
            mat::joint_matrix_store(sg, acc, sycl::address_space_cast<sycl::access::address_space::global_space,
                                                                      sycl::access::decorated::no>(C), TN, mat::layout::row_major);
        });
    }).wait();
    return C[0] == 32.f ? 0 : 1;
}
img.0 (SPIR-V, base64; base64 -d > img.0)
AwIjBwAEAQAOAAYATQAAAAAAAAARAAIABAAAABEAAgAFAAAAEQACAAYAAAARAAIACAAAABEAAgALAAAAEQACACYAAAARAAIAJwAA
ABEAAgA6AAAAEQACAP0VAAARAAIAhhcAAAoACABTUFZfS0hSX2Nvb3BlcmF0aXZlX21hdHJpeAAACgAHAFNQVl9LSFJfZXhwZWN0
X2Fzc3VtZQAAAAsABQABAAAAT3BlbkNMLnN0ZAAADgADAAIAAAACAAAADwAWAAYAAAAOAAAAX1pUU1paNG1haW5FTktVbFJONHN5
Y2wzX1YxN2hhbmRsZXJFRV9jbEVTMl9FVWxOUzBfN25kX2l0ZW1JTGkxRUVFRV8AAAAABQAAABAABAAOAAAAIwAAABAAAAADAAMA
BAAAAKCGAQAFAAsABQAAAF9fc3BpcnZfQnVpbHRJbkxvY2FsSW52b2NhdGlvbklkAAAAAAUADQAHAAAAY2xhc3Muc3ljbDo6X1Yx
OjpkZXRhaWw6OmhhbGZfaW1wbDo6aGFsZgAAAAAFAAUAFgAAAC5wcmVoZWFkZXIAAEcADQAFAAAAKQAAAF9fc3BpcnZfQnVpbHRJ
bkxvY2FsSW52b2NhdGlvbklkAAAAAAEAAABHAAMABQAAABYAAABHAAQABQAAAAsAAAAbAAAARwAEAA8AAAAsAAAAAgAAAEcABAAQ
AAAALAAAAAIAAABHAAQAEQAAACwAAAACAAAARwAEABIAAAAsAAAAAgAAAEcABAATAAAALAAAAAQAAABHAAMAIwAAAHURAABHAAMA
IwAAAHYRAABHAAMAMAAAAHURAABHAAMAMAAAAHYRAAAVAAQAAgAAAEAAAAAAAAAAFQAEAAsAAAAIAAAAAAAAABUABAApAAAAIAAA
AAAAAAArAAUAAgAAAB8AAAAAAAAAAAAAACsABQACAAAAIQAAAAAAAIAAAAAAKwAFAAIAAAAlAAAAgAAAAAAAAAArAAQAKQAAACoA
AAAAAAAAKwAFAAIAAAAuAAAAEAAAAAAAAAArAAUAAgAAADIAAAAAAQAAAAAAACsABAApAAAAOwAAAAIAAAArAAQAKQAAADwAAACQ
AwAAKwAEACkAAAA/AAAAAwAAACsABAApAAAAQAAAAAgAAAArAAQAKQAAAEEAAAAQAAAAKwAEACkAAABGAAAAAQAAABcABAADAAAA
AgAAAAMAAAAgAAQABAAAAAEAAAADAAAAEwACAAYAAAAWAAMACAAAABAAAAAeAAMABwAAAAgAAAAgAAQACQAAAAQAAAAHAAAAIAAE
AAoAAAAFAAAACAAAACAABAAMAAAABQAAAAsAAAAhAAgADQAAAAYAAAAJAAAACgAAAAkAAAAKAAAADAAAABQAAgAdAAAAIAAEACsA
AAAEAAAACAAAACAABAA5AAAACAAAAAsAAAAWAAMAPQAAACAAAABoEQcAQgAAAD0AAAA/AAAAQAAAAEEAAAA7AAAAaBEHAEQAAAAI
AAAAPwAAAEAAAABBAAAAKgAAAGgRBwBHAAAACAAAAD8AAABBAAAAQQAAAEYAAAAgAAQASwAAAAUAAAA9AAAAOwAEAAQAAAAFAAAA
AQAAACkAAwAdAAAAHgAAACsABAA9AAAAPgAAAAAAAAA2AAUABgAAAA4AAAAAAAAADQAAADcAAwAJAAAADwAAADcAAwAKAAAAEAAA
ADcAAwAJAAAAEQAAADcAAwAKAAAAEgAAADcAAwAMAAAAEwAAAPgAAgAUAAAAPQAGAAMAAAAbAAAABQAAAAIAAAAgAAAAUQAFAAIA
AAAcAAAAGwAAAAAAAACpAAYAAgAAACAAAAAeAAAAHAAAAB8AAACwAAUAHQAAACIAAAAgAAAAIQAAAP4VAgAiAAAA+QACABUAAAD4
AAIAFQAAAPUABwACAAAAJAAAACAAAAAUAAAAIwAAABcAAACwAAUAHQAAACYAAAAkAAAAJQAAAPoABAAmAAAAFwAAABYAAAD4AAIA
FgAAAPkAAgAYAAAA+AACABcAAABGAAUACQAAACcAAAAPAAAAJAAAAEYABQAKAAAAKAAAABAAAAAkAAAARgAGACsAAAAsAAAAJwAA
AB8AAAAqAAAAPQAGAAgAAAAtAAAAKAAAAAIAAAACAAAAPgAFACwAAAAtAAAAAgAAAAIAAACAAAUAAgAAACMAAAAkAAAALgAAAPkA
AgAVAAAA+AACABgAAAD1AAcAAgAAADEAAAAwAAAAGQAAACAAAAAWAAAAsAAFAB0AAAAzAAAAMQAAADIAAAD6AAQAMwAAABkAAAAa
AAAA+AACABkAAABGAAUACQAAADQAAAARAAAAMQAAAEYABQAKAAAANQAAABIAAAAxAAAARgAGACsAAAA2AAAANAAAAB8AAAAqAAAA
PQAGAAgAAAA3AAAANQAAAAIAAAACAAAAPgAFADYAAAA3AAAAAgAAAAIAAACAAAUAAgAAADAAAAAxAAAALgAAAPkAAgAYAAAA+AAC
ABoAAAB5AAQAOQAAADoAAAATAAAA4AAEADsAAAA7AAAAPAAAAFAABABCAAAAQwAAAD4AAABpEQcARAAAAEUAAAAPAAAAKgAAAC4A
AAAAAAAAaREHAEcAAABIAAAAEQAAACoAAAAuAAAAAAAAAGsRBwBCAAAASQAAAEUAAABIAAAAQwAAAAAAAAB7AAUADAAAAEoAAAA6
AAAABQAAAHwABABLAAAATAAAAEoAAABqEQYATAAAAEkAAAAqAAAALgAAAAAAAAD9AAEAOAABAA==
Versions
  • intel-igc-core-2 / intel-igc-opencl-2 2.41.5; intel-ocloc and intel-opencl-icd 26.35.39758.10
  • Intel oneAPI DPC++/C++ Compiler 2026.1.1 (2026.1.1.20260724)
  • Ubuntu 24.04; host Core Ultra 7 265K (Arrow Lake) with 2x Arc Pro B70
Where it was found

In the Mach X inference engine: 12 of its 168 SYCL device images use the matrix extension, and compiling them for the box's integrated GPU (-device 12.70.4) crashed the same way. The engine now checks aspect::ext_intel_matrix and takes non-matrix kernels there; Arc A-series (DG2) reports the aspect, so a fp16 matrix kernel still reaches the compiler on that platform.

主要言語
C++
スター
725
フォーク
200
PR マージ指標
30日以内にマージされた PR はありません

環境構築

はじめの一歩

  1. issue を最後まで読み、次にプロジェクトのコントリビューションガイドを読みます。
  2. 着手することを issue にコメントします — 二人が同じ作業をするのを防げます。
  3. リポジトリをフォークし、ブランチを切って変更します。
  4. issue 番号を参照したプルリクエストを送ります。

intel/intel-graphics-compiler のほかの issue

intel/intel-graphics-compiler の issue をすべて見る

似ている issue

C++ の issue をもっと見る

新しい issue をメールで受け取る

初心者向けの GitHub issue を短くまとめたダイジェスト。