GemmSYCL.h

Go to the documentation of this file.

// ----------------------------------------------------------------------------
// -                        Open3D: www.open3d.org                            -
// ----------------------------------------------------------------------------
// Copyright (c) 2018-2026 www.open3d.org
// SPDX-License-Identifier: MIT
// ----------------------------------------------------------------------------
 
// Shared column-major GEMM shim for the SYCL conv ops (ContinuousConv,
// SparseConv and their Backprop/Transpose variants), used exactly like the
// CUDA path's `cutlass::gemm::device::Gemm<float, ColumnMajor, ...>`:
//     D = alpha * op(A) * op(B) + beta * C          (all matrices ColumnMajor)
//
// Declaration only: the sycl-tla (CUTLASS-for-SYCL) template machinery is
// instantiated once in GemmSYCL.cpp, so the eight conv op translation units
// that call this do not each recompile a full set of GEMM kernels through
// both the host and SPIR-V device passes. See GemmSYCL.cpp for the backend
// design notes (TF32/XMX vs. IEEE fp32 paths, the row/column-major transpose
// duality used to expose a column-major output, and tile selection).
#pragma once
 
#include <cutlass/layout/matrix.h>
 
#include <cstdint>
#include <sycl/sycl.hpp>
#include <vector>
 
namespace open3d {
namespace ml {
namespace impl {
 
template <class LayoutA = cutlass::layout::ColumnMajor,
          class LayoutB = cutlass::layout::ColumnMajor>
sycl::event GemmColumnMajorSYCL(sycl::queue& queue,
                                int m,
                                int n,
                                int k,
                                float alpha,
                                const float* A,
                                int64_t lda,
                                const float* B,
                                int64_t ldb,
                                float beta,
                                float* C,
                                int64_t ldc,
                                bool allow_tf32 = false,
                                const std::vector<sycl::event>& deps = {});
 
// Explicit instantiation declarations; definitions live in GemmSYCL.cpp.
extern template sycl::event
GemmColumnMajorSYCL<cutlass::layout::ColumnMajor, cutlass::layout::ColumnMajor>(
        sycl::queue&,
        int,
        int,
        int,
        float,
        const float*,
        int64_t,
        const float*,
        int64_t,
        float,
        float*,
        int64_t,
        bool,
        const std::vector<sycl::event>&);
 
extern template sycl::event
GemmColumnMajorSYCL<cutlass::layout::ColumnMajor, cutlass::layout::RowMajor>(
        sycl::queue&,
        int,
        int,
        int,
        float,
        const float*,
        int64_t,
        const float*,
        int64_t,
        float,
        float*,
        int64_t,
        bool,
        const std::vector<sycl::event>&);
 
}  // namespace impl
}  // namespace ml
}  // namespace open3d