leejet · Nov 1, 2026 · Nov 1, 2026 · Nov 1, 2026
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -378,7 +378,7 @@ if (NOT SD_USE_UPSTREAM_GGML AND (SD_CUDA OR SD_HIPBLAS OR SD_MUSA))
 endif()
 if (SD_VULKAN AND NOT SD_USE_UPSTREAM_GGML)
     add_subdirectory(src/core/vulkan)
-    list(APPEND SD_BACKEND_REGISTRATIONS sd_ggml_register_vulkan_rope sd_ggml_register_vulkan_i8)
+    list(APPEND SD_BACKEND_REGISTRATIONS sd_ggml_register_vulkan_rope sd_ggml_register_vulkan_i8 sd_ggml_register_vulkan_modulate)
 endif()
 if(SD_BACKEND_REGISTRATIONS)
     string(REPLACE ";" ",\n    " SD_BACKEND_REGISTRATIONS "${SD_BACKEND_REGISTRATIONS}")

diff --git a/cmake/backend_extensions.h.in b/cmake/backend_extensions.h.in
@@ -5,6 +5,7 @@
 
 void sd_ggml_register_vulkan_rope(ggml_backend_dev_t device);
 void sd_ggml_register_vulkan_i8(ggml_backend_dev_t device);
+void sd_ggml_register_vulkan_modulate(ggml_backend_dev_t device);
 
 static void (*const sd_backend_extensions[])(ggml_backend_dev_t) = {
     @SD_BACKEND_REGISTRATIONS@

diff --git a/docs/ggml_extensions.md b/docs/ggml_extensions.md
@@ -135,6 +135,25 @@ CPU kernels and fusion rules live in `src/core/cpu/`, with registration collecte
 in `cpu/registration.cpp`. CUDA-style and Vulkan implementations live in their
 respective backend directories.
 
+`ggml_ext_modulate(ctx, x, shift, scale)` emits `sd.modulate.v1` for F32 inputs
+when the runner context uses CPU, CUDA/HIP/MUSA, or Vulkan.
+Its explicit inputs are `x`, `scale`, and optional `shift`; it computes
+`(x + x * scale) + shift` with separate FP32 multiply/add rounding. The output
+is contiguous and out of place. Inputs may be strided and use ggml's
+dimension-wise repetition rules. Callers prepare token, batch, or frame layouts
+before constructing the operation; normalization and residual gates remain
+separate operations. CPU, CUDA/HIP/MUSA, and Vulkan have standalone kernels.
+Other backends (including Metal, SYCL, and OpenCL), upstream ggml builds, and
+other input types retain the original MUL/ADD graph.
+
+ConvRot fusion matches the Modulate node directly, including scale-only
+modulation, instead of recognizing its former MUL/ADD/ADD decomposition.
+CUDA accepts contiguous per-channel parameters for `Modulate -> ConvRot` and
+`RMSNorm -> Mul(weight) -> Modulate -> ConvRot`. Vulkan limits the first fusion
+to widths 4096/6144/5376 and the RMSNorm fusion to width 5376.
+Other layouts execute the standalone operation followed by ConvRot. The
+gated-residual MUL/ADD fusion remains available for residual branches.
+
 `src/core/cpu/rope.cpp` registers the CPU implementation of
 `sd.rope_apply.v1` and a
 `RMSNorm -> Mul -> RoPE` fusion. The fusion accepts a one-dimensional weight
@@ -269,8 +288,9 @@ unsupported layouts retain ordinary operator execution. Fusion and force-cuBLAS
 disable settings are honored.
 
 `src/core/cuda/convrot.cu` registers CUDA fusion for `Scale -> ConvRot`,
-`Norm -> Mul -> ConvRot`, `RMSNorm -> Mul(weight) -> Mul(modulation) -> Add ->
-Add(shift) -> ConvRot`, `SwiGLU -> ConvRot`, and `SiLU/Sigmoid -> Mul -> ConvRot`.
+`Norm -> Mul -> ConvRot`, `Modulate -> ConvRot`,
+`RMSNorm -> Mul(weight) -> Modulate -> ConvRot`, `SwiGLU -> ConvRot`, and
+`SiLU/Sigmoid -> Mul -> ConvRot`.
 ConvRot outputs are custom nodes. SwiGLU accepts packed,
 swapped, or separate gate/up inputs. Independent modulation and up-branch
 producers are moved ahead of their consumers through the safe graph-delay API.

diff --git a/src/core/cpu/modulate.cpp b/src/core/cpu/modulate.cpp
@@ -0,0 +1,82 @@
+#include "core/cpu/registration.h"
+#include "core/ggml_extensions.h"
+
+#ifndef SD_USE_UPSTREAM_GGML
+#include <cstring>
+
+namespace sd_cpu {
+    static bool modulate_supports(ggml_backend_dev_t, const ggml_tensor* node, void*) {
+        return sd_ggml_modulate_describe(node);
+    }
+
+    static float modulate_value(const ggml_tensor* t, const int64_t* index) {
+        size_t offset = 0;
+        for (int d = 0; d < GGML_MAX_DIMS; ++d) {
+            offset += (index[d] % t->ne[d]) * t->nb[d];
+        }
+        return *reinterpret_cast<const float*>(static_cast<const char*>(t->data) + offset);
+    }
+
+    static void modulate_compute(ggml_tensor* dst, int ith, int nth, void*) {
+        auto* output       = static_cast<float*>(dst->data);
+        const int64_t rows = ggml_nrows(dst);
+        for (int64_t row = ith; row < rows; row += nth) {
+            int64_t index[] = {0, row % dst->ne[1], (row / dst->ne[1]) % dst->ne[2], row / (dst->ne[1] * dst->ne[2])};
+            for (int64_t col = 0; col < dst->ne[0]; ++col) {
+                index[0]      = col;
+                const float x = modulate_value(dst->src[0], index);
+                // Match separate ggml MUL/ADD rounding even with contraction enabled.
+                volatile float product         = x * modulate_value(dst->src[1], index);
+                volatile float scaled          = x + product;
+                output[row * dst->ne[0] + col] = dst->src[2] ? scaled + modulate_value(dst->src[2], index) : scaled;
+            }
+        }
+    }
+
+    static ggml_status modulate_prepare(ggml_backend_t, ggml_tensor* const* nodes, int, void**, size_t* bytes, void*) {
+        *bytes = sd_ggml_modulate_workspace_size(nodes[0]);
+        return GGML_STATUS_SUCCESS;
+    }
+
+    static void modulate_destroy(void*, void*) {}
+
+    static ggml_status modulate_execute(const ggml_ext_execution_context* context, ggml_tensor* const* nodes, int, void*, void*) {
+        auto reg = ggml_backend_dev_backend_reg(ggml_backend_get_device(context->backend));
+        auto run = reinterpret_cast<ggml_ext_cpu_parallel_fn>(ggml_backend_reg_get_proc_address(reg, "ggml_ext_cpu_parallel_v1"));
+        GGML_ASSERT(run);
+        auto* dst          = nodes[0];
+        const auto address = reinterpret_cast<uintptr_t>(dst->data);
+        bool alias         = false;
+        for (const auto* input : dst->src) {
+            if (input) {
+                const auto src = reinterpret_cast<uintptr_t>(input->data);
+                alias |= address < src + ggml_nbytes(input) && src < address + ggml_nbytes(dst);
+            }
+        }
+        if (!alias) {
+            return run(context->backend, dst, modulate_compute, nullptr);
+        }
+        ggml_tensor temporary = *dst;
+        temporary.data        = ggml_backend_buffer_get_base(context->workspace);
+        const auto status     = run(context->backend, &temporary, modulate_compute, nullptr);
+        if (status == GGML_STATUS_SUCCESS) {
+            std::memcpy(dst->data, temporary.data, ggml_nbytes(dst));
+        }
+        return status;
+    }
+
+    void register_modulate(ggml_backend_dev_t device) {
+        ggml_ext_kernel_desc desc = {};
+        desc.abi_version          = GGML_EXT_ABI_VERSION;
+        desc.name                 = "sd.modulate.cpu.v1";
+        desc.op                   = GGML_OP_CUSTOM;
+        desc.custom_op            = sd_ggml_modulate_op();
+        desc.priority             = 100;
+        desc.supports             = modulate_supports;
+        desc.prepare              = modulate_prepare;
+        desc.destroy              = modulate_destroy;
+        desc.execute              = modulate_execute;
+        GGML_ASSERT(ggml_ext_register_kernel(device, &desc));
+    }
+}
+#endif
diff --git a/src/core/cpu/registration.cpp b/src/core/cpu/registration.cpp
@@ -6,6 +6,7 @@ namespace sd_cpu {
     void register_extensions(ggml_backend_dev_t device) {
         register_i8(device);
         register_rope(device);
+        register_modulate(device);
     }
 
 }

diff --git a/src/core/cpu/registration.h b/src/core/cpu/registration.h
@@ -7,6 +7,7 @@
 namespace sd_cpu {
     void register_extensions(ggml_backend_dev_t device);
     void register_rope(ggml_backend_dev_t device);
+    void register_modulate(ggml_backend_dev_t device);
     void register_i8(ggml_backend_dev_t device);
 }
 #endif

diff --git a/src/core/cuda/CMakeLists.txt b/src/core/cuda/CMakeLists.txt
@@ -2,7 +2,7 @@ cmake_policy(SET CMP0079 NEW)
 
 function(sd_cuda_sources VENDOR LANGUAGE)
     set(SOURCES)
-    foreach(OP rope int8 convrot affine registration)
+    foreach(OP rope int8 convrot affine modulate registration)
         set(SOURCE "${CMAKE_CURRENT_BINARY_DIR}/${OP}_${VENDOR}.cu")
         string(TOUPPER "${VENDOR}" UPPER_VENDOR)
         file(GENERATE OUTPUT "${SOURCE}" CONTENT "#define SD_CUDA_VENDOR_${UPPER_VENDOR} 1\n#include \"core/cuda/${OP}.cu\"\n")

diff --git a/src/core/cuda/convrot.cu b/src/core/cuda/convrot.cu
@@ -27,6 +27,7 @@ namespace SD_CUDA_NAMESPACE {
         }
         enum Kind { SCALE,
                     NORM,
+                    MOD,
                     RMS_MOD,
                     SWIGLU,
                     SILU,
@@ -65,16 +66,20 @@ namespace SD_CUDA_NAMESPACE {
                 op.eps        = param_f32(first, 0);
                 if (op.k < 1024 || !vector(op.modulation, op.k))
                     return false;
-            } else if (count == 6 && first->op == GGML_OP_RMS_NORM && nodes[1]->op == GGML_OP_MUL && nodes[1]->src[0] == first &&
-                       nodes[2]->op == GGML_OP_MUL && nodes[2]->src[0] == nodes[1] &&
-                       nodes[3]->op == GGML_OP_ADD && nodes[3]->src[0] == nodes[1] && nodes[3]->src[1] == nodes[2] &&
-                       nodes[4]->op == GGML_OP_ADD && nodes[4]->src[0] == nodes[3] && value == nodes[4]) {
+            } else if (count == 2 && sd_ggml_modulate_describe(first) && value == first) {
+                op.kind       = MOD;
+                op.modulation = first->src[1];
+                op.shift      = first->src[2];
+                if (!vector(op.modulation, op.k) || (op.shift && !vector(op.shift, op.k)))
+                    return false;
+            } else if (count == 4 && first->op == GGML_OP_RMS_NORM && nodes[1]->op == GGML_OP_MUL && nodes[1]->src[0] == first &&
+                       sd_ggml_modulate_describe(nodes[2]) && nodes[2]->src[0] == nodes[1] && value == nodes[2]) {
                 op.kind       = RMS_MOD;
                 op.weight     = nodes[1]->src[1];
                 op.modulation = nodes[2]->src[1];
-                op.shift      = nodes[4]->src[1];
+                op.shift      = nodes[2]->src[2];
                 op.eps        = param_f32(first, 0);
-                if (op.k < 1024 || !vector(op.weight, op.k) || !vector(op.modulation, op.k) || !vector(op.shift, op.k))
+                if (op.k < 1024 || !vector(op.weight, op.k) || !vector(op.modulation, op.k) || (op.shift && !vector(op.shift, op.k)))
                     return false;
             } else if (count == 2 && first->op == GGML_OP_GLU && ggml_get_glu_op(first) == GGML_GLU_OP_SWIGLU && value == first) {
                 op.kind = SWIGLU;
@@ -119,7 +124,7 @@ namespace SD_CUDA_NAMESPACE {
             auto* g             = const_cast<ggml_cgraph*>(graph);
             const int available = ggml_graph_n_nodes(g) - first;
             const auto* node    = ggml_graph_node(g, first);
-            const int count     = node->op == GGML_OP_RMS_NORM ? 6 : node->op == GGML_OP_NORM || node->op == GGML_OP_UNARY ? 3
+            const int count     = node->op == GGML_OP_RMS_NORM ? 4 : node->op == GGML_OP_NORM || node->op == GGML_OP_UNARY ? 3
                                                                                                                            : 2;
             if (available < count)
                 return 0;
@@ -132,25 +137,26 @@ namespace SD_CUDA_NAMESPACE {
         static void optimize(ggml_backend_dev_t device, ggml_cgraph* graph, void*) {
             if (fusion_disabled())
                 return;
-            for (int i = 2; i < ggml_graph_n_nodes(graph); ++i) {
+            for (int i = 1; i < ggml_graph_n_nodes(graph); ++i) {
                 auto* quant = ggml_graph_node(graph, i);
                 int group;
                 if (!sd_ggml_convrot_describe(quant, &group) || group != 256)
                     continue;
                 ggml_tensor* chain[6] = {};
                 int count             = 0;
                 auto* value           = quant->src[0];
-                if (value->op == GGML_OP_ADD && value->src[0]->op == GGML_OP_ADD) {
-                    auto* add    = value->src[0];
-                    auto* weight = add->src[0];
-                    if (weight->op == GGML_OP_MUL) {
+                if (sd_ggml_modulate_describe(value)) {
+                    auto* weight = value->src[0];
+                    if (weight->op == GGML_OP_MUL && weight->src[0]->op == GGML_OP_RMS_NORM) {
                         chain[0] = weight->src[0];
                         chain[1] = weight;
-                        chain[2] = add->src[1];
-                        chain[3] = add;
-                        chain[4] = value;
-                        chain[5] = quant;
-                        count    = 6;
+                        chain[2] = value;
+                        chain[3] = quant;
+                        count    = 4;
+                    } else {
+                        chain[0] = value;
+                        chain[1] = quant;
+                        count    = 2;
                     }
                 } else if (value->op == GGML_OP_MUL) {
                     for (int src = 0; src < 2; ++src)
@@ -210,13 +216,13 @@ namespace SD_CUDA_NAMESPACE {
             const auto x = reinterpret_cast<uintptr_t>(a->data), y = reinterpret_cast<uintptr_t>(b->data);
             return x <= y ? y - x < ggml_nbytes(a) : x - y < ggml_nbytes(b);
         }
-        template <bool Swiglu = false, bool Normalize = false, bool Scaled = false, bool Sigmoid = false, bool RMS = false>
+        template <bool Swiglu = false, bool Normalize = false, bool Scaled = false, bool Sigmoid = false, bool RMS = false, bool Modulate = false>
         static CUDAError launch(Plan& p, const Operation& op, int8_t* output, CUDAStream stream) {
-            constexpr bool warp = Scaled || RMS;
+            constexpr bool warp = Scaled || RMS || Modulate;
             constexpr int extra = Normalize ? 384 : RMS ? 256
                                                         : 128;
             if (!p.configured) {
-                auto status = cuda_set_shared_memory(quantize<Swiglu, false, Normalize, Scaled, warp, float, Sigmoid, RMS>,
+                auto status = cuda_set_shared_memory(quantize<Swiglu, false, Normalize, Scaled, warp, float, Sigmoid, RMS, Modulate>,
                                                      int(p.device.shared_memory - extra));
                 if (status != cuda_success)
                     return status;
@@ -229,7 +235,7 @@ namespace SD_CUDA_NAMESPACE {
                                                                                         : 256;
             const int64_t padded = (op.rows + 3) / 4 * 4;
             auto data            = [](const ggml_tensor* t) { return t ? static_cast<const float*>(t->data) : nullptr; };
-            quantize<Swiglu, false, Normalize, Scaled, warp, float, Sigmoid, RMS><<<unsigned(op.rows), threads, size_t(op.k) * sizeof(float), stream>>>(
+            quantize<Swiglu, false, Normalize, Scaled, warp, float, Sigmoid, RMS, Modulate><<<unsigned(op.rows), threads, size_t(op.k) * sizeof(float), stream>>>(
                 data(op.input) + op.gate_offset, output, reinterpret_cast<float*>(output + op.k * padded), op.k, op.rows, op.k,
                 op.up ? data(op.up) + op.up_offset : nullptr, op.stride, op.up_stride, nullptr, nullptr, data(op.modulation), op.eps, op.scale, op.bias, data(op.weight), data(op.shift));
             return cuda_last_error();
@@ -260,6 +266,9 @@ namespace SD_CUDA_NAMESPACE {
                 case RMS_MOD:
                     error = launch<false, false, false, false, true>(plan, op, output, stream);
                     break;
+                case MOD:
+                    error = launch<false, false, false, false, false, true>(plan, op, output, stream);
+                    break;
                 case SWIGLU:
                 case SILU:
                     error = launch<true>(plan, op, output, stream);
@@ -277,12 +286,13 @@ namespace SD_CUDA_NAMESPACE {
     }
 
     void register_convrot(ggml_backend_dev_t device) {
-        for (auto op : {GGML_OP_SCALE, GGML_OP_NORM, GGML_OP_RMS_NORM, GGML_OP_GLU, GGML_OP_UNARY}) {
+        for (auto op : {GGML_OP_SCALE, GGML_OP_NORM, GGML_OP_RMS_NORM, GGML_OP_GLU, GGML_OP_UNARY, GGML_OP_CUSTOM}) {
             ggml_ext_kernel_desc desc = {};
             desc.abi_version          = GGML_EXT_ABI_VERSION;
             const std::string name    = std::string("sd.convrot_fusion.") + ggml_op_name(op) + "." SD_CUDA_NAME ".v1";
             desc.name                 = name.c_str();
             desc.op                   = op;
+            desc.custom_op            = op == GGML_OP_CUSTOM ? sd_ggml_modulate_op() : nullptr;
             desc.priority             = 225;
             desc.optimize             = op == GGML_OP_SCALE ? sd_convrot::optimize : nullptr;
             desc.match                = sd_convrot::match;

diff --git a/src/core/cuda/convrot_kernel.h b/src/core/cuda/convrot_kernel.h
@@ -40,7 +40,7 @@ namespace SD_CUDA_NAMESPACE {
             return __fadd_rn(__fadd_rn(__fadd_rn(a, b), c), d);
         }
 
-        template <bool swiglu, bool dequantize = false, bool normalize = false, bool scaled = false, bool warp_rotation = false, typename T = float, bool sigmoid_gate = false, bool rms_modulation = false>
+        template <bool swiglu, bool dequantize = false, bool normalize = false, bool scaled = false, bool warp_rotation = false, typename T = float, bool sigmoid_gate = false, bool rms_modulation = false, bool modulate = false>
         __global__ void quantize(
             const T* src,
             int8_t* dst,
@@ -103,7 +103,12 @@ namespace SD_CUDA_NAMESPACE {
                             float value = src_row[i];
                             if constexpr (rms_modulation) {
                                 value = __fmul_rn(__fmul_rn(value, rms_scale), rms_weight[i]);
-                                value = __fadd_rn(__fadd_rn(value, __fmul_rn(value, modulation[i])), rms_shift[i]);
+                            }
+                            if constexpr (rms_modulation || modulate) {
+                                value = __fadd_rn(value, __fmul_rn(value, modulation[i]));
+                                if (rms_shift) {
+                                    value = __fadd_rn(value, rms_shift[i]);
+                                }
                             }
                             if constexpr (scaled) {
                                 value = __fmaf_rn(value, input_scale, input_bias);