Fix device-agnostic fp32 GEMM compile on oneAPI 2025.2+

[ssheorey]

Context

We recently added support for sparse convolutions and continuous convolutions in Open3D for 3D point cloud data using sycl-tla. We ran into this issue. Currently, Open3D maintains this PR as a patch.

Description

Device-agnostic fp32 GemmUniversalAdapter::run() fails to compile on oneAPI 2025.2+ (reproduced on 2026.0):

error: unscoped enum 'Kind' requires fixed underlying type

This patch:

1. Gives fixed underlying types to Kind enums used in kernel metadata (scale_type.h, cache_operation.h, mma.h).
2. Uses sycl::detail::auto_name on the work-group-scratch launch path (#803 only fixed the other branch).
3. Adds __SYCL_DEVICE_ONLY__ fallbacks in arch/memory.h for global load/store (replacing CUDA-only / invalid paths on Intel GPU).

---

Type

• Bug

---

Testing

• Tests pass (sycl-tla CI — not run by PR author)
• A770
• Panther Lake

What was verified (author machine)

Check	Without patch	With patch
mre_compile_only.cpp device compile	Fails (Kind error, exit 1)	OK (exit 0)
mre_device_agnostic_fp32_gemm.cpp device compile	Fails (same error, exit 1)	OK (exit 0)
mre_device_agnostic_fp32_gemm.cpp link + run	Not run (no build)	OK; prints OK: device-agnostic fp32 GEMM completed (exit 0)

Environment: oneAPI 2026.0 (icpx 2026.0.0.20260331), sycl-tla main @ 2db1b7c, Intel GPU 0xb080, runtime via ONEAPI_DEVICE_SELECTOR=level_zero:gpu only.

Not verified

• oneAPI 2025.2 / 2025.3 (only 2026.0 retested end-to-end)
• OpenCL SYCL backend for the runtime MRE
• Xe20 hardware or sycl-tla CI

How to reproduce

From a sycl-tla repository root, save the two sources below as mre_compile_only.cpp and mre_device_agnostic_fp32_gemm.cpp, then:

source /opt/intel/oneapi/setvars.sh   # or oneapi-vars.sh for your install

ICXX_FLAGS="-std=c++17 -fsycl -DCUTLASS_ENABLE_SYCL=1 -DSYCL_INTEL_TARGET=1 \
  -I include -I tools/util/include"

# Before applying this PR's header changes:
icpx $ICXX_FLAGS -c mre_compile_only.cpp                    # expect Kind error
icpx $ICXX_FLAGS -c mre_device_agnostic_fp32_gemm.cpp      # expect Kind error

# After applying the PR:
icpx $ICXX_FLAGS -c mre_compile_only.cpp                    # expect success
icpx $ICXX_FLAGS mre_device_agnostic_fp32_gemm.cpp -o mre_gemm
ONEAPI_DEVICE_SELECTOR=level_zero:gpu ./mre_gemm                   # expect OK line above

mre_compile_only.cpp (compile-only repro)

// SPDX-License-Identifier: BSD-3-Clause
#include <cutlass/epilogue/collective/collective_builder.hpp>
#include <cutlass/gemm/collective/collective_builder.hpp>
#include <cutlass/gemm/device/gemm_universal_adapter.h>
#include <cutlass/layout/matrix.h>
#include <cute/tensor.hpp>
#include <sycl/sycl.hpp>

using Element = float;
using LayoutA = cutlass::layout::ColumnMajor;
using LayoutB = cutlass::layout::RowMajor;
using LayoutC = cutlass::layout::RowMajor;
using TileShape = cute::Shape<cute::_64, cute::_64, cute::_16>;

using CollectiveMainloop = typename cutlass::gemm::collective::CollectiveBuilder<
    cutlass::arch::Agnostic, cutlass::arch::OpMultiplyAdd,
    Element, LayoutA, 4, Element, LayoutB, 4, Element, TileShape,
    cute::Shape<cute::_1, cute::_1, cute::_1>,
    cutlass::gemm::collective::StageCountAuto,
    cutlass::gemm::collective::KernelScheduleAuto>::CollectiveOp;

using EpilogueOp =
    cutlass::epilogue::fusion::LinearCombination<Element, Element, Element, Element>;
using CollectiveEpilogue = typename cutlass::epilogue::collective::CollectiveBuilder<
    cutlass::arch::Agnostic, cutlass::arch::OpMultiplyAdd, TileShape,
    cute::Shape<cute::_1, cute::_1, cute::_1>,
    cutlass::epilogue::collective::EpilogueTileAuto, Element, Element,
    Element, LayoutC, 4, Element, LayoutC, 4,
    cutlass::epilogue::collective::EpilogueScheduleAuto, EpilogueOp>::CollectiveOp;

using GemmKernel = cutlass::gemm::kernel::GemmUniversal<
    cute::Shape<int, int, int, int>, CollectiveMainloop, CollectiveEpilogue>;
using Gemm = cutlass::gemm::device::GemmUniversalAdapter<GemmKernel>;

int main() {
    sycl::queue q;
    Gemm gemm;
    typename Gemm::Arguments args{};
    gemm.initialize(args, nullptr, &q);
    return gemm.run(&q) == cutlass::Status::kSuccess ? 0 : 1;
}

mre_device_agnostic_fp32_gemm.cpp (compile + GPU runtime smoke)

// SPDX-License-Identifier: BSD-3-Clause
#include <cutlass/epilogue/collective/collective_builder.hpp>
#include <cutlass/gemm/collective/collective_builder.hpp>
#include <cutlass/gemm/device/gemm_universal.h>
#include <cutlass/gemm/device/gemm_universal_adapter.h>
#include <cutlass/kernel_hardware_info.h>
#include <cutlass/layout/matrix.h>
#include <cutlass/util/packed_stride.hpp>

#include <cute/tensor.hpp>
#include <cstdlib>
#include <iostream>
#include <sycl/sycl.hpp>
#include <vector>

namespace {

using Element = float;
using LayoutA = cutlass::layout::ColumnMajor;
using LayoutB = cutlass::layout::RowMajor;
using LayoutC = cutlass::layout::RowMajor;
using TileShape = cute::Shape<cute::_64, cute::_64, cute::_16>;

using CollectiveMainloop =
    typename cutlass::gemm::collective::CollectiveBuilder<
        cutlass::arch::Agnostic, cutlass::arch::OpMultiplyAdd, Element,
        LayoutA, 4, Element, LayoutB, 4, Element, TileShape,
        cute::Shape<cute::_1, cute::_1, cute::_1>,
        cutlass::gemm::collective::StageCountAuto,
        cutlass::gemm::collective::KernelScheduleAuto>::CollectiveOp;

using EpilogueOp = cutlass::epilogue::fusion::LinearCombination<
    Element, Element, Element, Element>;

using CollectiveEpilogue =
    typename cutlass::epilogue::collective::CollectiveBuilder<
        cutlass::arch::Agnostic, cutlass::arch::OpMultiplyAdd, TileShape,
        cute::Shape<cute::_1, cute::_1, cute::_1>,
        cutlass::epilogue::collective::EpilogueTileAuto, Element, Element,
        Element, LayoutC, 4, Element, LayoutC, 4,
        cutlass::epilogue::collective::EpilogueScheduleAuto,
        EpilogueOp>::CollectiveOp;

using GemmKernel = cutlass::gemm::kernel::GemmUniversal<
    cute::Shape<int, int, int, int>, CollectiveMainloop, CollectiveEpilogue>;

using Gemm = cutlass::gemm::device::GemmUniversalAdapter<GemmKernel>;

}  // namespace

int main() {
    sycl::queue q(sycl::gpu_selector_v);
    const int m = 64, n = 64, k = 16;
    const float alpha = 1.f, beta = 0.f;

    std::vector<float> h_a(m * k, 1.f), h_b(k * n, 2.f), h_c(m * n, 0.f),
        h_d(m * n, 0.f);
    float *d_a = sycl::malloc_device<float>(h_a.size(), q);
    float *d_b = sycl::malloc_device<float>(h_b.size(), q);
    float *d_c = sycl::malloc_device<float>(h_c.size(), q);
    float *d_d = sycl::malloc_device<float>(h_d.size(), q);
    q.memcpy(d_a, h_a.data(), h_a.size() * sizeof(float)).wait();
    q.memcpy(d_b, h_b.data(), h_b.size() * sizeof(float)).wait();
    q.memcpy(d_c, h_c.data(), h_c.size() * sizeof(float)).wait();

    using StrideA = typename Gemm::GemmKernel::StrideA;
    using StrideB = typename Gemm::GemmKernel::StrideB;
    using StrideC = typename Gemm::GemmKernel::StrideC;
    using StrideD = typename Gemm::GemmKernel::StrideD;

    const int batch = 1;
    StrideA stride_a =
        cutlass::make_cute_packed_stride(StrideA{}, cute::make_shape(m, k, batch));
    StrideB stride_b =
        cutlass::make_cute_packed_stride(StrideB{}, cute::make_shape(n, k, batch));
    StrideC stride_c =
        cutlass::make_cute_packed_stride(StrideC{}, cute::make_shape(m, n, batch));
    StrideD stride_d = stride_c;

    cutlass::KernelHardwareInfo hw_info;
    hw_info.sm_count = 1;

    typename Gemm::GemmKernel::Arguments args{
        cutlass::gemm::GemmUniversalMode::kGemm, {m, n, k, 1},
        {d_a, stride_a, d_b, stride_b},
        {{alpha, beta}, d_c, stride_c, d_d, stride_d}, hw_info};

    Gemm gemm;
    void *workspace = nullptr;
    size_t ws = Gemm::get_workspace_size(args);
    if (ws) workspace = sycl::malloc_device(ws, q);

    if (gemm.can_implement(args) != cutlass::Status::kSuccess ||
        gemm.initialize(args, workspace, &q) != cutlass::Status::kSuccess ||
        gemm.run(&q) != cutlass::Status::kSuccess) {
        std::cerr << "GEMM failed\n";
        return EXIT_FAILURE;
    }
    q.wait();

    q.memcpy(h_d.data(), d_d, h_d.size() * sizeof(float)).wait();
    const float expected = alpha * 1.f * 2.f * static_cast<float>(k);
    if (std::abs(h_d[0] - expected) > 1e-3f) {
        std::cerr << "Mismatch: got " << h_d[0] << " expected " << expected << "\n";
        return EXIT_FAILURE;
    }

    std::cout << "OK: device-agnostic fp32 GEMM completed\n";
    sycl::free(d_a, q);
    sycl::free(d_b, q);
    sycl::free(d_c, q);
    sycl::free(d_d, q);
    if (workspace) sycl::free(workspace, q);
    return EXIT_SUCCESS;
}

---

Performance

Metric	Before	After
Compile (device-agnostic fp32 GEMM)	Fails	OK
Runtime	—	No intended change

---

References

Fixes #

Related: #803

---

Checklist

• Copyright
• Co-pilot Review
• Deprecated APIs not used

[taozha2]

I saw compile error here, this API maybe different for different compiler, please check this.