minimax_music3: performance pack — native rope/SwiGLU lowering + opt-in CFG reuse, chunk hop, batched ensemble takes, q4_k depth

CMakeLists.txt

@@ -2363,6 +2363,35 @@ if (ENGINE_BUILD_TESTS)
         COMMAND gguf_tensor_source_test
     )
 
+    add_engine_unittest(minimax_music3_lm_head_test tests/unittests/test_minimax_music3_lm_head.cpp)
+    target_include_directories(minimax_music3_lm_head_test PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/tests/unittests)
+
+    add_test(
+        NAME minimax_music3_lm_head_test
+        COMMAND minimax_music3_lm_head_test
+    )
+
+    add_engine_unittest(
+        minimax_music3_pipeline_buffers_test
+        tests/unittests/test_minimax_music3_pipeline_buffers.cpp)
+    target_include_directories(
+        minimax_music3_pipeline_buffers_test
+        PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/tests/unittests)
+
+    add_test(
+        NAME minimax_music3_pipeline_buffers_test
+        COMMAND minimax_music3_pipeline_buffers_test
+    )
+
+    add_engine_unittest(
+        minimax_music3_graph_release_policy_test
+        tests/unittests/test_minimax_music3_graph_release_policy.cpp)
+
+    add_test(
+        NAME minimax_music3_graph_release_policy_test
+        COMMAND minimax_music3_graph_release_policy_test
+    )
+
     add_engine_unittest(model_spec_system_test tests/unittests/test_model_spec_system.cpp)
     target_include_directories(model_spec_system_test PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/tests/unittests)
 

docs/community_models/minimax_music3.md

@@ -110,6 +110,42 @@ Common session options:
 | `minimax_music3.weight_context_mb` | `32` | Weight context size in MiB. |
 | `minimax_music3.mem_saver` | `true` | Load large generation stages only while needed to reduce peak VRAM. |
 
+## Performance options
+
+Everything below is opt-in; the defaults reproduce the reference trajectory
+exactly. The two cheapest wins that need no extra GGUFs are
+`flow_uncond_interval=3` and `flow_chunk_hop_frames=150`.
+
+Request options:
+
+| Option | Default | Notes |
+|---|---:|---|
+| `flow_uncond_interval` | `1` | Evaluate the flow unconditional CFG branch only every N-th step and reuse the cached guidance delta in between. `2`-`3` measures flow -25..-40% at ~0.3 dB mel-L1. `1` keeps the exact reference trajectory. |
+| `flow_uncond_warmup` | `2` | Steps at the start of each chunk that always evaluate both CFG branches before delta reuse kicks in. |
+| `flow_chunk_hop_frames` | `0` | Flow chunk hop in AR frames (~25/sec). `0` keeps the model config (100, 50% chunk overlap); `150` measures flow ~-35% with consistently rederived crops and carry. |
+| `ensemble_takes` | `1` | Generate K takes that share one batched AR pass and return them as named outputs (`take_01`..). Amortizes the AR stage across takes. |
+| `ensemble_prefix_frames` | `0` | With ensembles: share a master AR prefix of this many frames across takes (intro-lock), then let takes diverge. |
+
+Session options:
+
+| Option | Default | Notes |
+|---|---:|---|
+| `minimax_music3.pipeline_overlap` | `false` | Overlap AR decoding with per-chunk denoise/vocode on a second, lower-priority CUDA stream. Requires `mem_saver=false` (all stages stay resident). |
+
+Optional component GGUFs: `rvq_depth_decoder_q4_k.gguf` (the depth stage is
+weight-bandwidth-bound; measures depth about -24% against `q8_0` with no
+audible difference in our listening) can be selected with
+`minimax_music3.rvq_depth_decoder_gguf` as shown above.
+
+Environment variables (supported):
+
+| Variable | Effect |
+|---|---|
+| `MM3_LEGACY_ROPE=1` | Use the original split-table rope path instead of the fused native NEOX rope. |
+| `MM3_LEGACY_GLU=1` | Use the original slice+silu+mul GLU instead of the fused SwiGLU kernel. |
+| `MM3_DEPTH_GPU_SAMPLE=1` | Device philox top-k sampler for the depth stage (byte-exact, perf-neutral option). |
+| `MM3_DEPTH_GPU_FRAME=1` | Full GPU depth frame with device residual ids and hidden accumulation (byte-exact, perf-neutral option). |
+
 ## Notes
 
 The default package is intended to make the model practical on common GPUs,

