diff --git a/server/CMakeLists.txt b/server/CMakeLists.txt index 0da29ea33..cfa7d8fbf 100644 --- a/server/CMakeLists.txt +++ b/server/CMakeLists.txt @@ -1,12 +1,44 @@ cmake_minimum_required(VERSION 3.21) # 3.21 adds first-class HIP language support (project(LANGUAGES ... HIP)) set(DFLASH27B_GPU_BACKEND "cuda" CACHE STRING "GPU backend to build: cuda or hip") set_property(CACHE DFLASH27B_GPU_BACKEND PROPERTY STRINGS cuda hip) +option(DFLASH27B_ENABLE_MIXED_CUDA_HIP + "Primary GPU runtime plus an isolated in-process CUDA/HIP peer (Linux only)" + OFF) string(TOLOWER "${DFLASH27B_GPU_BACKEND}" DFLASH27B_GPU_BACKEND) +if(DFLASH27B_ENABLE_MIXED_CUDA_HIP) + # Reject unsupported targets before project() tries to discover both GPU + # toolchains. Toolchain files define CMAKE_SYSTEM_NAME for cross builds; + # otherwise the host system is the native target. + if(CMAKE_SYSTEM_NAME) + set(_dflash_mixed_target_system "${CMAKE_SYSTEM_NAME}") + else() + set(_dflash_mixed_target_system "${CMAKE_HOST_SYSTEM_NAME}") + endif() + if(NOT _dflash_mixed_target_system STREQUAL "Linux") + message(FATAL_ERROR + "The in-process CUDA+HIP runtime is currently supported on Linux only") + endif() + unset(_dflash_mixed_target_system) +endif() +# These are internal build-shape switches rather than user-selectable backend +# policy. Reset both on every configure so reusing a build directory after +# changing the primary backend cannot leave the old peer as a module. +set(GGML_CUDA_MODULE OFF CACHE BOOL "Build CUDA as a runtime-loadable module" FORCE) +set(GGML_HIP_MODULE OFF CACHE BOOL "Build HIP as a runtime-loadable module" FORCE) + if(DFLASH27B_GPU_BACKEND STREQUAL "cuda") set(DFLASH27B_USER_CUDA_ARCHITECTURES "${CMAKE_CUDA_ARCHITECTURES}") project(dflash LANGUAGES C CXX CUDA) elseif(DFLASH27B_GPU_BACKEND STREQUAL "hip") - project(dflash LANGUAGES C CXX HIP) + if(DFLASH27B_ENABLE_MIXED_CUDA_HIP) + set(DFLASH27B_CUDA_ARCHITECTURES "86" CACHE STRING + "Secondary CUDA GPU targets, e.g. 86 for RTX 3090") + set(CMAKE_CUDA_ARCHITECTURES "${DFLASH27B_CUDA_ARCHITECTURES}" + CACHE STRING "" FORCE) + project(dflash LANGUAGES C CXX HIP CUDA) + else() + project(dflash LANGUAGES C CXX HIP) + endif() else() message(FATAL_ERROR "DFLASH27B_GPU_BACKEND must be 'cuda' or 'hip', got '${DFLASH27B_GPU_BACKEND}'") endif() @@ -42,8 +74,9 @@ endif() # If we do not set this, ggml will output to bin and DLLs will not load set(CMAKE_RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}") -# ROCm root for HIP builds (rpath + rocwmma header discovery). -if(DFLASH27B_GPU_BACKEND STREQUAL "hip") +# ROCm root for HIP and mixed CUDA+HIP builds (rpath + header discovery). +if(DFLASH27B_GPU_BACKEND STREQUAL "hip" OR + DFLASH27B_ENABLE_MIXED_CUDA_HIP) if(DEFINED ROCM_PATH) set(_dflash_rocm_root "${ROCM_PATH}") elseif(DEFINED ENV{ROCM_PATH}) @@ -57,7 +90,8 @@ endif() # Bake portable rpath into all executables so bundled ggml backend libs / libggml-base # are found regardless of LD_LIBRARY_PATH or stale /usr/local/lib (closes #31). set(CMAKE_INSTALL_RPATH "$ORIGIN/deps/llama.cpp/ggml/src;$ORIGIN/deps/llama.cpp/ggml/src/ggml-cuda;$ORIGIN/deps/llama.cpp/ggml/src/ggml-hip;$ORIGIN/../deps/llama.cpp/ggml/src;$ORIGIN/../deps/llama.cpp/ggml/src/ggml-cuda;$ORIGIN/../deps/llama.cpp/ggml/src/ggml-hip") -if(DFLASH27B_GPU_BACKEND STREQUAL "hip" AND _dflash_rocm_root) +if((DFLASH27B_GPU_BACKEND STREQUAL "hip" OR + DFLASH27B_ENABLE_MIXED_CUDA_HIP) AND _dflash_rocm_root) list(APPEND CMAKE_BUILD_RPATH "${_dflash_rocm_root}/lib" "${_dflash_rocm_root}/lib64") set(CMAKE_INSTALL_RPATH "${CMAKE_INSTALL_RPATH};${_dflash_rocm_root}/lib;${_dflash_rocm_root}/lib64") endif() @@ -78,20 +112,52 @@ endif() if(DFLASH27B_GPU_BACKEND STREQUAL "cuda") set(GGML_CUDA ON CACHE BOOL "" FORCE) - set(GGML_HIP OFF CACHE BOOL "" FORCE) + if(DFLASH27B_ENABLE_MIXED_CUDA_HIP) + set(GGML_HIP ON CACHE BOOL "" FORCE) + set(GGML_HIP_RCCL OFF CACHE BOOL "" FORCE) + # Keep CUDA/CPU linked normally so every existing code path is + # unchanged. Build only HIP as an RTLD_LOCAL module: both GPU + # implementations use the historical ggml_backend_cuda_* symbol + # family and therefore cannot safely share the global symbol scope. + set(GGML_BACKEND_DL OFF CACHE BOOL "" FORCE) + set(GGML_HIP_MODULE ON CACHE BOOL + "Build HIP as an isolated runtime-loadable backend" FORCE) + # The peer module must link to a shared ggml core. Keep this as a + # scoped build requirement; do not overwrite the user's cached + # BUILD_SHARED_LIBS choice for later non-mixed reconfiguration. + set(DFLASH27B_MIXED_GGML_SHARED ON) + set(DFLASH27B_GGML_BACKEND_TARGET ggml-cuda) + set(DFLASH27B_HIP_ARCHITECTURES "" CACHE STRING + "Secondary HIP GPU targets, e.g. gfx1151") + else() + set(GGML_HIP OFF CACHE BOOL "" FORCE) + set(GGML_BACKEND_DL OFF CACHE BOOL "" FORCE) + set(DFLASH27B_GGML_BACKEND_TARGET ggml-cuda) + endif() set(GGML_CUDA_GRAPHS ON CACHE BOOL "Enable CUDA graphs for AR-decode replay (lucebox)" FORCE) - set(DFLASH27B_GGML_BACKEND_TARGET ggml-cuda) elseif(DFLASH27B_GPU_BACKEND STREQUAL "hip") - set(GGML_CUDA OFF CACHE BOOL "" FORCE) + if(DFLASH27B_ENABLE_MIXED_CUDA_HIP) + set(GGML_CUDA ON CACHE BOOL "" FORCE) + set(GGML_CUDA_MODULE ON CACHE BOOL + "Build CUDA as an isolated runtime-loadable backend" FORCE) + set(DFLASH27B_MIXED_GGML_SHARED ON) + else() + set(GGML_CUDA OFF CACHE BOOL "" FORCE) + set(GGML_CUDA_MODULE OFF CACHE BOOL "" FORCE) + endif() set(GGML_HIP ON CACHE BOOL "" FORCE) set(GGML_HIP_RCCL OFF CACHE BOOL "" FORCE) set(DFLASH27B_GGML_BACKEND_TARGET ggml-hip) set(DFLASH27B_HIP_ARCHITECTURES "" CACHE STRING "HIP GPU targets, e.g. gfx906;gfx1100") + set(GGML_BACKEND_DL OFF CACHE BOOL "" FORCE) if(DFLASH27B_HIP_ARCHITECTURES AND NOT CMAKE_HIP_ARCHITECTURES) set(CMAKE_HIP_ARCHITECTURES "${DFLASH27B_HIP_ARCHITECTURES}" CACHE STRING "" FORCE) endif() + if(DFLASH27B_ENABLE_MIXED_CUDA_HIP) + set(GGML_CUDA_GRAPHS ON CACHE BOOL + "Enable CUDA graphs in the secondary CUDA backend" FORCE) + endif() endif() -set(GGML_BACKEND_DL OFF CACHE BOOL "" FORCE) set(GGML_METAL OFF CACHE BOOL "" FORCE) set(GGML_VULKAN OFF CACHE BOOL "" FORCE) set(GGML_BLAS OFF CACHE BOOL "" FORCE) @@ -147,6 +213,16 @@ if(DFLASH27B_GPU_BACKEND STREQUAL "cuda") # which triggers massive first-request PTX JIT on newer GPUs even though # dflash_common itself is compiled for the intended arch set. set(CMAKE_CUDA_ARCHITECTURES "${_dflash_archs}" CACHE STRING "" FORCE) + if(DFLASH27B_ENABLE_MIXED_CUDA_HIP) + if(DFLASH27B_HIP_ARCHITECTURES) + set(_dflash_mixed_hip_archs "${DFLASH27B_HIP_ARCHITECTURES}") + elseif(AMDGPU_TARGETS) + set(_dflash_mixed_hip_archs "${AMDGPU_TARGETS}") + else() + set(_dflash_mixed_hip_archs "gfx1151") + endif() + set(CMAKE_HIP_ARCHITECTURES "${_dflash_mixed_hip_archs}" CACHE STRING "" FORCE) + endif() elseif(DFLASH27B_GPU_BACKEND STREQUAL "hip") # User override precedence: -DDFLASH27B_HIP_ARCHITECTURES → -DAMDGPU_TARGETS # → gfx1151 default (Strix Halo). @@ -200,10 +276,44 @@ if(WIN32 AND NOT CMAKE_ASM_COMPILER) set(CMAKE_ASM_COMPILER "${CMAKE_CXX_COMPILER}" CACHE FILEPATH "" FORCE) endif() -# Use only the ggml subtree of llama.cpp (skip libllama). +# Resolve GPU runtime packages before creating the ggml targets so their build +# rpaths include non-system toolkit installations. This matters when a peer +# module is loaded with dlopen and cannot inherit link-time search paths from +# the primary executable. +if(DFLASH27B_GPU_BACKEND STREQUAL "cuda" OR + DFLASH27B_ENABLE_MIXED_CUDA_HIP) + find_package(CUDAToolkit REQUIRED) + if(UNIX AND CUDAToolkit_LIBRARY_DIR) + list(APPEND CMAKE_BUILD_RPATH "${CUDAToolkit_LIBRARY_DIR}") + list(APPEND CMAKE_INSTALL_RPATH "${CUDAToolkit_LIBRARY_DIR}") + endif() +endif() +# Use only the ggml subtree of llama.cpp (skip libllama). Mixed builds need a +# shared ggml core for the isolated peer module, but that requirement belongs +# only to this sub-build. Restoring the previous value prevents a reused build +# directory from silently changing ordinary builds after mixed mode is turned +# off. +if(DFLASH27B_MIXED_GGML_SHARED) + set(_dflash_build_shared_libs_was_defined OFF) + if(DEFINED BUILD_SHARED_LIBS) + set(_dflash_build_shared_libs_was_defined ON) + set(_dflash_saved_build_shared_libs "${BUILD_SHARED_LIBS}") + endif() + set(BUILD_SHARED_LIBS ON) +endif() add_subdirectory(deps/llama.cpp/ggml EXCLUDE_FROM_ALL) +if(DFLASH27B_MIXED_GGML_SHARED) + if(_dflash_build_shared_libs_was_defined) + set(BUILD_SHARED_LIBS "${_dflash_saved_build_shared_libs}") + else() + unset(BUILD_SHARED_LIBS) + endif() + unset(_dflash_saved_build_shared_libs) + unset(_dflash_build_shared_libs_was_defined) +endif() -if(DFLASH27B_GPU_BACKEND STREQUAL "hip") +if(DFLASH27B_GPU_BACKEND STREQUAL "hip" OR + DFLASH27B_ENABLE_MIXED_CUDA_HIP) # The vendored ggml HIP shim still uses a few CUDA spellings that are not # mapped in this upstream snapshot. Keep the compatibility layer in this # repo so the build stays reproducible from a clean checkout. @@ -217,11 +327,22 @@ if(DFLASH27B_GPU_BACKEND STREQUAL "hip") ${CMAKE_CURRENT_SOURCE_DIR}/src/hip_compat) endif() -if(DFLASH27B_GPU_BACKEND STREQUAL "cuda") - # The CUDA-only sources include directly, so the toolkit - # headers must be available when compiling the library. - find_package(CUDAToolkit REQUIRED) -elseif(DFLASH27B_GPU_BACKEND STREQUAL "hip") +if(DFLASH27B_ENABLE_MIXED_CUDA_HIP) + # CUDA and HIP are compiled from the same ggml-cuda implementation and + # therefore share C++ typeinfo, vtable, template, and data-symbol names. + # Bind every definition in the secondary module locally: limiting this to + # functions still permits the primary runtime's pool/vtable state to + # interpose and hand an allocation to the wrong vendor kernel. + if(DFLASH27B_GPU_BACKEND STREQUAL "cuda") + target_link_options(ggml-hip PRIVATE "LINKER:-Bsymbolic") + else() + target_link_options(ggml-cuda PRIVATE "LINKER:-Bsymbolic") + endif() +endif() + +if(DFLASH27B_GPU_BACKEND STREQUAL "hip" OR + DFLASH27B_ENABLE_MIXED_CUDA_HIP) + # The ggml HIP subdirectory establishes ROCm's package search roots. find_package(hip REQUIRED) endif() @@ -281,6 +402,7 @@ add_library(dflash_common STATIC src/laguna/laguna_layer_split_adapter.cpp src/laguna/laguna_dflash_target.cpp src/common/backend_ipc.cpp + src/common/dynamic_backend.cpp src/common/domino_head.cpp src/common/dspark_head.cpp src/common/target_shard_ipc.cpp @@ -386,11 +508,19 @@ if(DFLASH27B_GPU_BACKEND STREQUAL "cuda") DFLASH27B_BACKEND_CUDA=1 DFLASH27B_CUDA_MIN_SM=${_dflash_cuda_min_sm} DFLASH27B_MIN_SM=${_dflash_cuda_min_sm}) + if(DFLASH27B_ENABLE_MIXED_CUDA_HIP) + target_compile_definitions(dflash_common PRIVATE + DFLASH27B_BACKEND_MIXED=1) + endif() elseif(DFLASH27B_GPU_BACKEND STREQUAL "hip") target_sources(dflash_common PRIVATE src/deepseek4/deepseek4_hc_cuda.cu) set_source_files_properties(src/deepseek4/deepseek4_hc_cuda.cu PROPERTIES LANGUAGE HIP) set_target_properties(dflash_common PROPERTIES HIP_ARCHITECTURES "${_dflash_archs}") target_compile_definitions(dflash_common PRIVATE DFLASH27B_BACKEND_HIP=1 GGML_USE_HIP) + if(DFLASH27B_ENABLE_MIXED_CUDA_HIP) + target_compile_definitions(dflash_common PRIVATE + DFLASH27B_BACKEND_MIXED=1) + endif() # hip_compat shim is needed by ALL dflash_common sources (peer_access.cpp, # dflash_feature_ring.cpp, flashprefill.cpp), not just the SM80_EQUIV path. target_include_directories(dflash_common PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/hip_compat) @@ -694,6 +824,14 @@ if(DFLASH27B_TESTS) add_test(NAME cuda_comm_api COMMAND test_cuda_comm_api) endif() + if(DFLASH27B_ENABLE_MIXED_CUDA_HIP) + add_executable(test_mixed_cuda_hip test/test_mixed_cuda_hip.cpp) + target_include_directories(test_mixed_cuda_hip PRIVATE + ${DFLASH27B_SRC_INCLUDE_DIRS}) + target_link_libraries(test_mixed_cuda_hip PRIVATE dflash_common) + add_test(NAME mixed_cuda_hip COMMAND test_mixed_cuda_hip) + endif() + add_executable(test_rocmfp4 deps/llama.cpp/ggml/rocmfp4/test_rocmfp4.c) target_link_libraries(test_rocmfp4 PRIVATE ggml-base) if(UNIX) @@ -1307,10 +1445,18 @@ if(DFLASH27B_SERVER) target_include_directories(dflash_server PRIVATE ${DFLASH27B_SRC_INCLUDE_DIRS}) if(DFLASH27B_GPU_BACKEND STREQUAL "hip") target_compile_definitions(dflash_server PRIVATE DFLASH27B_BACKEND_HIP=1 GGML_USE_HIP) + if(DFLASH27B_ENABLE_MIXED_CUDA_HIP) + target_compile_definitions(dflash_server PRIVATE + DFLASH27B_BACKEND_MIXED=1) + endif() else() target_compile_definitions(dflash_server PRIVATE DFLASH27B_BACKEND_CUDA=1 DFLASH27B_CUDA_MIN_SM=${_dflash_cuda_min_sm}) + if(DFLASH27B_ENABLE_MIXED_CUDA_HIP) + target_compile_definitions(dflash_server PRIVATE + DFLASH27B_BACKEND_MIXED=1) + endif() endif() if(NOT WIN32) target_link_libraries(dflash_server PRIVATE dflash_common ggml ${DFLASH27B_GGML_BACKEND_TARGET} pthread) @@ -1351,10 +1497,18 @@ if(DFLASH27B_SERVER) target_include_directories(backend_ipc_daemon PRIVATE ${DFLASH27B_SRC_INCLUDE_DIRS}) if(DFLASH27B_GPU_BACKEND STREQUAL "hip") target_compile_definitions(backend_ipc_daemon PRIVATE DFLASH27B_BACKEND_HIP=1 GGML_USE_HIP) + if(DFLASH27B_ENABLE_MIXED_CUDA_HIP) + target_compile_definitions(backend_ipc_daemon PRIVATE + DFLASH27B_BACKEND_MIXED=1) + endif() else() target_compile_definitions(backend_ipc_daemon PRIVATE DFLASH27B_BACKEND_CUDA=1 DFLASH27B_CUDA_MIN_SM=${_dflash_cuda_min_sm}) + if(DFLASH27B_ENABLE_MIXED_CUDA_HIP) + target_compile_definitions(backend_ipc_daemon PRIVATE + DFLASH27B_BACKEND_MIXED=1) + endif() endif() if(NOT WIN32) target_link_libraries(backend_ipc_daemon PRIVATE dflash_common ggml ${DFLASH27B_GGML_BACKEND_TARGET} pthread) diff --git a/server/deps/llama.cpp/ggml/include/ggml-backend.h b/server/deps/llama.cpp/ggml/include/ggml-backend.h index 834a5ee7c..1a943481e 100644 --- a/server/deps/llama.cpp/ggml/include/ggml-backend.h +++ b/server/deps/llama.cpp/ggml/include/ggml-backend.h @@ -377,8 +377,9 @@ extern "C" { ggml_backend_sched_t sched, bool enabled); - // Inputs in one split share a destination backend and copy generation, so - // they may reuse one generation wait before their ordered copies. + // Inputs in one split share a destination and copy generation. Batch + // unlike-runtime fallbacks through per-backend host arenas, synchronize + // each producer once, then enqueue the consumer transfers together. GGML_API void ggml_backend_sched_set_batch_split_copies( ggml_backend_sched_t sched, bool enabled); diff --git a/server/deps/llama.cpp/ggml/src/CMakeLists.txt b/server/deps/llama.cpp/ggml/src/CMakeLists.txt index 315eafe2e..cfd2d7c09 100644 --- a/server/deps/llama.cpp/ggml/src/CMakeLists.txt +++ b/server/deps/llama.cpp/ggml/src/CMakeLists.txt @@ -250,7 +250,9 @@ if (CMAKE_SYSTEM_NAME MATCHES "Linux") endif() function(ggml_add_backend_library backend) - if (GGML_BACKEND_DL) + if (GGML_BACKEND_DL OR + (GGML_HIP_MODULE AND backend STREQUAL "ggml-hip") OR + (GGML_CUDA_MODULE AND backend STREQUAL "ggml-cuda")) add_library(${backend} MODULE ${ARGN}) # write the shared library to the output directory set_target_properties(${backend} PROPERTIES LIBRARY_OUTPUT_DIRECTORY ${CMAKE_RUNTIME_OUTPUT_DIRECTORY}) @@ -302,7 +304,9 @@ function(ggml_add_backend backend) string(TOLOWER "ggml-${backend}" backend_target) add_subdirectory(${backend_target}) message(STATUS "Including ${backend} backend") - if (NOT GGML_BACKEND_DL) + if (NOT GGML_BACKEND_DL AND + NOT (GGML_HIP_MODULE AND backend STREQUAL "HIP") AND + NOT (GGML_CUDA_MODULE AND backend STREQUAL "CUDA")) string(TOUPPER "GGML_USE_${backend}" backend_use) target_compile_definitions(ggml PUBLIC ${backend_use}) endif() diff --git a/server/deps/llama.cpp/ggml/src/ggml-backend.cpp b/server/deps/llama.cpp/ggml/src/ggml-backend.cpp index c7ba5e473..f0aeaa73d 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-backend.cpp +++ b/server/deps/llama.cpp/ggml/src/ggml-backend.cpp @@ -20,6 +20,7 @@ #include #include #include +#include #include #ifdef __APPLE__ @@ -819,6 +820,11 @@ struct ggml_backend_sched_deferred_peer_copy { struct ggml_backend_sched { bool is_reset; // true if the scheduler has been reset since the last graph split bool is_alloc; + // True only after every backend stream has been synchronized and before + // another graph is submitted. In that state scheduler copy destinations + // are already reusable, so queueing the same generation event before + // every split is redundant (and particularly expensive across HIP/CUDA). + bool backends_synchronized; int n_backends; @@ -876,6 +882,27 @@ struct ggml_backend_sched { int deferred_peer_copies_capacity; bool split_deferred_peer_copies; bool batch_split_copies; + ggml_backend_buffer_t batch_staging_buffers[GGML_SCHED_MAX_BACKENDS]; + uint8_t * batch_staging_bases[GGML_SCHED_MAX_BACKENDS]; + size_t batch_staging_size; + bool profile_requested; + int profile_min_splits; + + struct { + bool active; + uint64_t loop_us; + uint64_t destination_wait_us; + uint64_t d2h_submit_us; + uint64_t host_relay_us; + uint64_t h2d_submit_us; + uint64_t compute_submit_us; + uint64_t staged_bytes; + int staged_copies; + uint64_t source_wait_us[GGML_SCHED_MAX_BACKENDS]; + uint64_t final_wait_us[GGML_SCHED_MAX_BACKENDS]; + int source_waits[GGML_SCHED_MAX_BACKENDS]; + int split_counts[GGML_SCHED_MAX_BACKENDS]; + } profile; int debug; @@ -1748,13 +1775,333 @@ static bool ggml_backend_sched_alloc_splits(ggml_backend_sched_t sched) { return true; } +static bool ggml_backend_sched_prepare_batch_staging( + ggml_backend_sched_t sched) { + if (!sched->batch_split_copies) { + return true; + } + + constexpr size_t staging_alignment = 64; + size_t required = 0; + for (int split_id = 0; split_id < sched->n_splits; ++split_id) { + const struct ggml_backend_sched_split & split = + sched->splits[split_id]; + for (int input_id = 0; input_id < split.n_inputs; ++input_id) { + const size_t bytes = ggml_nbytes(split.inputs[input_id]); + if (required > SIZE_MAX - (staging_alignment - 1)) { + return false; + } + required = (required + staging_alignment - 1) & + ~(staging_alignment - 1); + if (bytes > SIZE_MAX - required) { + return false; + } + required += bytes; + } + } + bool arenas_ready = required > 0; + for (int i = 0; arenas_ready && i < sched->n_backends; ++i) { + arenas_ready = sched->batch_staging_bases[i] != nullptr; + } + if (required <= sched->batch_staging_size && arenas_ready) { + return true; + } + if (required == 0) { + return true; + } + + ggml_backend_buffer_t new_buffers[GGML_SCHED_MAX_BACKENDS] = {}; + uint8_t * new_bases[GGML_SCHED_MAX_BACKENDS] = {}; + for (int i = 0; i < sched->n_backends; ++i) { + ggml_backend_dev_t device = ggml_backend_get_device(sched->backends[i]); + ggml_backend_buffer_type_t host_buft = device + ? ggml_backend_dev_host_buffer_type(device) : nullptr; + if (host_buft) { + new_buffers[i] = ggml_backend_buft_alloc_buffer(host_buft, required); + if (new_buffers[i]) { + new_bases[i] = static_cast( + ggml_backend_buffer_get_base(new_buffers[i])); + } + } + if (!new_bases[i]) { + if (new_buffers[i]) { + ggml_backend_buffer_free(new_buffers[i]); + new_buffers[i] = nullptr; + } + new_bases[i] = static_cast(malloc(required)); + } + if (!new_bases[i]) { + for (int j = 0; j <= i; ++j) { + if (new_buffers[j]) { + ggml_backend_buffer_free(new_buffers[j]); + } else { + free(new_bases[j]); + } + } + return false; + } + } + + // A graph may be reallocated after an earlier execution. Quiesce every + // runtime before releasing host pages that an asynchronous H2D transfer + // could still reference. + for (int i = 0; i < sched->n_backends; ++i) { + ggml_backend_synchronize(sched->backends[i]); + } + sched->backends_synchronized = true; + for (int i = 0; i < sched->n_backends; ++i) { + if (sched->batch_staging_buffers[i]) { + ggml_backend_buffer_free(sched->batch_staging_buffers[i]); + } else { + free(sched->batch_staging_bases[i]); + } + sched->batch_staging_buffers[i] = new_buffers[i]; + sched->batch_staging_bases[i] = new_bases[i]; + } + sched->batch_staging_size = required; + + for (int i = 0; i < sched->n_backends; ++i) { + GGML_LOG_DEBUG( + "%s: backend=%s size=%zu memory=%s\n", __func__, + ggml_backend_name(sched->backends[i]), required, + sched->batch_staging_buffers[i] ? "pinned" : "pageable"); + } + return true; +} + +static void ggml_backend_sched_free_batch_staging( + ggml_backend_sched_t sched) { + if (!sched) { + return; + } + for (int i = 0; i < sched->n_backends; ++i) { + if (sched->batch_staging_buffers[i]) { + ggml_backend_buffer_free(sched->batch_staging_buffers[i]); + } else { + free(sched->batch_staging_bases[i]); + } + sched->batch_staging_buffers[i] = nullptr; + sched->batch_staging_bases[i] = nullptr; + } + sched->batch_staging_size = 0; +} + +using ggml_backend_sched_profile_clock = std::chrono::steady_clock; + +static uint64_t ggml_backend_sched_elapsed_us( + ggml_backend_sched_profile_clock::time_point start, + ggml_backend_sched_profile_clock::time_point end) { + return (uint64_t) std::chrono::duration_cast( + end - start).count(); +} + +struct ggml_backend_sched_staged_copy { + ggml_backend_t source_backend; + int source_backend_id; + struct ggml_tensor * destination; + size_t offset; + size_t bytes; +}; + +// Owns copy synchronization and host-staging state for one scheduler split. +// Keeping this policy separate from graph submission makes the execution loop +// readable and gives every early/fallback path one synchronization contract. +struct ggml_backend_sched_split_copy_state { + ggml_backend_sched_t sched; + int destination_backend_id; + ggml_backend_t destination_backend; + bool staging_arena_reusable; + bool destination_generation_ready; + size_t & staging_cursor; + bool destination_host_ready = false; + ggml_backend_t synchronized_sources[GGML_SCHED_MAX_BACKENDS] = {}; + int n_synchronized_sources = 0; + ggml_backend_sched_staged_copy staged[GGML_SCHED_MAX_SPLIT_INPUTS] = {}; + int n_staged = 0; + + bool profiling() const { + return sched->profile.active; + } + + void mark_destination_ready() { + destination_generation_ready = true; + destination_host_ready = true; + } + + void wait_for_destination_generation() { + if (sched->batch_split_copies && destination_generation_ready) { + return; + } + const auto start = profiling() + ? ggml_backend_sched_profile_clock::now() + : ggml_backend_sched_profile_clock::time_point{}; + if (sched->events[destination_backend_id][sched->cur_copy]) { + ggml_backend_event_wait( + destination_backend, + sched->events[destination_backend_id][sched->cur_copy]); + } else { + ggml_backend_synchronize(destination_backend); + destination_host_ready = true; + } + if (profiling()) { + sched->profile.destination_wait_us += + ggml_backend_sched_elapsed_us( + start, ggml_backend_sched_profile_clock::now()); + } + destination_generation_ready = true; + } + + void synchronize_source(ggml_backend_t source_backend) { + const auto end = synchronized_sources + n_synchronized_sources; + if (sched->batch_split_copies && + std::find(synchronized_sources, end, source_backend) != end) { + return; + } + const auto start = profiling() + ? ggml_backend_sched_profile_clock::now() + : ggml_backend_sched_profile_clock::time_point{}; + ggml_backend_synchronize(source_backend); + if (profiling()) { + for (int backend_id = 0; backend_id < sched->n_backends; + ++backend_id) { + if (sched->backends[backend_id] == source_backend) { + sched->profile.source_wait_us[backend_id] += + ggml_backend_sched_elapsed_us( + start, ggml_backend_sched_profile_clock::now()); + sched->profile.source_waits[backend_id]++; + break; + } + } + } + if (sched->batch_split_copies) { + GGML_ASSERT(n_synchronized_sources < GGML_SCHED_MAX_BACKENDS); + synchronized_sources[n_synchronized_sources++] = source_backend; + } + } + + bool can_stage(int source_backend_id) const { + return sched->batch_split_copies && + staging_arena_reusable && + source_backend_id >= 0 && + sched->batch_staging_bases[source_backend_id] && + sched->batch_staging_bases[destination_backend_id] && + sched->backends[source_backend_id]->iface.get_tensor_async && + destination_backend->iface.set_tensor_async; + } + + void stage(ggml_backend_t source_backend, + int source_backend_id, + const struct ggml_tensor * source, + struct ggml_tensor * destination) { + constexpr size_t staging_alignment = 64; + staging_cursor = (staging_cursor + staging_alignment - 1) & + ~(staging_alignment - 1); + const size_t bytes = ggml_nbytes(source); + GGML_ASSERT(bytes <= sched->batch_staging_size); + GGML_ASSERT(staging_cursor <= sched->batch_staging_size - bytes); + + const auto start = profiling() + ? ggml_backend_sched_profile_clock::now() + : ggml_backend_sched_profile_clock::time_point{}; + ggml_backend_tensor_get_async( + source_backend, source, + sched->batch_staging_bases[source_backend_id] + staging_cursor, + 0, bytes); + if (profiling()) { + sched->profile.d2h_submit_us += ggml_backend_sched_elapsed_us( + start, ggml_backend_sched_profile_clock::now()); + sched->profile.staged_bytes += bytes; + sched->profile.staged_copies++; + } + + GGML_ASSERT(n_staged < GGML_SCHED_MAX_SPLIT_INPUTS); + staged[n_staged++] = { + source_backend, source_backend_id, destination, + staging_cursor, bytes}; + staging_cursor += bytes; + } + + void prepare_blocking_copy() { + if (sched->events[destination_backend_id][sched->cur_copy]) { + ggml_backend_event_synchronize( + sched->events[destination_backend_id][sched->cur_copy]); + destination_host_ready = true; + } else if (!destination_host_ready) { + ggml_backend_synchronize(destination_backend); + destination_host_ready = true; + } + } + + void flush_staged() { + for (int copy_id = 0; copy_id < n_staged; ++copy_id) { + synchronize_source(staged[copy_id].source_backend); + } + for (int copy_id = 0; copy_id < n_staged; ++copy_id) { + const ggml_backend_sched_staged_copy & copy = staged[copy_id]; + auto start = profiling() + ? ggml_backend_sched_profile_clock::now() + : ggml_backend_sched_profile_clock::time_point{}; + memcpy( + sched->batch_staging_bases[destination_backend_id] + + copy.offset, + sched->batch_staging_bases[copy.source_backend_id] + + copy.offset, + copy.bytes); + if (profiling()) { + sched->profile.host_relay_us += + ggml_backend_sched_elapsed_us( + start, ggml_backend_sched_profile_clock::now()); + } + + start = profiling() + ? ggml_backend_sched_profile_clock::now() + : ggml_backend_sched_profile_clock::time_point{}; + ggml_backend_tensor_set_async( + destination_backend, copy.destination, + sched->batch_staging_bases[destination_backend_id] + + copy.offset, + 0, copy.bytes); + if (profiling()) { + sched->profile.h2d_submit_us += + ggml_backend_sched_elapsed_us( + start, ggml_backend_sched_profile_clock::now()); + } + } + } +}; + +static ggml_backend_sched_profile_clock::time_point +ggml_backend_sched_begin_profile(ggml_backend_sched_t sched) { + const bool enabled = sched->profile_requested && + sched->n_splits >= sched->profile_min_splits; + sched->profile = {}; + sched->profile.active = enabled; + if (!enabled) { + return {}; + } + for (int split_id = 0; split_id < sched->n_splits; ++split_id) { + const int backend_id = sched->splits[split_id].backend_id; + if (backend_id >= 0 && backend_id < sched->n_backends) { + sched->profile.split_counts[backend_id]++; + } + } + return ggml_backend_sched_profile_clock::now(); +} + static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t sched) { GGML_ASSERT(sched); struct ggml_backend_sched_split * splits = sched->splits; + const auto profile_loop_start = ggml_backend_sched_begin_profile(sched); + ggml_tensor * prev_ids_tensor = nullptr; std::vector ids; std::vector used_ids; + size_t batch_staging_cursor = 0; + const bool copy_destinations_ready = sched->backends_synchronized; + // From this point onward an early return must remain conservative: a + // backend may have accepted work even if a later split fails. + sched->backends_synchronized = false; for (int split_id = 0; split_id < sched->n_splits; split_id++) { struct ggml_backend_sched_split * split = &splits[split_id]; @@ -1771,26 +2118,10 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s } } - // All copied inputs for this split use the same destination backend - // and scheduler copy generation. Waiting for that generation once is - // sufficient before overwriting any of its input buffers. The legacy - // loop waited again before every tensor; without scheduler events each - // wait synchronizes the whole destination stream and serializes a - // multi-input peer handoff. - bool split_copy_generation_ready = false; - auto wait_for_split_copy_generation = [&]() { - if (sched->batch_split_copies && split_copy_generation_ready) { - return; - } - if (sched->events[split_backend_id][sched->cur_copy] != NULL) { - ggml_backend_event_wait( - split_backend, - sched->events[split_backend_id][sched->cur_copy]); - } else { - ggml_backend_synchronize(split_backend); - } - split_copy_generation_ready = true; - }; + ggml_backend_sched_split_copy_state copy_state{ + sched, split_backend_id, split_backend, + copy_destinations_ready, copy_destinations_ready, + batch_staging_cursor}; // copy the input tensors to the split backend for (int input_id = 0; input_id < split->n_inputs; input_id++) { @@ -1806,11 +2137,11 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s } else { ggml_backend_synchronize(split_backend); } - split_copy_generation_ready = true; + copy_state.mark_destination_ready(); ggml_backend_tensor_copy(input, input_cpy); } else { // wait for the split backend to finish using the input before overwriting it - wait_for_split_copy_generation(); + copy_state.wait_for_destination_generation(); // when offloading MoE weights, we can reduce the amount of data copied by copying only the experts that are used ggml_tensor * node = split->graph.nodes[0]; @@ -1901,20 +2232,51 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s // try async copy, but if not possible, we can still use a sync copy without synchronizing the dst backend, since we handle the synchronization here with multiple copies and events // TODO: add public function to facilitate this, since applications do not have direct access to the backend interface if (!split_backend->iface.cpy_tensor_async || !split_backend->iface.cpy_tensor_async(input_backend, split_backend, input, input_cpy)) { - ggml_backend_synchronize(input_backend); - if (sched->events[split_backend_id][sched->cur_copy] != NULL) { - ggml_backend_event_synchronize(sched->events[split_backend_id][sched->cur_copy]); - } else { - ggml_backend_synchronize(split_backend); + const int input_backend_id = + ggml_backend_sched_backend_id(sched, input_backend); + if (copy_state.can_stage(input_backend_id)) { + copy_state.stage( + input_backend, input_backend_id, + input, input_cpy); + continue; } + + // All non-input copies are gathered before this split + // launches. Once a source backend is synchronized, its + // other inputs for the same split are ready as well. + // Cross-runtime MoE joins commonly carry activation, + // route-ID, and route-weight tensors together; waiting + // on the same source for each tensor serialized dozens + // of redundant stream round trips per verifier step. + copy_state.synchronize_source(input_backend); + + // With one scheduler copy there is no event object and + // wait_for_destination_generation() already established + // host-visible destination quiescence. Parallel-copy + // schedulers retain the explicit event synchronization + // before a blocking host-staged overwrite. + copy_state.prepare_blocking_copy(); ggml_backend_tensor_copy(input, input_cpy); } } } } + // Queue D2H for every source, wait once per source, then relay and + // enqueue all H2D transfers immediately before the consumer graph. + copy_state.flush_staged(); + if (!sched->callback_eval) { + const auto submit_start = sched->profile.active + ? ggml_backend_sched_profile_clock::now() + : ggml_backend_sched_profile_clock::time_point{}; enum ggml_status ec = ggml_backend_graph_compute_async(split_backend, &split->graph); + if (sched->profile.active) { + sched->profile.compute_submit_us += + ggml_backend_sched_elapsed_us( + submit_start, + ggml_backend_sched_profile_clock::now()); + } if (ec != GGML_STATUS_SUCCESS) { return ec; } @@ -1952,7 +2314,7 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s } } - // Publish cold-owner completion before a later main-backend graph + // Publish producer completion before a later consumer-backend graph // reaches its in-graph event wait. Recording is asynchronous and does // not block the host from immediately enqueueing independent work. for (int i = 0; i < sched->n_deferred_peer_copies; ++i) { @@ -1972,6 +2334,12 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s } } + if (sched->profile.active) { + sched->profile.loop_us = ggml_backend_sched_elapsed_us( + profile_loop_start, + ggml_backend_sched_profile_clock::now()); + } + return GGML_STATUS_SUCCESS; } @@ -1998,6 +2366,13 @@ ggml_backend_sched_t ggml_backend_sched_new( const char * GGML_SCHED_DEBUG_REALLOC = getenv("GGML_SCHED_DEBUG_REALLOC"); sched->debug_realloc = GGML_SCHED_DEBUG_REALLOC ? atoi(GGML_SCHED_DEBUG_REALLOC) : sched->debug_realloc; + const char * profile_raw = getenv("GGML_SCHED_PROFILE"); + const char * profile_min_raw = getenv("GGML_SCHED_PROFILE_MIN_SPLITS"); + sched->profile_requested = + profile_raw && profile_raw[0] && strcmp(profile_raw, "0") != 0; + sched->profile_min_splits = profile_min_raw + ? std::max(1, atoi(profile_min_raw)) : 1; + sched->n_backends = n_backends; sched->n_copies = parallel ? GGML_SCHED_MAX_COPIES : 1; @@ -2048,6 +2423,13 @@ void ggml_backend_sched_free(ggml_backend_sched_t sched) { if (sched == NULL) { return; } + // graph_compute_async callers may tear down without an explicit wait. + // Quiesce every backend before releasing graph allocations, native events, + // or pinned staging pages still referenced by queued transfers. + for (int b = 0; b < sched->n_backends; ++b) { + ggml_backend_synchronize(sched->backends[b]); + } + ggml_backend_sched_free_batch_staging(sched); for (int b = 0; b < sched->n_backends; b++) { for (int c = 0; c < sched->n_copies; c++) { ggml_backend_event_free(sched->events[b][c]); @@ -2080,12 +2462,14 @@ void ggml_backend_sched_reset(ggml_backend_sched_t sched) { // reset explicitly discards that graph, so carrying either list into the // next allocation would dereference stale metadata (and leak the events). sched->n_late_cross_input_split_nodes = 0; - if (sched->n_deferred_peer_copies > 0) { - // A caller may reset after an asynchronous submission. Do not destroy - // producer events while either backend can still reference them. + if (!sched->backends_synchronized) { + // A caller may reset after an asynchronous submission. Graph + // allocations, deferred events, and host-staging pages all remain + // reachable by queued work until every scheduler backend is idle. for (int backend_id = 0; backend_id < sched->n_backends; ++backend_id) { ggml_backend_synchronize(sched->backends[backend_id]); } + sched->backends_synchronized = true; } for (int i = 0; i < sched->n_deferred_peer_copies; ++i) { ggml_backend_event_free(sched->deferred_peer_copies[i].event); @@ -2143,6 +2527,10 @@ bool ggml_backend_sched_alloc_graph(ggml_backend_sched_t sched, struct ggml_cgra ggml_backend_sched_split_graph(sched, graph); + if (!ggml_backend_sched_prepare_batch_staging(sched)) { + return false; + } + if (!ggml_backend_sched_alloc_splits(sched)) { return false; } @@ -2170,12 +2558,6 @@ bool ggml_backend_sched_alloc_graph(ggml_backend_sched_t sched, struct ggml_cgra return true; } -enum ggml_status ggml_backend_sched_graph_compute(ggml_backend_sched_t sched, struct ggml_cgraph * graph) { - enum ggml_status err = ggml_backend_sched_graph_compute_async(sched, graph); - ggml_backend_sched_synchronize(sched); - return err; -} - enum ggml_status ggml_backend_sched_graph_compute_async(ggml_backend_sched_t sched, struct ggml_cgraph * graph) { GGML_ASSERT(sched); if (!sched->is_reset && !sched->is_alloc) { @@ -2191,10 +2573,52 @@ enum ggml_status ggml_backend_sched_graph_compute_async(ggml_backend_sched_t sch return ggml_backend_sched_compute_splits(sched); } +enum ggml_status ggml_backend_sched_graph_compute(ggml_backend_sched_t sched, struct ggml_cgraph * graph) { + enum ggml_status err = ggml_backend_sched_graph_compute_async(sched, graph); + ggml_backend_sched_synchronize(sched); + return err; +} + void ggml_backend_sched_synchronize(ggml_backend_sched_t sched) { GGML_ASSERT(sched); + using profile_clock = std::chrono::steady_clock; for (int i = 0; i < sched->n_backends; i++) { + const profile_clock::time_point wait_start = sched->profile.active + ? profile_clock::now() : profile_clock::time_point{}; ggml_backend_synchronize(sched->backends[i]); + if (sched->profile.active) { + sched->profile.final_wait_us[i] = + (uint64_t) std::chrono::duration_cast( + profile_clock::now() - wait_start).count(); + } + } + sched->backends_synchronized = true; + if (sched->profile.active) { + GGML_LOG_INFO( + "[ggml-sched-profile] splits=%d loop=%lluus dst_wait=%lluus " + "d2h_submit=%lluus host_relay=%lluus h2d_submit=%lluus " + "compute_submit=%lluus " + "staged_copies=%d staged_bytes=%llu\n", + sched->n_splits, + (unsigned long long) sched->profile.loop_us, + (unsigned long long) sched->profile.destination_wait_us, + (unsigned long long) sched->profile.d2h_submit_us, + (unsigned long long) sched->profile.host_relay_us, + (unsigned long long) sched->profile.h2d_submit_us, + (unsigned long long) sched->profile.compute_submit_us, + sched->profile.staged_copies, + (unsigned long long) sched->profile.staged_bytes); + for (int i = 0; i < sched->n_backends; ++i) { + GGML_LOG_INFO( + "[ggml-sched-profile] backend=%s splits=%d " + "source_waits=%d source_wait=%lluus final_wait=%lluus\n", + ggml_backend_name(sched->backends[i]), + sched->profile.split_counts[i], + sched->profile.source_waits[i], + (unsigned long long) sched->profile.source_wait_us[i], + (unsigned long long) sched->profile.final_wait_us[i]); + } + sched->profile.active = false; } if (!sched->is_alloc) { // if the graph is not already allocated, always use copy 0 after a synchronization diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu index 18a49f735..9b466fc3a 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu @@ -3417,7 +3417,12 @@ static thread_local ggml_cuda_pending_peer_copy_batch static bool ggml_cuda_batch_peer_copies_enabled() { static const bool enabled = [] { - const char * value = getenv("GGML_CUDA_BATCH_PEER_COPIES"); + const char * value = getenv("GGML_BATCH_PEER_COPIES"); + if (!value || !*value) { + // Compatibility with profiles created before the policy became + // runtime-neutral. + value = getenv("GGML_CUDA_BATCH_PEER_COPIES"); + } return value && *value && strcmp(value, "0") != 0; }(); return enabled; @@ -5178,7 +5183,14 @@ static const ggml_backend_i ggml_backend_cuda_interface = { }; static ggml_guid_t ggml_backend_cuda_guid() { +#if defined(GGML_USE_HIP) + // CUDA and HIP may coexist in one process as separate backend modules. + // A distinct GUID prevents either implementation from casting the other + // vendor's backend context and attempting an invalid peer copy. + static ggml_guid guid = { 0x87, 0x4b, 0x09, 0xd6, 0x72, 0x3f, 0x4e, 0x91, 0xa4, 0x66, 0x2d, 0x6e, 0x0a, 0x7f, 0x31, 0xbc }; +#else static ggml_guid guid = { 0x2c, 0xdd, 0xe8, 0x1c, 0x65, 0xb3, 0x65, 0x73, 0x6a, 0x12, 0x88, 0x61, 0x1c, 0xc9, 0xdc, 0x25 }; +#endif return &guid; } diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cu index 8f1f06a77..7bebd1b3a 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cu +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cu @@ -173,13 +173,32 @@ static void ggml_cuda_mul_mat_q_impl( const bool use_native_mxfp4 = blackwell_mma_available(cc) && src0->type == GGML_TYPE_MXFP4; const bool grouped_src = !ids && ggml_mul_mat_is_grouped_src(dst); const ggml_tensor * grouped_physical = grouped_src ? src1->view_src : nullptr; + int64_t grouped_width = 0; + int64_t grouped_row_stride = 0; + int64_t grouped_plane_stride = 0; if (grouped_src) { - GGML_ASSERT(grouped_physical && grouped_physical->type == GGML_TYPE_F32); - GGML_ASSERT(grouped_physical->ne[2] == - ggml_mul_mat_grouped_src_groups(dst)); - GGML_ASSERT(grouped_physical->ne[0] * grouped_physical->ne[2] == ne10); - GGML_ASSERT(grouped_physical->ne[1] == ne11); - GGML_ASSERT(grouped_physical->ne[3] == 1); + const int64_t groups = ggml_mul_mat_grouped_src_groups(dst); + GGML_ASSERT(groups > 1 && ne10 % groups == 0); + grouped_width = ne10 / groups; + + if (grouped_physical) { + GGML_ASSERT(grouped_physical->type == GGML_TYPE_F32); + GGML_ASSERT(grouped_physical->ne[0] == grouped_width); + GGML_ASSERT(grouped_physical->ne[1] == ne11); + GGML_ASSERT(grouped_physical->ne[2] == groups); + GGML_ASSERT(grouped_physical->ne[3] == 1); + grouped_row_stride = grouped_physical->nb[1] / ts_src1; + grouped_plane_stride = grouped_physical->nb[2] / ts_src1; + } else { + // A scheduler copy between unlike backends materializes the raw + // bytes of the logical 2-D view as a standalone tensor. The bytes + // retain the grouped physical ordering, while view_src metadata + // cannot cross the allocation boundary. ggml_mul_mat_grouped_src + // guarantees a contiguous [width, rows, groups] physical source, + // so reconstruct those two strides from the op's group count. + grouped_row_stride = grouped_width; + grouped_plane_stride = grouped_width * ne11; + } GGML_ASSERT(!use_native_mxfp4); } @@ -195,9 +214,8 @@ static void ggml_cuda_mul_mat_q_impl( if (grouped_src) { quantize_mmq_q8_1_grouped_cuda( src1_d, src1_q8_1.get(), src0->type, - ne10, grouped_physical->ne[0], - grouped_physical->nb[1] / ts_src1, - grouped_physical->nb[2] / ts_src1, + ne10, grouped_width, + grouped_row_stride, grouped_plane_stride, ne10_padded, ne11, stream); } else if (use_native_mxfp4) { static_assert(sizeof(block_fp4_mmq) == 4 * sizeof(block_q8_1)); diff --git a/server/deps/llama.cpp/ggml/src/ggml-hip/CMakeLists.txt b/server/deps/llama.cpp/ggml/src/ggml-hip/CMakeLists.txt index a7d4e0ea2..46e6a2928 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-hip/CMakeLists.txt +++ b/server/deps/llama.cpp/ggml/src/ggml-hip/CMakeLists.txt @@ -92,7 +92,13 @@ if (NOT GGML_BACKEND_DL) target_compile_definitions(ggml PUBLIC GGML_USE_CUDA) endif() -add_compile_definitions(GGML_USE_HIP) +if (GGML_HIP_MODULE) + # Do not leak HIP runtime spellings into the linked CUDA application in a + # mixed-vendor build. The module is the only target that needs this flag. + target_compile_definitions(ggml-hip PRIVATE GGML_USE_HIP) +else() + add_compile_definitions(GGML_USE_HIP) +endif() if (GGML_CUDA_FORCE_MMQ) add_compile_definitions(GGML_CUDA_FORCE_MMQ) diff --git a/server/docs/DS4.md b/server/docs/DS4.md index e0fc2dba7..3c1787421 100644 --- a/server/docs/DS4.md +++ b/server/docs/DS4.md @@ -124,6 +124,56 @@ decode and 415.52 tok/s median sparse prefill. Those numbers require the full qualified manifest, including the burn-in kernel switches; they are not a claim for the minimal activation example above. +#### CUDA 3090 + Strix Halo in one process + +The mixed-vendor build links the selected target runtime normally and loads +the other vendor as an isolated backend module. CUDA and HIP device indices +have separate namespaces, so `cuda:0` and `hip:0` are a valid pair. +Cross-vendor activations are staged in host memory inside the process; native +peer access is intentionally not attempted. + +```bash +cmake -S server -B server/build-cuda-hip \ + -DDFLASH27B_GPU_BACKEND=hip \ + -DDFLASH27B_ENABLE_MIXED_CUDA_HIP=ON \ + -DDFLASH27B_CUDA_ARCHITECTURES=86 \ + -DDFLASH27B_HIP_ARCHITECTURES=gfx1151 \ + -DCMAKE_BUILD_TYPE=Release +cmake --build server/build-cuda-hip -j +ctest --test-dir server/build-cuda-hip -R mixed_cuda_hip --output-on-failure +``` + +```bash +export DFLASH_DS4_MOE_TP=1 +export DFLASH_DS4_MOE_TP_INPROC=1 +export DFLASH_DS4_MOE_TP_BACKEND=cuda +export DFLASH_DS4_MOE_TP_GPU=0 # cuda:0 (RTX 3090) +export DFLASH_DS4_MOE_TP_CONCENTRATE_COLD=1 +export DFLASH_DS4_TP_SCHEDULE_BRANCHES=1 +export DFLASH_DS4_TP_TARGETED_JOIN_SPLIT=1 +export GGML_BATCH_PEER_COPIES=1 +# Start conservatively and tune from the startup placement and memory logs; +# the usable budget depends on the model, placement policy, and free VRAM. +export DFLASH_EXPERT_BUDGET_MB=85000 +export DFLASH_DS4_DRAFT=/path/to/dspark-draft.gguf +export DFLASH_DS4_DRAFT_BACKEND=cuda +export DFLASH_DS4_DRAFT_GPU=0 + +./server/build-cuda-hip/dflash_server /path/to/deepseek4-target.gguf \ + --target-device hip:0 \ + --ds4-expert-top-k 4 \ + --ds4-prefill sparse +``` + +The peer module is normally found beside the executable. Set +`DFLASH_CUDA_BACKEND_PATH` or `DFLASH_HIP_BACKEND_PATH` only when packaging it +elsewhere. Sparse/approximate DeepSeek4 prefill remains restricted to a HIP +target; CUDA-primary ROCmFP2 execution is not yet qualified. The mixed path is +burn-in functionality. On the qualified 3090 + Strix machine, the tuned top-4 +performance profile held 48.1 tok/s median on the deterministic 128-token +workload. The all-6-expert reference-exact mode is a correctness profile, not +a throughput profile. + ### Local single-shard If the adapter decides all 43 layers fit on one CUDA GPU, it loads a single shard locally and no IPC daemon is involved. @@ -182,13 +232,22 @@ The runtime logs the chosen split with a `[deepseek4-split] auto-split:` banner. | `DFLASH_DS4_CUDA_LAYERS` | Override the auto-split heuristic and pin the first `N` DeepSeek4 layers to CUDA. The remaining `43 - N` layers run on the Halo shard. | | `DFLASH_DS4_TIMING` | Enable DS4 timing logs for the layer-split parent and target-shard daemon. Useful for profiling prefill/decode breakdowns; leave unset for normal runs. | | `DFLASH_DS4_SPEC` / `DFLASH_DS4_DRAFT` | Enable DSpark and select its GGUF. | -| `DFLASH_DS4_DRAFT_GPU` | HIP device for the in-process drafter. | +| `DFLASH_DS4_DRAFT_BACKEND` / `DFLASH_DS4_DRAFT_GPU` | Backend and device for the in-process drafter. | | `DFLASH_DS4_MOE_TP` | Enable routed-expert partitioning. | -| `DFLASH_DS4_MOE_TP_INPROC` | Use two local HIP backends instead of an expert IPC worker. | -| `DFLASH_DS4_MOE_TP_GPU` | HIP device that owns the cold expert stack. | +| `DFLASH_DS4_MOE_TP_INPROC` | Use two local GPU backends instead of an expert IPC worker. | +| `DFLASH_DS4_MOE_TP_BACKEND` | Cold expert backend (`cuda` or `hip`); mixed builds default to the peer runtime. | +| `DFLASH_DS4_MOE_TP_GPU` | Device index within the cold expert backend. | +| `DFLASH_DS4_MOE_TP_CONCENTRATE_COLD` | Cross-vendor burn-in mode: place complete cold expert layers on the peer to reduce joins. | +| `DFLASH_DS4_MOE_TP_PEER_HOT` | With a routing profile, place its hottest experts on the secondary owner. | +| `DFLASH_DS4_CROSS_VENDOR_OWNER_SUMS` | Reduce each owner's routed outputs locally before the final cross-vendor add. This changes floating-point association and is not the byte-identity mode. | +| `DFLASH_DS4_TP_SCHEDULE_BRANCHES` | Submit the two owner branches independently through the mixed scheduler. | +| `DFLASH_DS4_TP_TARGETED_JOIN_SPLIT` | Gather the peer result at the join without an extra peer fence per layer. | +| `DFLASH_DS4_COMP_PAD_STRIDE` | Exact compressed-KV padding bucket; wider buckets trade small masked work for fewer verifier graph captures. | +| `DFLASH_DS4_DISABLE_GROUPED_OUTPUT_PROJECTION` | Diagnostic fallback for runtimes that cannot preserve grouped projection metadata across a scheduler copy. | +| `DFLASH_CUDA_BACKEND_PATH` / `DFLASH_HIP_BACKEND_PATH` | Optional explicit peer backend module path. | | `DFLASH_EXPERT_BUDGET_MB` | Main-GPU memory budget for hot experts. | | `DFLASH_DS4_HOTNESS_CSV` | Optional per-layer routing profile for hot placement. | -| `GGML_CUDA_BATCH_PEER_COPIES` | Batch ordered peer copies behind one dependency. | +| `GGML_BATCH_PEER_COPIES` | Batch peer-runtime copies and unlike-runtime pinned-host staging with one source wait per split. The old `GGML_CUDA_BATCH_PEER_COPIES` spelling remains an alias. | | `DFLASH_MOE_PREFILL_PERSISTENT_OWNER_ALLOC` | Long-prefill arena kill switch; set `0` to restore per-layer owner allocation. | `DFLASH_DS4_TIMING` enables the existing timing banners: @@ -249,11 +308,13 @@ whole-model GPU graph uses stable padded reduction shapes, so near-tied greedy logits can select a different token than the normal causal verifier even at temperature 0. Leave it unset when comparing against the normal verifier, or set `DFLASH_DS4_SEQ_VERIFY=1` for the slower token-at-a-time verification -diagnostic. Neither fused verification nor the separate +diagnostic. `DFLASH_DS4_SPEC_REFERENCE_EXACT=1` combines sequential target +verification with full rollback snapshots for byte-identity checks. Neither +fused verification nor the separate `--ds4-expert-top-k 4` approximation should be presented as byte-identical AR. DSpark can verify against in-process heterogeneous expert placement. The -drafter remains local to its selected HIP backend; a failed draft load is +drafter remains local to its selected GPU backend; a failed draft load is reported and falls back to normal autoregressive decode. The target cache and sampler stay on the main backend while routed target experts execute on their configured owners. `--ds4-expert-top-k 4` remains a separate approximate diff --git a/server/docs/ENVIRONMENT.md b/server/docs/ENVIRONMENT.md index a85e1fa65..960a59841 100644 --- a/server/docs/ENVIRONMENT.md +++ b/server/docs/ENVIRONMENT.md @@ -27,8 +27,19 @@ consolidation of this list into CLI flags is tracked as follow-up work. | `DFLASH_MMID_GROUPED_TYPES` | 7 | Grouped-kernel type mask; bit 3 (`8`) opts ROCmFP2/ROCmFP3 into the path. | | `DFLASH_MMID_GROUPED_DEVICE` | -1 | Optional zero-based device restriction; unset/-1 applies to every eligible device. | | `DFLASH_DS4_MOE_TP` / `DFLASH_DS4_MOE_TP_INPROC` | unset | BURN-IN: enable DeepSeek4 route-owner expert parallelism in one process. | -| `DFLASH_DS4_MOE_TP_GPU` | auto | HIP device for the cold DeepSeek4 expert owner. | -| `GGML_CUDA_BATCH_PEER_COPIES` | unset | BURN-IN: publish ordered HIP peer copies with one cross-device dependency per source/destination pair. | +| `DFLASH_DS4_MOE_TP_BACKEND` / `DFLASH_MOE_TP_BACKEND` | peer runtime in a mixed build; compiled runtime otherwise | Select the in-process cold expert owner backend. | +| `DFLASH_DS4_MOE_TP_GPU` | peer backend device 0 in a mixed build; other local device otherwise | Device index within the cold DeepSeek4 expert backend. | +| `DFLASH_DS4_MOE_TP_CONCENTRATE_COLD` | unset | BURN-IN: use complete peer-owned expert layers to reduce cross-runtime joins; falls back when the placement would exceed the target budget. | +| `DFLASH_DS4_MOE_TP_PEER_HOT` | unset | BURN-IN: with `DFLASH_DS4_HOTNESS_CSV`, reserve the profile's hottest experts for the secondary owner. | +| `DFLASH_DS4_CROSS_VENDOR_OWNER_SUMS` | unset | BURN-IN: reduce each mixed-vendor owner's routed outputs locally before the final owner add. This changes floating-point association and is not the byte-identity mode. | +| `DFLASH_DS4_TP_SCHEDULE_BRANCHES` | unset | BURN-IN: expose independent mixed-vendor expert branches to the common multi-backend scheduler. | +| `DFLASH_DS4_TP_TARGETED_JOIN_SPLIT` / `DFLASH_MOE_TP_TARGETED_JOIN_SPLIT` | unset | BURN-IN: start a main-GPU split only at each peer-result join, avoiding an extra peer fence per MoE layer. | +| `DFLASH_DS4_COMP_PAD_STRIDE` | 16 | BURN-IN: compressed-KV padding bucket (`16`, `32`, `64`, or `128`); wider exact-masked buckets reduce verifier graph recapture churn. | +| `DFLASH_DS4_DISABLE_GROUPED_OUTPUT_PROJECTION` | unset | DEBUG: restore the materialized output projection when diagnosing grouped-view copies across unlike runtimes. | +| `DFLASH_DS4_DRAFT_BACKEND` / `DFLASH_DS4_DRAFT_GPU` | compiled backend / target device | Select the in-process DSpark backend and device. | +| `DFLASH_CUDA_BACKEND_PATH` / `DFLASH_HIP_BACKEND_PATH` | auto-discovered beside the executable | Explicit peer module file path for a mixed CUDA+HIP build. | +| `GGML_BATCH_PEER_COPIES` | unset | BURN-IN: batch peer-runtime copies and unlike-runtime host staging with one source wait per split. `GGML_CUDA_BATCH_PEER_COPIES` remains a compatibility alias. | +| `GGML_SCHED_PROFILE` / `GGML_SCHED_PROFILE_MIN_SPLITS` | unset / 1 | DEBUG: report scheduler splits, copy volume, submission time, and source/destination synchronization time. | | `DFLASH_MOE_PREFILL_PERSISTENT_OWNER_ALLOC` | 1 for qualified long heterogeneous prefill | KILL SWITCH: =0 restores per-layer route/owner scratch allocation. | | `DFLASH_MOE_TP_*` / `DFLASH_MOE_HYBRID_PREFILL_EAGER` | unset | BURN-IN: model-neutral names for common heterogeneous-MoE scheduling and kernel policy. Existing `DFLASH_DS4_*` names remain compatibility aliases. | | `DFLASH_MMID_TELEMETRY` | unset | DEBUG: report MUL_MAT_ID dispatch, MMVQ variant, and per-node graph compatibility. | @@ -37,7 +48,7 @@ consolidation of this list into CLI flags is tracked as follow-up work. | `DFLASH_PREFILL_CACHE_SLOTS` | 0 | Container-entrypoint equivalent of `--prefill-cache-slots`; not read directly by the native binary. | | `DFLASH_SPLIT_FAST_ROLLBACK` | unset | OPT-IN: exact F32 checkpoints and replay-free rollback for local qwen35 target layer splits. Prefer `--target-split-fast-rollback`; adds checkpoint VRAM (~1.65 GiB for the measured Qwen3.6-27B q=16 split). | | `DFLASH_STALL_TOOL_PREFIX` | unset | OPT-IN: recover a stalled tool call by injecting the prepared tool prefix when generation stops after an action suffix. | -| `DFLASH_DS4_SPEC` / `DFLASH_DS4_DRAFT` / `DFLASH_DS4_DRAFT_GPU` | unset | OPT-IN: enable DeepSeek4 DSpark, select its draft GGUF, and optionally select the local drafter GPU. See `DS4.md`. | +| `DFLASH_DS4_SPEC` / `DFLASH_DS4_DRAFT` / `DFLASH_DS4_DRAFT_BACKEND` / `DFLASH_DS4_DRAFT_GPU` | unset | OPT-IN: enable DeepSeek4 DSpark, select its draft GGUF, and optionally select the local drafter backend/device. See `DS4.md`. | | `DFLASH_DS4_CUDA_LAYERS` | auto | Override the DeepSeek4 heterogeneous layer-split heuristic. See `DS4.md`. | ## Full inventory (generated) @@ -59,6 +70,7 @@ consolidation of this list into CLI flags is tracked as follow-up work. - `DFLASH_ADAPTIVE_WIDTH_MIN` - adaptive_verify_width.h - `DFLASH_ADAPTIVE_WIDTH_THETA` - adaptive_verify_width.h - `DFLASH_COLD_THREADS` - moe_expert_compute_cpu.cpp +- `DFLASH_CUDA_BACKEND_PATH` - dynamic_backend.cpp - `DFLASH_CUDA_MMVQ_MOE_ALIGN_SHARED_IDS` - moe_hybrid_ffn_eval.cpp - `DFLASH_CUDA_MMVQ_MOE_KERNEL` - moe_hybrid_ffn_eval.cpp - `DFLASH_DISABLE_DRAFT_ATTN` - draft_graph.cpp @@ -73,24 +85,34 @@ consolidation of this list into CLI flags is tracked as follow-up work. - `DFLASH_DRAFT_PERSIST` - laguna_backend.cpp - `DFLASH_DROP_COLD` - qwen35moe_backend.cpp, qwen35moe_pipelined_decode.cpp - `DFLASH_DS4_ADAPTIVE_WIDTH` - deepseek4_dspark_spec.cpp +- `DFLASH_DS4_COMP_PAD_STRIDE` - deepseek4_graph.cpp +- `DFLASH_DS4_CROSS_VENDOR_OWNER_SUMS` - deepseek4_fused_verify.inc - `DFLASH_DS4_CUDA_LAYERS` - deepseek4_layer_split_adapter.cpp - `DFLASH_DS4_DENSE_TP_MASK` - deepseek4_loader.cpp - `DFLASH_DS4_DENSE_TP_STRIX_FRACTION` - deepseek4_loader.cpp +- `DFLASH_DS4_DISABLE_GROUPED_OUTPUT_PROJECTION` - deepseek4_graph.cpp - `DFLASH_DS4_DRAFT` - deepseek4_backend.cpp +- `DFLASH_DS4_DRAFT_BACKEND` - deepseek4_backend.cpp - `DFLASH_DS4_DRAFT_GPU` - deepseek4_backend.cpp - `DFLASH_DS4_DSPARK_DEBUG` - deepseek4_graph.cpp - `DFLASH_DS4_FUSED_VERIFY` - deepseek4_dspark_spec.cpp, deepseek4_loader.cpp - `DFLASH_DS4_HOTNESS_CSV` - deepseek4_backend.cpp - `DFLASH_DS4_MOE_TP` - deepseek4_backend.cpp +- `DFLASH_DS4_MOE_TP_BACKEND` - deepseek4_backend.cpp +- `DFLASH_DS4_MOE_TP_CONCENTRATE_COLD` - deepseek4_backend.cpp - `DFLASH_DS4_MOE_TP_GPU` - deepseek4_backend.cpp - `DFLASH_DS4_MOE_TP_INPROC` - deepseek4_backend.cpp +- `DFLASH_DS4_MOE_TP_PEER_HOT` - deepseek4_backend.cpp - `DFLASH_DS4_ROUTING_STATS_OUT` - deepseek4_backend.cpp - `DFLASH_DS4_SEQ_VERIFY` - deepseek4_dspark_spec.cpp - `DFLASH_DS4_SPEC` - deepseek4_backend.cpp +- `DFLASH_DS4_SPEC_REFERENCE_EXACT` - deepseek4_dspark_spec.cpp - `DFLASH_DS4_SPEC_Q` - deepseek4_dspark_spec.cpp - `DFLASH_DS4_TIMING` - deepseek4_backend.cpp, deepseek4_target_shard_ipc_daemon.cpp - `DFLASH_DS4_TP_CAPTURE_CACHE_SLOTS` - deepseek4_fused_verify.inc - `DFLASH_DS4_TP_FUSED_CACHE_SLOTS` - deepseek4_fused_verify.inc +- `DFLASH_DS4_TP_SCHEDULE_BRANCHES` - deepseek4_fused_verify.inc +- `DFLASH_DS4_TP_TARGETED_JOIN_SPLIT` - deepseek4_fused_verify.inc, moe_hybrid_ffn_eval.cpp - `DFLASH_DS4_TOPK` - deepseek4_graph.cpp - `DFLASH_EXPERT_BUDGET_MB` - deepseek4_backend.cpp, laguna_backend.cpp, qwen35moe_backend.cpp - `DFLASH_EXPERT_BUDGET_PCT` - laguna_backend.cpp @@ -112,6 +134,7 @@ consolidation of this list into CLI flags is tracked as follow-up work. - `DFLASH_GPU_DRAFT_TOPK` - qwen35_dflash_target.cpp - `DFLASH_GPU_SAMPLE` - geometric_sampler_cuda.cu - `DFLASH_GPU_VERIFY_ARGMAX` - qwen35_dflash_target.cpp +- `DFLASH_HIP_BACKEND_PATH` - dynamic_backend.cpp - `DFLASH_IGNORE_EOS` - laguna_backend.cpp - `DFLASH_KVFLASH` - gemma4_backend.cpp, gemma4_layer_split_adapter.cpp, kvflash_pager.h, laguna_backend.cpp, laguna_layer_split_adapter.cpp, qwen35_backend.cpp, qwen35_layer_split_adapter.cpp - `DFLASH_KVFLASH_DRAFTER` - kvflash_pager.h @@ -125,6 +148,7 @@ consolidation of this list into CLI flags is tracked as follow-up work. - `DFLASH_LAGUNA_DSPARK_CONFIDENCE_THRESHOLD` - laguna_backend.cpp - `DFLASH_LAGUNA_DSPARK_TREE` - laguna_backend.cpp - `DFLASH_LAGUNA_EXPERT_CACHE` - moe_hybrid_ffn_eval.cpp +- `DFLASH_MOE_TP_TARGETED_JOIN_SPLIT` - moe_hybrid_ffn_eval.cpp, deepseek4_fused_verify.inc - `DFLASH_LAGUNA_FUSED_DOMINO` - laguna_backend.cpp - `DFLASH_LAGUNA_FUSED_DSPARK` - laguna_backend.cpp - `DFLASH_LAGUNA_FUSED_QK` - laguna_target_loader.cpp @@ -177,6 +201,7 @@ consolidation of this list into CLI flags is tracked as follow-up work. - `DFLASH_MOE_PREFILL_HOT_SUB_BATCH` - moe_hybrid_ffn_eval.cpp - `DFLASH_MOE_PREFILL_MASKED_COLD` - moe_hybrid_ffn_eval.cpp - `DFLASH_MOE_PREFILL_PERSISTENT_OWNER_ALLOC` - deepseek4_graph.cpp +- `DFLASH_MOE_TP_BACKEND` - deepseek4_backend.cpp - `DFLASH_NO_MASK` - laguna_backend.cpp - `DFLASH_NO_MOE_ROUTER_FUSE` - qwen35moe_ffn.cpp - `DFLASH_NO_MOE_SWIGLU_FUSE` - qwen35moe_ffn.cpp @@ -212,7 +237,6 @@ consolidation of this list into CLI flags is tracked as follow-up work. - `DFLASH_VERIFY_WIDTH` - qwen35moe_backend.cpp - `FAST_ROLLBACK_DIAG` - qwen35_dflash_target.cpp - `HOME` - spark_corpus.cpp -- `LUCE_CUDA_I32_REPEAT` - moe_hybrid_ffn_eval.cpp - `LUCE_MMVQ_MAX_NCOLS` - deepseek4_backend.cpp - `LUCE_QK_FUSE_LAYERS` - laguna_target_graph.cpp - `LUCE_QK_FUSE_MODE` - laguna_target_graph.cpp diff --git a/server/src/common/dynamic_backend.cpp b/server/src/common/dynamic_backend.cpp new file mode 100644 index 000000000..bf8293e35 --- /dev/null +++ b/server/src/common/dynamic_backend.cpp @@ -0,0 +1,189 @@ +#include "dynamic_backend.h" + +#include "ggml-cuda.h" + +#include +#include +#include +#include +#include + +#if defined(__linux__) +#include +#endif + +namespace dflash::common { +namespace { + +namespace fs = std::filesystem; + +static std::string executable_directory() { +#if defined(__linux__) + std::vector path(1024); + for (;;) { + const ssize_t size = ::readlink("/proc/self/exe", path.data(), path.size()); + if (size < 0) return {}; + if ((size_t)size < path.size()) { + return fs::path(std::string(path.data(), (size_t)size)) + .parent_path().string(); + } + path.resize(path.size() * 2); + } +#else + return {}; +#endif +} + +#if defined(DFLASH27B_BACKEND_MIXED) +static ggml_backend_reg_t load_peer_registry(PlacementBackend backend, + std::string * error) { + static std::mutex mutex; + static ggml_backend_reg_t cuda_registry = nullptr; + static ggml_backend_reg_t hip_registry = nullptr; + std::lock_guard lock(mutex); + + ggml_backend_reg_t & registry = backend == PlacementBackend::Cuda + ? cuda_registry : hip_registry; + if (registry) return registry; + const char * registry_name = backend == PlacementBackend::Cuda + ? "CUDA" : "ROCm"; + const char * path_variable = backend == PlacementBackend::Cuda + ? "DFLASH_CUDA_BACKEND_PATH" : "DFLASH_HIP_BACKEND_PATH"; + const char * module_name = backend == PlacementBackend::Cuda + ? "libggml-cuda.so" : "libggml-hip.so"; + + registry = ggml_backend_reg_by_name(registry_name); + if (registry) return registry; + + std::vector candidates; + if (const char * explicit_path = std::getenv(path_variable)) { + if (*explicit_path) candidates.emplace_back(explicit_path); + } + if (candidates.empty()) { + const std::string exe_dir = executable_directory(); + if (!exe_dir.empty()) { + candidates.emplace_back(fs::path(exe_dir) / module_name); + } + } + + std::string attempted; + for (const fs::path & candidate : candidates) { + std::error_code ec; + if (!fs::is_regular_file(candidate, ec)) { + if (!attempted.empty()) attempted += ", "; + attempted += candidate.string(); + continue; + } + ggml_backend_reg_t loaded = ggml_backend_load(candidate.string().c_str()); + if (loaded && + std::string(ggml_backend_reg_name(loaded)) == registry_name) { + registry = loaded; + return registry; + } + if (loaded) ggml_backend_unload(loaded); + if (!attempted.empty()) attempted += ", "; + attempted += candidate.string(); + } + + if (error) { + *error = "could not load the "; + *error += registry_name; + *error += " backend module"; + if (!attempted.empty()) *error += " (tried " + attempted + ")"; + *error += "; set "; + *error += path_variable; + *error += " to the full path of "; + *error += module_name; + } + return nullptr; +} +#endif + +} // namespace + +ggml_backend_t init_placement_backend(PlacementBackend backend, + int device, + std::string * error) { + if (backend == PlacementBackend::Auto) backend = compiled_placement_backend(); + if (device < 0) { + if (error) *error = "GPU device index must be non-negative"; + return nullptr; + } + + if (backend == compiled_placement_backend()) { + ggml_backend_t result = ggml_backend_cuda_init(device); + if (!result && error) { + *error = "failed to initialize "; + *error += placement_backend_name(backend); + *error += " device " + std::to_string(device); + } + return result; + } + +#if defined(DFLASH27B_BACKEND_MIXED) + if (backend == PlacementBackend::Cuda || backend == PlacementBackend::Hip) { + ggml_backend_reg_t registry = load_peer_registry(backend, error); + if (!registry) return nullptr; + const size_t count = ggml_backend_reg_dev_count(registry); + if ((size_t)device >= count) { + if (error) { + *error = std::string(placement_backend_name(backend)) + + " device " + std::to_string(device) + + " is out of range (found " + + std::to_string(count) + ")"; + } + return nullptr; + } + ggml_backend_t result = ggml_backend_dev_init( + ggml_backend_reg_dev_get(registry, (size_t)device), nullptr); + if (!result && error) { + *error = "failed to initialize "; + *error += placement_backend_name(backend); + *error += " device " + std::to_string(device); + } + return result; + } +#endif + + if (error) { + *error = "this binary does not contain the requested GPU backend"; + } + return nullptr; +} + +PlacementBackend placement_backend_of(ggml_backend_t backend) { + if (!backend) return PlacementBackend::Auto; + ggml_backend_dev_t device = ggml_backend_get_device(backend); + if (!device) return PlacementBackend::Auto; + ggml_backend_reg_t registry = ggml_backend_dev_backend_reg(device); + if (!registry) return PlacementBackend::Auto; + const char * name = ggml_backend_reg_name(registry); + if (!name) return PlacementBackend::Auto; + const std::string value(name); + if (value == "CUDA") return PlacementBackend::Cuda; + if (value == "ROCm") return PlacementBackend::Hip; + return PlacementBackend::Auto; +} + +BackendPairCapabilities backend_pair_capabilities(ggml_backend_t first, + ggml_backend_t second) { + BackendPairCapabilities result; + if (!first || !second) return result; + + result.same_runtime = ggml_guid_matches( + ggml_backend_guid(first), ggml_backend_guid(second)); + if (!result.same_runtime) return result; + + const auto is_gpu = [](ggml_backend_t backend) { + ggml_backend_dev_t device = ggml_backend_get_device(backend); + if (!device) return false; + const enum ggml_backend_dev_type type = + ggml_backend_dev_type(device); + return type == GGML_BACKEND_DEVICE_TYPE_GPU || + type == GGML_BACKEND_DEVICE_TYPE_IGPU; + }; + result.native_gpu_handoff = is_gpu(first) && is_gpu(second); + return result; +} + +} // namespace dflash::common diff --git a/server/src/common/dynamic_backend.h b/server/src/common/dynamic_backend.h new file mode 100644 index 000000000..77a6186b9 --- /dev/null +++ b/server/src/common/dynamic_backend.h @@ -0,0 +1,38 @@ +// Runtime backend selection for configurations that contain more than one +// GPU vendor in a single process. + +#pragma once + +#include "placement/placement_backend.h" + +#include "ggml-backend.h" + +#include + +namespace dflash::common { + +struct BackendPairCapabilities { + // Backends built by the same runtime implementation have compatible + // device pointers, streams, and events. Different vendors must exchange + // tensors through backend-neutral staging. + bool same_runtime = false; + bool native_gpu_handoff = false; +}; + +// Initialize one device from the requested backend. Ordinary builds keep +// using their linked backend; mixed builds load the isolated peer module on +// first use. +ggml_backend_t init_placement_backend(PlacementBackend backend, + int device, + std::string * error = nullptr); + +// Resolve the vendor represented by an initialized ggml backend. +PlacementBackend placement_backend_of(ggml_backend_t backend); + +// Describe the operations that are safe between two initialized backends. +// This deliberately keys off backend identity rather than vendor names so the +// scheduling code also works for future runtime modules. +BackendPairCapabilities backend_pair_capabilities(ggml_backend_t first, + ggml_backend_t second); + +} // namespace dflash::common diff --git a/server/src/common/feature_gate.cpp b/server/src/common/feature_gate.cpp index c7d66348c..0ada44496 100644 --- a/server/src/common/feature_gate.cpp +++ b/server/src/common/feature_gate.cpp @@ -169,14 +169,19 @@ std::string check_feature_compatibility( arch + "')"; } - // Approximate prefill and the fused decode options are implemented only - // in the monolithic HIP DeepSeek4 backend; the layer-split adapter and - // the CUDA path have no equivalent. + // Approximate prefill and fused decode are implemented only in the + // monolithic HIP DeepSeek4 backend. Expert top-k is model policy handled + // by either monolithic backend, but the layer-split adapter does not yet + // propagate it. const bool monolithic_ds4 = arch == "deepseek4" && target_backend == PlacementBackend::Hip && !args.device.is_layer_split() && !args.remote_target_shard.enabled(); + const bool local_ds4 = + arch == "deepseek4" && + !args.device.is_layer_split() && + !args.remote_target_shard.enabled(); // ── approximate --ds4-prefill × placement if (arch == "deepseek4" && @@ -188,11 +193,16 @@ std::string check_feature_compatibility( "--ds4-prefill exact for split, remote, or CUDA placement"; } - // ── --ds4-fused-decode / --ds4-expert-top-k × placement - if ((args.ds4_fused_decode || args.ds4_expert_top_k != 0) && - !monolithic_ds4) { - return "--ds4-fused-decode and --ds4-expert-top-k currently require " - "single-device HIP DeepSeek4"; + // ── --ds4-fused-decode × placement + if (args.ds4_fused_decode && !monolithic_ds4) { + return "--ds4-fused-decode currently requires single-device HIP " + "DeepSeek4"; + } + + // ── --ds4-expert-top-k × architecture/adapter + if (args.ds4_expert_top_k != 0 && !local_ds4) { + return "--ds4-expert-top-k currently requires a single local " + "DeepSeek4 backend"; } return {}; diff --git a/server/src/common/moe_hybrid_ffn_eval.cpp b/server/src/common/moe_hybrid_ffn_eval.cpp index 06be22801..8df318642 100644 --- a/server/src/common/moe_hybrid_ffn_eval.cpp +++ b/server/src/common/moe_hybrid_ffn_eval.cpp @@ -3,7 +3,6 @@ #include "ggml-alloc.h" #include "ggml-backend.h" -#include "ggml-cuda.h" #include #include @@ -71,117 +70,63 @@ static uint64_t elapsed_us(HybridClock::time_point start, HybridClock::time_poin return (uint64_t) std::chrono::duration_cast(end - start).count(); } -static bool compact_materialized_experts_enabled() { - static const bool enabled = [] { - const char * raw = std::getenv("DFLASH_MOE_COMPACT_MATERIALIZED"); - return raw && *raw && std::strcmp(raw, "0") != 0; - }(); - return enabled; +static bool backend_is_gpu(ggml_backend_t backend) { + if (!backend) return false; + ggml_backend_dev_t device = ggml_backend_get_device(backend); + if (!device) return false; + const enum ggml_backend_dev_type type = ggml_backend_dev_type(device); + return type == GGML_BACKEND_DEVICE_TYPE_GPU || + type == GGML_BACKEND_DEVICE_TYPE_IGPU; } -// The legacy ROCmFP2 verifier graph expands a q-token routed FFN into q -// independent single-token subgraphs. That was required before the CUDA/HIP -// backend gained its grouped MUL_MAT_ID MMVQ kernel, but it multiplies graph -// nodes, scheduler copies, and launches by the verify width. Keep the old -// lowering as the default while the grouped path is qualified on each ROCm -// architecture; opt in with DFLASH_MOE_TP_GROUPED_MMVQ=1. The original DS4 -// variable remains a compatibility alias. -static bool grouped_mmvq_moe_enabled() { - static const bool enabled = [] { - const char * raw = moe_policy_env( - "DFLASH_MOE_TP_GROUPED_MMVQ", "DFLASH_DS4_TP_GROUPED_MMVQ"); - return raw && *raw && std::strcmp(raw, "0") != 0; - }(); - return enabled; -} - -static bool gpu_i32_repeat_enabled() { +static bool compact_materialized_experts_enabled() { static const bool enabled = [] { - const char * raw = std::getenv("LUCE_CUDA_I32_REPEAT"); + const char * raw = std::getenv("DFLASH_MOE_COMPACT_MATERIALIZED"); return raw && *raw && std::strcmp(raw, "0") != 0; }(); return enabled; } -// The regular DeepSeek graph already uses ggml_laguna_moe_combine for the -// route-weighted expert reduction. Keep the heterogeneous graph A/B-able -// while replacing its MUL + shape-only REPEAT_BACK sequence with the same -// exact owner-local kernel. -static bool fused_moe_combine_enabled() { - static const bool enabled = [] { - const char * raw = std::getenv("DFLASH_MOE_FUSED_COMBINE"); - return raw && *raw && std::strcmp(raw, "0") != 0; - }(); - return enabled; +static bool moe_policy_flag(const char * name, const char * legacy_name = nullptr) { + const char * raw = moe_policy_env(name, legacy_name); + return raw && *raw && std::strcmp(raw, "0") != 0; } -// DeepSeek V4 stores routed gate and up projections as one tensor with the -// output rows concatenated. The generic graph computes that full tensor, -// materializes two contiguous views, clamps both, and only then runs SwiGLU. -// Present the two weight halves as views instead: the CUDA/HIP graph optimizer -// can fuse both MUL_MAT_ID operations and DS4 SwiGLU into one MMVQ launch. -// This is exact when the external gate/up scale is one (the ROCmFP checkpoint -// used by the heterogeneous path); other scale values retain the old graph. -static bool fused_gate_up_mmvq_enabled() { - static const bool enabled = [] { - const char * raw = moe_policy_env( +const MoeHybridGraphPolicy & moe_hybrid_graph_policy() { + static const MoeHybridGraphPolicy policy = [] { + MoeHybridGraphPolicy result; + result.grouped_mmvq = moe_policy_flag( + "DFLASH_MOE_TP_GROUPED_MMVQ", "DFLASH_DS4_TP_GROUPED_MMVQ"); + result.fused_combine = moe_policy_flag("DFLASH_MOE_FUSED_COMBINE"); + result.fused_gate_up = moe_policy_flag( "DFLASH_MOE_TP_FUSED_GATE_UP", "DFLASH_DS4_TP_FUSED_GATE_UP"); - return raw && *raw && std::strcmp(raw, "0") != 0; - }(); - return enabled; -} - -static bool coarse_owner_op_enabled() { - static const bool enabled = [] { - const char * raw = moe_policy_env( + result.coarse_owner = moe_policy_flag( "DFLASH_MOE_TP_COARSE_OWNER", "DFLASH_DS4_TP_COARSE_OWNER"); - return raw && *raw && std::strcmp(raw, "0") != 0; - }(); - return enabled; -} - -static bool coarse_owner_split_op_enabled() { - static const bool enabled = [] { - const char * raw = moe_policy_env( - "DFLASH_MOE_TP_COARSE_OWNER_SPLIT", "DFLASH_DS4_TP_COARSE_OWNER_SPLIT"); - return raw && *raw && std::strcmp(raw, "0") != 0; - }(); - return enabled; -} + result.coarse_owner_split = moe_policy_flag( + "DFLASH_MOE_TP_COARSE_OWNER_SPLIT", + "DFLASH_DS4_TP_COARSE_OWNER_SPLIT"); + result.device_join = moe_policy_flag( + "DFLASH_MOE_TP_DEVICE_JOIN", "DFLASH_DS4_TP_DEVICE_JOIN"); + result.route_prefork = moe_policy_flag( + "DFLASH_MOE_TP_ROUTE_PREFORK", "DFLASH_DS4_TP_ROUTE_PREFORK"); + result.targeted_join_split = moe_policy_flag( + "DFLASH_MOE_TP_TARGETED_JOIN_SPLIT", + "DFLASH_DS4_TP_TARGETED_JOIN_SPLIT"); -static bool align_shared_moe_ids_enabled() { - static const bool enabled = [] { - const char * raw = std::getenv("DFLASH_CUDA_MMVQ_MOE_ALIGN_SHARED_IDS"); - const bool requested = raw && *raw && std::strcmp(raw, "0") != 0; + const bool align_requested = moe_policy_flag( + "DFLASH_CUDA_MMVQ_MOE_ALIGN_SHARED_IDS"); const char * kernel = std::getenv("DFLASH_CUDA_MMVQ_MOE_KERNEL"); const bool dedicated_kernel = !kernel || !*kernel || std::strcmp(kernel, "0") != 0; - if (requested && !dedicated_kernel) { + result.align_shared_ids = align_requested && dedicated_kernel; + if (align_requested && !dedicated_kernel) { std::fprintf(stderr, "[moe-hybrid] shared-ID alignment disabled because the dedicated " "MMVQ MoE kernel is disabled\n"); } - return requested && dedicated_kernel; - }(); - return enabled; -} - -static bool device_join_enabled() { - static const bool enabled = [] { - const char * raw = moe_policy_env( - "DFLASH_MOE_TP_DEVICE_JOIN", "DFLASH_DS4_TP_DEVICE_JOIN"); - return raw && *raw && std::strcmp(raw, "0") != 0; + return result; }(); - return enabled; -} - -static bool route_prefork_enabled() { - static const bool enabled = [] { - const char * raw = moe_policy_env( - "DFLASH_MOE_TP_ROUTE_PREFORK", "DFLASH_DS4_TP_ROUTE_PREFORK"); - return raw && *raw && std::strcmp(raw, "0") != 0; - }(); - return enabled; + return policy; } static void add_hybrid_telemetry(MoeHybridFfnTelemetry & dst, @@ -631,7 +576,8 @@ static bool build_batched_routed_graph( bool tokenwise = false, std::vector * backend_nodes = nullptr, bool allow_fused_combine = false, - bool force_fused_combine = false) + bool force_fused_combine = false, + bool defer_route_reduction = false) { const auto track = [&](ggml_tensor * t) -> ggml_tensor * { if (backend_nodes && t) backend_nodes->push_back(t); @@ -653,7 +599,8 @@ static bool build_batched_routed_graph( inp_col, sel_col, wts_col, n_embd, n_ff_exp, n_used, 1, swiglu_clamp, &routed_col, false, backend_nodes, - allow_fused_combine, force_fused_combine)) { + allow_fused_combine, force_fused_combine, + defer_route_reduction)) { return false; } joined = joined ? track(ggml_concat(ctx, joined, routed_col, 1)) @@ -666,7 +613,9 @@ static bool build_batched_routed_graph( ggml_tensor * cur_3d = ggml_reshape_3d(ctx, inp, n_embd, 1, n_tokens); ggml_tensor * gu = nullptr; const bool coarse_split_requested = - coarse_owner_op_enabled() && coarse_owner_split_op_enabled(); + !defer_route_reduction && + moe_hybrid_graph_policy().coarse_owner && + moe_hybrid_graph_policy().coarse_owner_split; const bool coarse_split_eligible = gate_tensor && up_tensor && gate_tensor->type == GGML_TYPE_Q2_0_ROCMFP2 && @@ -694,7 +643,8 @@ static bool build_batched_routed_graph( n_ff_exp, swiglu_clamp, gate_scale, up_scale, down_scale)); return *out_routed != nullptr; - } else if (coarse_owner_op_enabled() && + } else if (!defer_route_reduction && + moe_hybrid_graph_policy().coarse_owner && gate_up_tensor && gate_up_scale == 1.0f && gate_up_tensor->type == GGML_TYPE_Q2_0_ROCMFP2 && @@ -704,7 +654,7 @@ static bool build_batched_routed_graph( n_ff_exp, swiglu_clamp, down_scale)); return *out_routed != nullptr; } else if (gate_up_tensor && - fused_gate_up_mmvq_enabled() && + moe_hybrid_graph_policy().fused_gate_up && gate_up_scale == 1.0f) { GGML_ASSERT(gate_up_tensor->ne[1] == 2 * n_ff_exp); ggml_tensor * gate_w = ggml_view_3d( @@ -748,8 +698,8 @@ static bool build_batched_routed_graph( ggml_mul_mat_id(ctx, down_tensor, gu, sel), down_scale)); // Weight and sum over experts: [n_embd, n_used, n_tokens] * [1, n_used, n_tokens] - if (allow_fused_combine && - (force_fused_combine || fused_moe_combine_enabled())) { + if (!defer_route_reduction && allow_fused_combine && + (force_fused_combine || moe_hybrid_graph_policy().fused_combine)) { *out_routed = track(ggml_laguna_moe_combine(ctx, experts, wts)); return *out_routed != nullptr; } @@ -757,6 +707,11 @@ static bool build_batched_routed_graph( ggml_tensor * w_view = ggml_reshape_3d(ctx, wts, 1, n_used, n_tokens); experts = track(ggml_mul(ctx, experts, w_view)); + if (defer_route_reduction) { + *out_routed = experts; + return true; + } + // repeat_back uses this tensor for shape only, but the scheduler still // treats it as a leaf. Keep it on the branch backend; otherwise every MoE // branch acquires a tiny CPU split solely for an uninitialized shape leaf. @@ -767,6 +722,248 @@ static bool build_batched_routed_graph( return true; } +// One routed-expert owner. The established hot/cold storage names describe +// primary and secondary ownership respectively; this view keeps the graph +// construction logic independent of whether the secondary owner is a CPU, +// another same-runtime GPU, or a dynamically loaded peer runtime. +struct MoeOwnerGraphSpec { + const std::vector * local_by_global = nullptr; + ggml_tensor * gate = nullptr; + ggml_tensor * up = nullptr; + ggml_tensor * down = nullptr; + ggml_tensor * gate_up = nullptr; + ggml_tensor ** local_lut = nullptr; + ggml_tensor ** valid_lut = nullptr; + std::vector * remap_nodes = nullptr; + std::vector * branch_nodes = nullptr; + ggml_tensor * local_ids = nullptr; + ggml_tensor * masked_weights = nullptr; + ggml_tensor * output = nullptr; + + bool available() const { + return (gate_up || (gate && up)) && down; + } +}; + +static bool build_moe_owner_remap( + ggml_context * ctx, + const MoeHybridConfig & cfg, + ggml_tensor * global_ids, + ggml_tensor * router_weights, + int n_tokens, + MoeOwnerGraphSpec & owner) { + if (!owner.local_by_global || + (int) owner.local_by_global->size() != cfg.n_expert || + !owner.local_lut || !owner.valid_lut) { + return false; + } + const auto track = [&owner](ggml_tensor * tensor) { + if (tensor && owner.remap_nodes) { + owner.remap_nodes->push_back(tensor); + } + return tensor; + }; + if (!*owner.local_lut) { + *owner.local_lut = ggml_new_tensor_4d( + ctx, GGML_TYPE_I32, 1, cfg.n_expert, n_tokens, 1); + ggml_set_input(*owner.local_lut); + // These inputs are consumed late in a whole-model graph. Preserve + // their allocation from graph start so activation scratch cannot + // reuse the small buffer before its layer executes. + ggml_set_output(*owner.local_lut); + } + if (!*owner.valid_lut) { + *owner.valid_lut = ggml_new_tensor_4d( + ctx, GGML_TYPE_F32, 1, cfg.n_expert, n_tokens, 1); + ggml_set_input(*owner.valid_lut); + ggml_set_output(*owner.valid_lut); + } + + // Store immutable q-replicated lookup rows as graph inputs instead of + // running owner-local REPEAT kernels in every layer and verifier step. + ggml_tensor * mapped = track(ggml_get_rows( + ctx, *owner.local_lut, global_ids)); + mapped = track(ggml_reshape_2d( + ctx, mapped, cfg.n_expert_used, n_tokens)); + owner.local_ids = track(ggml_cont(ctx, mapped)); + ggml_tensor * valid = track(ggml_get_rows( + ctx, *owner.valid_lut, global_ids)); + valid = track(ggml_reshape_2d( + ctx, valid, cfg.n_expert_used, n_tokens)); + owner.masked_weights = track(ggml_mul(ctx, router_weights, valid)); + return true; +} + +static bool prepare_moe_owner_branch( + ggml_context * ctx, + const MoeHybridConfig & cfg, + ggml_tensor * global_ids, + ggml_tensor * router_weights, + int n_tokens, + MoeOwnerGraphSpec & owner) { + return !owner.available() || build_moe_owner_remap( + ctx, cfg, global_ids, router_weights, n_tokens, owner); +} + +static void align_moe_owner_routes( + ggml_context * ctx, + int n_tokens, + MoeOwnerGraphSpec & owner) { + if (!owner.available() || !owner.local_ids || n_tokens <= 1 || + !moe_hybrid_graph_policy().align_shared_ids) { + return; + } + owner.local_ids = ggml_ds4_moe_align_ids(ctx, owner.local_ids); + if (owner.local_ids && owner.remap_nodes) { + owner.remap_nodes->push_back(owner.local_ids); + } +} + +static bool build_moe_owner_branch( + ggml_context * ctx, + const MoeHybridConfig & cfg, + const MoeLayerDesc & desc, + ggml_tensor * inp, + int n_tokens, + bool canonical_route_join, + bool allow_fused_combine, + MoeOwnerGraphSpec & owner) { + if (!owner.available()) { + return true; + } + const MoeHybridGraphPolicy & policy = moe_hybrid_graph_policy(); + const ggml_tensor * dispatch_weights = owner.gate_up + ? owner.gate_up : owner.gate; + const bool tokenwise = !canonical_route_join && + dispatch_weights->type == GGML_TYPE_Q2_0_ROCMFP2 && + !(n_tokens > 1 && policy.grouped_mmvq); + return build_batched_routed_graph( + ctx, owner.gate, owner.up, owner.down, owner.gate_up, + desc.ffn_gate_exps_s, desc.ffn_up_exps_s, + desc.ffn_down_exps_s, desc.ffn_gate_up_exps_s, + inp, owner.local_ids, owner.masked_weights, + cfg.n_embd, cfg.n_ff_exp, cfg.n_expert_used, n_tokens, + cfg.swiglu_clamp, &owner.output, tokenwise, + owner.branch_nodes, allow_fused_combine, + /*force_fused_combine=*/false, canonical_route_join); +} + +static ggml_tensor * build_moe_owner_join( + ggml_context * ctx, + ggml_cgraph * schedule_graph, + const MoeHybridConfig & cfg, + const MoeLayerDesc & desc, + ggml_tensor * inp, + ggml_tensor * global_ids, + ggml_tensor * router_weights, + int n_tokens, + bool include_shared, + bool canonical_route_join, + ggml_tensor * primary, + ggml_tensor * secondary, + MoeHybridGraphInputs & out) { + // Canonical reduction consumes the secondary tensor below, but branch + // scheduling still needs the original producer as an independent graph + // root. Keep that dependency separate from the owner-level add state. + ggml_tensor * secondary_branch = secondary; + if (canonical_route_join) { + ggml_tensor * routes = nullptr; + if (primary && secondary) { + routes = ggml_add(ctx, primary, secondary); + out.join_nodes.push_back(routes); + } else { + routes = primary ? primary : secondary; + } + if (!routes) return nullptr; + + ggml_tensor * sum_shape = ggml_new_tensor_3d( + ctx, GGML_TYPE_F32, cfg.n_embd, 1, n_tokens); + ggml_tensor * route_sum = ggml_repeat_back(ctx, routes, sum_shape); + primary = ggml_reshape_2d(ctx, route_sum, cfg.n_embd, n_tokens); + out.join_nodes.push_back(sum_shape); + out.join_nodes.push_back(route_sum); + out.join_nodes.push_back(primary); + // The secondary contribution is consumed by the canonical route + // reduction and must not be added again as an owner-level partial. + secondary = nullptr; + } + + ggml_tensor * shared = include_shared + ? build_shared_expert_subgraph(ctx, desc, inp, cfg.swiglu_clamp) + : nullptr; + if (shared) { + primary = primary ? ggml_add(ctx, primary, shared) : shared; + } + + const MoeHybridGraphPolicy & policy = moe_hybrid_graph_policy(); + if (schedule_graph && secondary_branch && policy.route_prefork) { + // Materialize shared route IDs and weights before either owner branch. + // This prevents a late cross-runtime dependency from synchronizing the + // secondary stream before independent primary work can be submitted. + out.route_prefork_nodes.push_back(global_ids); + out.route_prefork_nodes.push_back(router_weights); + ggml_build_forward_expand(schedule_graph, global_ids); + ggml_build_forward_expand(schedule_graph, router_weights); + } + + if (canonical_route_join) { + if (schedule_graph && secondary_branch) { + // Queue secondary expert work before the final graph traversal + // reaches the primary branch and canonical route-order join. + ggml_build_forward_expand(schedule_graph, secondary_branch); + } + return primary; + } + + if (schedule_graph && secondary && policy.device_join) { + // Submit secondary and primary work independently, then perform the + // peer wait/copy at its exact position in the primary consumer graph. + ggml_set_output(secondary); + ggml_build_forward_expand(schedule_graph, secondary); + if (primary) { + ggml_build_forward_expand(schedule_graph, primary); + } + ggml_tensor * secondary_ready = + ggml_ds4_deferred_peer_copy(ctx, secondary); + // Diagnostic host-copy paths may prefill this tensor before its split + // launches, so reserve a stable allocation for the full graph. + ggml_set_input(secondary_ready); + ggml_set_output(secondary_ready); + out.deferred_peer_copy_nodes.push_back(secondary_ready); + out.main_output = primary; + out.peer_output = secondary_ready; + return primary ? ggml_add(ctx, secondary_ready, primary) + : secondary_ready; + } + + if (schedule_graph && secondary && policy.targeted_join_split) { + // The scheduler starts a new primary split at this join, allowing both + // owner branches to be enqueued before unlike-runtime host staging. + ggml_build_forward_expand(schedule_graph, secondary); + ggml_tensor * combined = primary + ? ggml_add(ctx, secondary, primary) : secondary; + if (primary) { + out.join_nodes.push_back(combined); + } + return combined; + } + + if (schedule_graph && secondary) { + ggml_build_forward_expand(schedule_graph, secondary); + ggml_tensor * secondary_fence = ggml_cont(ctx, secondary); + out.cold_nodes.push_back(secondary_fence); + return primary ? ggml_add(ctx, primary, secondary_fence) + : secondary_fence; + } + + if (secondary && primary) { + ggml_tensor * combined = ggml_add(ctx, secondary, primary); + out.join_nodes.push_back(combined); + return combined; + } + return secondary ? secondary : primary; +} + bool build_moe_hybrid_ffn_graph( ggml_context * ctx, ggml_cgraph * schedule_graph, @@ -779,7 +976,8 @@ bool build_moe_hybrid_ffn_graph( int n_tokens, MoeHybridGraphInputs & out, bool include_shared, - bool allow_fused_combine) { + bool allow_fused_combine, + MoeHybridJoinMode join_mode) { out.output = nullptr; out.main_output = nullptr; @@ -790,211 +988,48 @@ bool build_moe_hybrid_ffn_graph( return false; } - const int n_used = cfg.n_expert_used; + const bool canonical_route_join = + join_mode == MoeHybridJoinMode::CanonicalRouteOrder; // Both owner remaps consume the same normalized top-k route weights. - // Expose the canonical tensor so the heterogeneous scheduler can keep it - // on the main GPU. Otherwise expanding the cold branch first lets backend - // assignment migrate this shared dependency to Strix, forcing the hot - // R9700 branch to wait for a reverse peer copy before it can launch. + // Expose the canonical tensor so the scheduler can keep it on the primary + // backend rather than discovering it late through the secondary branch. out.router_weights = router_weights; - auto build_remap = [&](const std::vector & local_by_global, - ggml_tensor ** local_lut, - ggml_tensor ** valid_lut, - ggml_tensor ** local_ids, - ggml_tensor ** masked_weights, - std::vector * backend_nodes) { - auto track = [backend_nodes](ggml_tensor * tensor) { - if (tensor && backend_nodes) backend_nodes->push_back(tensor); - return tensor; - }; - if (!*local_lut) { - *local_lut = ggml_new_tensor_2d( - ctx, GGML_TYPE_I32, 1, cfg.n_expert); - ggml_set_input(*local_lut); - // These inputs are consumed late in a whole-model graph. Preserve - // their allocation from graph start; otherwise gallocr may reuse - // the tiny buffer as activation scratch before its layer executes. - ggml_set_output(*local_lut); - } - if (!*valid_lut) { - *valid_lut = ggml_new_tensor_2d( - ctx, GGML_TYPE_F32, 1, cfg.n_expert); - ggml_set_input(*valid_lut); - ggml_set_output(*valid_lut); - } - - // q1 can address the 2-D LUT directly. Historically q>1 left its tiny - // I32 repeat unpinned for CPU fallback. With the exact GPU I32 repeat - // enabled, track it with the rest of the owner-local remap nodes. - ggml_tensor * local_lut_batched = *local_lut; - if (n_tokens > 1) { - ggml_tensor * repeated = ggml_repeat_4d( - ctx, *local_lut, 1, cfg.n_expert, n_tokens, 1); - local_lut_batched = - gpu_i32_repeat_enabled() ? track(repeated) : repeated; - } - ggml_tensor * mapped = track(ggml_get_rows( - ctx, local_lut_batched, global_ids)); - mapped = track(ggml_reshape_2d(ctx, mapped, n_used, n_tokens)); - *local_ids = track(ggml_cont(ctx, mapped)); - ggml_tensor * valid_lut_batched = *valid_lut; - if (n_tokens > 1) { - valid_lut_batched = track(ggml_repeat_4d( - ctx, *valid_lut, 1, cfg.n_expert, n_tokens, 1)); - } - ggml_tensor * valid = track(ggml_get_rows( - ctx, valid_lut_batched, global_ids)); - valid = track(ggml_reshape_2d(ctx, valid, n_used, n_tokens)); - *masked_weights = track(ggml_mul(ctx, router_weights, valid)); - return (int)local_by_global.size() == cfg.n_expert; - }; - - ggml_tensor * hot_ids = nullptr; - ggml_tensor * hot_weights = nullptr; - if (!build_remap(storage.hot_local_by_global, - &out.hot_local_lut, &out.hot_valid_lut, - &hot_ids, &hot_weights, &out.hot_remap_nodes)) { + MoeOwnerGraphSpec primary_owner{ + &storage.hot_local_by_global, + storage.gate_hot, storage.up_hot, storage.down_hot, + storage.gate_up_hot, + &out.hot_local_lut, &out.hot_valid_lut, + &out.hot_remap_nodes, &out.hot_nodes}; + MoeOwnerGraphSpec secondary_owner{ + &storage.cold_local_by_global, + storage.gate_cold, storage.up_cold, storage.down_cold, + storage.gate_up_cold, + &out.cold_local_lut, &out.cold_valid_lut, + &out.cold_remap_nodes, &out.cold_nodes}; + + // Keep graph construction order stable: both remaps, then both optional ID + // alignments, then both expert branches. + if (!prepare_moe_owner_branch( + ctx, cfg, global_ids, router_weights, n_tokens, primary_owner) || + !prepare_moe_owner_branch( + ctx, cfg, global_ids, router_weights, n_tokens, secondary_owner)) { return false; } - - ggml_tensor * cold_ids = nullptr; - ggml_tensor * cold_weights = nullptr; - if (!build_remap(storage.cold_local_by_global, - &out.cold_local_lut, &out.cold_valid_lut, - &cold_ids, &cold_weights, &out.cold_remap_nodes)) { + align_moe_owner_routes(ctx, n_tokens, primary_owner); + align_moe_owner_routes(ctx, n_tokens, secondary_owner); + if (!build_moe_owner_branch( + ctx, cfg, desc, inp, n_tokens, canonical_route_join, + allow_fused_combine, primary_owner) || + !build_moe_owner_branch( + ctx, cfg, desc, inp, n_tokens, canonical_route_join, + allow_fused_combine, secondary_owner)) { return false; } - // q-token verification often routes adjacent tokens to the same expert - // at different top-k ranks. Align those owner-local IDs before MMVQ so - // equal weights are consumed by warps in one block. The encoded original - // route slot is decoded by the dedicated MoE kernel, which scatters every - // result back before the unchanged weighted reduction. - if (n_tokens > 1 && align_shared_moe_ids_enabled()) { - hot_ids = ggml_ds4_moe_align_ids(ctx, hot_ids); - cold_ids = ggml_ds4_moe_align_ids(ctx, cold_ids); - out.hot_remap_nodes.push_back(hot_ids); - out.cold_remap_nodes.push_back(cold_ids); - } - - ggml_tensor * hot = nullptr; - if ((storage.gate_up_hot || (storage.gate_hot && storage.up_hot)) && - storage.down_hot) { - const ggml_tensor * hot_gate = storage.gate_up_hot - ? storage.gate_up_hot : storage.gate_hot; - const bool tokenwise = - hot_gate->type == GGML_TYPE_Q2_0_ROCMFP2 && - !(n_tokens > 1 && grouped_mmvq_moe_enabled()); - if (!build_batched_routed_graph( - ctx, - storage.gate_hot, storage.up_hot, storage.down_hot, - storage.gate_up_hot, - desc.ffn_gate_exps_s, desc.ffn_up_exps_s, - desc.ffn_down_exps_s, desc.ffn_gate_up_exps_s, - inp, hot_ids, hot_weights, - cfg.n_embd, cfg.n_ff_exp, n_used, n_tokens, - cfg.swiglu_clamp, &hot, tokenwise, - &out.hot_nodes, allow_fused_combine)) { - return false; - } - } - - ggml_tensor * cold = nullptr; - if ((storage.gate_up_cold || (storage.gate_cold && storage.up_cold)) && - storage.down_cold) { - const ggml_tensor * cold_gate = storage.gate_up_cold - ? storage.gate_up_cold : storage.gate_cold; - const bool tokenwise = - cold_gate->type == GGML_TYPE_Q2_0_ROCMFP2 && - !(n_tokens > 1 && grouped_mmvq_moe_enabled()); - if (!build_batched_routed_graph( - ctx, - storage.gate_cold, storage.up_cold, storage.down_cold, - storage.gate_up_cold, - desc.ffn_gate_exps_s, desc.ffn_up_exps_s, - desc.ffn_down_exps_s, desc.ffn_gate_up_exps_s, - inp, cold_ids, cold_weights, - cfg.n_embd, cfg.n_ff_exp, - n_used, n_tokens, - cfg.swiglu_clamp, &cold, tokenwise, - &out.cold_nodes, allow_fused_combine)) { - return false; - } - } - - ggml_tensor * main_branch = hot; - ggml_tensor * shared = include_shared - ? build_shared_expert_subgraph(ctx, desc, inp, cfg.swiglu_clamp) - : nullptr; - if (shared) main_branch = main_branch ? ggml_add(ctx, main_branch, shared) : shared; - - // The generic scheduler copies every cross-backend input before launching - // any node in a split. If hot/shared and the final add are one contiguous - // main-backend split, the add's cold input makes that whole split wait for - // the peer, serializing the nominally parallel branches. - // - // Expand cold now, then visit hot/shared, then a new peer-owned CONT fence, - // and finally the main-backend add. This yields: - // peer cold compute -> main hot/shared compute -> peer fence -> main join - // The scheduler can enqueue cold first and hot/shared second; the fence - // separates the final join so its event wait is inserted after hot/shared. - ggml_tensor * combined = nullptr; - if (schedule_graph && cold && device_join_enabled()) { - // Materialize both route IDs and normalized route weights before the - // cold split is expanded. Cross-device copies are not guaranteed to - // remain asynchronous on this heterogeneous ROCm pair. If the tiny - // weight copy is discovered after cold expert execution, the fallback - // copy synchronizes the Strix stream and prevents the host from - // enqueueing independent R9700 hot work until cold has completed. - // - // q4 verification calls this builder twice (4 routes + padded 2), so - // retain every derived route tensor rather than only the canonical - // six-wide routing output. - if (route_prefork_enabled()) { - out.route_prefork_nodes.push_back(global_ids); - out.route_prefork_nodes.push_back(router_weights); - ggml_build_forward_expand(schedule_graph, global_ids); - ggml_build_forward_expand(schedule_graph, router_weights); - } - // Enforce fork order without adding another backend graph: - // cold owner -> hot/shared -> in-graph event wait/copy -> add. - // The deferred copy remains in the same main-backend split as the - // hot branch, so its wait is reached only after useful main work. - // Keep the peer result live explicitly. The generic allocator normally - // derives lifetime from children on the same execution backend; this - // custom foreign-buffer edge intentionally bypasses that copy path. - ggml_set_output(cold); - ggml_build_forward_expand(schedule_graph, cold); - if (main_branch) { - ggml_build_forward_expand(schedule_graph, main_branch); - } - ggml_tensor * cold_ready = - ggml_ds4_deferred_peer_copy(ctx, cold); - // The scheduler may prefill this tensor in the host-copy diagnostic - // before its containing main split launches. Reserve a stable buffer - // from graph start so earlier hot-branch scratch cannot alias it. - ggml_set_input(cold_ready); - ggml_set_output(cold_ready); - out.deferred_peer_copy_nodes.push_back(cold_ready); - out.main_output = main_branch; - out.peer_output = cold_ready; - combined = main_branch ? ggml_add(ctx, cold_ready, main_branch) - : cold_ready; - } else if (schedule_graph && cold) { - ggml_build_forward_expand(schedule_graph, cold); - ggml_tensor * cold_fence = ggml_cont(ctx, cold); - out.cold_nodes.push_back(cold_fence); - combined = main_branch ? ggml_add(ctx, main_branch, cold_fence) - : cold_fence; - } else { - // Preserve the established default dependency order exactly. - if (cold && main_branch) { - combined = ggml_add(ctx, cold, main_branch); - out.join_nodes.push_back(combined); - } else { - combined = cold ? cold : main_branch; - } - } + ggml_tensor * combined = build_moe_owner_join( + ctx, schedule_graph, cfg, desc, inp, global_ids, router_weights, + n_tokens, include_shared, canonical_route_join, + primary_owner.output, secondary_owner.output, out); if (!combined) return false; out.output = ggml_cont(ctx, combined); @@ -1238,7 +1273,7 @@ bool build_cached_hot_batched_graph( desc.ffn_gate_exps_s, desc.ffn_up_exps_s, desc.ffn_down_exps_s, desc.ffn_gate_up_exps_s, out.inp, out.sel, out.wts, n_embd, n_ff_exp, n_used, n_tokens, cfg.swiglu_clamp, &routed, false, nullptr, - ggml_backend_is_cuda(gpu_backend)); + backend_is_gpu(gpu_backend)); } // Shared expert (always on GPU) @@ -1300,7 +1335,7 @@ static bool build_cached_cold_batched_graph( desc.ffn_gate_exps_s, desc.ffn_up_exps_s, desc.ffn_down_exps_s, desc.ffn_gate_up_exps_s, out.inp, out.sel, out.wts, n_embd, n_ff_exp, n_used, n_tokens, cfg.swiglu_clamp, &routed, false, nullptr, - ggml_backend_is_cuda(cpu_backend)); + backend_is_gpu(cpu_backend)); if (!routed) { out.free(); return false; } out.output = routed; @@ -1820,11 +1855,11 @@ static bool eval_moe_hybrid_ffn_batched_core( auto set_cur_input = [&](ggml_tensor * dst, ggml_backend_t dst_backend) { if (cur_backend) { - // Use the backend-aware copy path. On heterogeneous HIP this + // Use the backend-aware copy path. On heterogeneous runtimes this // submits on the producer stream and publishes an event to the // consumer stream; the generic buffer copy uses a per-thread - // stream with no cross-device dependency and can leave the Strix - // destination containing zeros. + // stream with no cross-device dependency and can leave the + // secondary destination containing zeros. ggml_backend_tensor_copy_async( gpu_backend, dst_backend, cur_backend, dst); } else { @@ -2002,7 +2037,7 @@ static bool eval_moe_hybrid_ffn_batched_core( if (!logged) { std::fprintf(stderr, "[hybrid-ffn] masked cold prefill routes active; " - "Strix skips R9700-owned expert GEMMs " + "secondary owner skips primary-owned expert GEMMs " "max_expert_rows=%d\n", (int)max_cold_expert_rows); logged = true; @@ -2050,7 +2085,7 @@ static bool eval_moe_hybrid_ffn_batched_core( desc.ffn_gate_exps_s, desc.ffn_up_exps_s, desc.ffn_down_exps_s, desc.ffn_gate_up_exps_s, inp, sel, wts, n_embd, n_ff_exp, n_used, n_tokens, cfg.swiglu_clamp, &routed, false, nullptr, - ggml_backend_is_cuda(gpu_backend)); + backend_is_gpu(gpu_backend)); } // Shared expert (always on GPU) @@ -2163,7 +2198,7 @@ static bool eval_moe_hybrid_ffn_batched_core( desc.ffn_gate_exps_s, desc.ffn_up_exps_s, desc.ffn_down_exps_s, desc.ffn_gate_up_exps_s, inp, sel, wts, n_embd, n_ff_exp, n_used, n_tokens, cfg.swiglu_clamp, &cold_routed, false, nullptr, - ggml_backend_is_cuda(cold_backend), + backend_is_gpu(cold_backend), /*force_fused_combine=*/mask_skipped_cold); ggml_cgraph * cold_gf = ggml_new_graph_custom(cold_ctx, 4096, false); @@ -2458,8 +2493,8 @@ static bool eval_moe_owner_expert_major_batched( (size_t)counts[(size_t)e]; } std::vector cursor(offsets.begin(), offsets.end() - 1); - // Device input keeps the normalized activation on the R9700 and gathers - // the owner rows on-device. The host fallback remains for decode and A/B. + // Device input keeps the normalized activation on the primary owner and + // gathers rows on-device. The host fallback remains for decode and A/B. std::vector packed_input; if (!cur_backend) { packed_input.resize(n_pairs * (size_t)n_embd); @@ -2871,7 +2906,7 @@ bool eval_moe_hot_only_batched( desc.ffn_gate_exps_s, desc.ffn_up_exps_s, desc.ffn_down_exps_s, desc.ffn_gate_up_exps_s, inp, sel, wts, n_embd, n_ff_exp, n_used, n_tokens, cfg.swiglu_clamp, &routed, false, nullptr, - ggml_backend_is_cuda(gpu_backend)); + backend_is_gpu(gpu_backend)); // Shared expert (always on GPU) ggml_tensor * combined = routed; @@ -3047,11 +3082,11 @@ bool eval_moe_hybrid_ffn_batched( return true; } - // A prefill-only full cold stack keeps Strix on its proven full-stack - // MUL_MAT_ID path. Do not pair it with hundreds of q<=4 hot sub-batches: - // pack the R9700 routes by expert and launch one owner graph instead. The - // hot-local map takes priority in eval_moe_hybrid_ffn_batched_core(), so - // skip_hot makes the duplicated Strix stack evaluate cold routes only. + // A prefill-only full secondary stack keeps that owner on its qualified + // full-stack MUL_MAT_ID path. Do not pair it with hundreds of q<=4 primary + // sub-batches: pack the primary routes by expert and launch one owner graph + // instead. The primary-local map takes priority in the core evaluator, so + // skip_hot makes the duplicated secondary stack evaluate its routes only. const bool inprocess_full_cold_hot_expert_major = !expert_compute && expert_major_prefill_enabled(n_tokens) && cold_on_gpu && storage.cold_backend && @@ -3063,8 +3098,9 @@ bool eval_moe_hybrid_ffn_batched( static bool logged = false; if (!logged) { std::fprintf(stderr, - "[hybrid-ffn] full-cold Strix MMID batch + expert-major " - "R9700 prefill active tokens=%d hot_stack=%d\n", + "[hybrid-ffn] full-secondary MMID batch + " + "expert-major primary prefill active tokens=%d " + "primary_stack=%d\n", n_tokens, n_hot_stack); logged = true; } @@ -3178,12 +3214,12 @@ bool eval_moe_hybrid_ffn_batched( const bool remote_cold_full_batch = expert_compute && expert_layer && parse_moe_expert_compute_ipc_mode() == MoeExpertComputeIpcMode::Batched; - // The in-process heterogeneous path owns a nearly-full cold stack on - // Strix and a small hot stack on the R9700. A full reduced-stack MMQ - // is pathological for the 24-expert hot side, but is stable once the - // cold stack covers most experts and the prompt supplies enough rows. - // Run that one cold batch on Strix while retaining the proven q<=4 - // MMVQ slices for R9700 hot/shared work. + // The in-process heterogeneous path may own a nearly full secondary + // stack and a small primary stack. A full reduced-stack MMQ is + // pathological for the small side, but stable once the secondary + // covers most experts and the prompt supplies enough rows. Run that + // secondary batch while retaining q<=4 MMVQ slices for primary/shared + // work. const bool inprocess_cold_full_batch = !expert_compute && cold_on_gpu && storage.cold_backend && storage.cold_backend != gpu_backend && diff --git a/server/src/common/moe_hybrid_ffn_eval.h b/server/src/common/moe_hybrid_ffn_eval.h index 5c5bc5125..e5213b922 100644 --- a/server/src/common/moe_hybrid_ffn_eval.h +++ b/server/src/common/moe_hybrid_ffn_eval.h @@ -123,6 +123,9 @@ struct MoeHybridGraphInputs { // main->peer copy in the middle of cold execution, which synchronizes the // peer stream before the hot branch can be submitted. std::vector route_prefork_nodes; + // [1, n_expert, q, 1] immutable per-owner lookup rows. Keeping the q + // replicas in the input avoids per-step GPU REPEAT kernels and the split + // boundaries they introduce in a heterogeneous graph. ggml_tensor * hot_local_lut = nullptr; ggml_tensor * hot_valid_lut = nullptr; ggml_tensor * cold_local_lut = nullptr; @@ -149,14 +152,43 @@ struct MoeHybridGraphInputs { std::vector deferred_peer_copy_nodes; }; +enum class MoeHybridJoinMode { + // Reduce each owner's routes locally, then add the two partial sums. This + // minimizes transfer size and is the fast path for one GPU runtime. + OwnerPartialSums, + // Preserve the model's route order across owners and perform one final + // reduction on the main backend. Cross-runtime execution uses this mode + // to avoid changing floating-point association at the owner boundary. + CanonicalRouteOrder, +}; + +// Process-wide graph policy parsed once from the model-neutral environment +// variables. Legacy DS4 spellings remain accepted by the implementation, but +// graph builders and scheduler setup consume this typed view instead of +// independently re-reading configuration in hot paths. +struct MoeHybridGraphPolicy { + bool grouped_mmvq = false; + bool fused_combine = false; + bool fused_gate_up = false; + bool coarse_owner = false; + bool coarse_owner_split = false; + bool align_shared_ids = false; + bool device_join = false; + bool route_prefork = false; + bool targeted_join_split = false; +}; + +const MoeHybridGraphPolicy & moe_hybrid_graph_policy(); + // Append a device-resident hot+cold+shared MoE FFN to an existing graph. // `global_ids` and `router_weights` are [n_expert_used, n_tokens]. Weight // tensors in `storage` determine scheduler placement on the two GPU backends. -// When `schedule_graph` is non-null, the cold branch is expanded immediately -// and a peer-owned fence is inserted before the final main-backend join. This -// forces three scheduler splits (cold, hot/shared, join), preventing the join's -// cold-result copy from blocking hot/shared launch. Consumers may use -// main_output + peer_output to fuse the exact final add into their next op. +// When `schedule_graph` is non-null, the cold branch is expanded immediately. +// The default path inserts a peer-owned fence before the final main-backend +// join; targeted-join scheduling can instead mark the join itself as a fresh +// split. Both forms submit cold and hot/shared independently before gathering +// the peer result. Consumers may use main_output + peer_output to fuse the +// exact final add into their next op. bool build_moe_hybrid_ffn_graph( ggml_context * ctx, ggml_cgraph * schedule_graph, @@ -169,7 +201,9 @@ bool build_moe_hybrid_ffn_graph( int n_tokens, MoeHybridGraphInputs & out, bool include_shared = true, - bool allow_fused_combine = false); + bool allow_fused_combine = false, + MoeHybridJoinMode join_mode = + MoeHybridJoinMode::OwnerPartialSums); int moe_hybrid_expert_compute_batch_limit(); int moe_hybrid_expert_compute_ipc_batch_limit(int n_tokens); diff --git a/server/src/deepseek4/deepseek4_backend.cpp b/server/src/deepseek4/deepseek4_backend.cpp index e6702903b..75c0c0e08 100644 --- a/server/src/deepseek4/deepseek4_backend.cpp +++ b/server/src/deepseek4/deepseek4_backend.cpp @@ -2,6 +2,7 @@ #include "deepseek4_backend.h" #include "deepseek4_internal.h" +#include "common/dynamic_backend.h" #include "common/peer_access.h" #include "common/sampler.h" @@ -89,17 +90,71 @@ static void configure_gfx1201_hybrid_sub_batch_default(int gpu) { #endif } -static bool ds4_inprocess_moe_tp_enabled() { - return env_flag_enabled("DFLASH_DS4_MOE_TP_INPROC"); +struct Ds4MoeTpConfig { + bool requested = false; + bool in_process = false; + bool backend_valid = true; + PlacementBackend secondary_backend = PlacementBackend::Auto; + int secondary_gpu = 0; + bool all_on_secondary = false; + bool concentrate_secondary = false; + bool profile_hot_on_secondary = false; +}; + +static Ds4MoeTpConfig ds4_moe_tp_config(int local_gpu) { + Ds4MoeTpConfig result; + result.requested = env_flag_enabled("DFLASH_DS4_MOE_TP"); + result.in_process = result.requested && + env_flag_enabled("DFLASH_DS4_MOE_TP_INPROC"); + result.all_on_secondary = result.requested && + env_flag_enabled("DFLASH_DS4_MOE_TP_ALL_COLD"); + result.concentrate_secondary = result.requested && + env_flag_enabled("DFLASH_DS4_MOE_TP_CONCENTRATE_COLD"); + result.profile_hot_on_secondary = result.in_process && + env_flag_enabled("DFLASH_DS4_MOE_TP_PEER_HOT"); + + const char * raw = std::getenv("DFLASH_DS4_MOE_TP_BACKEND"); + if (!raw || !*raw) raw = std::getenv("DFLASH_MOE_TP_BACKEND"); + if (!raw || !*raw) { +#if defined(DFLASH27B_BACKEND_MIXED) + result.secondary_backend = + compiled_placement_backend() == PlacementBackend::Cuda + ? PlacementBackend::Hip : PlacementBackend::Cuda; +#else + result.secondary_backend = compiled_placement_backend(); +#endif + } else { + result.backend_valid = parse_placement_backend( + raw, result.secondary_backend) && + result.secondary_backend != PlacementBackend::Auto; + } + + const char * gpu_raw = std::getenv("DFLASH_DS4_MOE_TP_GPU"); + if (!gpu_raw || !*gpu_raw) { + gpu_raw = std::getenv("DFLASH_MOE_EXPERT_COMPUTE_IPC_GPU"); + } + if (gpu_raw && *gpu_raw) { + result.secondary_gpu = std::max(0, std::atoi(gpu_raw)); + } else if (result.backend_valid && + result.secondary_backend != compiled_placement_backend()) { + // CUDA and HIP have independent device namespaces. The first device + // in the peer runtime is therefore backend:0 even when the target is + // also device zero in its own runtime. + result.secondary_gpu = 0; + } else { + result.secondary_gpu = local_gpu == 0 ? 1 : 0; + } + return result; } -static int ds4_moe_tp_gpu(int local_gpu) { - const char * raw = std::getenv("DFLASH_DS4_MOE_TP_GPU"); +static bool ds4_draft_backend(PlacementBackend & out) { + const char * raw = std::getenv("DFLASH_DS4_DRAFT_BACKEND"); if (!raw || !*raw) { - raw = std::getenv("DFLASH_MOE_EXPERT_COMPUTE_IPC_GPU"); + out = compiled_placement_backend(); + return true; } - if (raw && *raw) return std::max(0, std::atoi(raw)); - return local_gpu == 0 ? 1 : 0; + return parse_placement_backend(raw, out) && + out != PlacementBackend::Auto; } static double gib(uint64_t bytes) { @@ -310,9 +365,45 @@ static void fill_prefix_hot_placement(const DeepSeek4Weights & w, } } +// Cross-runtime joins are much more expensive than native peer handoffs. Keep +// approximately the same expert residency as the uniform placement, but +// concentrate the cold owner into complete layers. A partial cold layer costs +// another cross-runtime join and some CUDA prefill paths require a complete +// expert stack, so retain the small remainder on the target backend. +static int fill_concentrated_cold_placement(const DeepSeek4Weights & w, + int hot_per_layer, + MoeHybridPlacement & out) { + out = {}; + out.n_layer = w.n_layer; + out.n_expert = w.n_expert; + out.n_expert_used = w.n_expert_used; + out.hot_counts.assign((size_t) w.n_layer, w.n_expert); + out.hot_expert_ids.resize((size_t) w.n_layer); + + const int requested_cold = + w.n_layer * std::max(0, w.n_expert - hot_per_layer); + int cold_remaining = w.n_expert > 0 + ? requested_cold / w.n_expert * w.n_expert : 0; + const int retained_local = requested_cold - cold_remaining; + for (int il = w.n_layer - 1; il >= 0; --il) { + const int cold = std::min(w.n_expert, cold_remaining); + const int hot = w.n_expert - cold; + out.hot_counts[(size_t) il] = hot; + auto & ids = out.hot_expert_ids[(size_t) il]; + ids.reserve((size_t) hot); + for (int ie = 0; ie < hot; ++ie) { + ids.push_back((int32_t) ie); + } + out.total_hot += hot; + cold_remaining -= cold; + } + return retained_local; +} + static bool fill_profiled_hot_placement(const DeepSeek4Weights & w, int hot_per_layer, const char * profile_path, + bool profile_hot_on_secondary, MoeHybridPlacement & out, std::string * err) { MoeHybridRoutingStats stats; @@ -334,9 +425,41 @@ static bool fill_profiled_hot_placement(const DeepSeek4Weights & w, out.hot_expert_ids.resize((size_t)w.n_layer); out.total_hot = hot_per_layer * w.n_layer; for (int il = 0; il < w.n_layer; ++il) { - std::vector ranked = stats.hot_experts(il, hot_per_layer); auto & ids = out.hot_expert_ids[(size_t)il]; - ids.assign(ranked.begin(), ranked.end()); + if (!profile_hot_on_secondary) { + std::vector ranked = stats.hot_experts(il, hot_per_layer); + ids.assign(ranked.begin(), ranked.end()); + continue; + } + + // `hot` is the primary-backend side of MoeHybridPlacement. On a + // memory-rich iGPU paired with a smaller, faster dGPU, filling that + // primary side with the most frequently routed experts starves the + // dGPU of useful work. Reserve the peer-sized complement for the + // hottest experts and keep every other expert on the primary. This + // changes ownership only; route order and reduction semantics stay + // unchanged. + const int peer_count = w.n_expert - hot_per_layer; + const std::vector ranked_peer = + stats.hot_experts(il, peer_count); + std::vector on_peer((size_t)w.n_expert, 0); + for (int expert : ranked_peer) { + if (expert >= 0 && expert < w.n_expert) { + on_peer[(size_t)expert] = 1; + } + } + ids.reserve((size_t)hot_per_layer); + for (int expert = 0; expert < w.n_expert; ++expert) { + if (!on_peer[(size_t)expert]) { + ids.push_back((int32_t)expert); + } + } + if ((int)ids.size() != hot_per_layer) { + if (err) { + *err = "routing profile did not yield a complete expert ranking"; + } + return false; + } } return true; } @@ -345,18 +468,23 @@ static bool fill_profiled_hot_placement(const DeepSeek4Weights & w, // but distribute those slots across layers to minimize the predicted owner // critical path. Uniform expert counts are a poor fit for heterogeneous EP: // routing skew varies substantially by layer, while every layer joins on the -// slower of its R9700 hot/shared and Strix cold branches. +// slower of its primary/shared and secondary expert branches. // // The cost model intentionally uses measured bandwidth rather than advertised // peak bandwidth. It is only an allocation objective; actual placement still // uses authoritative router statistics and evaluates every selected expert. static bool compute_ds4_hybrid_budget_info(const DeepSeek4Weights & w, - int gpu, + ggml_backend_t backend, int max_ctx, Ds4HybridBudgetInfo & out, std::string * err) { out = {}; - ggml_backend_cuda_get_device_memory(gpu, &out.gpu_free, &out.gpu_total); + if (!backend || !ggml_backend_get_device(backend)) { + if (err) *err = "target backend has no device"; + return false; + } + ggml_backend_dev_memory( + ggml_backend_get_device(backend), &out.gpu_free, &out.gpu_total); if (out.gpu_total == 0) { if (err) *err = "could not query GPU memory"; return false; @@ -472,7 +600,7 @@ bool DeepSeek4Backend::load_model() { // Fused decode and layer-major prefill normally require monolithic expert // residency. Heterogeneous TP is the exception: its fused graph owns the - // routed experts across two HIP backends, so forcing a full load would + // routed experts across two local GPU backends, so forcing a full load would // disable the requested split before the TP runtime can initialize. const bool force_full = env_flag_enabled("DFLASH_DS4_FORCE_FULL_LOAD"); const bool heterogeneous_tp = env_flag_enabled("DFLASH_DS4_MOE_TP"); @@ -498,9 +626,9 @@ bool DeepSeek4Backend::load_model() { cfg_.model_path); return false; } - } else if (target_backend == PlacementBackend::Hip) { + } else if (target_backend == PlacementBackend::Hip || heterogeneous_tp) { std::fprintf(stderr, - "[deepseek4] HIP target detected; using hybrid expert load path\n"); + "[deepseek4] heterogeneous target detected; using hybrid expert load path\n"); if (!init_hybrid_model()) { std::fprintf(stderr, "[deepseek4] hybrid mode failed: %s\n", cfg_.model_path); return false; @@ -544,23 +672,37 @@ bool DeepSeek4Backend::load_spec_drafter() { } const bool separate_draft_stream = env_flag_enabled("DFLASH_DS4_DRAFT_SEPARATE_STREAM"); - if (draft_gpu != cfg_.device.gpu || separate_draft_stream) { - spec_backend_ = ggml_backend_cuda_init(draft_gpu); + PlacementBackend draft_kind = PlacementBackend::Auto; + if (!ds4_draft_backend(draft_kind)) { + std::fprintf(stderr, + "[deepseek4] invalid DFLASH_DS4_DRAFT_BACKEND; " + "expected cuda or hip\n"); + return false; + } + const PlacementBackend target_kind = placement_backend_of(backend_); + if (draft_kind != target_kind || draft_gpu != cfg_.device.gpu || + separate_draft_stream) { + std::string backend_error; + spec_backend_ = init_placement_backend( + draft_kind, draft_gpu, &backend_error); if (!spec_backend_) { std::fprintf(stderr, - "[deepseek4] failed to initialize DSpark GPU %d\n", - draft_gpu); + "[deepseek4] failed to initialize DSpark %s:%d: %s\n", + placement_backend_name(draft_kind), draft_gpu, + backend_error.c_str()); return false; } draft_backend = spec_backend_; const bool low_priority = separate_draft_stream && env_flag_enabled("DFLASH_DS4_DRAFT_LOW_PRIORITY"); const bool priority_configured = low_priority && + backend_pair_capabilities(backend_, spec_backend_).same_runtime && ggml_backend_cuda_set_low_priority_stream(spec_backend_); std::fprintf(stderr, - "[deepseek4] DSpark backend gpu=%d target_gpu=%d " + "[deepseek4] DSpark backend=%s:%d target=%s:%d " "separate_stream=%d low_priority=%d\n", - draft_gpu, cfg_.device.gpu, + placement_backend_name(draft_kind), draft_gpu, + placement_backend_name(target_kind), cfg_.device.gpu, (int) separate_draft_stream, (int) priority_configured); } @@ -727,7 +869,8 @@ bool DeepSeek4Backend::init_moe_tensor_parallel() { return false; } - if (ds4_inprocess_moe_tp_enabled()) { + const Ds4MoeTpConfig tp = ds4_moe_tp_config(cfg_.device.gpu); + if (tp.in_process) { if (!expert_backend_ || !moe_hybrid_->materialized_cold_experts || moe_hybrid_->cold_backend != expert_backend_) { std::fprintf(stderr, @@ -735,10 +878,16 @@ bool DeepSeek4Backend::init_moe_tensor_parallel() { return false; } expert_runtime_.reset(); + const PlacementBackend local_kind = + cfg_.device.backend == PlacementBackend::Auto + ? compiled_placement_backend() : cfg_.device.backend; std::fprintf(stderr, - "[deepseek4-moe-tp] enabled mode=in-process local_gpu=%d " - "expert_gpu=%d local_experts=%d remote_experts=%d\n", - cfg_.device.gpu, ds4_moe_tp_gpu(cfg_.device.gpu), + "[deepseek4-moe-tp] enabled mode=in-process local=%s:%d " + "secondary=%s:%d primary_experts=%d " + "secondary_experts=%d\n", + placement_backend_name(local_kind), cfg_.device.gpu, + placement_backend_name(tp.secondary_backend), + tp.secondary_gpu, moe_placement_.total_hot, w_.n_layer * w_.n_expert - moe_placement_.total_hot); return true; @@ -780,25 +929,46 @@ bool DeepSeek4Backend::compute_uniform_hybrid_placement(const DeepSeek4Weights & MoeHybridPlacement & out, std::string * err) const { Ds4HybridBudgetInfo budget; - if (!compute_ds4_hybrid_budget_info(w, cfg_.device.gpu, max_ctx, budget, err)) { + if (!compute_ds4_hybrid_budget_info(w, backend_, max_ctx, budget, err)) { return false; } - const bool all_cold = env_flag_enabled("DFLASH_DS4_MOE_TP_ALL_COLD"); - int hot_per_layer = all_cold ? 0 : budget.max_hot_per_layer; - if (all_cold) { + const Ds4MoeTpConfig tp = ds4_moe_tp_config(cfg_.device.gpu); + int hot_per_layer = tp.all_on_secondary ? 0 : budget.max_hot_per_layer; + if (tp.all_on_secondary) { std::fprintf(stderr, - "[deepseek4-moe-tp] all routed experts assigned to the cold backend\n"); - } - if (const char * profile_path = std::getenv("DFLASH_DS4_HOTNESS_CSV")) { + "[deepseek4-moe-tp] all routed experts assigned to the " + "secondary backend\n"); + } + const bool concentrate_requested = tp.concentrate_secondary; + bool concentrated = false; + int retained_local = 0; + const int requested_cold = + w.n_layer * std::max(0, w.n_expert - hot_per_layer); + if (concentrate_requested && requested_cold >= w.n_expert) { + retained_local = + fill_concentrated_cold_placement(w, hot_per_layer, out); + concentrated = true; + } else if (concentrate_requested) { + std::fprintf(stderr, + "[deepseek4] concentrated secondary placement needs at least " + "one complete layer; using uniform placement\n"); + fill_prefix_hot_placement(w, hot_per_layer, out); + } else if (const char * profile_path = std::getenv("DFLASH_DS4_HOTNESS_CSV")) { if (*profile_path) { + const bool profile_hot_on_secondary = + tp.in_process && tp.profile_hot_on_secondary; if (!fill_profiled_hot_placement( - w, hot_per_layer, profile_path, out, err)) { + w, hot_per_layer, profile_path, + profile_hot_on_secondary, + out, err)) { return false; } std::fprintf(stderr, - "[deepseek4] hybrid placement profile=%s\n", - profile_path); + "[deepseek4] hybrid placement profile=%s%s\n", + profile_path, + profile_hot_on_secondary + ? " profile-hot-owner=secondary" : ""); } else { fill_prefix_hot_placement(w, hot_per_layer, out); } @@ -810,6 +980,28 @@ bool DeepSeek4Backend::compute_uniform_hybrid_placement(const DeepSeek4Weights & if (!compute_ds4_expert_memory_info(w, &out, placed_mem, err)) { return false; } + if (concentrated && placed_mem.hot_bytes > budget.expert_budget) { + std::fprintf(stderr, + "[deepseek4] concentrated secondary placement exceeds the " + "primary expert budget; using uniform placement\n"); + fill_prefix_hot_placement(w, hot_per_layer, out); + if (!compute_ds4_expert_memory_info(w, &out, placed_mem, err)) { + return false; + } + concentrated = false; + } + if (concentrated) { + const int cold_layers = + w.n_expert > 0 + ? (w.n_layer * w.n_expert - out.total_hot) / w.n_expert : 0; + std::fprintf(stderr, + "[deepseek4] concentrated secondary placement: " + "cross-owner layers=%d primary_experts=%d " + "secondary_experts=%d retained_primary=%d\n", + cold_layers, out.total_hot, + w.n_layer * w.n_expert - out.total_hot, + retained_local); + } std::fprintf(stderr, "[deepseek4] hybrid placement: gpu_total=%.2f GiB gpu_free=%.2f GiB core=%.2f GiB kv=%.2f GiB warm=%.2f GiB safety=%.2f GiB expert_budget=%.2f GiB hot/layer=%d\n", @@ -853,26 +1045,47 @@ bool DeepSeek4Backend::init_hybrid_model() { auto hybrid = std::make_shared(); MoeHybridConfig hybrid_cfg = make_ds4_parent_worker_cfg(w_); - const bool inprocess_tp = - env_flag_enabled("DFLASH_DS4_MOE_TP") && ds4_inprocess_moe_tp_enabled(); + const Ds4MoeTpConfig tp = ds4_moe_tp_config(cfg_.device.gpu); + const bool inprocess_tp = tp.requested && tp.in_process; if (inprocess_tp) { - const int expert_gpu = ds4_moe_tp_gpu(cfg_.device.gpu); - if (expert_gpu == cfg_.device.gpu) { + const int expert_gpu = tp.secondary_gpu; + const PlacementBackend expert_kind = tp.secondary_backend; + if (!tp.backend_valid) { + std::fprintf(stderr, + "[deepseek4-moe-tp] invalid DFLASH_DS4_MOE_TP_BACKEND; " + "expected cuda or hip\n"); + return false; + } + const PlacementBackend local_kind = + cfg_.device.backend == PlacementBackend::Auto + ? compiled_placement_backend() : cfg_.device.backend; + if (expert_kind == local_kind && expert_gpu == cfg_.device.gpu) { std::fprintf(stderr, - "[deepseek4-moe-tp] in-process expert GPU must differ from local GPU\n"); + "[deepseek4-moe-tp] in-process secondary device must " + "differ from the primary device\n"); return false; } - if (g_peer_access_opt_in) { + if (expert_kind == local_kind && g_peer_access_opt_in) { const bool peer_ok = enable_peer_access_pair(cfg_.device.gpu, expert_gpu); std::fprintf(stderr, - "[deepseek4-moe-tp] peer access GPU %d <-> GPU %d: %s\n", - cfg_.device.gpu, expert_gpu, peer_ok ? "enabled" : "unavailable"); + "[deepseek4-moe-tp] peer access %s:%d <-> %s:%d: %s\n", + placement_backend_name(local_kind), cfg_.device.gpu, + placement_backend_name(expert_kind), expert_gpu, + peer_ok ? "enabled" : "unavailable"); + } else if (expert_kind != local_kind) { + std::fprintf(stderr, + "[deepseek4-moe-tp] cross-vendor owner join %s:%d <-> %s:%d " + "uses in-process host staging\n", + placement_backend_name(local_kind), cfg_.device.gpu, + placement_backend_name(expert_kind), expert_gpu); } - expert_backend_ = ggml_backend_cuda_init(expert_gpu); + expert_backend_ = init_placement_backend(expert_kind, expert_gpu, &err); if (!expert_backend_) { std::fprintf(stderr, - "[deepseek4-moe-tp] failed to initialize in-process expert GPU %d\n", - expert_gpu); + "[deepseek4-moe-tp] failed to initialize in-process " + "secondary backend %s:%d: %s\n", + placement_backend_name(expert_kind), expert_gpu, + err.c_str()); return false; } hybrid_cfg.materialize_cold_experts = true; diff --git a/server/src/deepseek4/deepseek4_dspark_spec.cpp b/server/src/deepseek4/deepseek4_dspark_spec.cpp index 198301062..21596aa6c 100644 --- a/server/src/deepseek4/deepseek4_dspark_spec.cpp +++ b/server/src/deepseek4/deepseek4_dspark_spec.cpp @@ -67,15 +67,18 @@ class DeepSeek4DFlashTarget : public DFlashTarget { n, n > 0 ? tokens[0] : -1, n > 1 ? tokens[1] : -1, w_.n_vocab); return false; } - // Sequential verify (measurement mode): q single-token forwards through - // the legacy AR decode path. Causal by construction, compressor fed every - // token. Slow; used to measure the drafter's token-at-a-time accept rate. - // It is not a bit-exact oracle: graph shape can change floating-point - // reduction order around near-tied logits. Enable: DFLASH_DS4_SEQ_VERIFY=1 - // (pair with DFLASH_DS4_FULL_SNAP=1 so rollback/replay stay exact). + // Sequential verify: q single-token forwards through the same cached + // graph as ordinary AR decode. This preserves target arithmetic; exact + // rollback still requires a full snapshot and replay after rejection. + // DFLASH_DS4_SEQ_VERIFY is a diagnostic. The supported reference mode, + // DFLASH_DS4_SPEC_REFERENCE_EXACT, enables both requirements together. static const bool seq_verify = [] { - const char * v = std::getenv("DFLASH_DS4_SEQ_VERIFY"); - return v && *v && *v != '0'; + const char * exact = + std::getenv("DFLASH_DS4_SPEC_REFERENCE_EXACT"); + const char * sequential = + std::getenv("DFLASH_DS4_SEQ_VERIFY"); + return (exact && *exact && *exact != '0') || + (sequential && *sequential && *sequential != '0'); }(); if (seq_verify) { std::vector am_all; @@ -91,7 +94,7 @@ class DeepSeek4DFlashTarget : public DFlashTarget { tokens.data() + t, 1, base_pos + t, am1, keep_logits_ ? &logits1 : nullptr, feat1, telemetry_, - /*allow_graph_reuse=*/false, + /*allow_graph_reuse=*/true, moe_hybrid_, expert_runtime_, routing_stats_)) { return false; @@ -109,13 +112,13 @@ class DeepSeek4DFlashTarget : public DFlashTarget { return true; } std::vector am; - // n==1 must take the dynamic (non-reuse) path: the reused decode graph - // skips the capture/all-logits hooks (backend HC), which this needs. + // Reuse the normal cached graph for q==1 so reference verification has + // exactly the same target arithmetic as ordinary AR decode. if (!deepseek4_dspark_verify_forward(backend_, device_, w_, cache_, capture_ids_, embed_buf_.data(), tokens.data(), n, base_pos, am, keep_logits_ ? &verify_logits_ : nullptr, verify_features_, telemetry_, - /*allow_graph_reuse=*/n > 1, + /*allow_graph_reuse=*/true, moe_hybrid_, expert_runtime_, routing_stats_)) { return false; @@ -607,6 +610,13 @@ bool deepseek4_dspark_verify_forward(ggml_backend_t backend, argmax_out = std::move(gpu_argmax); return true; } + if (n_tokens == 1 && (int) all_logits.size() < w.n_vocab && + (int) last_logits.size() >= w.n_vocab) { + // The reference-exact q1 path reuses the normal AR graph. That graph + // returns its logits through the regular output vector rather than + // the verifier's q-wide hook. + all_logits = last_logits; + } if ((int) all_logits.size() < w.n_vocab * n_tokens) { std::fprintf(stderr, "[ds4-verify] all_logits too small: got=%zu need=%d (cap=%zu)\n", all_logits.size(), w.n_vocab * n_tokens, capture_out.size()); @@ -648,14 +658,23 @@ bool run_deepseek4_dspark_spec_decode( const bool debug = spec_env_flag("DFLASH_DS4_DSPARK_DEBUG"); const bool timing = spec_env_flag("DFLASH_DS4_TIMING"); - const bool full_snap = spec_env_flag("DFLASH_DS4_FULL_SNAP"); - const bool seq_verify_mode = spec_env_flag("DFLASH_DS4_SEQ_VERIFY"); + const bool reference_exact = + spec_env_flag("DFLASH_DS4_SPEC_REFERENCE_EXACT"); + const bool full_snap = reference_exact || + spec_env_flag("DFLASH_DS4_FULL_SNAP"); + const bool seq_verify_mode = reference_exact || + spec_env_flag("DFLASH_DS4_SEQ_VERIFY"); const bool async_rollback = spec_env_flag("DFLASH_DS4_ASYNC_ROLLBACK"); const bool pinned_rollback = spec_env_flag("DFLASH_DS4_PINNED_ROLLBACK"); const bool draft_overlap_probe = spec_env_flag("DFLASH_DS4_DRAFT_OVERLAP_PROBE"); const bool draft_overlap_reuse_context = spec_env_flag("DFLASH_DS4_DRAFT_OVERLAP_REUSE_CONTEXT"); + if (reference_exact) { + std::fprintf(stderr, + "[ds4-spec] reference-exact verifier: sequential target replay " + "with full rollback snapshots\n"); + } ggml_backend_t drafter_backend = drafter.core.backend ? drafter.core.backend : backend; const bool draft_overlap_probe_active = @@ -1001,7 +1020,7 @@ bool run_deepseek4_dspark_spec_decode( // The bonus token is DEFERRED: it becomes the next step's seed, whose // KV is written then. t0 = SpecClock::now(); - if (full_snap) { + if (full_snap && accept < q) { // Legacy: full restore + replay the committed tokens through the // target so ring/compressor/n_comp advance exactly. std::vector kv_toks; @@ -1019,7 +1038,7 @@ bool run_deepseek4_dspark_spec_decode( ok = false; break; } - } else if (accept < q) { + } else if (!full_snap && accept < q) { // The prev-half flush is bad only if the boundary sits at-or-past // the commit point (its chunk then contains rejected tokens). const bool restore_prev = boundary_crossed && first_boundary >= commit_pos; diff --git a/server/src/deepseek4/deepseek4_fused_verify.inc b/server/src/deepseek4/deepseek4_fused_verify.inc index b6aebc6f1..046a23bbc 100644 --- a/server/src/deepseek4/deepseek4_fused_verify.inc +++ b/server/src/deepseek4/deepseek4_fused_verify.inc @@ -13,6 +13,71 @@ static void ds4_fv_set(ggml_tensor * t, const void * data, size_t nbytes) { if (t && t->buffer) ggml_backend_tensor_set(t, data, 0, nbytes); } + +struct Ds4MixedMoePolicy { + bool owner_local_reduction = false; + bool direct_device_join = false; + bool schedule_branches = false; + bool native_route_width = false; + bool fused_hc_join = false; + bool late_join_split = false; + bool deferred_join_split = false; + bool batch_peer_copies = false; + bool report_split_count = false; + bool scheduler_trace = false; + bool pin_route_weights = false; + bool preserve_routes_for_diagnostics = false; +}; + +static const Ds4MixedMoePolicy & ds4_mixed_moe_policy() { + static const Ds4MixedMoePolicy policy = [] { + Ds4MixedMoePolicy result; + result.owner_local_reduction = + ds4_env_flag("DFLASH_DS4_CROSS_VENDOR_OWNER_SUMS"); + result.direct_device_join = + ds4_env_flag("DFLASH_DS4_TP_PEER_FENCE") || + ds4_env_flag("DFLASH_DS4_TP_DEVICE_JOIN"); + result.schedule_branches = + ds4_env_flag("DFLASH_DS4_TP_SCHEDULE_BRANCHES"); + result.native_route_width = + ds4_env_flag("DFLASH_DS4_TP_NATIVE_ROUTE_WIDTH"); + result.fused_hc_join = + ds4_env_flag("DFLASH_DS4_TP_FUSED_HC_JOIN"); + result.late_join_split = + ds4_env_flag("DFLASH_DS4_TP_LATE_JOIN_SPLIT"); + result.deferred_join_split = + ds4_env_flag("DFLASH_DS4_TP_DEVICE_JOIN_SPLIT"); + result.batch_peer_copies = + ds4_env_flag("GGML_BATCH_PEER_COPIES") || + ds4_env_flag("GGML_CUDA_BATCH_PEER_COPIES"); + result.report_split_count = + ds4_env_flag("DFLASH_DS4_TP_SPLIT_COUNT"); + result.scheduler_trace = + ds4_env_flag("DFLASH_DS4_TP_SCHED_TRACE"); + result.pin_route_weights = + ds4_env_flag("DFLASH_DS4_TP_MAIN_ROUTE_WEIGHTS"); + result.preserve_routes_for_diagnostics = + ds4_env_flag("DFLASH_DS4_TP_ROUTE_STATS") || + ds4_env_flag("DFLASH_DS4_ROUTING_STATS_OUT") || + ds4_env_flag("DFLASH_DS4_TP_CACHE_AUDIT"); + return result; + }(); + return policy; +} + +template +static void ds4_fv_set_repeated_rows( + ggml_tensor * tensor, + const std::vector & row) { + if (!tensor || row.empty()) return; + const size_t elements = (size_t) ggml_nelements(tensor); + GGML_ASSERT(elements % row.size() == 0); + std::vector repeated(elements); + for (size_t offset = 0; offset < elements; offset += row.size()) { + std::copy(row.begin(), row.end(), repeated.begin() + offset); + } + ds4_fv_set(tensor, repeated.data(), repeated.size() * sizeof(T)); +} static bool ds4_fused_verify_enabled() { static int enabled = -1; if (enabled < 0) { @@ -117,14 +182,10 @@ static void ds4_fused_verify_refresh_hybrid_luts( cold_valid[(size_t) ie] = 1.0f; } } - ds4_fv_set(inputs.hot_local_lut, hot_lut.data(), - sizeof(int32_t) * hot_lut.size()); - ds4_fv_set(inputs.hot_valid_lut, hot_valid.data(), - sizeof(float) * hot_valid.size()); - ds4_fv_set(inputs.cold_local_lut, cold_lut.data(), - sizeof(int32_t) * cold_lut.size()); - ds4_fv_set(inputs.cold_valid_lut, cold_valid.data(), - sizeof(float) * cold_valid.size()); + ds4_fv_set_repeated_rows(inputs.hot_local_lut, hot_lut); + ds4_fv_set_repeated_rows(inputs.hot_valid_lut, hot_valid); + ds4_fv_set_repeated_rows(inputs.cold_local_lut, cold_lut); + ds4_fv_set_repeated_rows(inputs.cold_valid_lut, cold_valid); } } @@ -297,6 +358,10 @@ static bool ds4_build_fused_verify_graph( } const int n_embd = w.n_embd; const int n_hc = w.n_hc; + const BackendPairCapabilities pair_capabilities = hybrid + ? backend_pair_capabilities(backend, hybrid->cold_backend) + : BackendPairCapabilities{true, true}; + const bool same_gpu_runtime = pair_capabilities.same_runtime; const size_t arena_size = 256u * 1024 * 1024; if (fg.sg.meta_arena.size() < arena_size) fg.sg.meta_arena.resize(arena_size); @@ -517,9 +582,34 @@ static bool ds4_build_fused_verify_graph( } if (hybrid) { + const Ds4MixedMoePolicy & mixed_policy = + ds4_mixed_moe_policy(); + const bool cross_vendor_owner_sums = + !same_gpu_runtime && + mixed_policy.owner_local_reduction; const bool allow_fused_combine = - ggml_backend_is_cuda(backend) && - ggml_backend_is_cuda(hybrid->cold_backend); + same_gpu_runtime || cross_vendor_owner_sums; + const bool allow_direct_device_join = + pair_capabilities.native_gpu_handoff && + mixed_policy.direct_device_join; + const bool schedule_host_staged_branches = + !same_gpu_runtime && + mixed_policy.schedule_branches; + ggml_cgraph * hybrid_schedule_graph = + (allow_direct_device_join || schedule_host_staged_branches) + ? gf : nullptr; + // Cross-vendor staging normally preserves canonical route order + // so it matches the single-owner reduction as closely as + // possible. For performance qualification, allow each owner to + // reduce its routed experts locally and stage one [n_embd, q] + // partial instead of the full [n_embd, n_used, q] route tensor. + // This is the same join shape used by the qualified same-runtime + // heterogeneous path, but remains opt-in until exact-output and + // mixed-vendor burn-in have passed. + const MoeHybridJoinMode join_mode = + same_gpu_runtime || cross_vendor_owner_sums + ? MoeHybridJoinMode::OwnerPartialSums + : MoeHybridJoinMode::CanonicalRouteOrder; MoeHybridConfig hybrid_cfg = make_ds4_moe_hybrid_config(w); hybrid_cfg.n_expert_used = (int) selected->ne[0]; MoeLayerDesc desc = make_ds4_moe_layer_desc(L); @@ -540,7 +630,8 @@ static bool ds4_build_fused_verify_graph( // six routes in one owner graph: the legacy 4 + padded-2 split // computes two duplicate expert paths whose weights are zero. const bool native_route_width = - ds4_env_flag("DFLASH_DS4_TP_NATIVE_ROUTE_WIDTH"); + !same_gpu_runtime || + mixed_policy.native_route_width; if (lane_q > 1 && hybrid_cfg.n_expert_used > 4 && has_hot_routed && !native_route_width) { const int route_width = hybrid_cfg.n_expert_used; @@ -584,41 +675,39 @@ static bool ds4_build_fused_verify_graph( } if (!build_moe_hybrid_ffn_graph( ctx, - (ds4_env_flag("DFLASH_DS4_TP_PEER_FENCE") || - ds4_env_flag("DFLASH_DS4_TP_DEVICE_JOIN")) ? gf : nullptr, + hybrid_schedule_graph, hybrid_cfg, desc, hybrid->layers[(size_t) il], ffn_normed, first_ids, first_weights, lane_q, inputs, true, - allow_fused_combine)) { + allow_fused_combine, join_mode)) { return false; } ffn_out = inputs.output; if (!build_moe_hybrid_ffn_graph( ctx, - (ds4_env_flag("DFLASH_DS4_TP_PEER_FENCE") || - ds4_env_flag("DFLASH_DS4_TP_DEVICE_JOIN")) ? gf : nullptr, + hybrid_schedule_graph, hybrid_cfg, desc, hybrid->layers[(size_t) il], ffn_normed, padded_ids, padded_weights, lane_q, inputs, false, - allow_fused_combine)) { + allow_fused_combine, join_mode)) { return false; } ffn_out = ggml_add(ctx, ffn_out, inputs.output); } else { if (!build_moe_hybrid_ffn_graph( ctx, - (ds4_env_flag("DFLASH_DS4_TP_PEER_FENCE") || - ds4_env_flag("DFLASH_DS4_TP_DEVICE_JOIN")) ? gf : nullptr, + hybrid_schedule_graph, hybrid_cfg, desc, hybrid->layers[(size_t) il], ffn_normed, selected, router_weights, - lane_q, inputs, true, allow_fused_combine)) { + lane_q, inputs, true, allow_fused_combine, + join_mode)) { return false; } ffn_out = inputs.output; if (lane_q > 1 && native_route_width && - ds4_env_flag("DFLASH_DS4_TP_FUSED_HC_JOIN") && + mixed_policy.fused_hc_join && inputs.main_output && inputs.peer_output) { fused_hc_join_inputs = &inputs; } @@ -753,6 +842,17 @@ static bool ds4_build_fused_verify_graph( "[ds4-fused-verify] scheduler CPU fallback unavailable\n"); return false; } + const Ds4MixedMoePolicy & mixed_policy = + ds4_mixed_moe_policy(); + if (!same_gpu_runtime && mixed_policy.direct_device_join) { + static bool warned_cross_vendor_join = false; + if (!warned_cross_vendor_join) { + std::fprintf(stderr, + "[ds4-fused-verify] direct peer join disabled across GPU " + "vendors; using scheduler host staging\n"); + warned_cross_vendor_join = true; + } + } ggml_backend_t backends[3] = { backend, peer, hybrid->cpu_backend}; fg.sched = ggml_backend_sched_new( @@ -762,25 +862,21 @@ static bool ds4_build_fused_verify_graph( "[ds4-fused-verify] scheduler creation failed\n"); return false; } - const bool late_join_split = - ds4_env_flag("DFLASH_DS4_TP_LATE_JOIN_SPLIT"); + const bool late_join_split = mixed_policy.late_join_split; + const MoeHybridGraphPolicy & moe_policy = + moe_hybrid_graph_policy(); const bool targeted_join_split = - ds4_env_flag("DFLASH_DS4_TP_TARGETED_JOIN_SPLIT"); - const bool device_join_split = - ds4_env_flag("DFLASH_DS4_TP_DEVICE_JOIN_SPLIT"); - const bool batch_split_copies = - ds4_env_flag("GGML_CUDA_BATCH_PEER_COPIES"); - const bool gpu_i32_repeat = - ds4_env_flag("LUCE_CUDA_I32_REPEAT"); - const bool report_split_count = - ds4_env_flag("DFLASH_DS4_TP_SPLIT_COUNT"); + moe_policy.targeted_join_split; + const bool device_join_split = mixed_policy.deferred_join_split; + const bool batch_split_copies = mixed_policy.batch_peer_copies; + const bool report_split_count = mixed_policy.report_split_count; ggml_backend_sched_set_late_cross_input_split( fg.sched, late_join_split); ggml_backend_sched_set_deferred_peer_copy_split( fg.sched, device_join_split); ggml_backend_sched_set_batch_split_copies( fg.sched, batch_split_copies); - if (ds4_env_flag("DFLASH_DS4_TP_SCHED_TRACE")) { + if (mixed_policy.scheduler_trace) { ggml_backend_sched_set_eval_callback( fg.sched, ds4_fused_verify_trace_node, nullptr); } @@ -807,26 +903,22 @@ static bool ds4_build_fused_verify_graph( pin_main(fg.mask_bundle); for (ggml_tensor * hids : fg.hash_ids) pin_main(hids); for (const MoeHybridGraphInputs & inputs : fg.hybrid_inputs) { - if (ds4_env_flag("DFLASH_DS4_TP_MAIN_ROUTE_WEIGHTS")) { + if (mixed_policy.pin_route_weights) { for (ggml_tensor * node : inputs.router_nodes) { pin_main(node); } pin_main(inputs.router_weights); } - if (ds4_env_flag("DFLASH_DS4_TP_ROUTE_PREFORK")) { + if (moe_policy.route_prefork) { for (ggml_tensor * node : inputs.route_prefork_nodes) { pin_main(node); } } pin_main(inputs.hot_local_lut); pin_main(inputs.hot_valid_lut); - // q1 has no repeat. The opt-in exact I32 GPU repeat lets q>1 keep - // both the LUT and its repeated output on the cold owner as well. - if (q == 1 || gpu_i32_repeat) { - pin_peer(inputs.cold_local_lut); - } else { - pin_main(inputs.cold_local_lut); - } + // The LUT is already q-batched and consumed only by the cold + // remap, so keep it with that owner for every verifier width. + pin_peer(inputs.cold_local_lut); pin_peer(inputs.cold_valid_lut); for (ggml_tensor * node : inputs.hot_remap_nodes) { pin_main(node); @@ -860,9 +952,7 @@ static bool ds4_build_fused_verify_graph( pin_main(route.selected); pin_main(route.weights); } - if (ds4_env_flag("DFLASH_DS4_TP_ROUTE_STATS") || - ds4_env_flag("DFLASH_DS4_ROUTING_STATS_OUT") || - ds4_env_flag("DFLASH_DS4_TP_CACHE_AUDIT")) { + if (mixed_policy.preserve_routes_for_diagnostics) { // Preserve the unsplit authoritative route and weight matrices. // Reading the pre-fork list is insufficient in the legacy 4+2 // lowering because its second build overwrites the per-layer diff --git a/server/src/deepseek4/deepseek4_graph.cpp b/server/src/deepseek4/deepseek4_graph.cpp index ee212bb5b..f99134fec 100644 --- a/server/src/deepseek4/deepseek4_graph.cpp +++ b/server/src/deepseek4/deepseek4_graph.cpp @@ -13,6 +13,7 @@ #include "internal.h" #include "../common/step_graph.h" #include "../common/cuda_graph_overrides.h" +#include "../common/dynamic_backend.h" #include "../common/moe_expert_compute.h" #include "../common/moe_hybrid_ffn_eval.h" #include "../common/moe_hybrid_routing_stats.h" @@ -1375,10 +1376,33 @@ static int ds4_comp_rows_used(const ggml_tensor * comp_cache, int n_cached, int // rows in [n_comp, padded) are masked to -1e30 in the score matrix, which // underflows to exactly 0 in softmax, so a padded read is bit-identical to an // unpadded read of the first n_comp rows. -static constexpr int DS4_COMP_PAD_STRIDE = 16; +static int ds4_comp_pad_stride() { + static const int stride = [] { + constexpr int default_stride = 16; + const char * raw = std::getenv("DFLASH_DS4_COMP_PAD_STRIDE"); + if (!raw || !*raw) return default_stride; + const int requested = std::atoi(raw); + switch (requested) { + case 16: + case 32: + case 64: + case 128: + return requested; + default: + std::fprintf(stderr, + "[deepseek4] invalid DFLASH_DS4_COMP_PAD_STRIDE=%s; " + "using %d\n", + raw, default_stride); + return default_stride; + } + }(); + return stride; +} + static int ds4_padded_comp_rows(int n_comp, int cap) { if (n_comp <= 0) return 0; - const int padded = ((n_comp + DS4_COMP_PAD_STRIDE - 1) / DS4_COMP_PAD_STRIDE) * DS4_COMP_PAD_STRIDE; + const int stride = ds4_comp_pad_stride(); + const int padded = ((n_comp + stride - 1) / stride) * stride; return padded < cap ? padded : cap; } @@ -2153,13 +2177,19 @@ static ggml_tensor * build_mla_attention( ggml_tensor * attn_low = ggml_mul_mat(ctx, out_a_3d, attn_out); // attn_low: [n_lora_o, n_tokens, n_out_group] ggml_tensor * out = nullptr; - if (n_tokens > 1) { + const bool grouped_output_projection = + n_tokens > 1 && + !ds4_env_flag("DFLASH_DS4_DISABLE_GROUPED_OUTPUT_PROJECTION"); + if (grouped_output_projection) { // Batched ROCmFPX MMQ consumes src1's channel stride directly. This // avoids materializing both permutations (~256 MiB/layer at 2K). out = ggml_mul_mat_grouped_src(ctx, L.attn_output_b, attn_low); } else { - // Preserve the established single-token graph and its numerical - // behavior. Decode is intentionally outside the prefill fast path. + // Preserve the established single-token graph and provide an exact + // fallback for heterogeneous runtimes that cannot retain grouped-view + // metadata across a scheduler copy. At verifier widths (q <= 4), this + // materializes at most 128 KiB per layer rather than the long-prefill + // volume avoided by the grouped path. attn_low = ggml_cont(ctx, ggml_permute(ctx, attn_low, 0, 2, 1, 3)); attn_low = ggml_reshape_2d( ctx, attn_low, n_lora_o * n_out_group, n_tokens); @@ -2885,7 +2915,8 @@ static bool eval_ds4_hybrid( ggml_tensor * ffn_normed_backend = nullptr, const MoeHybridDeviceOutputs * device_outputs = nullptr) { const auto ffn_t0 = Ds4TimingClock::now(); - if (!storage.down_cold && !storage.gate_up_cold && + if (!storage.cold_expert_ids.empty() && + !storage.down_cold && !storage.gate_up_cold && !(expert_compute && expert_layer)) { if (!hybrid_owner || !stream_engine || !stream_engine->is_ready() || !hybrid_owner->has_mmap() || @@ -4444,10 +4475,13 @@ bool deepseek4_step( // while a variant recurs, which is what the ggml-cuda/HIP graph cache keys // on, enabling graph replay for the bulk of decode steps. -static bool ds4_fused_decode_enabled() { - static const bool enabled = +static bool ds4_fused_decode_enabled(const DeepSeek4Weights & w) { + // The supported control is --ds4-fused-decode, propagated through the + // loaded weights. Keep the old environment spelling as a compatibility + // fallback for existing launch scripts. + static const bool legacy_env_enabled = ds4_env_flag("DFLASH_DS4_FUSED_DECODE"); - return enabled; + return w.fused_decode || legacy_env_enabled; } struct DeepSeek4FusedDecodeGraph { @@ -6760,7 +6794,8 @@ bool deepseek4_step_layer_range( if (!moe_hybrid && n_tokens == 1 && allow_decode_graph_reuse && layer_begin == 0 && is_last_shard && !(verify_hooks && verify_hooks->capture_layer_ids && verify_hooks->capture_out) && - out_logits && ds4_backend_is_gpu(backend) && ds4_fused_decode_enabled()) { + out_logits && ds4_backend_is_gpu(backend) && + ds4_fused_decode_enabled(w)) { const int rc = ds4_try_fused_decode_step( fused_decode_graph_cache, backend, w, cache, hc_layer_weights_range, hc_output_weights_range, hash_routing_tables_range, scratch.hash_expert_ids, @@ -6848,6 +6883,38 @@ bool deepseek4_step_layer_range( hc_state.data(), 0, sizeof(float) * hc_state.size()); hc_state_backend = cached_decode_hc_post_graph.residual_hc; } + const auto capture_requested = [&](int layer) { + if (!verify_hooks || !verify_hooks->capture_layer_ids || + !verify_hooks->capture_out) { + return false; + } + const std::vector & ids = *verify_hooks->capture_layer_ids; + return std::find(ids.begin(), ids.end(), layer) != ids.end(); + }; + const auto capture_hc_layer = [&](int layer, const float * state) { + if (!state || !capture_requested(layer)) return; + const std::vector & ids = *verify_hooks->capture_layer_ids; + std::vector & capture = *verify_hooks->capture_out; + if ((int) capture.size() != (int) ids.size() * n_embd * n_tokens) { + capture.assign( + (size_t) ids.size() * n_embd * n_tokens, 0.0f); + } + for (size_t ci = 0; ci < ids.size(); ++ci) { + if (ids[ci] != layer) continue; + for (int t = 0; t < n_tokens; ++t) { + float * dst = capture.data() + + (size_t) t * ids.size() * n_embd + ci * n_embd; + const float * hs = state + (size_t) t * hc_dim; + for (int d = 0; d < n_embd; ++d) { + float sum = 0.0f; + for (int h = 0; h < n_hc; ++h) { + sum += hs[(size_t) h * n_embd + d]; + } + dst[d] = sum / (float) n_hc; + } + } + } + }; for (int il = layer_begin; il < layer_end; ++il) { const DeepSeek4Layer & L = w.layers[(size_t)il]; DeepSeek4LayerCache & lc = cache.layers[(size_t)il]; @@ -7518,25 +7585,15 @@ bool deepseek4_step_layer_range( n_hc); std::memcpy(hc_state.data(), next_hc.data(), next_hc.size() * sizeof(float)); if (telemetry) telemetry->hc_post_ffn_us += ds4_elapsed_us(hc_post_ffn_t0, Ds4TimingClock::now()); - if (verify_hooks && verify_hooks->capture_layer_ids && verify_hooks->capture_out) { - const std::vector & _ids = *verify_hooks->capture_layer_ids; - for (size_t _ci = 0; _ci < _ids.size(); ++_ci) { - if (_ids[_ci] != il) continue; - const int _ncap = (int) _ids.size(); - std::vector & _cap = *verify_hooks->capture_out; - if ((int) _cap.size() != _ncap * n_embd * n_tokens) - _cap.assign((size_t) _ncap * n_embd * n_tokens, 0.0f); - for (int _t = 0; _t < n_tokens; ++_t) { - float * _dst = _cap.data() + (size_t) _t * _ncap * n_embd + (size_t) _ci * n_embd; - const float * _hs = hc_state.data() + (size_t) _t * hc_dim; - for (int _d = 0; _d < n_embd; ++_d) { - float _acc = 0.0f; - for (int _h = 0; _h < n_hc; ++_h) _acc += _hs[(size_t) _h * n_embd + _d]; - _dst[_d] = _acc / (float) n_hc; - } - } - } - } + capture_hc_layer(il, hc_state.data()); + } + if ((use_backend_prefill_hc || use_backend_decode_hc_graph || + use_backend_decode_hc_direct) && + hc_state_backend && capture_requested(il)) { + ggml_backend_tensor_get( + hc_state_backend, hc_state.data(), 0, + sizeof(float) * hc_state.size()); + capture_hc_layer(il, hc_state.data()); } } } @@ -7585,7 +7642,9 @@ bool deepseek4_step_layer_range( ggml_context * ctx = ggml_init(params); if (!ctx) return false; - const bool last_only = n_tokens > 1; + const bool need_all_logits = + verify_hooks && verify_hooks->all_logits_out; + const bool last_only = n_tokens > 1 && !need_all_logits; const int output_tokens = last_only ? 1 : n_tokens; ggml_tensor * inp = ggml_new_tensor_2d( ctx, GGML_TYPE_F32, n_embd, output_tokens); diff --git a/server/src/placement/placement_backend.h b/server/src/placement/placement_backend.h index 0dfa053c4..4a0b39666 100644 --- a/server/src/placement/placement_backend.h +++ b/server/src/placement/placement_backend.h @@ -46,9 +46,4 @@ inline PlacementBackend compiled_placement_backend() { #endif } -inline bool placement_backend_supported(PlacementBackend backend) { - return backend == PlacementBackend::Auto || - backend == compiled_placement_backend(); -} - } // namespace dflash::common diff --git a/server/test/test_feature_gate.cpp b/server/test/test_feature_gate.cpp index 63091244b..071b8d44c 100644 --- a/server/test/test_feature_gate.cpp +++ b/server/test/test_feature_gate.cpp @@ -288,6 +288,19 @@ static void test_feature_gate_ds4_decode_options_require_monolithic_hip() { topk, "qwen35", PlacementBackend::Hip).empty()); TEST_ASSERT(gate_result( topk, "deepseek4", PlacementBackend::Hip).empty()); + + // Top-k is a model policy in the monolithic backend and is independent of + // the GPU vendor. Unlike fused decode, mixed CUDA-primary expert + // placement can therefore use it. + BackendArgs cuda_topk = topk; + cuda_topk.device.backend = PlacementBackend::Cuda; + TEST_ASSERT(gate_result( + cuda_topk, "deepseek4", PlacementBackend::Cuda).empty()); + + BackendArgs split_topk = topk; + split_topk.device.layer_split_gpus = {0, 1}; + TEST_ASSERT(!gate_result( + split_topk, "deepseek4", PlacementBackend::Hip).empty()); } static void test_feature_gate_remote_draft_requires_supported_arch() { diff --git a/server/test/test_mixed_cuda_hip.cpp b/server/test/test_mixed_cuda_hip.cpp new file mode 100644 index 000000000..8b74050d8 --- /dev/null +++ b/server/test/test_mixed_cuda_hip.cpp @@ -0,0 +1,394 @@ +#include "common/dynamic_backend.h" + +#include "ggml.h" +#include "ggml-alloc.h" +#include "ggml-backend.h" +#include "ggml-cpu.h" + +#include +#include +#include +#include + +using dflash::common::PlacementBackend; +using dflash::common::backend_pair_capabilities; +using dflash::common::init_placement_backend; +using dflash::common::placement_backend_of; + +namespace { + +bool run_scale(ggml_backend_t backend, const char * label) { + constexpr int64_t n = 4096; + ggml_init_params params{}; + params.mem_size = 4 * 1024 * 1024; + params.no_alloc = true; + ggml_context * ctx = ggml_init(params); + if (!ctx) return false; + + ggml_tensor * input = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, n); + ggml_set_input(input); + ggml_tensor * output = ggml_scale(ctx, input, 2.0f); + ggml_set_output(output); + ggml_cgraph * graph = ggml_new_graph(ctx); + ggml_build_forward_expand(graph, output); + ggml_gallocr_t alloc = ggml_gallocr_new( + ggml_backend_get_default_buffer_type(backend)); + if (!alloc || !ggml_gallocr_alloc_graph(alloc, graph)) { + if (alloc) ggml_gallocr_free(alloc); + ggml_free(ctx); + return false; + } + + std::vector host((size_t)n); + for (int64_t i = 0; i < n; ++i) host[(size_t)i] = (float)i / 128.0f; + ggml_backend_tensor_set(input, host.data(), 0, ggml_nbytes(input)); + const enum ggml_status status = ggml_backend_graph_compute(backend, graph); + std::vector result((size_t)n); + if (status == GGML_STATUS_SUCCESS) { + ggml_backend_tensor_get(output, result.data(), 0, ggml_nbytes(output)); + } + + bool ok = status == GGML_STATUS_SUCCESS; + for (int64_t i = 0; ok && i < n; ++i) { + ok = std::fabs(result[(size_t)i] - 2.0f * host[(size_t)i]) < 1.0e-5f; + } + std::printf("mixed-backend %s scale: %s\n", label, ok ? "ok" : "FAILED"); + ggml_gallocr_free(alloc); + ggml_free(ctx); + return ok; +} + +bool run_cross_copy(ggml_backend_t src_backend, + ggml_backend_t dst_backend, + const char * label) { + constexpr int64_t n = 4096; + ggml_init_params params{}; + params.mem_size = 2 * 1024 * 1024; + params.no_alloc = true; + ggml_context * src_ctx = ggml_init(params); + ggml_context * dst_ctx = ggml_init(params); + if (!src_ctx || !dst_ctx) { + if (src_ctx) ggml_free(src_ctx); + if (dst_ctx) ggml_free(dst_ctx); + return false; + } + + ggml_tensor * src = ggml_new_tensor_1d(src_ctx, GGML_TYPE_F32, n); + ggml_tensor * dst = ggml_new_tensor_1d(dst_ctx, GGML_TYPE_F32, n); + ggml_backend_buffer_t src_buf = ggml_backend_alloc_ctx_tensors(src_ctx, src_backend); + ggml_backend_buffer_t dst_buf = ggml_backend_alloc_ctx_tensors(dst_ctx, dst_backend); + if (!src_buf || !dst_buf) { + if (src_buf) ggml_backend_buffer_free(src_buf); + if (dst_buf) ggml_backend_buffer_free(dst_buf); + ggml_free(src_ctx); + ggml_free(dst_ctx); + return false; + } + + std::vector input((size_t)n); + std::vector output((size_t)n, 0.0f); + for (int64_t i = 0; i < n; ++i) input[(size_t)i] = (float)(i * 17 - 31); + ggml_backend_tensor_set(src, input.data(), 0, ggml_nbytes(src)); + ggml_backend_tensor_copy_async(src_backend, dst_backend, src, dst); + ggml_backend_synchronize(dst_backend); + ggml_backend_tensor_get(dst, output.data(), 0, ggml_nbytes(dst)); + const bool ok = input == output; + std::printf("mixed-backend %s copy: %s\n", label, ok ? "ok" : "FAILED"); + + ggml_backend_buffer_free(src_buf); + ggml_backend_buffer_free(dst_buf); + ggml_free(src_ctx); + ggml_free(dst_ctx); + return ok; +} + +bool run_cross_graph(ggml_backend_t first, + ggml_backend_t second, + const char * label, + bool batch_split_copies) { + constexpr int64_t n = 4096; + ggml_init_params params{}; + params.mem_size = 4 * 1024 * 1024; + params.no_alloc = true; + ggml_context * ctx = ggml_init(params); + ggml_backend_t cpu = ggml_backend_cpu_init(); + if (!ctx || !cpu) { + if (cpu) ggml_backend_free(cpu); + if (ctx) ggml_free(ctx); + return false; + } + + ggml_tensor * input = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, n); + ggml_set_input(input); + ggml_tensor * first_head = ggml_scale(ctx, input, 2.0f); + ggml_tensor * first_sibling = ggml_scale(ctx, input, 3.0f); + ggml_tensor * second_middle = ggml_add(ctx, first_head, first_sibling); + ggml_tensor * first_tail = ggml_scale(ctx, second_middle, 4.0f); + ggml_set_output(first_tail); + ggml_cgraph * graph = ggml_new_graph(ctx); + ggml_build_forward_expand(graph, first_tail); + + ggml_backend_t backends[] = { first, second, cpu }; + ggml_backend_sched_t sched = ggml_backend_sched_new( + backends, nullptr, 3, 64, false, true); + bool ok = sched != nullptr; + if (ok) { + ggml_backend_sched_set_tensor_backend(sched, input, first); + ggml_backend_sched_set_tensor_backend(sched, first_head, first); + ggml_backend_sched_set_tensor_backend(sched, first_sibling, first); + ggml_backend_sched_set_tensor_backend(sched, second_middle, second); + ggml_backend_sched_set_tensor_backend(sched, first_tail, first); + ggml_backend_sched_set_batch_split_copies( + sched, batch_split_copies); + ok = ggml_backend_sched_alloc_graph(sched, graph); + } + + std::vector host((size_t)n); + std::vector result((size_t)n, 0.0f); + for (int iteration = 0; ok && iteration < 3; ++iteration) { + for (int64_t i = 0; i < n; ++i) { + host[(size_t)i] = ((float)i + 17.0f * iteration) / 128.0f; + } + ggml_backend_tensor_set(input, host.data(), 0, ggml_nbytes(input)); + ok = ggml_backend_sched_graph_compute(sched, graph) == + GGML_STATUS_SUCCESS; + if (ok) { + ggml_backend_tensor_get( + first_tail, result.data(), 0, ggml_nbytes(first_tail)); + } + for (int64_t i = 0; ok && i < n; ++i) { + ok = std::fabs(result[(size_t)i] - 20.0f * host[(size_t)i]) < + 1.0e-4f; + } + } + std::printf("mixed-backend %s graph batch=%d: %s\n", label, + batch_split_copies ? 1 : 0, + ok ? "ok" : "FAILED"); + + if (sched) ggml_backend_sched_free(sched); + ggml_backend_free(cpu); + ggml_free(ctx); + return ok; +} + +struct SchedulerGraphCase { + ggml_context * ctx = nullptr; + ggml_tensor * input = nullptr; + ggml_tensor * first = nullptr; + ggml_tensor * second = nullptr; + ggml_tensor * output = nullptr; + ggml_cgraph * graph = nullptr; + int64_t elements = 0; +}; + +SchedulerGraphCase make_resize_graph(int64_t elements) { + SchedulerGraphCase result; + result.elements = elements; + ggml_init_params params{}; + params.mem_size = 4 * 1024 * 1024; + params.no_alloc = true; + result.ctx = ggml_init(params); + if (!result.ctx) return result; + + result.input = ggml_new_tensor_1d( + result.ctx, GGML_TYPE_F32, elements); + ggml_set_input(result.input); + result.first = ggml_scale(result.ctx, result.input, 2.0f); + result.second = ggml_scale(result.ctx, result.first, 3.0f); + result.output = ggml_scale(result.ctx, result.second, 4.0f); + ggml_set_output(result.output); + result.graph = ggml_new_graph(result.ctx); + ggml_build_forward_expand(result.graph, result.output); + return result; +} + +bool pin_resize_graph(ggml_backend_sched_t sched, + const SchedulerGraphCase & graph, + ggml_backend_t first, + ggml_backend_t second) { + if (!sched || !graph.ctx || !graph.input || !graph.first || + !graph.second || !graph.output || !graph.graph) { + return false; + } + ggml_backend_sched_set_tensor_backend(sched, graph.input, first); + ggml_backend_sched_set_tensor_backend(sched, graph.first, first); + ggml_backend_sched_set_tensor_backend(sched, graph.second, second); + ggml_backend_sched_set_tensor_backend(sched, graph.output, first); + return ggml_backend_sched_alloc_graph(sched, graph.graph); +} + +bool set_and_check_resize_graph(ggml_backend_sched_t sched, + const SchedulerGraphCase & graph, + float bias, + bool async) { + std::vector input((size_t) graph.elements); + std::vector output((size_t) graph.elements, 0.0f); + for (int64_t i = 0; i < graph.elements; ++i) { + input[(size_t) i] = ((float) i + bias) / 256.0f; + } + ggml_backend_tensor_set( + graph.input, input.data(), 0, ggml_nbytes(graph.input)); + const enum ggml_status status = async + ? ggml_backend_sched_graph_compute_async(sched, graph.graph) + : ggml_backend_sched_graph_compute(sched, graph.graph); + if (status != GGML_STATUS_SUCCESS) return false; + if (async) return true; + + ggml_backend_tensor_get( + graph.output, output.data(), 0, ggml_nbytes(graph.output)); + for (int64_t i = 0; i < graph.elements; ++i) { + if (std::fabs(output[(size_t) i] - 24.0f * input[(size_t) i]) >= + 1.0e-4f) { + return false; + } + } + return true; +} + +bool run_async_reset_resize_and_free(ggml_backend_t first, + ggml_backend_t second, + const char * label) { + ggml_backend_t cpu = ggml_backend_cpu_init(); + SchedulerGraphCase small = make_resize_graph(1024); + SchedulerGraphCase large = make_resize_graph(32768); + ggml_backend_t backends[] = { first, second, cpu }; + ggml_backend_sched_t sched = cpu + ? ggml_backend_sched_new(backends, nullptr, 3, 64, false, true) + : nullptr; + bool ok = sched && small.ctx && large.ctx; + if (ok) { + ggml_backend_sched_set_batch_split_copies(sched, true); + ok = pin_resize_graph(sched, small, first, second) && + set_and_check_resize_graph( + sched, small, 7.0f, /*async=*/true); + } + if (ok) { + // Reset while the small graph is still in flight, then grow both the + // graph allocation and staging arena. Allocation must quiesce old host + // pages before replacing them. + ggml_backend_sched_reset(sched); + ok = pin_resize_graph(sched, large, first, second) && + set_and_check_resize_graph( + sched, large, 11.0f, /*async=*/false); + } + if (ok) { + // Leave one final submission in flight. Scheduler teardown owns the + // wait required before freeing graph buffers and staging pages. + ok = set_and_check_resize_graph( + sched, large, 19.0f, /*async=*/true); + } + if (sched) ggml_backend_sched_free(sched); + if (cpu) ggml_backend_free(cpu); + if (large.ctx) ggml_free(large.ctx); + if (small.ctx) ggml_free(small.ctx); + std::printf("mixed-backend %s async-reset-resize-free: %s\n", + label, ok ? "ok" : "FAILED"); + return ok; +} + +bool run_multi_source_graph(ggml_backend_t first, + ggml_backend_t second, + const char * label) { + constexpr int64_t n = 8192; + ggml_init_params params{}; + params.mem_size = 4 * 1024 * 1024; + params.no_alloc = true; + ggml_context * ctx = ggml_init(params); + ggml_backend_t cpu = ggml_backend_cpu_init(); + if (!ctx || !cpu) { + if (cpu) ggml_backend_free(cpu); + if (ctx) ggml_free(ctx); + return false; + } + + ggml_tensor * input = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, n); + ggml_set_input(input); + ggml_tensor * first_head = ggml_scale(ctx, input, 2.0f); + ggml_tensor * cpu_head = ggml_scale(ctx, input, 5.0f); + ggml_tensor * second_head = ggml_scale(ctx, input, 3.0f); + ggml_tensor * second_join = ggml_add(ctx, first_head, cpu_head); + ggml_tensor * output = ggml_add(ctx, second_join, second_head); + ggml_set_output(output); + ggml_cgraph * graph = ggml_new_graph(ctx); + ggml_build_forward_expand(graph, output); + + ggml_backend_t backends[] = { first, second, cpu }; + ggml_backend_sched_t sched = ggml_backend_sched_new( + backends, nullptr, 3, 64, false, true); + bool ok = sched != nullptr; + if (ok) { + ggml_backend_sched_set_tensor_backend(sched, input, first); + ggml_backend_sched_set_tensor_backend(sched, first_head, first); + ggml_backend_sched_set_tensor_backend(sched, cpu_head, cpu); + ggml_backend_sched_set_tensor_backend(sched, second_head, second); + ggml_backend_sched_set_tensor_backend(sched, second_join, second); + ggml_backend_sched_set_tensor_backend(sched, output, first); + ggml_backend_sched_set_batch_split_copies(sched, true); + ok = ggml_backend_sched_alloc_graph(sched, graph); + } + + std::vector host((size_t) n); + std::vector result((size_t) n, 0.0f); + for (int iteration = 0; ok && iteration < 3; ++iteration) { + for (int64_t i = 0; i < n; ++i) { + host[(size_t) i] = ((float) i + 13.0f * iteration) / 128.0f; + } + ggml_backend_tensor_set(input, host.data(), 0, ggml_nbytes(input)); + ok = ggml_backend_sched_graph_compute(sched, graph) == + GGML_STATUS_SUCCESS; + if (ok) { + ggml_backend_tensor_get( + output, result.data(), 0, ggml_nbytes(output)); + } + for (int64_t i = 0; ok && i < n; ++i) { + ok = std::fabs(result[(size_t) i] - 10.0f * host[(size_t) i]) < + 1.0e-4f; + } + } + std::printf("mixed-backend %s multi-source graph: %s\n", + label, ok ? "ok" : "FAILED"); + if (sched) ggml_backend_sched_free(sched); + ggml_backend_free(cpu); + ggml_free(ctx); + return ok; +} + +} // namespace + +int main() { + std::string error; + ggml_backend_t cuda = init_placement_backend(PlacementBackend::Cuda, 0, &error); + if (!cuda) { + std::fprintf(stderr, "CUDA initialization failed: %s\n", error.c_str()); + return 1; + } + ggml_backend_t hip = init_placement_backend(PlacementBackend::Hip, 0, &error); + if (!hip) { + std::fprintf(stderr, "HIP initialization failed: %s\n", error.c_str()); + ggml_backend_free(cuda); + return 1; + } + + const auto pair = backend_pair_capabilities(cuda, hip); + bool ok = placement_backend_of(cuda) == PlacementBackend::Cuda && + placement_backend_of(hip) == PlacementBackend::Hip && + !pair.same_runtime && !pair.native_gpu_handoff; + ok = run_scale(cuda, "CUDA") && ok; + ok = run_scale(hip, "HIP") && ok; + ok = run_cross_copy(cuda, hip, "CUDA->HIP") && ok; + ok = run_cross_copy(hip, cuda, "HIP->CUDA") && ok; + ok = run_cross_graph(cuda, hip, "CUDA->HIP->CUDA", false) && ok; + ok = run_cross_graph(cuda, hip, "CUDA->HIP->CUDA", true) && ok; + ok = run_cross_graph(hip, cuda, "HIP->CUDA->HIP", false) && ok; + ok = run_cross_graph(hip, cuda, "HIP->CUDA->HIP", true) && ok; + ok = run_multi_source_graph(cuda, hip, "CUDA+CPU->HIP") && ok; + ok = run_multi_source_graph(hip, cuda, "HIP+CPU->CUDA") && ok; + ok = run_async_reset_resize_and_free( + cuda, hip, "CUDA->HIP->CUDA") && ok; + ok = run_async_reset_resize_and_free( + hip, cuda, "HIP->CUDA->HIP") && ok; + + ggml_backend_free(hip); + ggml_backend_free(cuda); + return ok ? 0 : 1; +}