Commit 2090f60f0 for llama.cpp
commit 2090f60f0b60b67b3bcc96573c6e68a0c33f30c7
Author: Aman Gupta <amangupta052@gmail.com>
Date: Wed Sep 30 20:24:59 2026 +0800
ggml : add BF16 unary, GLU, binary and scale ops (CPU, CUDA) (#29675)
* ggml : add BF16 unary, GLU, binary and scale ops (CPU, CUDA)
* ggml-cpu : use per-op _bf16 functions for BF16 unary and GLU ops
Assisted-by: Claude Opus 5.5
* CUDA: use ggml_cuda_cast in binbcast and unary kernels to fix the HIP bf16 build
* ggml-openvino : reject BF16 SCALE and mixed-type BF16 ADD/MUL/SUB
diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp
index a6ccfe954..3e377298f 100644
--- a/ggml/src/ggml-cpu/ops.cpp
+++ b/ggml/src/ggml-cpu/ops.cpp
@@ -2106,6 +2106,46 @@ void ggml_compute_forward_concat(
}
}
+// row loop shared by the BF16 unary ops
+template <void (*vec_op)(const int, ggml_bf16_t *, const ggml_bf16_t *)>
+static void ggml_compute_forward_unary_bf16_rows(
+ const ggml_compute_params * params,
+ ggml_tensor * dst) {
+
+ const ggml_tensor * src0 = dst->src[0];
+
+ assert(ggml_is_contiguous_rows(src0));
+ assert(ggml_are_same_shape(src0, dst));
+
+ GGML_TENSOR_LOCALS(int64_t, ne0, src0, ne)
+ GGML_TENSOR_LOCALS(size_t, nb0, src0, nb)
+ GGML_TENSOR_LOCALS(int64_t, ne, dst, ne)
+ GGML_TENSOR_LOCALS(size_t, nb, dst, nb)
+
+ const int ith = params->ith;
+ const int nth = params->nth;
+
+ const int nc = src0->ne[0];
+ const int nr = ggml_nrows(src0);
+
+ // rows per thread
+ const int dr = (nr + nth - 1)/nth;
+
+ // row range for this thread
+ const int ir0 = dr*ith;
+ const int ir1 = MIN(ir0 + dr, nr);
+
+ for (int ir = ir0; ir < ir1; ++ir) {
+ const int i3 = ir/(ne02*ne01);
+ const int i2 = (ir - i3*ne02*ne01)/ne01;
+ const int i1 = (ir - i3*ne02*ne01 - i2*ne01);
+
+ vec_op(nc,
+ (ggml_bf16_t *) ((char *) dst->data + i3*nb3 + i2*nb2 + i1*nb1),
+ (ggml_bf16_t *) ((char *) src0->data + i3*nb03 + i2*nb02 + i1*nb01));
+ }
+}
+
// ggml_compute_forward_gelu
static void ggml_compute_forward_gelu_f32(
@@ -2203,6 +2243,12 @@ static void ggml_compute_forward_gelu_f16(
}
}
+static void ggml_compute_forward_gelu_bf16(
+ const ggml_compute_params * params,
+ ggml_tensor * dst) {
+ ggml_compute_forward_unary_bf16_rows<ggml_vec_gelu_bf16>(params, dst);
+}
+
static void ggml_compute_forward_gelu(
const ggml_compute_params * params,
ggml_tensor * dst) {
@@ -2218,6 +2264,10 @@ static void ggml_compute_forward_gelu(
{
ggml_compute_forward_gelu_f16(params, dst);
} break;
+ case GGML_TYPE_BF16:
+ {
+ ggml_compute_forward_gelu_bf16(params, dst);
+ } break;
default:
{
GGML_ABORT("fatal error");
@@ -2433,6 +2483,12 @@ static void ggml_compute_forward_gelu_erf_f16(
}
}
+static void ggml_compute_forward_gelu_erf_bf16(
+ const ggml_compute_params * params,
+ ggml_tensor * dst) {
+ ggml_compute_forward_unary_bf16_rows<ggml_vec_gelu_erf_bf16>(params, dst);
+}
+
static void ggml_compute_forward_gelu_erf(
const ggml_compute_params * params,
ggml_tensor * dst) {
@@ -2448,6 +2504,10 @@ static void ggml_compute_forward_gelu_erf(
{
ggml_compute_forward_gelu_erf_f16(params, dst);
} break;
+ case GGML_TYPE_BF16:
+ {
+ ggml_compute_forward_gelu_erf_bf16(params, dst);
+ } break;
default:
{
GGML_ABORT("fatal error");
@@ -2552,6 +2612,12 @@ static void ggml_compute_forward_gelu_quick_f16(
}
}
+static void ggml_compute_forward_gelu_quick_bf16(
+ const ggml_compute_params * params,
+ ggml_tensor * dst) {
+ ggml_compute_forward_unary_bf16_rows<ggml_vec_gelu_quick_bf16>(params, dst);
+}
+
static void ggml_compute_forward_gelu_quick(
const ggml_compute_params * params,
ggml_tensor * dst) {
@@ -2567,6 +2633,10 @@ static void ggml_compute_forward_gelu_quick(
{
ggml_compute_forward_gelu_quick_f16(params, dst);
} break;
+ case GGML_TYPE_BF16:
+ {
+ ggml_compute_forward_gelu_quick_bf16(params, dst);
+ } break;
default:
{
GGML_ABORT("fatal error");
@@ -2671,6 +2741,12 @@ static void ggml_compute_forward_silu_f16(
}
}
+static void ggml_compute_forward_silu_bf16(
+ const ggml_compute_params * params,
+ ggml_tensor * dst) {
+ ggml_compute_forward_unary_bf16_rows<ggml_vec_silu_bf16>(params, dst);
+}
+
static void ggml_compute_forward_silu(
const ggml_compute_params * params,
ggml_tensor * dst) {
@@ -2686,6 +2762,10 @@ static void ggml_compute_forward_silu(
{
ggml_compute_forward_silu_f16(params, dst);
} break;
+ case GGML_TYPE_BF16:
+ {
+ ggml_compute_forward_silu_bf16(params, dst);
+ } break;
default:
{
GGML_ABORT("fatal error");
@@ -2889,6 +2969,59 @@ void ggml_compute_forward_silu_back(
// ggml_compute_forward_reglu
+// row loop shared by the BF16 GLU ops
+template <void (*vec_op)(const int, ggml_bf16_t *, const ggml_bf16_t *, const ggml_bf16_t *)>
+static void ggml_compute_forward_glu_bf16_rows(
+ const ggml_compute_params * params,
+ ggml_tensor * dst) {
+
+ const ggml_tensor * src0 = dst->src[0];
+ const ggml_tensor * src1 = dst->src[1];
+ char * src0_d = (char *) src0->data;
+ char * src1_d = (char *) (src1 ? src1->data : src0->data);
+ const size_t src0_o = src0->nb[1];
+ const size_t src1_o = src1 ? src1->nb[1] : src0->nb[1];
+
+ GGML_ASSERT(ggml_is_contiguous_1(src0));
+ GGML_ASSERT(ggml_is_contiguous_1(dst));
+
+ if (src1) {
+ GGML_ASSERT(ggml_is_contiguous_1(src1));
+ GGML_ASSERT(src0->type == src1->type);
+ }
+
+ const int ith = params->ith;
+ const int nth = params->nth;
+
+ const int nc = src1 ? src0->ne[0] : src0->ne[0] / 2;
+ const int nr = ggml_nrows(src0);
+
+ GGML_ASSERT(dst->ne[0] == nc);
+ GGML_ASSERT(ggml_nrows(dst) == nr);
+
+ const int32_t swapped = ggml_get_op_params_i32(dst, 1);
+
+ // rows per thread
+ const int dr = (nr + nth - 1)/nth;
+
+ // row range for this thread
+ const int ir0 = dr*ith;
+ const int ir1 = MIN(ir0 + dr, nr);
+
+ for (int i1 = ir0; i1 < ir1; i1++) {
+ const ggml_bf16_t * src0_p = (const ggml_bf16_t *) (src0_d + i1*src0_o);
+ const ggml_bf16_t * src1_p = (const ggml_bf16_t *) (src1_d + i1*src1_o);
+ ggml_bf16_t * dst_p = (ggml_bf16_t *) ((char *) dst->data + i1*(dst->nb[1]));
+
+ if (!src1) {
+ src0_p += swapped ? nc : 0;
+ src1_p += swapped ? 0 : nc;
+ }
+
+ vec_op(nc, dst_p, src0_p, src1_p);
+ }
+}
+
static void ggml_compute_forward_reglu_f32(
const ggml_compute_params * params,
ggml_tensor * dst) {
@@ -3008,6 +3141,12 @@ static void ggml_compute_forward_reglu_f16(
}
}
+static void ggml_compute_forward_reglu_bf16(
+ const ggml_compute_params * params,
+ ggml_tensor * dst) {
+ ggml_compute_forward_glu_bf16_rows<ggml_vec_reglu_bf16>(params, dst);
+}
+
static void ggml_compute_forward_reglu(
const ggml_compute_params * params,
ggml_tensor * dst) {
@@ -3023,6 +3162,10 @@ static void ggml_compute_forward_reglu(
{
ggml_compute_forward_reglu_f16(params, dst);
} break;
+ case GGML_TYPE_BF16:
+ {
+ ggml_compute_forward_reglu_bf16(params, dst);
+ } break;
default:
{
GGML_ABORT("fatal error");
@@ -3151,6 +3294,12 @@ static void ggml_compute_forward_geglu_f16(
}
}
+static void ggml_compute_forward_geglu_bf16(
+ const ggml_compute_params * params,
+ ggml_tensor * dst) {
+ ggml_compute_forward_glu_bf16_rows<ggml_vec_geglu_bf16>(params, dst);
+}
+
static void ggml_compute_forward_geglu(
const ggml_compute_params * params,
ggml_tensor * dst) {
@@ -3166,6 +3315,10 @@ static void ggml_compute_forward_geglu(
{
ggml_compute_forward_geglu_f16(params, dst);
} break;
+ case GGML_TYPE_BF16:
+ {
+ ggml_compute_forward_geglu_bf16(params, dst);
+ } break;
default:
{
GGML_ABORT("fatal error");
@@ -3294,6 +3447,12 @@ static void ggml_compute_forward_swiglu_f16(
}
}
+static void ggml_compute_forward_swiglu_bf16(
+ const ggml_compute_params * params,
+ ggml_tensor * dst) {
+ ggml_compute_forward_glu_bf16_rows<ggml_vec_swiglu_bf16>(params, dst);
+}
+
static void ggml_compute_forward_swiglu(
const ggml_compute_params * params,
ggml_tensor * dst) {
@@ -3309,6 +3468,10 @@ static void ggml_compute_forward_swiglu(
{
ggml_compute_forward_swiglu_f16(params, dst);
} break;
+ case GGML_TYPE_BF16:
+ {
+ ggml_compute_forward_swiglu_bf16(params, dst);
+ } break;
default:
{
GGML_ABORT("fatal error");
@@ -3657,6 +3820,12 @@ static void ggml_compute_forward_geglu_erf_f16(
}
}
+static void ggml_compute_forward_geglu_erf_bf16(
+ const ggml_compute_params * params,
+ ggml_tensor * dst) {
+ ggml_compute_forward_glu_bf16_rows<ggml_vec_geglu_erf_bf16>(params, dst);
+}
+
static void ggml_compute_forward_geglu_erf(
const ggml_compute_params * params,
ggml_tensor * dst) {
@@ -3672,6 +3841,10 @@ static void ggml_compute_forward_geglu_erf(
{
ggml_compute_forward_geglu_erf_f16(params, dst);
} break;
+ case GGML_TYPE_BF16:
+ {
+ ggml_compute_forward_geglu_erf_bf16(params, dst);
+ } break;
default:
{
GGML_ABORT("fatal error");
@@ -3800,6 +3973,12 @@ static void ggml_compute_forward_geglu_quick_f16(
}
}
+static void ggml_compute_forward_geglu_quick_bf16(
+ const ggml_compute_params * params,
+ ggml_tensor * dst) {
+ ggml_compute_forward_glu_bf16_rows<ggml_vec_geglu_quick_bf16>(params, dst);
+}
+
static void ggml_compute_forward_geglu_quick(
const ggml_compute_params * params,
ggml_tensor * dst) {
@@ -3815,6 +3994,10 @@ static void ggml_compute_forward_geglu_quick(
{
ggml_compute_forward_geglu_quick_f16(params, dst);
} break;
+ case GGML_TYPE_BF16:
+ {
+ ggml_compute_forward_geglu_quick_bf16(params, dst);
+ } break;
default:
{
GGML_ABORT("fatal error");
@@ -4746,6 +4929,44 @@ static void ggml_compute_forward_scale_f32(
}
}
+static void ggml_compute_forward_scale_bf16(
+ const ggml_compute_params * params,
+ ggml_tensor * dst) {
+
+ const ggml_tensor * src0 = dst->src[0];
+
+ GGML_ASSERT(ggml_is_contiguous(src0));
+ GGML_ASSERT(ggml_is_contiguous(dst));
+ GGML_ASSERT(ggml_are_same_shape(src0, dst));
+
+ float s; // scale factor
+ float b; // bias
+
+ memcpy(&s, (float *) dst->op_params + 0, sizeof(float));
+ memcpy(&b, (float *) dst->op_params + 1, sizeof(float));
+
+ const int ith = params->ith;
+ const int nth = params->nth;
+
+ const int nc = src0->ne[0];
+ const int nr = ggml_nrows(src0);
+
+ // rows per thread
+ const int dr = (nr + nth - 1)/nth;
+
+ // row range for this thread
+ const int ir0 = dr*ith;
+ const int ir1 = MIN(ir0 + dr, nr);
+
+ for (int i1 = ir0; i1 < ir1; i1++) {
+ const ggml_bf16_t * x = (const ggml_bf16_t *) ((const char *) src0->data + i1*src0->nb[1]);
+ ggml_bf16_t * y = (ggml_bf16_t *) ((char *) dst->data + i1*dst->nb[1]);
+ for (int k = 0; k < nc; k++) {
+ y[k] = GGML_FP32_TO_BF16(s*GGML_BF16_TO_FP32(x[k]) + b);
+ }
+ }
+}
+
void ggml_compute_forward_scale(
const ggml_compute_params * params,
ggml_tensor * dst) {
@@ -4757,6 +4978,10 @@ void ggml_compute_forward_scale(
{
ggml_compute_forward_scale_f32(params, dst);
} break;
+ case GGML_TYPE_BF16:
+ {
+ ggml_compute_forward_scale_bf16(params, dst);
+ } break;
default:
{
GGML_ABORT("fatal error");
diff --git a/ggml/src/ggml-cpu/vec.h b/ggml/src/ggml-cpu/vec.h
index ec1f0a14f..59fa88f2c 100644
--- a/ggml/src/ggml-cpu/vec.h
+++ b/ggml/src/ggml-cpu/vec.h
@@ -976,6 +976,12 @@ inline static void ggml_vec_gelu_f16(const int n, ggml_fp16_t * y, const ggml_fp
}
}
+inline static void ggml_vec_gelu_bf16(const int n, ggml_bf16_t * y, const ggml_bf16_t * x) {
+ for (int i = 0; i < n; ++i) {
+ y[i] = GGML_FP32_TO_BF16(ggml_gelu_f32(GGML_BF16_TO_FP32(x[i])));
+ }
+}
+
inline static void ggml_vec_gelu_erf_f16(const int n, ggml_fp16_t * y, const ggml_fp16_t * x) {
for (int i = 0; i < n; ++i) {
float xi = GGML_CPU_FP16_TO_FP32(x[i]);
@@ -984,6 +990,13 @@ inline static void ggml_vec_gelu_erf_f16(const int n, ggml_fp16_t * y, const ggm
}
}
+inline static void ggml_vec_gelu_erf_bf16(const int n, ggml_bf16_t * y, const ggml_bf16_t * x) {
+ for (int i = 0; i < n; ++i) {
+ float xi = GGML_BF16_TO_FP32(x[i]);
+ y[i] = GGML_FP32_TO_BF16(0.5f*xi*(1.0f + erff(xi*SQRT_2_INV)));
+ }
+}
+
#ifdef GGML_GELU_FP16
inline static void ggml_vec_gelu_f32(const int n, float * y, const float * x) {
uint16_t t;
@@ -1025,6 +1038,12 @@ inline static void ggml_vec_gelu_quick_f16(const int n, ggml_fp16_t * y, const g
}
}
+inline static void ggml_vec_gelu_quick_bf16(const int n, ggml_bf16_t * y, const ggml_bf16_t * x) {
+ for (int i = 0; i < n; ++i) {
+ y[i] = GGML_FP32_TO_BF16(ggml_gelu_quick_f32(GGML_BF16_TO_FP32(x[i])));
+ }
+}
+
#ifdef GGML_GELU_QUICK_FP16
inline static void ggml_vec_gelu_quick_f32(const int n, float * y, const float * x) {
uint16_t t;
@@ -1375,6 +1394,12 @@ inline static void ggml_vec_silu_f16(const int n, ggml_fp16_t * y, const ggml_fp
}
}
+inline static void ggml_vec_silu_bf16(const int n, ggml_bf16_t * y, const ggml_bf16_t * x) {
+ for (int i = 0; i < n; ++i) {
+ y[i] = GGML_FP32_TO_BF16(ggml_silu_f32(GGML_BF16_TO_FP32(x[i])));
+ }
+}
+
inline static float ggml_silu_backward_f32(float x, float dy) {
const float s = 1.0f/(1.0f + expf(-x));
return dy*s*(1.0f + x*(1.0f - s));
@@ -1411,6 +1436,13 @@ inline static void ggml_vec_reglu_f16 (const int n, ggml_fp16_t * y, const ggml_
}
}
+inline static void ggml_vec_reglu_bf16(const int n, ggml_bf16_t * y, const ggml_bf16_t * x, const ggml_bf16_t * g) {
+ for (int i = 0; i < n; ++i) {
+ float v = GGML_BF16_TO_FP32(x[i]);
+ y[i] = GGML_FP32_TO_BF16((v > 0.f) ? v * GGML_BF16_TO_FP32(g[i]) : 0.f);
+ }
+}
+
#ifdef GGML_GELU_FP16
inline static void ggml_vec_geglu_f32(const int n, float * y, const float * x, const float * g) {
uint16_t t;
@@ -1442,6 +1474,12 @@ inline static void ggml_vec_geglu_f16(const int n, ggml_fp16_t * y, const ggml_f
}
}
+inline static void ggml_vec_geglu_bf16(const int n, ggml_bf16_t * y, const ggml_bf16_t * x, const ggml_bf16_t * g) {
+ for (int i = 0; i < n; ++i) {
+ y[i] = GGML_FP32_TO_BF16(ggml_gelu_f32(GGML_BF16_TO_FP32(x[i])) * GGML_BF16_TO_FP32(g[i]));
+ }
+}
+
void ggml_vec_swiglu_f32(const int n, float * y, const float * x, const float * g);
inline static void ggml_vec_swiglu_f16(const int n, ggml_fp16_t * y, const ggml_fp16_t * x, const ggml_fp16_t * g) {
@@ -1452,6 +1490,14 @@ inline static void ggml_vec_swiglu_f16(const int n, ggml_fp16_t * y, const ggml_
}
}
+inline static void ggml_vec_swiglu_bf16(const int n, ggml_bf16_t * y, const ggml_bf16_t * x, const ggml_bf16_t * g) {
+ for (int i = 0; i < n; ++i) {
+ float xi = GGML_BF16_TO_FP32(x[i]);
+ float gi = GGML_BF16_TO_FP32(g[i]);
+ y[i] = GGML_FP32_TO_BF16((xi/(1.0f + expf(-xi))) * gi);
+ }
+}
+
inline static void ggml_vec_geglu_erf_f32(const int n, float * y, const float * x, const float * g) {
for (int i = 0; i < n; ++i) {
float xi = x[i];
@@ -1467,6 +1513,14 @@ inline static void ggml_vec_geglu_erf_f16(const int n, ggml_fp16_t * y, const gg
}
}
+inline static void ggml_vec_geglu_erf_bf16(const int n, ggml_bf16_t * y, const ggml_bf16_t * x, const ggml_bf16_t * g) {
+ for (int i = 0; i < n; ++i) {
+ float xi = GGML_BF16_TO_FP32(x[i]);
+ float gi = GGML_BF16_TO_FP32(g[i]);
+ y[i] = GGML_FP32_TO_BF16(0.5f * xi * (1.0f + erff(xi*SQRT_2_INV)) * gi);
+ }
+}
+
#ifdef GGML_GELU_QUICK_FP16
inline static void ggml_vec_geglu_quick_f32(const int n, float * y, const float * x, const float * g) {
uint16_t t;
@@ -1492,6 +1546,12 @@ inline static void ggml_vec_geglu_quick_f16(const int n, ggml_fp16_t * y, const
}
}
+inline static void ggml_vec_geglu_quick_bf16(const int n, ggml_bf16_t * y, const ggml_bf16_t * x, const ggml_bf16_t * g) {
+ for (int i = 0; i < n; ++i) {
+ y[i] = GGML_FP32_TO_BF16(ggml_gelu_quick_f32(GGML_BF16_TO_FP32(x[i])) * GGML_BF16_TO_FP32(g[i]));
+ }
+}
+
inline static void ggml_vec_sum_f32(const int n, float * s, const float * x) {
#ifndef GGML_USE_ACCELERATE
ggml_float sum = 0.0;
diff --git a/ggml/src/ggml-cuda/binbcast.cu b/ggml/src/ggml-cuda/binbcast.cu
index 2e38077bf..89ae2bc78 100644
--- a/ggml/src/ggml-cuda/binbcast.cu
+++ b/ggml/src/ggml-cuda/binbcast.cu
@@ -1,4 +1,5 @@
#include "binbcast.cuh"
+#include "convert.cuh"
#include <cstdint>
#include <utility>
@@ -82,14 +83,14 @@ static __global__ void k_bin_bcast(const src0_t * src0,
for (uint32_t i0 = i0s; i0 < ne0; i0 += s0) {
const uint32_t i10 = fastmodulo(i0, ne10);
- float result = src0_row ? (float) src0_row[size_t(i0)*s00] : 0.0f;
+ float result = src0_row ? ggml_cuda_cast<float>(src0_row[size_t(i0)*s00]) : 0.0f;
if constexpr (sizeof...(src1_ptrs) > 0) {
- result = (..., (result = bin_op(result, (float)src1s[i_src1 + size_t(i10)*s10])));
+ result = (..., (result = bin_op(result, ggml_cuda_cast<float>(src1s[i_src1 + size_t(i10)*s10]))));
} else {
- result = bin_op(result, (float)src1[i_src1 + size_t(i10)*s10]);
+ result = bin_op(result, ggml_cuda_cast<float>(src1[i_src1 + size_t(i10)*s10]));
}
- dst_row[i0] = (dst_t) result;
+ dst_row[i0] = ggml_cuda_cast<dst_t>(result);
// protect i0 from overflow
if (ne0 - i0 <= s0) {
@@ -154,14 +155,14 @@ static __global__ void k_bin_bcast_unravel(const src0_t * src0,
const uint32_t i10 = fastmodulo(i0, ne10);
ggml_cuda_pdl_sync();
- float result = src0_row ? (float) src0_row[size_t(i0)*s00] : 0.0f;
+ float result = src0_row ? ggml_cuda_cast<float>(src0_row[size_t(i0)*s00]) : 0.0f;
if constexpr (sizeof...(src1_ptrs) > 0) {
- result = (..., (result = bin_op(result, (float)src1s[i_src1 + size_t(i10)*s10])));
+ result = (..., (result = bin_op(result, ggml_cuda_cast<float>(src1s[i_src1 + size_t(i10)*s10]))));
} else {
- result = bin_op(result, (float)src1[i_src1 + size_t(i10)*s10]);
+ result = bin_op(result, ggml_cuda_cast<float>(src1[i_src1 + size_t(i10)*s10]));
}
- dst_row[i0] = (dst_t) result;
+ dst_row[i0] = ggml_cuda_cast<dst_t>(result);
}
template <float (*bin_op)(const float, const float), typename src0_t, typename src1_t, typename dst_t, size_t... I>
@@ -413,10 +414,14 @@ static void ggml_cuda_op_bin_bcast(
const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst,
const void * src0_dd, const void * src1_dd, void * dst_dd, cudaStream_t stream) {
- GGML_ASSERT(src1->type == GGML_TYPE_F32 || src1->type == GGML_TYPE_F16);
+ GGML_ASSERT(src1->type == GGML_TYPE_F32 || src1->type == GGML_TYPE_F16 || src1->type == GGML_TYPE_BF16);
- if (src0->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) {
+ if (src0->type == GGML_TYPE_F32 && src1->type != GGML_TYPE_BF16 && dst->type == GGML_TYPE_F32) {
op()(src0, src1, dst, (const float *)src0_dd, (const float *)src1_dd, (float *)dst_dd, stream);
+ } else if (src0->type == GGML_TYPE_BF16 && src1->type == GGML_TYPE_BF16 && dst->type == GGML_TYPE_BF16) {
+ op()(src0, src1, dst, (const nv_bfloat16 *) src0_dd, (const nv_bfloat16 *) src1_dd, (nv_bfloat16 *) dst_dd, stream);
+ } else if (src0->type == GGML_TYPE_BF16 && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_BF16) {
+ op()(src0, src1, dst, (const nv_bfloat16 *) src0_dd, (const float *) src1_dd, (nv_bfloat16 *) dst_dd, stream);
} else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F16 && dst->type == GGML_TYPE_F16) {
op()(src0, src1, dst, (const half *) src0_dd, (const half *)src1_dd, (half *) dst_dd, stream);
} else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F16) {
@@ -457,10 +462,18 @@ static void ggml_cuda_op_fused_binbcast_impl(ggml_backend_cuda_context & ctx, gg
const ggml_tensor * src0 = dst->src[0];
const ggml_tensor * src1 = dst->src[1];
- if (src0->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) {
+ if (src0->type == GGML_TYPE_F32 && src1->type != GGML_TYPE_BF16 && dst->type == GGML_TYPE_F32) {
launch_bin_bcast_pack<op, float, float, float>(src0, src1, dst,
(const float *) src0->data, (const float *) src1->data, (float *) dst->data,
stream, std::make_index_sequence<n_fuse>{});
+ } else if (src0->type == GGML_TYPE_BF16 && src1->type == GGML_TYPE_BF16 && dst->type == GGML_TYPE_BF16) {
+ launch_bin_bcast_pack<op, nv_bfloat16, nv_bfloat16, nv_bfloat16>(src0, src1, dst,
+ (const nv_bfloat16 *) src0->data, (const nv_bfloat16 *) src1->data, (nv_bfloat16 *) dst->data,
+ stream, std::make_index_sequence<n_fuse>{});
+ } else if (src0->type == GGML_TYPE_BF16 && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_BF16) {
+ launch_bin_bcast_pack<op, nv_bfloat16, float, nv_bfloat16>(src0, src1, dst,
+ (const nv_bfloat16 *) src0->data, (const float *) src1->data, (nv_bfloat16 *) dst->data,
+ stream, std::make_index_sequence<n_fuse>{});
} else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F16 && dst->type == GGML_TYPE_F16) {
launch_bin_bcast_pack<op, half, half, half>(src0, src1, dst,
(const half *) src0->data, (const half *) src1->data, (half *) dst->data,
diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
index 9afc2aa4a..ebbb5c106 100644
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -3371,7 +3371,7 @@ static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph,
return false;
}
- if (unary->type != GGML_TYPE_F32 && unary->type != GGML_TYPE_F16) {
+ if (unary->type != GGML_TYPE_F32 && unary->type != GGML_TYPE_F16 && unary->type != GGML_TYPE_BF16) {
return false;
}
@@ -3420,8 +3420,10 @@ static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph,
const ggml_tensor *tanh = cgraph->nodes[node_idx+1];
const ggml_tensor *scale2 = cgraph->nodes[node_idx+2];
- GGML_ASSERT(scale->src[0]->type == GGML_TYPE_F32);
- GGML_ASSERT(scale->type == GGML_TYPE_F32);
+ // the fused softcap kernel is F32 only
+ if (scale->src[0]->type != GGML_TYPE_F32 || scale->type != GGML_TYPE_F32) {
+ return false;
+ }
if (ggml_get_unary_op(tanh) != GGML_UNARY_OP_TANH) {
return false;
@@ -5182,6 +5184,9 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
case GGML_UNARY_OP_CEIL:
case GGML_UNARY_OP_ROUND:
case GGML_UNARY_OP_TRUNC:
+ if (op->src[0]->type == GGML_TYPE_BF16 && ggml_get_unary_op(op) == GGML_UNARY_OP_XIELU) {
+ return false;
+ }
// TODO: should become:
//return ggml_is_contiguous_rows(op->src[0]);
return ggml_is_contiguous(op->src[0]);
@@ -5198,6 +5203,10 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
case GGML_GLU_OP_GEGLU_ERF:
case GGML_GLU_OP_GEGLU_QUICK:
case GGML_GLU_OP_SWIGLU_CLAMP:
+ if (op->src[0]->type == GGML_TYPE_BF16 &&
+ (ggml_get_glu_op(op) == GGML_GLU_OP_SWIGLU_OAI || ggml_get_glu_op(op) == GGML_GLU_OP_SWIGLU_CLAMP)) {
+ return false;
+ }
return ggml_is_contiguous_1(op->src[0]);
default:
return false;
@@ -5467,7 +5476,6 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
case GGML_OP_TRANSPOSE:
case GGML_OP_ADD_ID:
case GGML_OP_ADD1:
- case GGML_OP_SCALE:
case GGML_OP_SQR:
case GGML_OP_SQRT:
case GGML_OP_SIN:
@@ -5475,10 +5483,16 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
case GGML_OP_CLAMP:
case GGML_OP_LOG:
return true;
+ case GGML_OP_SCALE:
+ return (op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_BF16) && op->type == op->src[0]->type;
case GGML_OP_ADD:
case GGML_OP_SUB:
case GGML_OP_MUL:
case GGML_OP_DIV:
+ if (op->src[0]->type == GGML_TYPE_BF16 || op->src[1]->type == GGML_TYPE_BF16 || op->type == GGML_TYPE_BF16) {
+ return op->src[0]->type == GGML_TYPE_BF16 && op->type == GGML_TYPE_BF16 &&
+ (op->src[1]->type == GGML_TYPE_BF16 || op->src[1]->type == GGML_TYPE_F32);
+ }
return (op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16) &&
(op->src[1]->type == GGML_TYPE_F32 || op->src[1]->type == GGML_TYPE_F16) &&
(op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16);
diff --git a/ggml/src/ggml-cuda/scale.cu b/ggml/src/ggml-cuda/scale.cu
index 7b2e59a43..5648b9261 100644
--- a/ggml/src/ggml-cuda/scale.cu
+++ b/ggml/src/ggml-cuda/scale.cu
@@ -1,37 +1,42 @@
#include "scale.cuh"
+#include "convert.cuh"
#define MAX_GRIDDIM_X 0x7FFFFFFF
-static __global__ void scale_f32(const float * x, float * dst, const float scale, const float bias, const int64_t nelements) {
+template <typename T>
+static __global__ void scale_f32(const T * x, T * dst, const float scale, const float bias, const int64_t nelements) {
ggml_cuda_pdl_lc();
int64_t tid = (int64_t)blockIdx.x * (int64_t)blockDim.x + (int64_t)threadIdx.x;
int64_t stride = (int64_t)blockDim.x * (int64_t)gridDim.x;
ggml_cuda_pdl_sync();
for (int64_t i = tid; i < nelements; i += stride) {
- dst[i] = scale * x[i] + bias;
+ dst[i] = ggml_cuda_cast<T>(scale * ggml_cuda_cast<float>(x[i]) + bias);
}
}
-static void scale_f32_cuda(const float * x, float * dst, const float scale, const float bias, const int64_t nelements, cudaStream_t stream) {
+template <typename T>
+static void scale_f32_cuda(const T * x, T * dst, const float scale, const float bias, const int64_t nelements, cudaStream_t stream) {
const int64_t num_blocks = (nelements + CUDA_SCALE_BLOCK_SIZE - 1) / CUDA_SCALE_BLOCK_SIZE;
const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(MIN(MAX_GRIDDIM_X, num_blocks), CUDA_SCALE_BLOCK_SIZE, 0, stream);
- ggml_cuda_kernel_launch(scale_f32, launch_params, x, dst, scale, bias, nelements);
+ ggml_cuda_kernel_launch(scale_f32<T>, launch_params, x, dst, scale, bias, nelements);
}
void ggml_cuda_op_scale(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
const ggml_tensor * src0 = dst->src[0];
- const float * src0_d = (const float *)src0->data;
- float * dst_d = (float *)dst->data;
cudaStream_t stream = ctx.stream();
- GGML_ASSERT(src0->type == GGML_TYPE_F32);
- GGML_ASSERT( dst->type == GGML_TYPE_F32);
+ GGML_ASSERT(src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_BF16);
+ GGML_ASSERT(dst->type == src0->type);
float scale;
float bias;
memcpy(&scale, (float *) dst->op_params + 0, sizeof(float));
memcpy(&bias, (float *) dst->op_params + 1, sizeof(float));
- scale_f32_cuda(src0_d, dst_d, scale, bias, ggml_nelements(src0), stream);
+ if (src0->type == GGML_TYPE_BF16) {
+ scale_f32_cuda((const nv_bfloat16 *) src0->data, (nv_bfloat16 *) dst->data, scale, bias, ggml_nelements(src0), stream);
+ } else {
+ scale_f32_cuda((const float *) src0->data, (float *) dst->data, scale, bias, ggml_nelements(src0), stream);
+ }
}
diff --git a/ggml/src/ggml-cuda/unary.cu b/ggml/src/ggml-cuda/unary.cu
index d3e594878..84788c5dc 100644
--- a/ggml/src/ggml-cuda/unary.cu
+++ b/ggml/src/ggml-cuda/unary.cu
@@ -124,7 +124,7 @@ static __global__ void unary_op_kernel(const T * x, T * dst, const int k) {
}
ggml_cuda_pdl_sync();
- dst[i] = (T)op((float)x[i]);
+ dst[i] = ggml_cuda_cast<T>(op(ggml_cuda_cast<float>(x[i])));
}
template <float (*op)(float), typename T>
@@ -143,12 +143,13 @@ void ggml_cuda_op_unary(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
GGML_ASSERT(ggml_is_contiguous(src0));
- GGML_ASSERT(src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_F16);
- GGML_ASSERT( dst->type == GGML_TYPE_F32 || dst->type == GGML_TYPE_F16);
+ GGML_ASSERT(src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_F16 || src0->type == GGML_TYPE_BF16);
GGML_ASSERT(src0->type == dst->type);
if (src0->type == GGML_TYPE_F16) {
unary_cuda<op>((const half *)src0_d, (half *)dst_d, ggml_nelements(src0), stream);
+ } else if (src0->type == GGML_TYPE_BF16) {
+ unary_cuda<op>((const nv_bfloat16 *)src0_d, (nv_bfloat16 *)dst_d, ggml_nelements(src0), stream);
} else {
unary_cuda<op>((const float *)src0_d, (float *)dst_d, ggml_nelements(src0), stream);
}
@@ -273,7 +274,7 @@ static __global__ void unary_gated_op_kernel(const T * x, const T * g, T * dst,
const int64_t j1 = o0 == o1 ? j0 : (i / n) * o1 + (i % n);
ggml_cuda_pdl_sync();
- dst[i] = (T)(op((float)x[j0]) * (float)g[j1]);
+ dst[i] = ggml_cuda_cast<T>(op(ggml_cuda_cast<float>(x[j0])) * ggml_cuda_cast<float>(g[j1]));
}
template <float (*op)(float), typename T>
@@ -299,8 +300,7 @@ void ggml_cuda_op_unary_gated(ggml_backend_cuda_context & ctx, ggml_tensor * dst
GGML_ASSERT(src0->nb[0] == ggml_element_size(src0));
GGML_ASSERT(ggml_is_contiguous(dst));
- GGML_ASSERT(src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_F16);
- GGML_ASSERT( dst->type == GGML_TYPE_F32 || dst->type == GGML_TYPE_F16);
+ GGML_ASSERT(src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_F16 || src0->type == GGML_TYPE_BF16);
GGML_ASSERT(src0->type == dst->type);
GGML_ASSERT(dst->ne[0] == nc);
GGML_ASSERT(ggml_nrows(dst) == ggml_nrows(src0));
@@ -324,6 +324,16 @@ void ggml_cuda_op_unary_gated(ggml_backend_cuda_context & ctx, ggml_tensor * dst
}
unary_gated_cuda<op>(src0_p, src1_p, (half *)dst_d, ggml_nelements(dst), nc, src0_o / sizeof(half), src1_o / sizeof(half), stream);
+ } else if (src0->type == GGML_TYPE_BF16) {
+ nv_bfloat16 * src0_p = (nv_bfloat16 *) src0_d;
+ nv_bfloat16 * src1_p = (nv_bfloat16 *) src1_d;
+
+ if (!src1) {
+ src0_p += swapped ? nc : 0;
+ src1_p += swapped ? 0 : nc;
+ }
+
+ unary_gated_cuda<op>(src0_p, src1_p, (nv_bfloat16 *)dst_d, ggml_nelements(dst), nc, src0_o / sizeof(nv_bfloat16), src1_o / sizeof(nv_bfloat16), stream);
} else {
float * src0_p = (float *) src0_d;
float * src1_p = (float *) src1_d;
@@ -664,7 +674,7 @@ static void ggml_cuda_op_unary_mul_impl(ggml_backend_cuda_context & ctx, ggml_te
GGML_ASSERT(other_src->nb[0] == ggml_element_size(other_src));
GGML_ASSERT(ggml_are_same_shape(unary_src, other_src));
- GGML_ASSERT(unary_src->type == GGML_TYPE_F32 || unary_src->type == GGML_TYPE_F16);
+ GGML_ASSERT(unary_src->type == GGML_TYPE_F32 || unary_src->type == GGML_TYPE_F16 || unary_src->type == GGML_TYPE_BF16);
GGML_ASSERT(unary_src->type == other_src->type);
GGML_ASSERT(unary_src->type == mul_node->type);
@@ -679,6 +689,10 @@ static void ggml_cuda_op_unary_mul_impl(ggml_backend_cuda_context & ctx, ggml_te
unary_gated_cuda<op>((const half *) unary_src->data, (const half *) other_src->data,
(half *) mul_node->data, k, nc,
unary_stride / sizeof(half), other_stride / sizeof(half), stream);
+ } else if (unary_src->type == GGML_TYPE_BF16) {
+ unary_gated_cuda<op>((const nv_bfloat16 *) unary_src->data, (const nv_bfloat16 *) other_src->data,
+ (nv_bfloat16 *) mul_node->data, k, nc,
+ unary_stride / sizeof(nv_bfloat16), other_stride / sizeof(nv_bfloat16), stream);
} else {
unary_gated_cuda<op>((const float *) unary_src->data, (const float *) other_src->data,
(float *) mul_node->data, k, nc,
diff --git a/ggml/src/ggml-openvino/ggml-openvino.cpp b/ggml/src/ggml-openvino/ggml-openvino.cpp
index e15f1c1ab..d5050bff7 100644
--- a/ggml/src/ggml-openvino/ggml-openvino.cpp
+++ b/ggml/src/ggml-openvino/ggml-openvino.cpp
@@ -1207,6 +1207,10 @@ static ggml_openvino_op_support is_op_supported_case(const ggml_tensor * op) {
if (op->src[1]->op == GGML_OP_PERMUTE) {
return {false, "ADD/MUL/SUB with PERMUTE src1 is not supported"};
}
+ if (op->src[0]->type != op->src[1]->type &&
+ (op->src[0]->type == GGML_TYPE_BF16 || op->src[1]->type == GGML_TYPE_BF16)) {
+ return {false, "ADD/MUL/SUB with BF16 and a different src1 type is not supported"};
+ }
// >8-expert MoE ReduceSum drifts past the 1e-7 tolerance (f32 order vs CPU); intermittent.
if (op->op == GGML_OP_ADD && is_moe_expert_sum_add(op) && op->src[1]->src[0]->ne[1] > 8) {
return {false, "MoE expert-plane sum with more than 8 experts is not supported"};
@@ -1220,6 +1224,12 @@ static ggml_openvino_op_support is_op_supported_case(const ggml_tensor * op) {
}
break;
}
+ case GGML_OP_SCALE: {
+ if (op->type == GGML_TYPE_BF16) {
+ return {false, "SCALE with BF16 type is not supported"};
+ }
+ break;
+ }
case GGML_OP_ADD_ID: {
// Keep support aligned with the CPU backend implementation, which only handles f32 inputs/output and i32 ids.
if (op->type != GGML_TYPE_F32 || op->src[0]->type != GGML_TYPE_F32 || op->src[1]->type != GGML_TYPE_F32 ||
diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index 23b23147d..5dc9cef3d 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -3295,11 +3295,16 @@ struct test_bin_bcast : public test_case {
int nf; // number of fused ops, nf == 1 -> single op (no fusion)
bool perm1; // permute src1?
bool src_overlap; // src0 and src1 are overlapping views of the same buffer
+ const ggml_type type_b; // type of src1, GGML_TYPE_COUNT -> same as src0
bool run_whole_graph() override { return nf > 1; }
std::string vars() override {
- return VARS_TO_STR6(type, ne, nr, nf, perm1, src_overlap);
+ std::string s = VARS_TO_STR6(type, ne, nr, nf, perm1, src_overlap);
+ if (type_b != GGML_TYPE_COUNT) {
+ s += "," + VAR_TO_STR(type_b);
+ }
+ return s;
}
size_t op_size(ggml_tensor * t) override {
@@ -3310,8 +3315,9 @@ struct test_bin_bcast : public test_case {
std::array<int64_t, 4> ne = {10, 10, 1, 1},
std::array<int, 4> nr = {1, 2, 1, 1},
int nf = 1,
- bool perm1 = false, bool src_overlap = false)
- : op(op), type(type), ne(ne), nr(nr), nf(nf), perm1(perm1), src_overlap(src_overlap) {}
+ bool perm1 = false, bool src_overlap = false,
+ ggml_type type_b = GGML_TYPE_COUNT)
+ : op(op), type(type), ne(ne), nr(nr), nf(nf), perm1(perm1), src_overlap(src_overlap), type_b(type_b) {}
ggml_tensor * build_graph(ggml_context * ctx) override {
GGML_ASSERT(nf <= 16);
@@ -3319,17 +3325,19 @@ struct test_bin_bcast : public test_case {
ggml_tensor * a = ggml_new_tensor_4d(ctx, type, ne[0]*nr[0], ne[1]*nr[1], ne[2]*nr[2], ne[3]*nr[3]);
ggml_set_name(a, "a");
+ const ggml_type tb = type_b == GGML_TYPE_COUNT ? type : type_b;
+
ggml_tensor * b[16];
for (int i = 0; i < nf; ++i) {
if (perm1) {
const int p[4] = { 1, 2, 0, 3 }; // hardcoded for now
- b[i] = ggml_new_tensor_4d(ctx, type, ne[p[0]], ne[p[1]], ne[p[2]], ne[p[3]]);
+ b[i] = ggml_new_tensor_4d(ctx, tb, ne[p[0]], ne[p[1]], ne[p[2]], ne[p[3]]);
b[i] = ggml_permute(ctx, b[i], p[0], p[1], p[2], p[3]);
} else if (src_overlap) {
b[i] = ggml_view_4d(ctx, a, ne[0], ne[1], ne[2], 2 * (ne[3] / 3), a->nb[1], a->nb[2], a->nb[3], (ne[3] / 3) * a->nb[3]);
} else {
- b[i] = ggml_new_tensor(ctx, type, 4, ne.data());
+ b[i] = ggml_new_tensor(ctx, tb, 4, ne.data());
}
ggml_set_name(b[i], (std::string("b") + std::to_string(i)).c_str());
}
@@ -3378,9 +3386,14 @@ struct test_bin_bcast : public test_case {
}
double max_nmse_err() override {
- if (op == ggml_add && type == GGML_TYPE_F16 && nf > 1) {
- // Fused ADDs can keep FP32 intermediates while the CPU rounds each ADD to FP16.
- return 1e-6;
+ if (op == ggml_add && nf > 1) {
+ // Fused ADDs can keep FP32 intermediates while the CPU rounds each ADD to FP16/BF16.
+ if (type == GGML_TYPE_F16) {
+ return 1e-6;
+ }
+ if (type == GGML_TYPE_BF16) {
+ return 1e-4;
+ }
}
return test_case::max_nmse_err();
}
@@ -4015,9 +4028,10 @@ struct test_unary_mul : public test_case {
// the fused kernel elides the rounding of the unary result that the CPU chain
// performs; relax the tolerance to match that drift
switch (type) {
- case GGML_TYPE_F16: return 5e-5;
+ case GGML_TYPE_F16: return 5e-5;
+ case GGML_TYPE_BF16: return 1e-4;
// gelu shader uses exp form, CPU uses tanhf
- default: return op == GGML_UNARY_OP_GELU ? 5e-7 : 1e-7;
+ default: return op == GGML_UNARY_OP_GELU ? 5e-7 : 1e-7;
}
}
@@ -9169,7 +9183,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
std::default_random_engine rng(0);
// unary ops
- for (ggml_type type : {GGML_TYPE_F16, GGML_TYPE_F32}) {
+ for (ggml_type type : {GGML_TYPE_F16, GGML_TYPE_F32, GGML_TYPE_BF16}) {
for (int v : {0, 1}) {
for (int op = 0; op < GGML_UNARY_OP_COUNT; op++) {
if (op == GGML_UNARY_OP_XIELU) {
@@ -9189,7 +9203,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
// fused unary + mul (gated activations that are not expressed as GGML_OP_GLU)
for (ggml_unary_op op : { GGML_UNARY_OP_GELU, GGML_UNARY_OP_SILU, GGML_UNARY_OP_SIGMOID, GGML_UNARY_OP_SOFTPLUS }) {
- for (ggml_type type : { GGML_TYPE_F16, GGML_TYPE_F32 }) {
+ for (ggml_type type : { GGML_TYPE_F16, GGML_TYPE_F32, GGML_TYPE_BF16 }) {
for (bool swap : { false, true }) {
test_cases.emplace_back(new test_unary_mul(op, type, { 128, 2, 2, 2 }, swap));
}
@@ -9254,7 +9268,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
test_cases.emplace_back(new test_dsv4_hc_post(31, 17, false, true));
// glu ops
- for (ggml_type type : {GGML_TYPE_F16, GGML_TYPE_F32}) {
+ for (ggml_type type : {GGML_TYPE_F16, GGML_TYPE_F32, GGML_TYPE_BF16}) {
for (int v : {0, 1}) {
for (int op = 0; op < GGML_GLU_OP_COUNT; op++) {
if (op == GGML_GLU_OP_SWIGLU_OAI || op == GGML_GLU_OP_SWIGLU_CLAMP) {
@@ -9864,7 +9878,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
test_cases.emplace_back(new test_bin_bcast(op, type, ne, nr, 1, perm1, src_overlap));
}
};
- for (ggml_type type : {GGML_TYPE_F16, GGML_TYPE_F32}) {
+ for (ggml_type type : {GGML_TYPE_F16, GGML_TYPE_F32, GGML_TYPE_BF16}) {
for (bool perm1 : {false, true}) {
add_test_bin_bcast(type, {1, 1, 8, 1}, {1, 1, 1, 1}, perm1);
add_test_bin_bcast(type, {1, 1, 1, 1}, {32, 1, 1, 1}, perm1);
@@ -9916,9 +9930,18 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
test_cases.emplace_back(new test_bin_bcast(ggml_sub_inplace, GGML_TYPE_F32, {16, 5, 4, 3}, {1, 1, 1, 1}, 16));
test_cases.emplace_back(new test_bin_bcast(ggml_div_inplace, GGML_TYPE_F32, {16, 5, 4, 3}, {1, 1, 1, 1}, 16));
+ // BF16 activations with F32 weights
+ for (auto op : {ggml_add, ggml_sub, ggml_mul, ggml_div}) {
+ test_cases.emplace_back(new test_bin_bcast(op, GGML_TYPE_BF16, {10, 5, 4, 3}, {1, 1, 1, 1}, 1, false, false, GGML_TYPE_F32));
+ test_cases.emplace_back(new test_bin_bcast(op, GGML_TYPE_BF16, {10, 5, 4, 3}, {2, 2, 2, 2}, 1, false, false, GGML_TYPE_F32));
+ test_cases.emplace_back(new test_bin_bcast(op, GGML_TYPE_BF16, {5120, 1, 1, 1}, {1, 32, 1, 1}, 1, false, false, GGML_TYPE_F32));
+ }
+
// fusion
test_cases.emplace_back(new test_bin_bcast(ggml_add, GGML_TYPE_F32, {10, 5, 4, 3}, {2, 1, 1, 1}, 2));
test_cases.emplace_back(new test_bin_bcast(ggml_add, GGML_TYPE_F16, {10, 5, 4, 3}, {2, 1, 1, 1}, 2));
+ test_cases.emplace_back(new test_bin_bcast(ggml_add, GGML_TYPE_BF16, {10, 5, 4, 3}, {2, 1, 1, 1}, 2));
+ test_cases.emplace_back(new test_bin_bcast(ggml_add, GGML_TYPE_BF16, {16, 5, 4, 3}, {1, 2, 1, 1}, 3));
test_cases.emplace_back(new test_bin_bcast(ggml_add, GGML_TYPE_F32, {16, 5, 4, 3}, {1, 1, 1, 1}, 2, true));
test_cases.emplace_back(new test_bin_bcast(ggml_add, GGML_TYPE_F32, {16, 5, 4, 3}, {1, 2, 1, 1}, 3));
test_cases.emplace_back(new test_bin_bcast(ggml_add, GGML_TYPE_F32, {10, 5, 4, 3}, {1, 1, 2, 1}, 4));
@@ -9932,6 +9955,9 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
test_cases.emplace_back(new test_scale(GGML_TYPE_F32, {10, 10, 10, 10}, 2.0f, 1.0f));
test_cases.emplace_back(new test_scale(GGML_TYPE_F32, {10, 10, 10, 10}, 2.0f, 1.0f, true)); // inplace test
test_cases.emplace_back(new test_scale(GGML_TYPE_F32, {100, 10, 10, 10}, 2.0f, 1.0f));
+ test_cases.emplace_back(new test_scale(GGML_TYPE_BF16, {10, 10, 10, 10}, 2.0f, 1.0f));
+ test_cases.emplace_back(new test_scale(GGML_TYPE_BF16, {10, 10, 10, 10}, 2.0f, 1.0f, true)); // inplace test
+ test_cases.emplace_back(new test_scale(GGML_TYPE_BF16, {100, 10, 10, 10}, 0.5f, 0.0f));
test_cases.emplace_back(new test_softcap(GGML_TYPE_F32, {10, 10, 10, 10}, 50.0f));
test_cases.emplace_back(new test_silu_back());
@@ -11356,7 +11382,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
// SWIGLU at a 27B-class FFN width, fused [gate|up] vs split operands
// note: same bytes either way, so a backend that indexes them differently shows it here
- for (ggml_type type : {GGML_TYPE_F16, GGML_TYPE_F32}) {
+ for (ggml_type type : {GGML_TYPE_F16, GGML_TYPE_F32, GGML_TYPE_BF16}) {
for (int64_t n_tokens : {512, 2048}) {
test_cases.emplace_back(new test_glu(GGML_GLU_OP_SWIGLU, type, { 2*17408, n_tokens, 1, 1 }, 0, false));
test_cases.emplace_back(new test_glu_split(GGML_GLU_OP_SWIGLU, type, { 17408, n_tokens, 1, 1 }, 0));