From fd66b6eccec135e397daf0067cf7a3d82b595b5f Mon Sep 17 00:00:00 2001 From: Aman Gupta Date: Wed, 30 Sep 2026 20:24:59 +0800 Subject: [PATCH] ggml : add BF16 unary, GLU, binary and scale ops (CPU, CUDA) (llama/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 --- ggml/src/ggml-cpu/ops.cpp | 225 +++++++++++++++++++++++ ggml/src/ggml-cpu/vec.h | 60 ++++++ ggml/src/ggml-cuda/binbcast.cu | 35 ++-- ggml/src/ggml-cuda/ggml-cuda.cu | 22 ++- ggml/src/ggml-cuda/scale.cu | 23 ++- ggml/src/ggml-cuda/unary.cu | 28 ++- ggml/src/ggml-openvino/ggml-openvino.cpp | 10 + 7 files changed, 372 insertions(+), 31 deletions(-) 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 +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(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(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(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(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 +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(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(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(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(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(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 #include @@ -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(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(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(src1[i_src1 + size_t(i10)*s10])); } - dst_row[i0] = (dst_t) result; + dst_row[i0] = ggml_cuda_cast(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(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(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(src1[i_src1 + size_t(i10)*s10])); } - dst_row[i0] = (dst_t) result; + dst_row[i0] = ggml_cuda_cast(result); } template @@ -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(src0, src1, dst, (const float *) src0->data, (const float *) src1->data, (float *) dst->data, stream, std::make_index_sequence{}); + } else if (src0->type == GGML_TYPE_BF16 && src1->type == GGML_TYPE_BF16 && dst->type == GGML_TYPE_BF16) { + launch_bin_bcast_pack(src0, src1, dst, + (const nv_bfloat16 *) src0->data, (const nv_bfloat16 *) src1->data, (nv_bfloat16 *) dst->data, + stream, std::make_index_sequence{}); + } else if (src0->type == GGML_TYPE_BF16 && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_BF16) { + launch_bin_bcast_pack(src0, src1, dst, + (const nv_bfloat16 *) src0->data, (const float *) src1->data, (nv_bfloat16 *) dst->data, + stream, std::make_index_sequence{}); } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F16 && dst->type == GGML_TYPE_F16) { launch_bin_bcast_pack(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 +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(scale * ggml_cuda_cast(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 +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, 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(op(ggml_cuda_cast(x[i]))); } template @@ -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((const half *)src0_d, (half *)dst_d, ggml_nelements(src0), stream); + } else if (src0->type == GGML_TYPE_BF16) { + unary_cuda((const nv_bfloat16 *)src0_d, (nv_bfloat16 *)dst_d, ggml_nelements(src0), stream); } else { unary_cuda((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(op(ggml_cuda_cast(x[j0])) * ggml_cuda_cast(g[j1])); } template @@ -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(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(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((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((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((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 ||