Implementation:Ggml org Ggml Sycl softmax

Knowledge Sources	GGML
Domains	ML_Infrastructure, GPU_Compute
Last Updated	2025-05-15 12:00 GMT

Overview

SYCL softmax kernel implementation with support for masking, ALiBi positional bias, attention sinks, and backward pass gradient computation.

Description

softmax.cpp implements the attention softmax operation critical for transformer inference on the SYCL backend. The main kernel soft_max_f32 is heavily templated for compile-time optimization:

• Template parameters: use_shared (whether to use shared memory for intermediate values), ncols_template (compile-time column count for common sizes, 0 for dynamic), block_size_template (work-group size), and T (mask type: float or sycl::half).

• Three-phase computation:
  • Phase 1 (Max): Each thread computes a partial maximum across its assigned columns, applying optional mask scaling (slope * mask[col]) for ALiBi bias and optional attention sink initialization. Warp-level reduction finds the row maximum.
  • Phase 2 (Exp + Sum): Each thread computes exp(x[col] * scale + mask_bias - max_val) for its columns, accumulating the sum. Cross-warp reduction via shared memory produces the total sum.
  • Phase 3 (Normalize): Each thread divides its exponential values by the total sum to produce normalized probabilities.

• ALiBi support: The get_alibi_slope function computes per-head linear bias slopes for Attention with Linear Biases, using the head index and n_head_log2 parameter.

• Attention sinks: When the sinks pointer is provided, the initial max value is seeded from the sink attention scores rather than -INFINITY, implementing attention sink optimization.

• Backward pass: soft_max_back computes the gradient of the softmax function for fine-tuning support.

The soft_max_params struct encapsulates all configuration including tensor dimensions, strides, scale, max_bias, and ALiBi slope parameters.

Usage

Called from the main SYCL backend when the compute graph contains GGML_OP_SOFT_MAX or GGML_OP_SOFT_MAX_BACK operations. This is invoked in every attention layer during transformer inference.

Code Reference

Source Location

• Repository: GGML
• File: src/ggml-sycl/softmax.cpp
• Lines: 426

Signatures

// Parameters structure
struct soft_max_params {
    int64_t nheads, ncols, nrows_x, nrows_y;
    int64_t ne00, ne01, ne02, ne03;
    int64_t nb11, nb12, nb13, ne12, ne13;
    uint32_t n_head_log2;
    float scale, max_bias, m0, m1;
};

// Main softmax kernel
template <bool use_shared, int ncols_template, int block_size_template, typename T>
static void soft_max_f32(const float * x, const T * mask, const float * sinks,
    float * dst, const soft_max_params p, uint8_t * dpct_local);

// Public dispatch functions
void ggml_sycl_op_soft_max(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
void ggml_sycl_op_soft_max_back(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; src[0] is input logits, src[1] is optional mask, src[2] is optional attention sinks

Outputs

Name	Type	Description
dst->data	void *	Normalized probability distribution (softmax output) in f32

Usage Examples

// Apply softmax to attention scores:
ggml_sycl_op_soft_max(sycl_ctx, softmax_output_tensor);

// Compute softmax gradient during backward pass:
ggml_sycl_op_soft_max_back(sycl_ctx, softmax_grad_tensor);

Related Pages

Implements Principle

• Principle:Ggml_org_Ggml_SYCL_GPU_Computation
• Environment:Ggml_org_Ggml_C_Cpp_Build_Environment