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)

Aperta
  • Qualcosa non va?
#475 0 commenti 0 reazioni 0 assegnatari Vedi su GitHub

Nessuno ha ancora preso questa issue.

Valutazione

Difficoltà
4/5
Tempo stimato
3-5 giorni
Idoneità per principianti
49/100
Tipo di issue
Bug
Chiarezza
Abbastanza chiara
Stato di attività
Attiva
Stack tecnologico
cpp
Ambito
compilers

Direzione di ricerca

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.

Scritto dal modello di indicizzazione a partire dal testo della issue.

Descrizione

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.

Lingua principale
C++
Stelle
725
Fork
200
Nessuna PR unita negli ultimi 30g

Preparare l'ambiente

Come iniziare

  1. Leggi tutta la issue e poi la guida ai contributi del progetto.
  2. Commenta sulla issue per dire che te ne occupi tu — evita che due persone facciano lo stesso lavoro.
  3. Fai un fork del repository e lavora su un branch.
  4. Apri una pull request che faccia riferimento al numero della issue.

Altre issue di intel/intel-graphics-compiler

Tutte le issue di intel/intel-graphics-compiler

Issue simili

Altre issue su C++

Ricevi le nuove issue nella tua casella

Un breve riepilogo di issue GitHub adatte ai principianti.