Repository navigation
oneMKL: transposed sparse gemv reads y when beta is zero #703
Description
On the GPU, oneapi::mkl::sparse::gemv with a CSR handle and transpose::trans or transpose::conjtrans reads y when beta is zero, so NaNs in an uninitialized output buffer end up in the result. The same call with transpose::nontrans ignores y, as do CSR gemm with any op and the CPU device. Prefilling y with 42 gives the correct result, so y is scaled by beta rather than skipped. BLAS convention is that y need not be set on input when beta is zero, and Julia's mul! relies on that: * passes uninitialized memory from similar.
In oneAPI.jl this made oneSparseMatrixCSC * x return NaNs depending on what the allocator handed out, because CSC matrices are stored as the CSR handle of their transpose and multiplied with a transposed op. That is the intermittent oneMKL.oneSparseMatrixCSC test failure in GPUArrays.jl's downstream CI. #702 works around it by zeroing the output before oneMKL calls that take a beta when beta is zero.
Reproducer
The calls mirror what our support library does (deps/src/onemkl.cpp): init_matrix_handle, set_csr_data with one-based indices, then gemv/gemm with USM pointers and no optimize_* call.
repro.cpp
// oneMKL sparse::gemv on a CSR handle with op = trans/conjtrans reads y when beta == 0
#include <sycl/sycl.hpp>
#include <oneapi/mkl.hpp>
#include <mkl.h>
#include <algorithm>
#include <cmath>
#include <cstdint>
#include <cstdio>
#include <utility>
using oneapi::mkl::transpose;
namespace sparse = oneapi::mkl::sparse;
int main() {
sycl::queue q{sycl::default_selector_v, sycl::property::queue::in_order()};
char ver[256];
mkl_get_version_string(ver, sizeof(ver));
auto dev = q.get_device();
printf("%s\n%s (driver %s)\n\n", ver, dev.get_info<sycl::info::device::name>().c_str(),
dev.get_info<sycl::info::device::driver_version>().c_str());
// 4x3 matrix in one-based CSR; row 2 and column 2 are empty
// [1 0 4]
// [0 0 0]
// [0 0 2]
// [3 0 0]
const int64_t m = 4, n = 3, nnz = 4;
int32_t *row_ptr = sycl::malloc_shared<int32_t>(m + 1, q);
int32_t *col_ind = sycl::malloc_shared<int32_t>(nnz, q);
float *val = sycl::malloc_shared<float>(nnz, q);
const int32_t rp[] = {1, 3, 3, 4, 5}, ci[] = {1, 3, 3, 1};
const float v[] = {1, 4, 2, 3};
std::copy(rp, rp + m + 1, row_ptr);
std::copy(ci, ci + nnz, col_ind);
std::copy(v, v + nnz, val);
sparse::matrix_handle_t A;
sparse::init_matrix_handle(&A);
sparse::set_csr_data(q, A, m, n, nnz, oneapi::mkl::index_base::one, row_ptr, col_ind, val).wait();
// y = 1 * op(A) * x + 0 * y with x = ones and y prefilled; gemm uses 2 columns
const float expect_N[] = {5, 0, 2, 3}, expect_T[] = {4, 0, 6};
float *x = sycl::malloc_shared<float>(2 * m, q), *y = sycl::malloc_shared<float>(2 * m, q);
const std::pair<transpose, const char *> ops[] = {
{transpose::nontrans, "nontrans"}, {transpose::trans, "trans"}, {transpose::conjtrans, "conjtrans"}};
for (int64_t ncols : {1, 2})
for (auto [op, opname] : ops)
for (float fill : {NAN, 42.0f}) {
const bool N = op == transpose::nontrans;
const int64_t nx = N ? n : m, ny = N ? m : n;
std::fill(x, x + nx * ncols, 1.0f);
std::fill(y, y + ny * ncols, fill);
if (ncols == 1)
sparse::gemv(q, op, 1.0f, A, x, 0.0f, y).wait();
else
sparse::gemm(q, oneapi::mkl::layout::col_major, op, transpose::nontrans, 1.0f, A,
x, ncols, nx, 0.0f, y, ny).wait();
bool ok = true;
printf("%s %-9s y_in=%-3g y =", ncols == 1 ? "gemv" : "gemm", opname, fill);
for (int64_t i = 0; i < ny * ncols; i++) {
printf(" %g", y[i]);
ok &= y[i] == (N ? expect_N : expect_T)[i % ny];
}
printf(" %s\n", ok ? "OK" : "WRONG");
}
sparse::release_matrix_handle(q, &A).wait();
for (void *p : {(void *)row_ptr, (void *)col_ind, (void *)val, (void *)x, (void *)y})
sycl::free(p, q);
}icpx -fsycl -qmkl=sequential -O2 repro.cpp -o repro
ONEAPI_DEVICE_SELECTOR=level_zero:gpu ./repro
Intel(R) oneAPI Math Kernel Library Version 2026.1-Product Build 20260612 for Intel(R) 64 architecture applications
Intel(R) Iris(R) Xe Graphics (driver 1.15.38308)
gemv nontrans y_in=nan y = 5 0 2 3 OK
gemv nontrans y_in=42 y = 5 0 2 3 OK
gemv trans y_in=nan y = nan nan nan WRONG
gemv trans y_in=42 y = 4 0 6 OK
gemv conjtrans y_in=nan y = nan nan nan WRONG
gemv conjtrans y_in=42 y = 4 0 6 OK
gemm nontrans y_in=nan y = 5 0 2 3 5 0 2 3 OK
gemm nontrans y_in=42 y = 5 0 2 3 5 0 2 3 OK
gemm trans y_in=nan y = 4 0 6 4 0 6 OK
gemm trans y_in=42 y = 4 0 6 4 0 6 OK
gemm conjtrans y_in=nan y = 4 0 6 4 0 6 OK
gemm conjtrans y_in=42 y = 4 0 6 4 0 6 OK
With ONEAPI_DEVICE_SELECTOR=opencl:cpu every case is OK.
Variations:
- Every output entry becomes NaN, not just the one for the empty column, and a matrix without empty rows or columns fails the same way.
- Zero-based indices fail the same way.
- Calling
sparse::optimize_gemv(q, op, A)before each gemv makes all cases correct. Calling it only fornontransand then running atransgemv still fails.
On the older oneMKL of the Aurora LTS stack, sparse gemm with beta = 0 also propagated NaN, for every op (first CI run of #702). I haven't been able to reproduce that on 2026.1.
Versions
- oneMKL 2026.1 (Product Build 20260612), the release
oneAPI_Support_jll0.11.0 is built from - Intel oneAPI DPC++/C++ Compiler 2026.1.0 (2026.1.0.20260617)
- Intel Iris Xe Graphics (TigerLake-LP GT2), Level Zero driver 1.15.38308 (NEO 26.18.38308 from
NEO_jll), Level Zero loader 1.29.0 - Debian 13
Float64 is untested: this GPU has no FP64 support.
To do: report this to Intel on the oneMKL community forum, and link the report here.
Activity
maleadt commentedon Oct 16, 2026
Reported upstream as https://community.intel.com/t/sparse-gemv-with-transpose-trans-reads-y-when-beta-is-zero-gpu-onemkl-2026-1/144080
maleadt commentedon Oct 23, 2026
Confirmed by upstream and will be fixed in 2026.3.