diff --git a/docs/models/breeze_tts.md b/docs/models/breeze_tts.md index e13c7cdf4..9aca60d01 100644 --- a/docs/models/breeze_tts.md +++ b/docs/models/breeze_tts.md @@ -67,3 +67,13 @@ audiocpp_cli \ | `--request-option top_p=` | `0..1` | `1.0` | Top-p sampling limit. | | `--request-option seed=` | integer >= 0 | `0` | Generation seed. | | `--session-option breeze_tts.reference_cache_slots=` | integer >= 0 | `1` | Prepared reference-audio cache slots. | +| `--session-option weight_type=` | `native`, `f32`, `f16`, `bf16`, `q8_0`, `q4_0`, `q4_k` | `native` | Weight storage type; quantized types convert at load time from the BF16 package. | + +Quantized weight storage is the largest measured speedup and applies to CUDA +and HIP alike: `q8_0` cut the fixed 100-token regression case from RTF ~1.5 to +~0.95 on gfx1151 and from ~0.77 to ~0.56 on an RTX 2080 Ti, and `q4_k` reached +~0.84 / ~0.49 respectively, with no audible quality regression in the Chinese +voice-design regression cases. Counter to intuition, fp32 is the one +configuration known to be *worse* for this model (mispronunciations and +runaway repetition), because the model is trained and tuned in bf16. + diff --git a/external/ggml/include/ggml.h b/external/ggml/include/ggml.h index 0de79aed5..1ba7dac76 100644 --- a/external/ggml/include/ggml.h +++ b/external/ggml/include/ggml.h @@ -615,6 +615,7 @@ extern "C" { GGML_UNARY_OP_CEIL, GGML_UNARY_OP_ROUND, GGML_UNARY_OP_TRUNC, + GGML_UNARY_OP_ROUND_BF16, GGML_UNARY_OP_COUNT, }; @@ -1258,6 +1259,12 @@ extern "C" { struct ggml_context * ctx, struct ggml_tensor * a); + // Rounds each element to bf16 precision, stored as f32. Equivalent to a + // cast f32 -> bf16 -> f32 round trip, but fused into a single op. + GGML_API struct ggml_tensor * ggml_round_bf16( + struct ggml_context * ctx, + struct ggml_tensor * a); + // xIELU activation function diff --git a/external/ggml/src/ggml-cpu/ggml-cpu.c b/external/ggml/src/ggml-cpu/ggml-cpu.c index d9ec09939..91e53b5e2 100644 --- a/external/ggml/src/ggml-cpu/ggml-cpu.c +++ b/external/ggml/src/ggml-cpu/ggml-cpu.c @@ -2260,6 +2260,7 @@ static int ggml_get_n_tasks(struct ggml_tensor * node, int n_threads) { case GGML_UNARY_OP_CEIL: case GGML_UNARY_OP_ROUND: case GGML_UNARY_OP_TRUNC: + case GGML_UNARY_OP_ROUND_BF16: { n_tasks = 1; } break; diff --git a/external/ggml/src/ggml-cpu/ops.cpp b/external/ggml/src/ggml-cpu/ops.cpp index 0f0f57399..4047e105f 100644 --- a/external/ggml/src/ggml-cpu/ops.cpp +++ b/external/ggml/src/ggml-cpu/ops.cpp @@ -10068,6 +10068,10 @@ void ggml_compute_forward_unary( { ggml_compute_forward_trunc(params, dst); } break; + case GGML_UNARY_OP_ROUND_BF16: + { + ggml_compute_forward_round_bf16(params, dst); + } break; case GGML_UNARY_OP_XIELU: { ggml_compute_forward_xielu(params, dst); diff --git a/external/ggml/src/ggml-cpu/unary-ops.cpp b/external/ggml/src/ggml-cpu/unary-ops.cpp index a82ffc260..4a813271c 100644 --- a/external/ggml/src/ggml-cpu/unary-ops.cpp +++ b/external/ggml/src/ggml-cpu/unary-ops.cpp @@ -97,6 +97,10 @@ static inline float op_trunc(float x) { return truncf(x); } +static inline float op_round_bf16(float x) { + return bf16_to_f32(f32_to_bf16(x)); +} + template static inline void vec_unary_op(int64_t n, dst_t * y, const src0_t * x) { constexpr auto src0_to_f32 = type_conversion_table::to_f32; @@ -322,6 +326,10 @@ void ggml_compute_forward_trunc(const ggml_compute_params * params, ggml_tensor unary_op(params, dst); } +void ggml_compute_forward_round_bf16(const ggml_compute_params * params, ggml_tensor * dst) { + unary_op(params, dst); +} + void ggml_compute_forward_xielu(const ggml_compute_params * params, ggml_tensor * dst) { const float alpha_n = ggml_get_op_params_f32(dst, 1); const float alpha_p = ggml_get_op_params_f32(dst, 2); diff --git a/external/ggml/src/ggml-cpu/unary-ops.h b/external/ggml/src/ggml-cpu/unary-ops.h index d75037369..06f3e1c37 100644 --- a/external/ggml/src/ggml-cpu/unary-ops.h +++ b/external/ggml/src/ggml-cpu/unary-ops.h @@ -28,6 +28,7 @@ void ggml_compute_forward_floor(const struct ggml_compute_params * params, struc void ggml_compute_forward_ceil(const struct ggml_compute_params * params, struct ggml_tensor * dst); void ggml_compute_forward_round(const struct ggml_compute_params * params, struct ggml_tensor * dst); void ggml_compute_forward_trunc(const struct ggml_compute_params * params, struct ggml_tensor * dst); +void ggml_compute_forward_round_bf16(const struct ggml_compute_params * params, struct ggml_tensor * dst); void ggml_compute_forward_xielu(const struct ggml_compute_params * params, struct ggml_tensor * dst); #ifdef __cplusplus diff --git a/external/ggml/src/ggml-cuda/ggml-cuda.cu b/external/ggml/src/ggml-cuda/ggml-cuda.cu index 8dc80a82d..278a49c43 100644 --- a/external/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/external/ggml/src/ggml-cuda/ggml-cuda.cu @@ -3065,6 +3065,9 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg case GGML_UNARY_OP_TRUNC: ggml_cuda_op_trunc(ctx, dst); break; + case GGML_UNARY_OP_ROUND_BF16: + ggml_cuda_op_round_bf16(ctx, dst); + break; case GGML_UNARY_OP_EXPM1: ggml_cuda_op_expm1(ctx, dst); break; @@ -4670,7 +4673,12 @@ static bool ggml_cuda_graph_set_enabled(ggml_backend_cuda_context * cuda_ctx, co ggml_cuda_graph * graph = cuda_ctx->cuda_graph(graph_key); if (graph->graph == nullptr) { - if (ggml_cuda_info().devices[cuda_ctx->device].cc < GGML_CUDA_CC_AMPERE) { + // CUDA graphs are disabled by default on pre-Ampere GPUs (matching + // upstream, where they regressed on some parts), but can be force + // enabled; decode loops made of many tiny kernels benefit even on + // Turing. + static const bool allow_pre_ampere = getenv("GGML_CUDA_GRAPHS_PRE_AMPERE") != nullptr; + if (!allow_pre_ampere && ggml_cuda_info().devices[cuda_ctx->device].cc < GGML_CUDA_CC_AMPERE) { if (!graph->disable_due_to_gpu_arch) { GGML_LOG_DEBUG("%s: disabling CUDA graphs due to GPU architecture\n", __func__); } @@ -5356,6 +5364,12 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g // TODO: should become: //return ggml_is_contiguous_rows(op->src[0]); return ggml_is_contiguous(op->src[0]); + case GGML_UNARY_OP_ROUND_BF16: + // f32/f16/bf16 src with contiguous rows, contiguous f32 dst. + return (op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16 || + op->src[0]->type == GGML_TYPE_BF16) && + op->type == GGML_TYPE_F32 && ggml_is_contiguous(op) && + ggml_is_contiguous_rows(op->src[0]); default: return false; } diff --git a/external/ggml/src/ggml-cuda/unary.cu b/external/ggml/src/ggml-cuda/unary.cu index fd3ac7550..213aeb4fe 100644 --- a/external/ggml/src/ggml-cuda/unary.cu +++ b/external/ggml/src/ggml-cuda/unary.cu @@ -114,6 +114,11 @@ static __device__ __forceinline__ float op_trunc(float x) { return trunc(x); } +static __device__ __forceinline__ float op_round_bf16(float x) { + // Matches the f32 -> bf16 -> f32 cpy round trip. + return __bfloat162float(__float2bfloat16(x)); +} + template static __global__ void unary_op_kernel(const T * x, T * dst, const int k) { const int i = blockDim.x*blockIdx.x + threadIdx.x; @@ -125,6 +130,75 @@ static __global__ void unary_op_kernel(const T * x, T * dst, const int k) { dst[i] = (T)op((float)x[i]); } +// Variant for a src with contiguous rows but arbitrary row strides; dst must be contiguous. +template +static __global__ void unary_op_kernel_strided( + const char * cx, T * dst, const int64_t k, + const int64_t ne0, const int64_t ne1, const int64_t ne2, + const int64_t nb01, const int64_t nb02, const int64_t nb03) { + const int64_t i = (int64_t)blockDim.x*blockIdx.x + threadIdx.x; + + if (i >= k) { + return; + } + + const int64_t i0 = i % ne0; + const int64_t i1 = (i / ne0) % ne1; + const int64_t i2 = (i / (ne0*ne1)) % ne2; + const int64_t i3 = i / (ne0*ne1*ne2); + + const T * x = (const T *) (cx + i1*nb01 + i2*nb02 + i3*nb03); + dst[i] = (T)op((float)x[i0]); +} + +// round-to-bf16 kernels: any of f32/f16/bf16 in, always f32 out. +template +static __global__ void round_bf16_kernel(const T * x, float * dst, const int64_t k) { + const int64_t i = (int64_t)blockDim.x*blockIdx.x + threadIdx.x; + + if (i >= k) { + return; + } + + dst[i] = op_round_bf16((float)x[i]); +} + +template +static __global__ void round_bf16_kernel_strided( + const char * cx, float * dst, const int64_t k, + const int64_t ne0, const int64_t ne1, const int64_t ne2, + const int64_t nb01, const int64_t nb02, const int64_t nb03) { + const int64_t i = (int64_t)blockDim.x*blockIdx.x + threadIdx.x; + + if (i >= k) { + return; + } + + const int64_t i0 = i % ne0; + const int64_t i1 = (i / ne0) % ne1; + const int64_t i2 = (i / (ne0*ne1)) % ne2; + const int64_t i3 = i / (ne0*ne1*ne2); + + const T * x = (const T *) (cx + i1*nb01 + i2*nb02 + i3*nb03); + dst[i] = op_round_bf16((float)x[i0]); +} + +template +static void round_bf16_cuda(const ggml_tensor * src0, float * dst, cudaStream_t stream) { + const int64_t k = ggml_nelements(src0); + const int64_t num_blocks = (k + CUDA_NEG_BLOCK_SIZE - 1) / CUDA_NEG_BLOCK_SIZE; + GGML_ASSERT(num_blocks < UINT_MAX); + + if (ggml_is_contiguous(src0)) { + round_bf16_kernel<<<(unsigned int) num_blocks, CUDA_NEG_BLOCK_SIZE, 0, stream>>>( + (const T *) src0->data, dst, k); + } else { + round_bf16_kernel_strided<<<(unsigned int) num_blocks, CUDA_NEG_BLOCK_SIZE, 0, stream>>>( + (const char *) src0->data, dst, k, src0->ne[0], src0->ne[1], src0->ne[2], + src0->nb[1], src0->nb[2], src0->nb[3]); + } +} + template static void unary_cuda(const T * x, T * dst, const int k, cudaStream_t stream) { const int num_blocks = (k + CUDA_NEG_BLOCK_SIZE - 1) / CUDA_NEG_BLOCK_SIZE; @@ -247,6 +321,31 @@ void ggml_cuda_op_trunc(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { ggml_cuda_op_unary(ctx, dst); } +void ggml_cuda_op_round_bf16(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { + const ggml_tensor * src0 = dst->src[0]; + + // ggml_round_bf16 always produces a contiguous f32 dst; src may be + // f32/f16/bf16 with contiguous rows. + GGML_ASSERT(dst->type == GGML_TYPE_F32); + GGML_ASSERT(ggml_is_contiguous(dst)); + GGML_ASSERT(ggml_is_contiguous_rows(src0)); + + cudaStream_t stream = ctx.stream(); + switch (src0->type) { + case GGML_TYPE_F32: + round_bf16_cuda(src0, (float *) dst->data, stream); + break; + case GGML_TYPE_F16: + round_bf16_cuda(src0, (float *) dst->data, stream); + break; + case GGML_TYPE_BF16: + round_bf16_cuda(src0, (float *) dst->data, stream); + break; + default: + GGML_ABORT("%s: unsupported src type %s", __func__, ggml_type_name(src0->type)); + } +} + void ggml_cuda_op_expm1(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { ggml_cuda_op_unary(ctx, dst); } diff --git a/external/ggml/src/ggml-cuda/unary.cuh b/external/ggml/src/ggml-cuda/unary.cuh index fbb229cb8..06a7e589f 100644 --- a/external/ggml/src/ggml-cuda/unary.cuh +++ b/external/ggml/src/ggml-cuda/unary.cuh @@ -75,6 +75,8 @@ void ggml_cuda_op_round(ggml_backend_cuda_context & ctx, ggml_tensor * dst); void ggml_cuda_op_trunc(ggml_backend_cuda_context & ctx, ggml_tensor * dst); +void ggml_cuda_op_round_bf16(ggml_backend_cuda_context & ctx, ggml_tensor * dst); + void ggml_cuda_op_reglu(ggml_backend_cuda_context & ctx, ggml_tensor * dst); void ggml_cuda_op_geglu(ggml_backend_cuda_context & ctx, ggml_tensor * dst); diff --git a/external/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/external/ggml/src/ggml-vulkan/ggml-vulkan.cpp index 334651a11..3e9a5217b 100644 --- a/external/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/external/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -774,8 +774,8 @@ struct vk_device_struct { vk_pipeline pipeline_pad_reflect_1d_f32; vk_pipeline pipeline_roll_f32; vk_pipeline pipeline_repeat_f32, pipeline_repeat_back_f32; - vk_pipeline pipeline_cpy_f32_f32, pipeline_cpy_f32_f16, pipeline_cpy_f16_f16, pipeline_cpy_f16_f32, pipeline_cpy_f32_bf16, pipeline_cpy_f32_i32, pipeline_cpy_i32_f32; - vk_pipeline pipeline_contig_cpy_f32_f32, pipeline_contig_cpy_f32_f16, pipeline_contig_cpy_f16_f16, pipeline_contig_cpy_f16_f32, pipeline_contig_cpy_f32_bf16, pipeline_contig_cpy_f32_i32, pipeline_contig_cpy_i32_f32; + vk_pipeline pipeline_cpy_f32_f32, pipeline_cpy_f32_f16, pipeline_cpy_f16_f16, pipeline_cpy_f16_f32, pipeline_cpy_f32_bf16, pipeline_cpy_bf16_f32, pipeline_cpy_f16_bf16, pipeline_cpy_bf16_f16, pipeline_cpy_f32_i32, pipeline_cpy_i32_f32; + vk_pipeline pipeline_contig_cpy_f32_f32, pipeline_contig_cpy_f32_f16, pipeline_contig_cpy_f16_f16, pipeline_contig_cpy_f16_f32, pipeline_contig_cpy_f32_bf16, pipeline_contig_cpy_bf16_f32, pipeline_contig_cpy_f16_bf16, pipeline_contig_cpy_bf16_f16, pipeline_contig_cpy_f32_i32, pipeline_contig_cpy_i32_f32; vk_pipeline pipeline_cpy_f32_quant[GGML_TYPE_COUNT]; vk_pipeline pipeline_cpy_quant_f32[GGML_TYPE_COUNT]; vk_pipeline pipeline_cpy_transpose_16, pipeline_cpy_transpose_32; @@ -4594,7 +4594,10 @@ static void ggml_vk_load_shaders(vk_device& device) { ggml_vk_create_pipeline(device, device->pipeline_cpy_f32_f16, "cpy_f32_f16", cpy_f32_f16_len, cpy_f32_f16_data, "main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_cpy_f16_f16, "cpy_f16_f16", cpy_f16_f16_len, cpy_f16_f16_data, "main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_cpy_f16_f32, "cpy_f16_f32", cpy_f16_f32_len, cpy_f16_f32_data, "main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1); - ggml_vk_create_pipeline(device, device->pipeline_cpy_f32_bf16,"cpy_f32_bf16",cpy_f32_bf16_len,cpy_f32_bf16_data,"main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1); + ggml_vk_create_pipeline(device, device->pipeline_cpy_f32_bf16,"cpy_f32_bf16",cpy_f32_bf16_len,cpy_f32_bf16_data,"main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1); + ggml_vk_create_pipeline(device, device->pipeline_cpy_bf16_f32,"cpy_bf16_f32",cpy_bf16_f32_len,cpy_bf16_f32_data,"main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1); + ggml_vk_create_pipeline(device, device->pipeline_cpy_f16_bf16,"cpy_f16_bf16",cpy_f16_bf16_len,cpy_f16_bf16_data,"main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1); + ggml_vk_create_pipeline(device, device->pipeline_cpy_bf16_f16,"cpy_bf16_f16",cpy_bf16_f16_len,cpy_bf16_f16_data,"main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_cpy_i32_f32, "cpy_i32_f32", cpy_i32_f32_len, cpy_i32_f32_data, "main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_cpy_f32_i32, "cpy_f32_i32", cpy_f32_i32_len, cpy_f32_i32_data, "main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1); @@ -4602,7 +4605,10 @@ static void ggml_vk_load_shaders(vk_device& device) { ggml_vk_create_pipeline(device, device->pipeline_contig_cpy_f32_f16, "contig_cpy_f32_f16", contig_cpy_f32_f16_len, contig_cpy_f32_f16_data, "main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_contig_cpy_f16_f16, "contig_cpy_f16_f16", contig_cpy_f16_f16_len, contig_cpy_f16_f16_data, "main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_contig_cpy_f16_f32, "contig_cpy_f16_f32", contig_cpy_f16_f32_len, contig_cpy_f16_f32_data, "main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1); - ggml_vk_create_pipeline(device, device->pipeline_contig_cpy_f32_bf16,"contig_cpy_f32_bf16",contig_cpy_f32_bf16_len,contig_cpy_f32_bf16_data,"main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1); + ggml_vk_create_pipeline(device, device->pipeline_contig_cpy_f32_bf16,"contig_cpy_f32_bf16",contig_cpy_f32_bf16_len,contig_cpy_f32_bf16_data,"main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1); + ggml_vk_create_pipeline(device, device->pipeline_contig_cpy_bf16_f32,"contig_cpy_bf16_f32",contig_cpy_bf16_f32_len,contig_cpy_bf16_f32_data,"main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1); + ggml_vk_create_pipeline(device, device->pipeline_contig_cpy_f16_bf16,"contig_cpy_f16_bf16",contig_cpy_f16_bf16_len,contig_cpy_f16_bf16_data,"main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1); + ggml_vk_create_pipeline(device, device->pipeline_contig_cpy_bf16_f16,"contig_cpy_bf16_f16",contig_cpy_bf16_f16_len,contig_cpy_bf16_f16_data,"main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_contig_cpy_i32_f32, "contig_cpy_i32_f32", contig_cpy_i32_f32_len, contig_cpy_i32_f32_data, "main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_contig_cpy_f32_i32, "contig_cpy_f32_i32", contig_cpy_f32_i32_len, contig_cpy_f32_i32_data, "main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1); @@ -7578,6 +7584,27 @@ static vk_pipeline ggml_vk_get_cpy_pipeline(ggml_backend_vk_context * ctx, const return ctx->device->pipeline_cpy_f32_bf16; } } + if (src->type == GGML_TYPE_BF16 && to == GGML_TYPE_F32) { + if (contig) { + return ctx->device->pipeline_contig_cpy_bf16_f32; + } else { + return ctx->device->pipeline_cpy_bf16_f32; + } + } + if (src->type == GGML_TYPE_F16 && to == GGML_TYPE_BF16) { + if (contig) { + return ctx->device->pipeline_contig_cpy_f16_bf16; + } else { + return ctx->device->pipeline_cpy_f16_bf16; + } + } + if (src->type == GGML_TYPE_BF16 && to == GGML_TYPE_F16) { + if (contig) { + return ctx->device->pipeline_contig_cpy_bf16_f16; + } else { + return ctx->device->pipeline_cpy_bf16_f16; + } + } if (src->type == GGML_TYPE_F32 && to == GGML_TYPE_I32) { if (contig) { return ctx->device->pipeline_contig_cpy_f32_i32; diff --git a/external/ggml/src/ggml-vulkan/vulkan-shaders/contig_copy.comp b/external/ggml/src/ggml-vulkan/vulkan-shaders/contig_copy.comp index 066b27ca0..cacddbdf7 100644 --- a/external/ggml/src/ggml-vulkan/vulkan-shaders/contig_copy.comp +++ b/external/ggml/src/ggml-vulkan/vulkan-shaders/contig_copy.comp @@ -19,7 +19,10 @@ void main() { if (idx + (num_iter-1)*num_threads < p.ne) { [[unroll]] for (uint i = 0; i < num_iter; ++i) { -#if defined(DATA_D_BF16) +#if defined(DATA_A_BF16) + float f = bf16_to_fp32(uint32_t(data_a[get_aoffset() + idx])); + data_d[get_doffset() + idx] = D_TYPE(f); +#elif defined(DATA_D_BF16) float f = float(data_a[get_aoffset() + idx]); data_d[get_doffset() + idx] = D_TYPE(fp32_to_bf16(f)); #elif !defined(OPTIMIZATION_ERROR_WORKAROUND) @@ -35,7 +38,10 @@ void main() { continue; } -#if defined(DATA_D_BF16) +#if defined(DATA_A_BF16) + float f = bf16_to_fp32(uint32_t(data_a[get_aoffset() + idx])); + data_d[get_doffset() + idx] = D_TYPE(f); +#elif defined(DATA_D_BF16) float f = float(data_a[get_aoffset() + idx]); data_d[get_doffset() + idx] = D_TYPE(fp32_to_bf16(f)); #elif !defined(OPTIMIZATION_ERROR_WORKAROUND) diff --git a/external/ggml/src/ggml-vulkan/vulkan-shaders/copy.comp b/external/ggml/src/ggml-vulkan/vulkan-shaders/copy.comp index a1ba96702..81cce864e 100644 --- a/external/ggml/src/ggml-vulkan/vulkan-shaders/copy.comp +++ b/external/ggml/src/ggml-vulkan/vulkan-shaders/copy.comp @@ -12,7 +12,10 @@ void main() { return; } -#if defined(DATA_D_BF16) +#if defined(DATA_A_BF16) + float f = bf16_to_fp32(uint32_t(data_a[get_aoffset() + src0_idx(idx)])); + data_d[get_doffset() + dst_idx(idx)] = D_TYPE(f); +#elif defined(DATA_D_BF16) float f = float(data_a[get_aoffset() + src0_idx(idx)]); data_d[get_doffset() + dst_idx(idx)] = D_TYPE(fp32_to_bf16(f)); #elif !defined(OPTIMIZATION_ERROR_WORKAROUND) diff --git a/external/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp b/external/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp index 7295ef107..a533af69f 100644 --- a/external/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp +++ b/external/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp @@ -730,14 +730,20 @@ void process_shaders() { string_to_spv("cpy_f32_f16", "copy.comp", {{"A_TYPE", "float"}, {"D_TYPE", "float16_t"}}); string_to_spv("cpy_f16_f16", "copy.comp", {{"A_TYPE", "float16_t"}, {"D_TYPE", "float16_t"}, {"OPTIMIZATION_ERROR_WORKAROUND", "1"}}); string_to_spv("cpy_f16_f32", "copy.comp", {{"A_TYPE", "float16_t"}, {"D_TYPE", "float"}, {"OPTIMIZATION_ERROR_WORKAROUND", "1"}}); - string_to_spv("cpy_f32_bf16","copy.comp", {{"A_TYPE", "float"}, {"D_TYPE", "uint16_t"}, {"DATA_D_BF16", "1"}}); + string_to_spv("cpy_f32_bf16","copy.comp", {{"A_TYPE", "float"}, {"D_TYPE", "uint16_t"}, {"DATA_D_BF16", "1"}}); + string_to_spv("cpy_bf16_f32","copy.comp", {{"A_TYPE", "uint16_t"}, {"D_TYPE", "float"}, {"DATA_A_BF16", "1"}}); + string_to_spv("cpy_f16_bf16","copy.comp", {{"A_TYPE", "float16_t"}, {"D_TYPE", "uint16_t"}, {"DATA_D_BF16", "1"}}); + string_to_spv("cpy_bf16_f16","copy.comp", {{"A_TYPE", "uint16_t"}, {"D_TYPE", "float16_t"}, {"DATA_A_BF16", "1"}}); string_to_spv("contig_cpy_f32_f32", "contig_copy.comp", {{"A_TYPE", "float"}, {"D_TYPE", "float"}}); string_to_spv("contig_cpy_f32_i32", "contig_copy.comp", {{"A_TYPE", "float"}, {"D_TYPE", "int"}}); string_to_spv("contig_cpy_i32_f32", "contig_copy.comp", {{"A_TYPE", "int"}, {"D_TYPE", "float"}}); string_to_spv("contig_cpy_f32_f16", "contig_copy.comp", {{"A_TYPE", "float"}, {"D_TYPE", "float16_t"}}); string_to_spv("contig_cpy_f16_f16", "contig_copy.comp", {{"A_TYPE", "float16_t"}, {"D_TYPE", "float16_t"}, {"OPTIMIZATION_ERROR_WORKAROUND", "1"}}); string_to_spv("contig_cpy_f16_f32", "contig_copy.comp", {{"A_TYPE", "float16_t"}, {"D_TYPE", "float"}, {"OPTIMIZATION_ERROR_WORKAROUND", "1"}}); - string_to_spv("contig_cpy_f32_bf16","contig_copy.comp",{{"A_TYPE", "float"}, {"D_TYPE", "uint16_t"}, {"DATA_D_BF16", "1"}}); + string_to_spv("contig_cpy_f32_bf16","contig_copy.comp",{{"A_TYPE", "float"}, {"D_TYPE", "uint16_t"}, {"DATA_D_BF16", "1"}}); + string_to_spv("contig_cpy_bf16_f32","contig_copy.comp",{{"A_TYPE", "uint16_t"}, {"D_TYPE", "float"}, {"DATA_A_BF16", "1"}}); + string_to_spv("contig_cpy_f16_bf16","contig_copy.comp",{{"A_TYPE", "float16_t"}, {"D_TYPE", "uint16_t"}, {"DATA_D_BF16", "1"}}); + string_to_spv("contig_cpy_bf16_f16","contig_copy.comp",{{"A_TYPE", "uint16_t"}, {"D_TYPE", "float16_t"}, {"DATA_A_BF16", "1"}}); string_to_spv("cpy_f32_i32", "copy.comp", {{"A_TYPE", "float"}, {"D_TYPE", "int"}}); string_to_spv("cpy_i32_f32", "copy.comp", {{"A_TYPE", "int"}, {"D_TYPE", "float"}}); diff --git a/external/ggml/src/ggml.c b/external/ggml/src/ggml.c index 0ba560368..dbad3596d 100644 --- a/external/ggml/src/ggml.c +++ b/external/ggml/src/ggml.c @@ -1229,9 +1229,10 @@ static const char * GGML_UNARY_OP_NAME[GGML_UNARY_OP_COUNT] = { "CEIL", "ROUND", "TRUNC", + "ROUND_BF16", }; -static_assert(GGML_UNARY_OP_COUNT == 22, "GGML_UNARY_OP_COUNT != 22"); +static_assert(GGML_UNARY_OP_COUNT == 23, "GGML_UNARY_OP_COUNT != 23"); static const char * GGML_GLU_OP_NAME[GGML_GLU_OP_COUNT] = { "REGLU", @@ -2959,6 +2960,26 @@ struct ggml_tensor * ggml_trunc_inplace( return ggml_unary_inplace(ctx, a, GGML_UNARY_OP_TRUNC); } +//ggml_round_bf16 + +struct ggml_tensor * ggml_round_bf16( + struct ggml_context * ctx, + struct ggml_tensor * a) { + GGML_ASSERT(a->type == GGML_TYPE_F32 || a->type == GGML_TYPE_F16 || a->type == GGML_TYPE_BF16); + GGML_ASSERT(ggml_is_contiguous_rows(a)); + + // Unlike ggml_unary, the result is always f32: bf16/f16 inputs are widened + // while rounding, matching an f32 -> bf16 -> f32 cast round trip. + struct ggml_tensor * result = ggml_new_tensor(ctx, GGML_TYPE_F32, GGML_MAX_DIMS, a->ne); + + ggml_set_op_params_i32(result, 0, (int32_t) GGML_UNARY_OP_ROUND_BF16); + + result->op = GGML_OP_UNARY; + result->src[0] = a; + + return result; +} + struct ggml_tensor * ggml_glu( struct ggml_context * ctx, struct ggml_tensor * a, diff --git a/include/engine/framework/modules/transformers/qwen_decoder.h b/include/engine/framework/modules/transformers/qwen_decoder.h index 3b43a744b..eaa80c12e 100644 --- a/include/engine/framework/modules/transformers/qwen_decoder.h +++ b/include/engine/framework/modules/transformers/qwen_decoder.h @@ -53,6 +53,10 @@ enum class QwenDecoderPositionEncoding { struct QwenDecoderActivationCastPolicy { bool enabled = false; ggml_type type = GGML_TYPE_BF16; + // Use the fused single-kernel round-to-bf16 op instead of a + // cast -> bf16 -> cast -> f32 round trip. Only valid on backends that + // implement GGML_UNARY_OP_ROUND_BF16 (CUDA/HIP, CPU fallback). + bool fused_round = false; bool after_input_norm = false; bool after_qkv_projection = false; bool after_qk_norm = false; diff --git a/src/framework/modules/conv_modules.cpp b/src/framework/modules/conv_modules.cpp index 185b66e35..e6f562b92 100644 --- a/src/framework/modules/conv_modules.cpp +++ b/src/framework/modules/conv_modules.cpp @@ -319,7 +319,8 @@ bool is_conv_transpose1d_col2im_fast_path_eligible( const core::ModuleBuildContext & ctx, const ConvTranspose1dConfig & config) noexcept { return (core::uses_ggml_cuda_or_hip_backend(ctx.backend_type) || - ctx.backend_type == core::BackendType::Metal) && + ctx.backend_type == core::BackendType::Metal || + ctx.backend_type == core::BackendType::Vulkan) && config.dilation == 1; } diff --git a/src/framework/modules/transformers/qwen_decoder.cpp b/src/framework/modules/transformers/qwen_decoder.cpp index a50b7191c..d7031270b 100644 --- a/src/framework/modules/transformers/qwen_decoder.cpp +++ b/src/framework/modules/transformers/qwen_decoder.cpp @@ -48,6 +48,9 @@ int64_t require_head_dim(const QwenDecoderLayerConfig & config) { config.activation_cast.type != GGML_TYPE_F16 && config.activation_cast.type != GGML_TYPE_BF16) { throw std::runtime_error("QwenDecoderLayerConfig activation cast supports only f32, f16, and bf16"); } + if (config.activation_cast.fused_round && config.activation_cast.type != GGML_TYPE_BF16) { + throw std::runtime_error("QwenDecoderLayerConfig fused activation rounding requires bf16"); + } return config.head_dim; } @@ -235,6 +238,13 @@ core::TensorValue activation_cast( if (policy.type == GGML_TYPE_F32) { return core::wrap_tensor(ggml_cast(ctx.ggml, input.tensor, GGML_TYPE_F32), input.shape, GGML_TYPE_F32); } + if (policy.fused_round && policy.type == GGML_TYPE_BF16 && ggml_is_contiguous_rows(input.tensor)) { + // Fused single-kernel round-to-bf16: f32/f16/bf16 in, always f32 out, + // values rounded to bf16. Numerically identical to the cast round trip + // below (bf16 input is already rounded, so rounding is a widening no-op), + // but avoids the intermediate bf16 tensor and one kernel launch. + return core::wrap_tensor(ggml_round_bf16(ctx.ggml, input.tensor), input.shape, GGML_TYPE_F32); + } auto rounded = core::wrap_tensor(ggml_cast(ctx.ggml, input.tensor, policy.type), input.shape, policy.type); return core::wrap_tensor(ggml_cast(ctx.ggml, rounded.tensor, GGML_TYPE_F32), input.shape, GGML_TYPE_F32); } diff --git a/src/models/breeze_tts/generator.cpp b/src/models/breeze_tts/generator.cpp index 7d8483877..83f149540 100644 --- a/src/models/breeze_tts/generator.cpp +++ b/src/models/breeze_tts/generator.cpp @@ -48,6 +48,37 @@ struct GgmlContextDeleter { } }; +// The official Breeze-TTS 2 inference runs the backbone and depth decoder with +// bf16 activations and a bf16 KV cache. Pure fp32 activations measurably drift +// into degenerate trajectories on some prompts (mispronunciations, repetition +// collapse), so match the reference bf16 behavior on GPU backends. +modules::QwenDecoderActivationCastPolicy breeze_bf16_activation_policy(core::BackendType backend_type) { + modules::QwenDecoderActivationCastPolicy policy; + if (backend_type != core::BackendType::Cuda && backend_type != core::BackendType::Hip && + backend_type != core::BackendType::Vulkan) { + return policy; + } + policy.enabled = true; + policy.type = GGML_TYPE_BF16; + // CUDA/HIP implement the fused round-to-bf16 unary op; Vulkan does not and + // keeps the cast round trip. + policy.fused_round = backend_type == core::BackendType::Cuda || backend_type == core::BackendType::Hip; + policy.after_input_norm = true; + policy.after_qkv_projection = true; + policy.after_qk_norm = true; + policy.after_rope = true; + policy.after_static_cache_update = true; + policy.after_attention = true; + policy.after_attention_output = true; + policy.after_residual = true; + policy.after_ffn_norm = true; + policy.after_mlp_projection = true; + policy.after_mlp_silu = true; + policy.after_mlp_mul = true; + policy.after_output = true; + return policy; +} + modules::QwenCausalDecodeRuntimeConfig backbone_config( const BreezeTTSConfig & config, core::BackendType backend_type, @@ -66,6 +97,8 @@ modules::QwenCausalDecodeRuntimeConfig backbone_config( out.decoder.stack.rope_theta = config.rope_theta; out.decoder.stack.rope_type = GGML_ROPE_TYPE_NEOX; out.decoder.stack.use_qk_norm = true; + out.decoder.stack.qkv_layout = modules::QwenDecoderQKVLayout::PackedQKV; + out.decoder.stack.runtime.mlp.mode = modules::QwenDecoderMLPMode::PackedGateUp; out.decoder.stack.attention_precision = GGML_PREC_F32; out.decoder.stack.projection_precision = GGML_PREC_DEFAULT; out.decoder.stack.runtime.attention.prefill_mode = modules::QwenDecoderAttentionMode::FlashGroupedViewKV; @@ -74,7 +107,12 @@ modules::QwenCausalDecodeRuntimeConfig backbone_config( out.decoder.stack.runtime.static_cache.set_rows_mode = modules::QwenDecoderStaticCacheSetRowsMode::BackendViewOptimized; if (backend_type == core::BackendType::Cuda || backend_type == core::BackendType::Hip || backend_type == core::BackendType::Vulkan) { - out.decoder.static_cache_type = GGML_TYPE_F16; + // BF16 KV cache matches the reference implementation, but flash + // attention only accelerates bf16 cache with native bf16 MMA + // (sm_80+); on older parts it is ~3x slower, so only HIP uses it. + out.decoder.static_cache_type = + backend_type == core::BackendType::Hip ? GGML_TYPE_BF16 : GGML_TYPE_F16; + out.decoder.stack.activation_cast = breeze_bf16_activation_policy(backend_type); } out.decoder.logits_size = config.lm_head_size; out.decoder.logits_mode = modules::QwenCausalDecoderLogitsMode::LastStep; @@ -106,6 +144,8 @@ modules::QwenCausalDecodeRuntimeConfig depth_config( out.decoder.stack.rope_theta = config.depth_rope_theta; out.decoder.stack.rope_type = GGML_ROPE_TYPE_NEOX; out.decoder.stack.use_qk_norm = false; + out.decoder.stack.qkv_layout = modules::QwenDecoderQKVLayout::PackedQKV; + out.decoder.stack.runtime.mlp.mode = modules::QwenDecoderMLPMode::PackedGateUp; out.decoder.stack.attention_precision = GGML_PREC_F32; out.decoder.stack.projection_precision = GGML_PREC_DEFAULT; out.decoder.stack.runtime.attention.prefill_mode = modules::QwenDecoderAttentionMode::FlashGroupedViewKV; @@ -114,7 +154,11 @@ modules::QwenCausalDecodeRuntimeConfig depth_config( out.decoder.stack.runtime.static_cache.set_rows_mode = modules::QwenDecoderStaticCacheSetRowsMode::BackendViewOptimized; if (backend_type == core::BackendType::Cuda || backend_type == core::BackendType::Hip || backend_type == core::BackendType::Vulkan) { - out.decoder.static_cache_type = GGML_TYPE_F16; + // See backbone_config: only HIP uses a bf16 KV cache; CUDA and Vulkan + // keep F16. + out.decoder.static_cache_type = + backend_type == core::BackendType::Hip ? GGML_TYPE_BF16 : GGML_TYPE_F16; + out.decoder.stack.activation_cast = breeze_bf16_activation_policy(backend_type); } out.decoder.logits_mode = modules::QwenCausalDecoderLogitsMode::LastStep; out.output_mode = modules::QwenCausalDecodeOutputMode::Hidden; @@ -160,6 +204,35 @@ std::vector llama3_rope_factors( return out; } +// Pack [a; b; ...] projection rows into a single tensor, matching the +// higgs_audio_tts loader: fewer, larger matmuls per layer. Parts may have +// different row counts (e.g. q vs k/v in GQA models). +core::TensorValue pack_projection_rows( + core::BackendWeightStore & store, + const assets::TensorSource & source, + const std::vector> & parts, + assets::TensorStorageType storage_type, + int64_t in_dim) { + std::vector packed; + int64_t total_out = 0; + ggml_type packed_type = GGML_TYPE_COUNT; + for (const auto & [name, out_dim] : parts) { + const auto part = source.require_tensor(name, storage_type, {out_dim, in_dim}); + if (packed_type == GGML_TYPE_COUNT) { + packed_type = part.type; + } else if (part.type != packed_type) { + throw std::runtime_error("BreezeTTS packed projection weights require matching storage types"); + } + packed.insert(packed.end(), part.bytes.begin(), part.bytes.end()); + total_out += out_dim; + } + return store.make_tensor( + core::TensorShape::from_dims({total_out, in_dim}), + packed_type, + packed.data(), + packed.size()); +} + modules::QwenDecoderLayerWeights load_backbone_layer( core::BackendWeightStore & store, const assets::TensorSource & source, @@ -168,17 +241,33 @@ modules::QwenDecoderLayerWeights load_backbone_layer( const std::optional & rope_factors, int64_t layer) { const std::string prefix = "backbone_model.layers." + std::to_string(layer); + const int64_t q_out = config.heads * config.head_dim; + const int64_t kv_out = config.kv_heads * config.head_dim; modules::QwenDecoderLayerWeights out; out.input_norm = binding::norm_weight_from_source(store, source, prefix + ".input_layernorm", config.hidden_size); - out.self_attention.q_weight = store.load_tensor(source, prefix + ".self_attn.q_proj.weight", storage_type, {config.heads * config.head_dim, config.hidden_size}); - out.self_attention.k_weight = store.load_tensor(source, prefix + ".self_attn.k_proj.weight", storage_type, {config.kv_heads * config.head_dim, config.hidden_size}); - out.self_attention.v_weight = store.load_tensor(source, prefix + ".self_attn.v_proj.weight", storage_type, {config.kv_heads * config.head_dim, config.hidden_size}); + // Packed layout is [q; k; v] with row counts q_out, kv_out, kv_out. + out.self_attention.qkv_weight = pack_projection_rows( + store, + source, + {{prefix + ".self_attn.q_proj.weight", q_out}, + {prefix + ".self_attn.k_proj.weight", kv_out}, + {prefix + ".self_attn.v_proj.weight", kv_out}}, + storage_type, + config.hidden_size); out.self_attention.out_weight = store.load_tensor(source, prefix + ".self_attn.o_proj.weight", storage_type, {config.hidden_size, config.heads * config.head_dim}); out.q_norm = binding::norm_weight_from_source(store, source, prefix + ".self_attn.q_norm", config.head_dim); out.k_norm = binding::norm_weight_from_source(store, source, prefix + ".self_attn.k_norm", config.head_dim); out.post_norm = binding::norm_weight_from_source(store, source, prefix + ".post_attention_layernorm", config.hidden_size); - out.mlp.gate_proj = binding::linear_from_source(store, source, prefix + ".mlp.gate_proj", storage_type, config.intermediate_size, config.hidden_size, false); - out.mlp.up_proj = binding::linear_from_source(store, source, prefix + ".mlp.up_proj", storage_type, config.intermediate_size, config.hidden_size, false); + out.mlp.gate_up_proj = modules::LinearWeights{ + pack_projection_rows( + store, + source, + {{prefix + ".mlp.gate_proj.weight", config.intermediate_size}, + {prefix + ".mlp.up_proj.weight", config.intermediate_size}}, + storage_type, + config.hidden_size), + std::nullopt, + }; out.mlp.down_proj = binding::linear_from_source(store, source, prefix + ".mlp.down_proj", storage_type, config.hidden_size, config.intermediate_size, false); out.rope_frequency_factors = rope_factors; return out; @@ -192,15 +281,31 @@ modules::QwenDecoderLayerWeights load_depth_layer( const std::optional & rope_factors, int64_t layer) { const std::string prefix = "depth_decoder.model.layers." + std::to_string(layer); + const int64_t q_out = config.depth_heads * config.depth_head_dim; + const int64_t kv_out = config.depth_kv_heads * config.depth_head_dim; modules::QwenDecoderLayerWeights out; out.input_norm = binding::norm_weight_from_source(store, source, prefix + ".input_layernorm", config.depth_hidden_size); - out.self_attention.q_weight = store.load_tensor(source, prefix + ".self_attn.q_proj.weight", storage_type, {config.depth_heads * config.depth_head_dim, config.depth_hidden_size}); - out.self_attention.k_weight = store.load_tensor(source, prefix + ".self_attn.k_proj.weight", storage_type, {config.depth_kv_heads * config.depth_head_dim, config.depth_hidden_size}); - out.self_attention.v_weight = store.load_tensor(source, prefix + ".self_attn.v_proj.weight", storage_type, {config.depth_kv_heads * config.depth_head_dim, config.depth_hidden_size}); + // Packed layout is [q; k; v] with row counts q_out, kv_out, kv_out. + out.self_attention.qkv_weight = pack_projection_rows( + store, + source, + {{prefix + ".self_attn.q_proj.weight", q_out}, + {prefix + ".self_attn.k_proj.weight", kv_out}, + {prefix + ".self_attn.v_proj.weight", kv_out}}, + storage_type, + config.depth_hidden_size); out.self_attention.out_weight = store.load_tensor(source, prefix + ".self_attn.o_proj.weight", storage_type, {config.depth_hidden_size, config.depth_heads * config.depth_head_dim}); out.post_norm = binding::norm_weight_from_source(store, source, prefix + ".post_attention_layernorm", config.depth_hidden_size); - out.mlp.gate_proj = binding::linear_from_source(store, source, prefix + ".mlp.gate_proj", storage_type, config.depth_intermediate_size, config.depth_hidden_size, false); - out.mlp.up_proj = binding::linear_from_source(store, source, prefix + ".mlp.up_proj", storage_type, config.depth_intermediate_size, config.depth_hidden_size, false); + out.mlp.gate_up_proj = modules::LinearWeights{ + pack_projection_rows( + store, + source, + {{prefix + ".mlp.gate_proj.weight", config.depth_intermediate_size}, + {prefix + ".mlp.up_proj.weight", config.depth_intermediate_size}}, + storage_type, + config.depth_hidden_size), + std::nullopt, + }; out.mlp.down_proj = binding::linear_from_source(store, source, prefix + ".mlp.down_proj", storage_type, config.depth_hidden_size, config.depth_intermediate_size, false); out.rope_frequency_factors = rope_factors; return out; @@ -489,6 +594,12 @@ class BreezeDepthProjectionRuntime { ggml_backend_synchronize(backend_); ggml_backend_tensor_get(graph.output, head_paired_staging_.data(), 0, head_paired_staging_.size() * sizeof(float)); const size_t vocab = static_cast(vocab_); + if (guidance_scale == 1.0F) { + // CFG is a no-op at scale 1: copy the conditional half directly, + // avoiding an inexact uncond + 1 * (cond - uncond) round trip. + std::memcpy(out, head_paired_staging_.data(), vocab * sizeof(float)); + return; + } for (int64_t token = 0; token < vocab_; ++token) { const size_t index = static_cast(token); out[index] = head_paired_staging_[vocab + index] + guidance_scale * (head_paired_staging_[index] - head_paired_staging_[vocab + index]); @@ -862,21 +973,30 @@ struct BreezeGeneratorRuntime::Impl { std::vector uncond_embeddings; int64_t cond_steps = 0; int64_t uncond_steps = 0; + // guidance_scale == 1 makes CFG a no-op (logits == cond), so skip the + // unconditional branch entirely and halve the backbone work. + const bool use_cfg = request.guidance_scale != 1.0F; const double prompt_ms = engine::debug::measure_ms([&] { if (!reference_codes.empty()) { if (request.reference_text.empty()) { throw std::runtime_error("BreezeTTS clone requires reference_text"); } cond_branch = tokenizer.build_clone(request.text, request.instruction, request.reference_text, reference_frames); - uncond_branch = tokenizer.build_clone_negative(request.text, request.reference_text, reference_frames); + if (use_cfg) { + uncond_branch = tokenizer.build_clone_negative(request.text, request.reference_text, reference_frames); + } } else { cond_branch = tokenizer.build_tts_instruction(request.text, request.instruction); - uncond_branch = tokenizer.build_tts_plain(request.text); + if (use_cfg) { + uncond_branch = tokenizer.build_tts_plain(request.text); + } } cond_embeddings = merge_prompt(cond_branch, reference_codes); - uncond_embeddings = merge_prompt(uncond_branch, reference_codes); cond_steps = static_cast(cond_branch.input_ids.size()); - uncond_steps = static_cast(uncond_branch.input_ids.size()); + if (use_cfg) { + uncond_embeddings = merge_prompt(uncond_branch, reference_codes); + uncond_steps = static_cast(uncond_branch.input_ids.size()); + } }); engine::debug::timing_log_scalar("breeze_tts.generate.prompt_ms", prompt_ms); text_encoder.release_runtime_graphs(); @@ -891,12 +1011,17 @@ struct BreezeGeneratorRuntime::Impl { backbone_cond_prefill_ms = engine::debug::measure_ms([&] { cond = backbone_cond->prefill_embeddings(cond_embeddings, cond_steps); }); - modules::QwenCausalPrefillResult uncond; - backbone_uncond_prefill_ms = engine::debug::measure_ms([&] { - uncond = backbone_uncond->prefill_embeddings(uncond_embeddings, uncond_steps); - }); + std::optional uncond; + if (use_cfg) { + uncond.emplace(); + backbone_uncond_prefill_ms = engine::debug::measure_ms([&] { + *uncond = backbone_uncond->prefill_embeddings(uncond_embeddings, uncond_steps); + }); + } backbone_cond->start_decode_embeddings(cond.state, cond_steps + request.max_tokens); - backbone_uncond->start_decode_embeddings(uncond.state, uncond_steps + request.max_tokens); + if (use_cfg) { + backbone_uncond->start_decode_embeddings(uncond->state, uncond_steps + request.max_tokens); + } sampling::HfSamplerScratch scratch; scratch.reserve_vocab(static_cast(config.lm_head_size)); @@ -913,12 +1038,17 @@ struct BreezeGeneratorRuntime::Impl { codes.reserve(static_cast(request.max_tokens * config.num_codebooks)); for (int64_t step = 0; step < request.max_tokens; ++step) { - if (cond.logits.size() != uncond.logits.size()) { + if (use_cfg && cond.logits.size() != uncond->logits.size()) { throw std::runtime_error("BreezeTTS CFG logits shape mismatch"); } - std::vector logits(cond.logits.size(), 0.0F); - for (size_t i = 0; i < logits.size(); ++i) { - logits[i] = uncond.logits[i] + request.guidance_scale * (cond.logits[i] - uncond.logits[i]); + std::vector logits; + if (use_cfg) { + logits.resize(cond.logits.size()); + for (size_t i = 0; i < logits.size(); ++i) { + logits[i] = uncond->logits[i] + request.guidance_scale * (cond.logits[i] - uncond->logits[i]); + } + } else { + logits = cond.logits; } suppress_reserved(logits, kCodecCodebookSize, config.vocab_size); const int32_t first_token = sample_logits( @@ -940,7 +1070,7 @@ struct BreezeGeneratorRuntime::Impl { } const auto frame = generate_frame( cond.hidden, - uncond.hidden, + use_cfg ? uncond->hidden : cond.hidden, first_token, request, scratch, @@ -959,14 +1089,16 @@ struct BreezeGeneratorRuntime::Impl { backbone_cond_decode_ms += engine::debug::measure_ms([&] { cond_step = backbone_cond->decode_embedding(embedded); }); - modules::QwenCausalDecodeStepResult uncond_step; - backbone_uncond_decode_ms += engine::debug::measure_ms([&] { - uncond_step = backbone_uncond->decode_embedding(embedded); - }); cond.logits = cond_step.logits; cond.hidden = cond_step.hidden; - uncond.logits = uncond_step.logits; - uncond.hidden = uncond_step.hidden; + if (use_cfg) { + modules::QwenCausalDecodeStepResult uncond_step; + backbone_uncond_decode_ms += engine::debug::measure_ms([&] { + uncond_step = backbone_uncond->decode_embedding(embedded); + }); + uncond->logits = uncond_step.logits; + uncond->hidden = uncond_step.hidden; + } } }); engine::debug::timing_log_scalar("breeze_tts.ar.total_ms", ar_ms); @@ -975,7 +1107,9 @@ struct BreezeGeneratorRuntime::Impl { engine::debug::timing_log_scalar("breeze_tts.ar.backbone_cond_decode_ms", backbone_cond_decode_ms); engine::debug::timing_log_scalar("breeze_tts.ar.backbone_uncond_decode_ms", backbone_uncond_decode_ms); backbone_cond->release_runtime_graphs(); - backbone_uncond->release_runtime_graphs(); + if (use_cfg) { + backbone_uncond->release_runtime_graphs(); + } depth_pair->release_runtime_graphs(); if (codes.empty()) { throw std::runtime_error("BreezeTTS generated no audio codes"); diff --git a/src/models/breeze_tts/session.cpp b/src/models/breeze_tts/session.cpp index 264d67bd1..10f7e5e86 100644 --- a/src/models/breeze_tts/session.cpp +++ b/src/models/breeze_tts/session.cpp @@ -315,7 +315,10 @@ BreezeGenerationRequest BreezeTTSSession::build_generation_request( generation.text = request.text_input->text; generation.instruction = runtime::find_option(request.options, {"instruction"}).value_or(""); generation.reference_text = runtime::find_option(request.options, {"reference_text"}).value_or(""); - generation.guidance_scale = runtime::parse_positive_finite_float_option(request.options, {"guidance_scale"}).value_or(generation.guidance_scale); + generation.guidance_scale = runtime::parse_finite_float_option(request.options, {"guidance_scale"}).value_or(generation.guidance_scale); + if (generation.guidance_scale < 0.0F) { + throw std::runtime_error("BreezeTTS guidance_scale must be non-negative"); + } generation.temperature = runtime::parse_positive_finite_float_option(request.options, {"temperature"}).value_or(generation.temperature); generation.depth_temperature = runtime::parse_positive_finite_float_option(request.options, {"depth_temperature"}).value_or(generation.depth_temperature); generation.top_k = runtime::parse_i64_option(request.options, {"top_k"}).value_or(generation.top_k);