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)
 

external/ggml/include/ggml-cuda.h

@@ -1,52 +1,55 @@
-#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);
+// 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

@@ -1451,7 +1451,16 @@ struct 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));
+            // GGML_CUDA_STREAM_PRIORITY (read at stream creation, not cached):
+            // lets a host create backend instances whose streams differ in
+            // scheduling priority (CUDA: numerically lower = higher priority).
+            const char * priority_env = getenv("GGML_CUDA_STREAM_PRIORITY");
+            if (priority_env != nullptr && atoi(priority_env) != 0) {
+                CUDA_CHECK(cudaStreamCreateWithPriority(
+                    &streams[device][stream], cudaStreamNonBlocking, atoi(priority_env)));
+            } else {
+                CUDA_CHECK(cudaStreamCreateWithFlags(&streams[device][stream], cudaStreamNonBlocking));
+            }
         }
         return streams[device][stream];
     }
@@ -1521,12 +1530,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;
 };