From 7552d8ec917e9cea67cf15dc1e9e2b64d95895c8 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 30 Sep 2026 23:56:35 -0700 Subject: [PATCH 1/4] MiniMax-H3: flash attention in the video VAE under --diffusion-fa; carry ggml-cuda speed patch The H3 video VAE decoder is a 36-layer ViT run per 16x16 latent tile. With only --diffusion-fa (what Studio passes) it used mul_mat + scale + softmax attention, materialising a 32 x L x L f32 score matrix per tile and layer. --diffusion-fa now also covers it; SD_H3_VAE_FLASH_ATTN=0 restores the old path. The text encoder is unchanged (that is what --fa would also switch, which moves the conditioning). scripts/unsloth/ggml-patches/ carries two ggml-cuda commits on top of the pinned submodule (BF16 cuBLAS for large-batch quantized matmul behind GGML_CUDA_QUANT_CUBLAS_MIN_BATCH, a cached device integrated flag, a vectorized row copy); the prebuilt workflow applies them after the submodule checkout. --- .github/workflows/unsloth-sd-prebuilt.yml | 6 + docs/minimax_h3.md | 4 + .../0001-ggml-cuda-h3-speed.patch | 255 ++++++++++++++++++ src/pipeline/diffusion_engine.cpp | 13 + 4 files changed, 278 insertions(+) create mode 100644 scripts/unsloth/ggml-patches/0001-ggml-cuda-h3-speed.patch diff --git a/.github/workflows/unsloth-sd-prebuilt.yml b/.github/workflows/unsloth-sd-prebuilt.yml index 291ee0f7ae..81708c3f11 100644 --- a/.github/workflows/unsloth-sd-prebuilt.yml +++ b/.github/workflows/unsloth-sd-prebuilt.yml @@ -222,6 +222,12 @@ jobs: # (we build with SD_SERVER_BUILD_FRONTEND / SD_WEBP / SD_WEBM off), so fetch only # ggml to keep the source tarball small and the build self-contained. git submodule update --init --recursive --depth 1 ggml + # ggml changes we carry on top of the pinned submodule commit (the submodule points at + # leejet/ggml, which we do not push to). A patch that no longer applies fails the build. + for p in scripts/unsloth/ggml-patches/*.patch; do + [ -e "$p" ] || continue + git -C ggml apply --verbose "../$p" + done COMMIT="$(git rev-parse HEAD)" SRC_ARTIFACT="sd-source-${TAG}" if [ "$EXISTS" != "true" ] || [ "${{ github.event_name }}" = "workflow_dispatch" ]; then diff --git a/docs/minimax_h3.md b/docs/minimax_h3.md index 82d0e2ffb4..0417d0011a 100644 --- a/docs/minimax_h3.md +++ b/docs/minimax_h3.md @@ -46,6 +46,10 @@ are detected from their weights. Omitting `--audio-vae` still runs the joint diffusion model but produces video without a decoded audio track. +`--diffusion-fa` also enables flash attention in the video VAE decoder (a ViT), which removes +most of its attention cost without touching the text encoder. Set `SD_H3_VAE_FLASH_ATTN=0` +to decode with the previous mul_mat + softmax attention. + ## First/last-frame conditioning Add `--init-img` for I2VA, or both `--init-img` and `--end-img` for FL2VA: diff --git a/scripts/unsloth/ggml-patches/0001-ggml-cuda-h3-speed.patch b/scripts/unsloth/ggml-patches/0001-ggml-cuda-h3-speed.patch new file mode 100644 index 0000000000..f8b379f189 --- /dev/null +++ b/scripts/unsloth/ggml-patches/0001-ggml-cuda-h3-speed.patch @@ -0,0 +1,255 @@ +From 22b560daa5e90bdb72c492f0b55667468ad3fd6b Mon Sep 17 00:00:00 2001 +From: Daniel Han +Date: Wed, 30 Sep 2026 22:28:21 -0700 +Subject: [PATCH 1/2] ggml-cuda: BF16 cuBLAS for large-batch quantized matmul, + cache device integrated flag + +GGML_CUDA_QUANT_CUBLAS_MIN_BATCH= dequantizes a quantized weight to BF16 +and runs a cuBLAS GEMM (FP32 accumulate, FP32 output) instead of MMQ when src1 +has at least that many rows; 0 or unset keeps MMQ everywhere. +GGML_CUDA_QUANT_CUBLAS_TYPE=f32 dequantizes to FP32 instead (accuracy oracle). +At diffusion-transformer sizes (MiniMax-H3 960x544x124: 19108 tokens per call) +MMQ reaches 130-240 TFLOP/s on a B200 while the BF16 GEMM reaches ~1300, and +the result is closer to an FP32 reference because activations stay BF16 +instead of being quantized to q8_1. + +ggml_backend_cuda_device_get_memory / get_type called cudaGetDeviceProperties +(~1-2 ms) on every call; stable-diffusion.cpp polls them per graph, 2798 times +in one tiled H3 VAE decode. The flag is now read once at registration. +--- + src/ggml-cuda/ggml-cuda.cu | 55 ++++++++++++++++++++++++++++++++------ + 1 file changed, 47 insertions(+), 8 deletions(-) + +diff --git a/src/ggml-cuda/ggml-cuda.cu b/src/ggml-cuda/ggml-cuda.cu +index e15fdc5e..517d0894 100644 +--- a/src/ggml-cuda/ggml-cuda.cu ++++ b/src/ggml-cuda/ggml-cuda.cu +@@ -2427,6 +2427,25 @@ static bool ggml_cuda_should_fuse_mul_mat_vec_q(const ggml_tensor * tensor) { + return use_mul_mat_vec_q; + } + ++// Large-batch matmul with a quantized weight: dequantize the weight to BF16 once and run a cuBLAS ++// GEMM (FP32 accumulate, FP32 output) instead of MMQ. MMQ quantizes the activations to q8_1 and runs ++// int8 MMA, which is the right call for LLM prompt sizes, but at diffusion-transformer sizes (10^4 ++// tokens per call) a tensor-core BF16 GEMM is faster on GPUs whose BF16 rate is not far below their ++// int8 rate, and it is more accurate (activations stay BF16 instead of 8-bit). ++// GGML_CUDA_QUANT_CUBLAS_MIN_BATCH=: use the BF16 path when src1 has at least this many rows; ++// 0 disables it (stock MMQ everywhere). Unset = the per-architecture default below. ++static int64_t ggml_cuda_quant_cublas_min_batch(int cc) { ++ static const int64_t env = [] { ++ const char * e = getenv("GGML_CUDA_QUANT_CUBLAS_MIN_BATCH"); ++ return e != nullptr ? (int64_t) atoll(e) : (int64_t) -1; ++ }(); ++ if (env >= 0) { ++ return env; ++ } ++ GGML_UNUSED(cc); ++ return 0; ++} ++ + static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { + GGML_TENSOR_BINARY_OP_LOCALS + +@@ -2492,6 +2511,23 @@ static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor + ggml_cuda_mul_mat_vec_q(ctx, src0, src1, nullptr, dst); + return; + } ++ if (ggml_is_quantized(src0->type) && ne11 > MMVQ_MAX_BATCH_SIZE && GGML_CUDA_CC_IS_NVIDIA(cc) && ++ bf16_mma_hardware_available(cc) && ggml_get_to_bf16_cuda(src0->type) != nullptr) { ++ const int64_t min_batch = ggml_cuda_quant_cublas_min_batch(cc); ++ if (min_batch > 0 && ne11*ne12*ne13 >= min_batch) { ++ // GGML_CUDA_QUANT_CUBLAS_TYPE=f32 dequantizes to FP32 instead: slow, for accuracy oracles only. ++ static const bool f32_oracle = [] { ++ const char * e = getenv("GGML_CUDA_QUANT_CUBLAS_TYPE"); ++ return e != nullptr && (strcmp(e, "f32") == 0 || strcmp(e, "fp32") == 0); ++ }(); ++ if (f32_oracle) { ++ ggml_cuda_mul_mat_cublas_impl(ctx, src0, src1, dst); ++ } else { ++ ggml_cuda_mul_mat_cublas_impl(ctx, src0, src1, dst); ++ } ++ return; ++ } ++ } + if (ggml_cuda_should_use_mmq(src0->type, cc, ne11, /*n_experts =*/ 0)) { + ggml_cuda_mul_mat_q(ctx, src0, src1, nullptr, dst); + return; +@@ -5200,6 +5236,10 @@ struct ggml_backend_cuda_device_context { + std::string description; + std::string pci_bus_id; + int op_offload_min_batch_size; ++ // cudaDeviceProp::integrated, read once at registration. cudaGetDeviceProperties costs ~1-2 ms per call, and ++ // get_memory / get_type are polled per graph by callers that budget VRAM (stable-diffusion.cpp asks thousands ++ // of times per tiled VAE decode, ~5 s of host time per H3 video on a B200 host). ++ bool integrated = false; + }; + + static const char * ggml_backend_cuda_device_get_name(ggml_backend_dev_t dev) { +@@ -5303,12 +5343,9 @@ static void ggml_backend_cuda_device_get_memory(ggml_backend_dev_t dev, size_t * + // ref: https://github.com/ggml-org/llama.cpp/pull/17368 + #if defined(__linux__) + // Check if this is a UMA (Unified Memory Architecture) system +- cudaDeviceProp prop; +- CUDA_CHECK(cudaGetDeviceProperties(&prop, ggml_cuda_get_physical_device(ctx->device))); +- + // Check if UMA is explicitly enabled via environment variable + bool uma_env = getenv("GGML_CUDA_ENABLE_UNIFIED_MEMORY") != nullptr; +- bool is_uma = prop.integrated > 0 || uma_env; ++ bool is_uma = ctx->integrated || uma_env; + + if (is_uma) { + // For UMA systems (like DGX Spark), use system memory info +@@ -5332,10 +5369,7 @@ static void ggml_backend_cuda_device_get_memory(ggml_backend_dev_t dev, size_t * + static enum ggml_backend_dev_type ggml_backend_cuda_device_get_type(ggml_backend_dev_t dev) { + ggml_backend_cuda_device_context * ctx = (ggml_backend_cuda_device_context *) dev->context; + +- cudaDeviceProp prop; +- CUDA_CHECK(cudaGetDeviceProperties(&prop, ggml_cuda_get_physical_device(ctx->device))); +- +- return prop.integrated ++ return ctx->integrated + ? GGML_BACKEND_DEVICE_TYPE_IGPU + : GGML_BACKEND_DEVICE_TYPE_GPU; + } +@@ -6095,6 +6129,11 @@ ggml_backend_reg_t ggml_backend_cuda_reg() { + dev_ctx->device = i; + dev_ctx->name = GGML_CUDA_NAME + std::to_string(i); + dev_ctx->description = ggml_cuda_device_description(i); ++ { ++ cudaDeviceProp prop; ++ CUDA_CHECK(cudaGetDeviceProperties(&prop, physical_id)); ++ dev_ctx->integrated = prop.integrated > 0; ++ } + + char pci_bus_id[32] = {}; + CUDA_CHECK(cudaDeviceGetPCIBusId(pci_bus_id, sizeof(pci_bus_id), physical_id)); +-- +2.43.0 + + +From 0bf02a4e324ef837685c63ebaeddf2b4a4c0f317 Mon Sep 17 00:00:00 2001 +From: Daniel Han +Date: Wed, 30 Sep 2026 22:54:01 -0700 +Subject: [PATCH 2/2] ggml-cuda: vectorized row copy for strided F32->F32/F16 + copies with contiguous rows + +cpy_scalar recomputes the 4-D source and destination index (64-bit div/mod) for +every element, so the permuted ggml_cont copies of a DiT's Q/K/V (rows of +head_dim floats) run at a few percent of memory bandwidth: a 411 MB +[128,42,19108] permute copy takes 1.67 ms on a B200, 0.30 ms as one float4 per +thread with the index math done once per row. Pure copy / same per-element +conversion, so the output is bit-identical; GGML_CUDA_CPY_ROWS=0 restores +cpy_scalar. +--- + src/ggml-cuda/cpy.cu | 83 ++++++++++++++++++++++++++++++++++++++++++++ + 1 file changed, 83 insertions(+) + +diff --git a/src/ggml-cuda/cpy.cu b/src/ggml-cuda/cpy.cu +index 4c6fb4a2..c26a081c 100644 +--- a/src/ggml-cuda/cpy.cu ++++ b/src/ggml-cuda/cpy.cu +@@ -1,4 +1,5 @@ + #include "cpy.cuh" ++#include + #include "dequantize.cuh" + #include "cpy-utils.cuh" + #include "fp8.cuh" +@@ -203,6 +204,84 @@ cudaStream_t stream) { + ggml_cuda_kernel_launch(cpy_scalar_contiguous, launch_params, cx, cdst, ne); + } + ++// Strided copy whose rows (dim 0) are contiguous on both sides and of equal length: one float4 per thread, the ++// 4-D index math done once per row instead of per element. cpy_scalar spends most of its time in the 64-bit ++// div/mod of that per-element math; the permuted ggml_cont copies of a DiT's Q/K/V (rows of head_dim floats) run ++// at a few percent of memory bandwidth through it. A pure copy (or the same per-element conversion), so the ++// result is bit-identical to cpy_scalar. GGML_CUDA_CPY_ROWS=0 disables it. ++template ++static __global__ void cpy_rows_vec4(const char * cx, char * cdst, const int64_t nrows, const int64_t ne00_4, ++ const int64_t ne01, const int64_t ne02, const int64_t nb01, const int64_t nb02, const int64_t nb03, ++ const int64_t ne11, const int64_t ne12, const int64_t nb11, const int64_t nb12, const int64_t nb13) { ++ ggml_cuda_pdl_lc(); ++ const int64_t row = (int64_t) blockIdx.x*blockDim.y + threadIdx.y; ++ if (row >= nrows) { ++ return; ++ } ++ const int64_t i01 = row % ne01; ++ const int64_t i02 = (row / ne01) % ne02; ++ const int64_t i03 = row / (ne01*ne02); ++ const int64_t i11 = row % ne11; ++ const int64_t i12 = (row / ne11) % ne12; ++ const int64_t i13 = row / (ne11*ne12); ++ const float4 * src = (const float4 *) (cx + i01*nb01 + i02*nb02 + i03*nb03); ++ char * dst_row = cdst + i11*nb11 + i12*nb12 + i13*nb13; ++ ggml_cuda_pdl_sync(); ++ for (int64_t c = threadIdx.x; c < ne00_4; c += blockDim.x) { ++ const float4 v = src[c]; ++ if constexpr (std::is_same_v) { ++ ((float4 *) dst_row)[c] = v; ++ } else { ++ dst_t * d = (dst_t *) dst_row + 4*c; ++ cpy_1_scalar((const char *) &v.x, (char *) (d + 0)); ++ cpy_1_scalar((const char *) &v.y, (char *) (d + 1)); ++ cpy_1_scalar((const char *) &v.z, (char *) (d + 2)); ++ cpy_1_scalar((const char *) &v.w, (char *) (d + 3)); ++ } ++ } ++} ++ ++static bool ggml_cuda_cpy_rows_enabled() { ++ static const bool enabled = [] { ++ const char * e = getenv("GGML_CUDA_CPY_ROWS"); ++ return e == nullptr || atoi(e) != 0; ++ }(); ++ return enabled; ++} ++ ++// True (and launched) when src0 -> src1 is a row copy cpy_rows_vec4 can do; false leaves the copy to the caller. ++template ++static bool ggml_cpy_rows_vec4_cuda(const ggml_tensor * src0, ggml_tensor * src1, cudaStream_t stream) { ++ if (!ggml_cuda_cpy_rows_enabled() || src0->type != GGML_TYPE_F32) { ++ return false; ++ } ++ const int64_t ne00 = src0->ne[0]; ++ if (src0->nb[0] != (int64_t) sizeof(float) || src1->nb[0] != (int64_t) sizeof(dst_t) || src1->ne[0] != ne00 || ne00 % 4 != 0) { ++ return false; ++ } ++ const size_t dst_align = std::is_same_v ? 16 : 2*sizeof(dst_t); ++ if (((uintptr_t) src0->data | src0->nb[1] | src0->nb[2] | src0->nb[3]) % 16 != 0 || ++ ((uintptr_t) src1->data | src1->nb[1] | src1->nb[2] | src1->nb[3]) % dst_align != 0) { ++ return false; ++ } ++ const int64_t nrows = ggml_nrows(src0); ++ if (nrows != ggml_nrows(src1) || src0->ne[1] == 0 || src0->ne[2] == 0 || src1->ne[1] == 0 || src1->ne[2] == 0) { ++ return false; ++ } ++ const int64_t ne00_4 = ne00 / 4; ++ const int threads_x = ne00_4 >= 32 ? 32 : (int) ne00_4; ++ const int threads_y = 256 / threads_x; ++ const int64_t blocks = (nrows + threads_y - 1) / threads_y; ++ if (blocks > INT_MAX) { ++ return false; ++ } ++ const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(dim3((unsigned) blocks), dim3(threads_x, threads_y), 0, stream); ++ ggml_cuda_kernel_launch(cpy_rows_vec4, launch_params, (const char *) src0->data, (char *) src1->data, nrows, ne00_4, ++ src0->ne[1], src0->ne[2], (int64_t) src0->nb[1], (int64_t) src0->nb[2], (int64_t) src0->nb[3], ++ src1->ne[1], src1->ne[2], (int64_t) src1->nb[1], (int64_t) src1->nb[2], (int64_t) src1->nb[3]); ++ return true; ++} ++ + template + static void ggml_cpy_scalar_cuda( + const char * cx, char * cdst, const int64_t ne, +@@ -509,6 +588,10 @@ void ggml_cuda_cpy(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, gg + ggml_cpy_scalar_cuda + (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream); + } ++ } else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F32 && ggml_cpy_rows_vec4_cuda(src0, src1, main_stream)) { ++ // row copy done ++ } else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F16 && !contiguous_srcs && ggml_cpy_rows_vec4_cuda(src0, src1, main_stream)) { ++ // row copy with conversion done + } else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F32) { + if (can_be_transposed) { + ggml_cpy_scalar_cuda +-- +2.43.0 + diff --git a/src/pipeline/diffusion_engine.cpp b/src/pipeline/diffusion_engine.cpp index e41c1151a8..eba4cca4da 100644 --- a/src/pipeline/diffusion_engine.cpp +++ b/src/pipeline/diffusion_engine.cpp @@ -4,6 +4,7 @@ #include #include #include +#include #include #include #include @@ -1177,6 +1178,18 @@ bool StableDiffusionGGML::validate_and_load_runners() { high_noise_diffusion_model->set_flash_attention_enabled(true); } } + // The MiniMax-H3 video VAE decoder is a 36-layer ViT (32 x 64 heads, ~1.5k tokens per 16x16 latent tile). + // Without flash attention every tile and layer materialises a 32 x L x L f32 score matrix (mul_mat + scale + + // softmax + mul_mat), which is most of the decode's GPU time. --diffusion-fa therefore also covers it, the + // way --fa would, without switching the text encoder's attention. SD_H3_VAE_FLASH_ATTN=0 restores the old path. + if (!sd_ctx_params->flash_attn && sd_ctx_params->diffusion_flash_attn && first_stage_model && + sd_version_is_minimax_h3(version)) { + const char* env = getenv("SD_H3_VAE_FLASH_ATTN"); + if (env == nullptr || strcmp(env, "0") != 0) { + LOG_INFO("Using flash attention in the MiniMax-H3 video VAE"); + first_stage_model->set_flash_attention_enabled(true); + } + } if (sd_ctx_params->sage_attn && !set_sage_attention_enabled(true)) { return false; } From b01506b950a67ec841d4eb02a0d49c170def2b48 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Thu, 1 Oct 2026 00:07:32 -0700 Subject: [PATCH 2/4] tiling: walk tile planes contiguously when splitting and merging process_tiles_2d copied and blended each pixel across all planes in the inner loop, striding a whole output plane (960x544 floats for H3) per element. A video tile has frames x channels planes, so every element was a cache and TLB miss: on MiniMax-H3 960x544x124 (105 tiles) this was about 10 s of single-threaded host time per decode. Loop planes outermost and precompute the per-column and per-row blend weights; each element still computes old + new * wy * wx in the same order, so the output is bitwise identical. --- src/runtime/tiling.cpp | 77 ++++++++++++++++++++++++++++-------------- 1 file changed, 51 insertions(+), 26 deletions(-) diff --git a/src/runtime/tiling.cpp b/src/runtime/tiling.cpp index f0a9dbc73b..767e9c4d19 100644 --- a/src/runtime/tiling.cpp +++ b/src/runtime/tiling.cpp @@ -74,12 +74,20 @@ static sd::Tensor sd_tensor_split_2d(const sd::Tensor& input, int int64_t input_plane = sd_tensor_plane_size(input); int64_t output_plane = sd_tensor_plane_size(output); int64_t plane_count = input.numel() / input_plane; - for (int iy = 0; iy < height; iy++) { - for (int ix = 0; ix < width; ix++) { - int64_t src_xy = (ix + x) % input_width + input_width * ((iy + y) % input_height); - int64_t dst_xy = ix + width * iy; - for (int64_t plane = 0; plane < plane_count; ++plane) { - output[plane * output_plane + dst_xy] = input[plane * input_plane + src_xy]; + // Plane-outer order: a plane is contiguous, so the inner loop streams instead of striding a + // whole plane per element (a video tile has frames x channels planes). + std::vector src_x(static_cast(width)); + for (int ix = 0; ix < width; ix++) { + src_x[ix] = (ix + x) % input_width; + } + for (int64_t plane = 0; plane < plane_count; ++plane) { + const float* src = input.data() + plane * input_plane; + float* dst = output.data() + plane * output_plane; + for (int iy = 0; iy < height; iy++) { + const float* src_row = src + input_width * ((iy + y) % input_height); + float* dst_row = dst + static_cast(width) * iy; + for (int ix = 0; ix < width; ix++) { + dst_row[ix] = src_row[src_x[ix]]; } } } @@ -112,26 +120,43 @@ static void sd_tensor_merge_2d(const sd::Tensor& input, return x * x * x * (x * (6.0f * x - 15.0f) + 10.0f); }; - for (int iy = y_skip; iy < height; iy++) { - for (int ix = x_skip; ix < width; ix++) { - int64_t src_xy = ix + width * iy; - int64_t ox = (x + ix) % img_width; - int64_t oy = (y + iy) % img_height; - int64_t dst_xy = ox + img_width * oy; - for (int64_t plane = 0; plane < plane_count; ++plane) { - float new_value = input[plane * input_plane + src_xy]; - if (overlap_x > 0 || overlap_y > 0) { - float old_value = (*output)[plane * output_plane + dst_xy]; - const float x_f_0 = (circular_x || (overlap_x > 0 && x > 0)) ? (ix - x_skip) / float(overlap_x) : 1.f; - const float x_f_1 = (circular_x || (overlap_x > 0 && x < (img_width - width))) ? (width - ix) / float(overlap_x) : 1.f; - const float y_f_0 = (circular_y || (overlap_y > 0 && y > 0)) ? (iy - y_skip) / float(overlap_y) : 1.f; - const float y_f_1 = (circular_y || (overlap_y > 0 && y < (img_height - height))) ? (height - iy) / float(overlap_y) : 1.f; - const float x_f = std::min(std::min(x_f_0, x_f_1), 1.f); - const float y_f = std::min(std::min(y_f_0, y_f_1), 1.f); - (*output)[plane * output_plane + dst_xy] = - old_value + new_value * smootherstep_f32(y_f) * smootherstep_f32(x_f); - } else { - (*output)[plane * output_plane + dst_xy] = new_value; + // Weights depend on ix or iy only; precompute them and walk each contiguous plane row by row. + // Same arithmetic per element as the per-pixel form, so the result is bit-identical. + const bool blend = overlap_x > 0 || overlap_y > 0; + std::vector wx, wy; + std::vector dst_x(static_cast(width)); + for (int64_t ix = x_skip; ix < width; ix++) { + dst_x[ix] = (x + ix) % img_width; + } + if (blend) { + wx.resize(static_cast(width)); + wy.resize(static_cast(height)); + for (int64_t ix = x_skip; ix < width; ix++) { + const float x_f_0 = (circular_x || (overlap_x > 0 && x > 0)) ? (ix - x_skip) / float(overlap_x) : 1.f; + const float x_f_1 = (circular_x || (overlap_x > 0 && x < (img_width - width))) ? (width - ix) / float(overlap_x) : 1.f; + wx[ix] = smootherstep_f32(std::min(std::min(x_f_0, x_f_1), 1.f)); + } + for (int64_t iy = y_skip; iy < height; iy++) { + const float y_f_0 = (circular_y || (overlap_y > 0 && y > 0)) ? (iy - y_skip) / float(overlap_y) : 1.f; + const float y_f_1 = (circular_y || (overlap_y > 0 && y < (img_height - height))) ? (height - iy) / float(overlap_y) : 1.f; + wy[iy] = smootherstep_f32(std::min(std::min(y_f_0, y_f_1), 1.f)); + } + } + for (int64_t plane = 0; plane < plane_count; ++plane) { + const float* src = input.data() + plane * input_plane; + float* dst = output->data() + plane * output_plane; + for (int64_t iy = y_skip; iy < height; iy++) { + const float* src_row = src + width * iy; + float* dst_row = dst + img_width * ((y + iy) % img_height); + if (blend) { + const float sy = wy[iy]; + for (int64_t ix = x_skip; ix < width; ix++) { + float& out = dst_row[dst_x[ix]]; + out = out + src_row[ix] * sy * wx[ix]; + } + } else { + for (int64_t ix = x_skip; ix < width; ix++) { + dst_row[dst_x[ix]] = src_row[ix]; } } } From ca92d95ddf041bf885a36a7fdafbffd88706c28e Mon Sep 17 00:00:00 2001 From: oobabooga <112222186+oobabooga@users.noreply.github.com> Date: Thu, 8 Oct 2026 00:15:39 -0300 Subject: [PATCH 3/4] Tighten H3 VAE and tiling comments --- src/pipeline/diffusion_engine.cpp | 6 ++---- src/runtime/tiling.cpp | 6 ++---- 2 files changed, 4 insertions(+), 8 deletions(-) diff --git a/src/pipeline/diffusion_engine.cpp b/src/pipeline/diffusion_engine.cpp index eba4cca4da..ec50c620cc 100644 --- a/src/pipeline/diffusion_engine.cpp +++ b/src/pipeline/diffusion_engine.cpp @@ -1178,10 +1178,8 @@ bool StableDiffusionGGML::validate_and_load_runners() { high_noise_diffusion_model->set_flash_attention_enabled(true); } } - // The MiniMax-H3 video VAE decoder is a 36-layer ViT (32 x 64 heads, ~1.5k tokens per 16x16 latent tile). - // Without flash attention every tile and layer materialises a 32 x L x L f32 score matrix (mul_mat + scale + - // softmax + mul_mat), which is most of the decode's GPU time. --diffusion-fa therefore also covers it, the - // way --fa would, without switching the text encoder's attention. SD_H3_VAE_FLASH_ATTN=0 restores the old path. + // The H3 video VAE decoder is a ViT whose f32 L x L attention scores dominate decode time, so --diffusion-fa + // covers it too, without switching the text encoder. SD_H3_VAE_FLASH_ATTN=0 restores the old path. if (!sd_ctx_params->flash_attn && sd_ctx_params->diffusion_flash_attn && first_stage_model && sd_version_is_minimax_h3(version)) { const char* env = getenv("SD_H3_VAE_FLASH_ATTN"); diff --git a/src/runtime/tiling.cpp b/src/runtime/tiling.cpp index 767e9c4d19..782dc4083d 100644 --- a/src/runtime/tiling.cpp +++ b/src/runtime/tiling.cpp @@ -74,8 +74,7 @@ static sd::Tensor sd_tensor_split_2d(const sd::Tensor& input, int int64_t input_plane = sd_tensor_plane_size(input); int64_t output_plane = sd_tensor_plane_size(output); int64_t plane_count = input.numel() / input_plane; - // Plane-outer order: a plane is contiguous, so the inner loop streams instead of striding a - // whole plane per element (a video tile has frames x channels planes). + // Plane-outer: planes are contiguous, so the inner loop streams instead of striding a plane per element. std::vector src_x(static_cast(width)); for (int ix = 0; ix < width; ix++) { src_x[ix] = (ix + x) % input_width; @@ -120,8 +119,7 @@ static void sd_tensor_merge_2d(const sd::Tensor& input, return x * x * x * (x * (6.0f * x - 15.0f) + 10.0f); }; - // Weights depend on ix or iy only; precompute them and walk each contiguous plane row by row. - // Same arithmetic per element as the per-pixel form, so the result is bit-identical. + // Same per-element arithmetic as the per-pixel form, so the result stays bit-identical. const bool blend = overlap_x > 0 || overlap_y > 0; std::vector wx, wy; std::vector dst_x(static_cast(width)); From ff941a8a853c951c4f77310d5ba204ea7bf9ac0e Mon Sep 17 00:00:00 2001 From: oobabooga <112222186+oobabooga@users.noreply.github.com> Date: Thu, 8 Oct 2026 22:51:20 -0300 Subject: [PATCH 4/4] Keep the H3 VAE on its old attention outside CUDA and ROCm --- docs/minimax_h3.md | 8 +++++--- src/pipeline/diffusion_engine.cpp | 11 +++++++---- 2 files changed, 12 insertions(+), 7 deletions(-) diff --git a/docs/minimax_h3.md b/docs/minimax_h3.md index 0417d0011a..bcceb73c47 100644 --- a/docs/minimax_h3.md +++ b/docs/minimax_h3.md @@ -46,9 +46,11 @@ are detected from their weights. Omitting `--audio-vae` still runs the joint diffusion model but produces video without a decoded audio track. -`--diffusion-fa` also enables flash attention in the video VAE decoder (a ViT), which removes -most of its attention cost without touching the text encoder. Set `SD_H3_VAE_FLASH_ATTN=0` -to decode with the previous mul_mat + softmax attention. +On CUDA and ROCm, `--diffusion-fa` also enables flash attention in the video VAE decoder (a ViT), +which removes most of its attention cost without touching the text encoder. Set +`SD_H3_VAE_FLASH_ATTN=0` to decode with the previous mul_mat + softmax attention. Other backends +keep that attention by default (on Vulkan the flash attention decode was slower); +`SD_H3_VAE_FLASH_ATTN=1` turns it on there. ## First/last-frame conditioning diff --git a/src/pipeline/diffusion_engine.cpp b/src/pipeline/diffusion_engine.cpp index ec50c620cc..db5800991c 100644 --- a/src/pipeline/diffusion_engine.cpp +++ b/src/pipeline/diffusion_engine.cpp @@ -1178,12 +1178,15 @@ bool StableDiffusionGGML::validate_and_load_runners() { high_noise_diffusion_model->set_flash_attention_enabled(true); } } - // The H3 video VAE decoder is a ViT whose f32 L x L attention scores dominate decode time, so --diffusion-fa - // covers it too, without switching the text encoder. SD_H3_VAE_FLASH_ATTN=0 restores the old path. + // The H3 video VAE decoder is a ViT whose f32 L x L attention scores dominate decode time, so on CUDA and ROCm + // --diffusion-fa covers it too, without switching the text encoder. It decoded slower on Vulkan, so other + // backends keep the old path unless SD_H3_VAE_FLASH_ATTN=1; SD_H3_VAE_FLASH_ATTN=0 restores it everywhere. if (!sd_ctx_params->flash_attn && sd_ctx_params->diffusion_flash_attn && first_stage_model && sd_version_is_minimax_h3(version)) { - const char* env = getenv("SD_H3_VAE_FLASH_ATTN"); - if (env == nullptr || strcmp(env, "0") != 0) { + const char* env = getenv("SD_H3_VAE_FLASH_ATTN"); + ggml_backend_t vae_backend = backend_for(SDBackendModule::VAE); + const bool fa_faster_by_default = sd_backend_is(vae_backend, "CUDA") || sd_backend_is(vae_backend, "ROCm"); + if (env != nullptr ? strcmp(env, "0") != 0 : fa_faster_by_default) { LOG_INFO("Using flash attention in the MiniMax-H3 video VAE"); first_stage_model->set_flash_attention_enabled(true); }