[SYCL][BMG] joint_matrix reports no matrix hardware on Battlemage G31 (matrix_combinations empty)

[Wirhoss]

Describe the bug

On Arc Pro B70 (BMG-G31, 8086:e223), joint_matrix kernels compile but throw at runtime: no matrix hardware on the target device, joint_matrix is not supported

matrix_combinations returns an empty vector and ext_intel_matrix is missing from the device's aspect list, on both the Level Zero and OpenCL backends, with JIT and with AOT (-device bmg_g31). The device reports architecture::intel_gpu_bmg_g31, which the extension spec lists as supported.

the same device advertises cl_intel_subgroup_matrix_multiply_accumulate, cl_intel_subgroup_matrix_multiply_accumulate_tf32 and cl_intel_subgroup_2d_block_io through OpenCL, and sycl::ext::intel::esimd::xmx::dpas (INT8, 8x16x32) runs correctly on it — the IGC dump shows dpas.s8.s8.8.8. So the DPAS hardware is reachable; only the cooperative-matrix path reports it as absent.

possibly related: #16851 (same message on Lunar Lake, closed, predates the release that added BMG/LNL/ARL-H to Joint Matrix).

To reproduce

#include <sycl/sycl.hpp>
#include
using namespace sycl;
namespace syclex = sycl::ext::oneapi::experimental;
namespace matrix = sycl::ext::oneapi::experimental::matrix;

constexpr int TM = 8, TN = 16, TK = 32, SG = 16;

int main() {
queue q;
auto dev = q.get_device();
std::cout << dev.get_infoinfo::device::name() << " / "
<< dev.get_infoinfo::device::driver_version() << "\n";
std::cout << "matrix_combinations: "
<< dev.get_infosyclex::info::device::matrix_combinations().size() << "\n";

auto *A = malloc_shared<int8_t>(TM * TK, q);
auto *B = malloc_shared<int8_t>(TK * TN, q);
auto *C = malloc_shared<int32_t>(TM * TN, q);
for (int i = 0; i < TM * TK; ++i) A[i] = 1;
for (int i = 0; i < TK * TN; ++i) B[i] = 1;

try {
q.submit([&](handler &h) {
h.parallel_for(nd_range<2>({1, SG}, {1, SG}),
[=](nd_item<2> it) [[sycl::reqd_sub_group_size(SG)]] {
auto sg = it.get_sub_group();
auto pA = address_space_cast<access::address_space::global_space,
access::decorated::no>(A);
auto pB = address_space_cast<access::address_space::global_space,
access::decorated::no>(B);
auto pC = address_space_cast<access::address_space::global_space,
access::decorated::no>(C);
matrix::joint_matrix<sub_group, int8_t, matrix::use::a, TM, TK,
matrix::layout::row_major> tA;
matrix::joint_matrix<sub_group, int8_t, matrix::use::b, TK, TN,
matrix::layout::ext_intel_packed> tB;
matrix::joint_matrix<sub_group, int32_t, matrix::use::accumulator, TM, TN> tC;
matrix::joint_matrix_fill(sg, tC, 0);
matrix::joint_matrix_load(sg, tA, pA, TK);
matrix::joint_matrix_load(sg, tB, pB, TN);
matrix::joint_matrix_mad(sg, tC, tA, tB, tC);
matrix::joint_matrix_store(sg, tC, pC, TN, matrix::layout::row_major);
});
}).wait();
std::cout << "ok, C[0]=" << C[0] << " (expected " << TK << ")\n";
} catch (exception const &e) {
std::cout << "EXCEPTION: " << e.what() << "\n";
}
free(A, q); free(B, q); free(C, q);
}

icpx -fsycl -O2 -o jm_repro jm_repro.cpp
icpx -fsycl -O2 -fsycl-targets=spir64_gen -Xsycl-target-backend "-device bmg_g31" -o jm_repro_aot jm_repro.cpp

Output (identical for the AOT binary, and for ONEAPI_DEVICE_SELECTOR=opencl:*):
Intel(R) Graphics [0xe223] / 1.14.37020+3
matrix_combinations: 0
EXCEPTION: no matrix hardware on the target device, joint_matrix is not supported

Environment

GPU: Intel Arc Pro B70, BMG-G31, 8086:e223, reported as intel_gpu_bmg_g31, sub-group sizes 16/32
Compiler: Intel oneAPI DPC++/C++ 2025.3.3 (2025.3.3.20260319)
Runtime under test: NEO 26.05.37020.3, Level Zero 1.14.37020+3
Also reproduced with compute runtime 26.22.38646.7 and 38308
Host: Ubuntu 26.04, kernel 7.0.0-28-generic, xe driver

Additional context

No response

[dkhaldi]

SYCL Joint Matrix is supported on https://github.com/intel/llvm/blob/sycl/sycl/source/detail/device_impl.hpp#L1526
But that support was added in end of 2025.
I see you are using a relatively old compiler ( Intel oneAPI DPC++/C++ 2025.3.3)
Can you use a more recent compiler and check again?
You can also get daily builds from https://github.com/intel/llvm/releases