external/ggml/include/ggml-cuda.h

@@ -1,52 +1,59 @@
-#pragma once
-
-#include "ggml.h"
-#include "ggml-backend.h"
-
-#ifdef  __cplusplus
-extern "C" {
-#endif
-
-#ifdef GGML_USE_HIP
-#define GGML_CUDA_NAME "ROCm"
-#define GGML_CUBLAS_NAME "hipBLAS"
-#elif defined(GGML_USE_MUSA)
-#define GGML_CUDA_NAME "MUSA"
-#define GGML_CUBLAS_NAME "muBLAS"
-#else
-#define GGML_CUDA_NAME "CUDA"
-#define GGML_CUBLAS_NAME "cuBLAS"
-#endif
-#define GGML_CUDA_MAX_DEVICES       16
-
-// backend API
-GGML_BACKEND_API ggml_backend_t ggml_backend_cuda_init(int device);
-
-GGML_BACKEND_API bool ggml_backend_is_cuda(ggml_backend_t backend);
+#pragma once
+
+#include "ggml.h"
+#include "ggml-backend.h"
+
+#ifdef  __cplusplus
+extern "C" {
+#endif
+
+#ifdef GGML_USE_HIP
+#define GGML_CUDA_NAME "ROCm"
+#define GGML_CUBLAS_NAME "hipBLAS"
+#elif defined(GGML_USE_MUSA)
+#define GGML_CUDA_NAME "MUSA"
+#define GGML_CUBLAS_NAME "muBLAS"
+#else
+#define GGML_CUDA_NAME "CUDA"
+#define GGML_CUBLAS_NAME "cuBLAS"
+#endif
+#define GGML_CUDA_MAX_DEVICES       16
+
+// backend API
+GGML_BACKEND_API ggml_backend_t ggml_backend_cuda_init(int device);
+
+GGML_BACKEND_API bool ggml_backend_is_cuda(ggml_backend_t backend);
 GGML_BACKEND_API void ggml_backend_cuda_trim_pools(ggml_backend_t backend);
+// Sets the CUDA scheduling priority used for this backend instance's lazily
+// created streams (lower = higher priority, 0 = default). Scoped to the
+// instance; already-created streams keep their priority.
+GGML_BACKEND_API void ggml_backend_cuda_set_stream_priority(ggml_backend_t backend, int priority);
+// Returns the backend's current compute CUDA stream (cudaStream_t) so host
+// code can enqueue its own kernels/copies ordered with graph computes.
+GGML_BACKEND_API void * ggml_backend_cuda_get_stream(ggml_backend_t backend);
 GGML_BACKEND_API void ggml_backend_cuda_clear_graph(ggml_backend_t backend, const struct ggml_cgraph * graph);
-
-// device buffer
-GGML_BACKEND_API ggml_backend_buffer_type_t ggml_backend_cuda_buffer_type(int device);
-
-// conduct allreduce operation between devices
-GGML_BACKEND_API bool ggml_backend_cuda_allreduce_tensor(ggml_backend_t * backends, struct ggml_tensor ** tensors, size_t n_backends);
-
-// split tensor buffer that splits matrices by rows across multiple devices
-GGML_BACKEND_API ggml_backend_buffer_type_t ggml_backend_cuda_split_buffer_type(int main_device, const float * tensor_split);
-
-// pinned host buffer for use with the CPU backend for faster copies between CPU and GPU
-GGML_BACKEND_API ggml_backend_buffer_type_t ggml_backend_cuda_host_buffer_type(void);
-
-GGML_BACKEND_API int  ggml_backend_cuda_get_device_count(void);
-GGML_BACKEND_API void ggml_backend_cuda_get_device_description(int device, char * description, size_t description_size);
-GGML_BACKEND_API void ggml_backend_cuda_get_device_memory(int device, size_t * free, size_t * total);
-
-GGML_BACKEND_API bool ggml_backend_cuda_register_host_buffer(void * buffer, size_t size);
-GGML_BACKEND_API void ggml_backend_cuda_unregister_host_buffer(void * buffer);
-
-GGML_BACKEND_API ggml_backend_reg_t ggml_backend_cuda_reg(void);
-
-#ifdef  __cplusplus
-}
-#endif
+
+// device buffer
+GGML_BACKEND_API ggml_backend_buffer_type_t ggml_backend_cuda_buffer_type(int device);
+
+// conduct allreduce operation between devices
+GGML_BACKEND_API bool ggml_backend_cuda_allreduce_tensor(ggml_backend_t * backends, struct ggml_tensor ** tensors, size_t n_backends);
+
+// split tensor buffer that splits matrices by rows across multiple devices
+GGML_BACKEND_API ggml_backend_buffer_type_t ggml_backend_cuda_split_buffer_type(int main_device, const float * tensor_split);
+
+// pinned host buffer for use with the CPU backend for faster copies between CPU and GPU
+GGML_BACKEND_API ggml_backend_buffer_type_t ggml_backend_cuda_host_buffer_type(void);
+
+GGML_BACKEND_API int  ggml_backend_cuda_get_device_count(void);
+GGML_BACKEND_API void ggml_backend_cuda_get_device_description(int device, char * description, size_t description_size);
+GGML_BACKEND_API void ggml_backend_cuda_get_device_memory(int device, size_t * free, size_t * total);
+
+GGML_BACKEND_API bool ggml_backend_cuda_register_host_buffer(void * buffer, size_t size);
+GGML_BACKEND_API void ggml_backend_cuda_unregister_host_buffer(void * buffer);
+
+GGML_BACKEND_API ggml_backend_reg_t ggml_backend_cuda_reg(void);
+
+#ifdef  __cplusplus
+}
+#endif

