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)
まだ誰も着手していません。
評価
調査の方向性
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 はありません
環境構築
- Dockerfile・Docker Compose ファイルなし
- プルリクエストのテンプレートなし
- コントリビューションガイドを読む
はじめの一歩
- issue を最後まで読み、次にプロジェクトのコントリビューションガイドを読みます。
- 着手することを issue にコメントします — 二人が同じ作業をするのを防げます。
- リポジトリをフォークし、ブランチを切って変更します。
- issue 番号を参照したプルリクエストを送ります。
intel/intel-graphics-compiler のほかの issue
-
functionControl internal option is documented but only the misspelled functonControl works再び着手できるかも このイシューのプルリクエストはマージされずにクローズされました。 オープン
難易度 2/5 1〜3時間 初心者へのやさしさ 75/100
intel/intel-graphics-compiler#422 · コメント 2 件 ·
-
add_strerror_r_to_error missing int overload対応中かも @cdecker08 が 425 日前に担当しました。 オープン
難易度 2/5 1〜3時間 初心者へのやさしさ 68/100
intel/intel-graphics-compiler#373 ·
-
CMake doesn't handle the case when multiple `opencl-c.h` are found対応中かも @pvelesko が 75 日前に担当しました。 オープン
難易度 2/5 1〜3時間 初心者へのやさしさ 68/100
intel/intel-graphics-compiler#314 · コメント 1 件 ·
-
Gen12LP: OWord spills of indirectly accessed variables use the wrong slot対応中かも @maleadt が 4 日前に担当しました。 オープン
難易度 4/5 3〜5日 初心者へのやさしさ 53/100
intel/intel-graphics-compiler#472 ·
-
Gen12LP: post-RA spill/fill propagation corrupts OWord spills対応中かも @maleadt が 4 日前に担当しました。 オープン
難易度 4/5 3〜5日 初心者へのやさしさ 58/100
intel/intel-graphics-compiler#471 ·
intel/intel-graphics-compiler の issue をすべて見る
似ている issue
-
runtime: lf_f128_mod is documented as fmodq (truncating) but implements floor-based modulo; unusedオープンinternal correctness low priority rt_lib
難易度 2/5 1〜3時間 初心者へのやさしさ 72/100
メンテナーはふだん 1 日以内に返信
-
難易度 1/5 1時間未満 初心者へのやさしさ 87/100
-
PriorityNotificationQueue::wakeUpAll() loses the wake-up for a thread that is not waiting yet対応中かも @matejk が今日担当しました。 オープンbug
難易度 2/5 1〜3時間 初心者へのやさしさ 78/100
pocoproject/poco#5674 ·
メンテナーはふだん 2 日以内に返信
-
bug docs GUI
難易度 2/5 1〜3時間 初心者へのやさしさ 62/100
sonic-pi-net/sonic-pi#3609 ·
-
bug
難易度 2/5 1〜3時間 初心者へのやさしさ 70/100
SanderMertens/flecs#2395 ·
メンテナーはふだん 1 日以内に返信