open3d::ml::impl::sycl_gemm_detail Namespace Reference

Data Structures
struct	GemmArgs
struct	TransposeLayout
struct	TransposeLayout< cutlass::layout::ColumnMajor >
struct	TransposeLayout< cutlass::layout::RowMajor >

Typedefs
using	Tile = cute::Shape< cute::_16, cute::_16, cute::_8 >

Functions
template<class Layout >
auto	MakeStrideA (int64_t ld)
template<class Layout >
auto	MakeStrideB (int64_t ld)
template<class TileShape , class LayoutA , class LayoutB >
std::pair< cutlass::Status, sycl::event >	RunGemmXmxTf32RowMajorOutput (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, const float *C, int64_t ldc, float *D, int64_t ldd, const std::vector< sycl::event > &deps={})
template<class TileShape , class LayoutA , class LayoutB >
std::pair< cutlass::Status, sycl::event >	RunGemmIeeeFp32RowMajorOutput (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, const float *C, int64_t ldc, float *D, int64_t ldd, const std::vector< sycl::event > &deps={})
template<class LayoutA , class LayoutB >
std::pair< cutlass::Status, sycl::event >	RunTf32 (const GemmArgs &a)
	Runs the TF32/XMX path.
template<class LayoutA , class LayoutB >
std::pair< cutlass::Status, sycl::event >	RunIeee (const GemmArgs &a)
	Runs the IEEE fp32 path.
bool	IsTf32Aligned (int64_t value)

Typedef Documentation

Tile

using open3d::ml::impl::sycl_gemm_detail::Tile = typedef cute::Shape<cute::_16, cute::_16, cute::_8>

Collective (work-group) tile shape. Small per-chunk GEMMs are typical of the conv ops (~32 output columns per run, see SparseConvSYCL.h/ContinuousConvSYCL.h), and can_implement's outcome does not depend on the tile shape (verified against sycl-tla's tile scheduler and epilogue/mainloop can_implement overloads – partial edge tiles are always predicated, not rejected), so a single tile per precision path is sufficient; there is no correctness reason to probe multiple tiles.

Function Documentation

IsTf32Aligned()

bool open3d::ml::impl::sycl_gemm_detail::IsTf32Aligned

(

int64_t

value

)

True if value is a multiple of the sycl-tla 128-bit (4-fp32-element) alignment requirement checked by the TF32/XMX path's can_implement (see the AlignmentA/B/C/D comment in RunGemmXmxTf32RowMajorOutput above).

MakeStrideA()

template<class Layout >

auto open3d::ml::impl::sycl_gemm_detail::MakeStrideA

(

int64_t

ld

)

Builds the CuTe stride for the A operand (canonical CuTe modes [M,K,L]) from a runtime leading dimension, matching cutlass::detail::TagToStrideA_t's convention (unit-stride dim is a compile-time cute::Int<1>): RowMajor -> (ld, 1, batch), ColumnMajor -> (1, ld, batch).

MakeStrideB()

template<class Layout >

auto open3d::ml::impl::sycl_gemm_detail::MakeStrideB

(

int64_t

ld

)

Builds the CuTe stride for the B operand. Note: cutlass::detail:: TagToStrideB_t is in canonical CuTe modes [N,K,L] — the opposite convention from A/C/D — so RowMajor/ColumnMajor map to the reverse stride pattern compared to MakeStrideA: RowMajor -> (1, ld, batch), ColumnMajor -> (ld, 1, batch). (Verified against cutlass/detail/layout.hpp; this asymmetry is intentional in CUTLASS/CuTe, reflecting that B's mathematical (K,N) shape is expressed as a (N,K) CuTe tensor.)

RunGemmIeeeFp32RowMajorOutput()

template<class TileShape , class LayoutA , class LayoutB >

std::pair< cutlass::Status, sycl::event > open3d::ml::impl::sycl_gemm_detail::RunGemmIeeeFp32RowMajorOutput	(	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,
		const float *	C,
		int64_t	ldc,
		float *	D,
		int64_t	ldd,
		const std::vector< sycl::event > &	deps = {}
	)

Runs IEEE float32 GEMM through sycl-tla's device-agnostic path. Returns {status, completion_event}; does not block (see RunGemmXmxTf32RowMajorOutput above for the barrier/workspace-free rationale).

RunGemmXmxTf32RowMajorOutput()

template<class TileShape , class LayoutA , class LayoutB >

std::pair< cutlass::Status, sycl::event > open3d::ml::impl::sycl_gemm_detail::RunGemmXmxTf32RowMajorOutput	(	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,
		const float *	C,
		int64_t	ldc,
		float *	D,
		int64_t	ldd,
		const std::vector< sycl::event > &	deps = {}
	)

Runs alpha * op(A) * op(B) + beta * C on the Intel Xe DPAS path via sycl-tla, producing a RowMajor (M x N) output D (D may alias C). A is (M x K, LayoutA), B is (K x N, LayoutB). GEMM element type is cutlass::tfloat32_t (see file header); accumulation/output stay float. Returns {status, completion_event}; does not block. The event is a barrier over all commands previously submitted to queue (see GemmColumnMajorSYCL's doc comment for why sycl-tla leaves no tighter option).

RunIeee()

template<class LayoutA , class LayoutB >

std::pair< cutlass::Status, sycl::event > open3d::ml::impl::sycl_gemm_detail::RunIeee

(

const GemmArgs &

a

)

Runs the IEEE fp32 path.

RunTf32()

template<class LayoutA , class LayoutB >

std::pair< cutlass::Status, sycl::event > open3d::ml::impl::sycl_gemm_detail::RunTf32

(

const GemmArgs &

a

)

Runs the TF32/XMX path.