external/ggml/src/ggml-cuda/common.cuh

@@ -1446,12 +1446,23 @@ struct ggml_backend_cuda_context {
 
     ggml_cuda_stream_context concurrent_stream_context;
 
+    // Scheduling priority for this context's lazily created streams
+    // (CUDA: numerically lower = higher priority, 0 = default). Set through
+    // ggml_backend_cuda_set_stream_priority(); scoped to this backend
+    // instance, unlike a process-wide environment variable.
+    int stream_priority = 0;
+
     ~ggml_backend_cuda_context();
 
     cudaStream_t stream(int device, int stream) {
         if (streams[device][stream] == nullptr) {
             ggml_cuda_set_device(device);
-            CUDA_CHECK(cudaStreamCreateWithFlags(&streams[device][stream], cudaStreamNonBlocking));
+            if (stream_priority != 0) {
+                CUDA_CHECK(cudaStreamCreateWithPriority(
+                    &streams[device][stream], cudaStreamNonBlocking, stream_priority));
+            } else {
+                CUDA_CHECK(cudaStreamCreateWithFlags(&streams[device][stream], cudaStreamNonBlocking));
+            }
         }
         return streams[device][stream];
     }
@@ -1521,12 +1532,12 @@ struct ggml_cuda_mm_fusion_args_host {
     const ggml_tensor * gate = nullptr;
     const ggml_tensor * gate_bias = nullptr;
     ggml_glu_op glu_op;
-    bool residual_only = false;
+    bool residual_only = false;
 };
 struct ggml_cuda_mm_fusion_args_device {
     const void * x_bias = nullptr;
     const void * gate = nullptr;
     const void * gate_bias = nullptr;
     ggml_glu_op glu_op;
-    bool residual_only = false;
+    bool residual_only = false;
 };