Implementation:Ggml org Ggml Sycl binary broadcast
| Knowledge Sources | |
|---|---|
| Domains | ML_Infrastructure, GPU_Compute |
| Last Updated | 2025-05-15 12:00 GMT |
Overview
SYCL binary broadcast operation kernels implementing element-wise arithmetic with NumPy-style broadcasting for the SYCL backend.
Description
binbcast.cpp provides GPU-accelerated binary element-wise operations (add, subtract, multiply, divide, repeat) with full multi-dimensional broadcasting support on SYCL devices. The file defines two core kernel templates -- k_bin_bcast and k_bin_bcast_unravel -- that compute per-element source indices using modular arithmetic to handle shape mismatches between operands. The bin_bcast_sycl struct template dispatches the appropriate kernel based on whether the total element count exceeds the maximum work-group range, selecting between the standard 3D-indexed kernel and the unraveled flat-index variant. Each binary operation (add, sub, mul, div) is instantiated through ggml_sycl_op_bin_bcast which extracts tensor metadata, determines strides and contiguity, and launches the kernel on the SYCL queue.
Usage
These functions are called from the main SYCL backend dispatcher (ggml-sycl.cpp) when the compute graph contains binary arithmetic operations (GGML_OP_ADD, GGML_OP_SUB, GGML_OP_MUL, GGML_OP_DIV, GGML_OP_REPEAT). The broadcasting logic automatically handles mismatched tensor shapes following standard broadcasting rules.
Code Reference
Source Location
- Repository: GGML
- File: src/ggml-sycl/binbcast.cpp
- Lines: 345
Signatures
// Core kernel template (3D indexed)
template<float (*bin_op)(const float, const float), typename src0_t, typename src1_t, typename dst_t>
static void k_bin_bcast(const src0_t * src0, const src1_t * src1, dst_t * dst,
int ne0, int ne1, int ne2, int ne3,
int ne10, int ne11, int ne12, int ne13,
int s1, int s2, int s3,
int s01, int s02, int s03,
int s11, int s12, int s13,
const sycl::nd_item<3> &item_ct1);
// Flat-index unraveled variant
template<float (*bin_op)(const float, const float), typename src0_t, typename src1_t, typename dst_t>
static void k_bin_bcast_unravel(const src0_t * src0, const src1_t * src1, dst_t * dst,
int ne0, int ne1, int ne2, int ne3, ...);
// Public entry points
void ggml_sycl_add(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
void ggml_sycl_sub(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
void ggml_sycl_mul(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
void ggml_sycl_div(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
void ggml_sycl_repeat(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
I/O Contract
Inputs
| Name | Type | Required | Description |
|---|---|---|---|
| ctx | ggml_backend_sycl_context & | Yes | SYCL backend context providing the device queue |
| dst | ggml_tensor * | Yes | Destination tensor whose src[0] and src[1] are the operands |
Outputs
| Name | Type | Description |
|---|---|---|
| dst->data | void * | Result tensor with broadcasted binary operation applied element-wise |
Usage Examples
// Called internally by the SYCL backend when dispatching a compute graph node:
// For an ADD operation with broadcasting (e.g., [1024, 768] + [1, 768])
ggml_sycl_add(sycl_ctx, add_tensor);
// For element-wise multiplication (e.g., [batch, seq, hidden] * [1, 1, hidden])
ggml_sycl_mul(sycl_ctx, mul_tensor);