7f59912916
The model code was spread across eight crates that had grown into each other:
ggml and cuda and mlx each owned part of a tensor runtime, llama and tts and
voice2 each owned part of a model, and libs/diffusion owned everything else.
They are now one tree with an explicit shape:
libs/ai/cuda — kernels and launch surface
libs/ai/metal — Metal shaders and the shim
libs/ai/llm — the language-model runtime (sessions, lanes, contexts,
the CUDA and Metal executors, the compiled Metal path)
libs/ai/models/ — common, flux, h3, music, paint, speech, stems, vision
libs/diffusion is not deleted but demoted: what remains is the VALIDATOR
crate — several dozen `*_validate.rs` oracles that check a native
implementation against a reference, which is where they belong now that the
implementations live next door.
The functional work inside the move is mostly in the LLM runtime: N lanes that
draft while one verify batch serves all of them, per-slot prefill over a shared
folded attention arena, speculation that survives batching, and a scheduler
that reports rather than publishes. And in the CUDA build: a machine without
usable CUDA must still LINK (and say so), the default kernel arch is the
building machine's GPU, `NO_CUDA` forces the stub even where the toolkit
exists, and kernels compile in parallel with progress.
libs/video_flow is new here: classical optical flow estimation and the `mkfl`
motion-field payload — a flow field measured from a clip without a model,
which is what drives free-rate bounce-looping playback and the uprez/tween
enhance pipe.
65 lines
2.4 KiB
Text
65 lines
2.4 KiB
Text
// SPDX-License-Identifier: MIT
|
|
// Copyright (c) 2023-2026 The ggml authors
|
|
//
|
|
// Substantial portions derived from ggml / llama.cpp
|
|
// (https://github.com/ggml-org/llama.cpp), MIT licensed.
|
|
// The original copyright notice and permission notice are retained.
|
|
// See libs/ai/NOTICE and, where present, LICENSE in this directory.
|
|
//
|
|
// Simplified API for asynchronous data loading.
|
|
|
|
#include "common.cuh"
|
|
|
|
|
|
static __device__ __forceinline__ unsigned int ggml_cuda_cvta_generic_to_shared(void * generic_ptr) {
|
|
#ifdef CP_ASYNC_AVAILABLE
|
|
return __cvta_generic_to_shared(generic_ptr);
|
|
#else
|
|
GGML_UNUSED(generic_ptr);
|
|
NO_DEVICE_CODE;
|
|
return 0;
|
|
#endif // CP_ASYNC_AVAILABLE
|
|
}
|
|
|
|
// Copies data from global to shared memory, cg == cache global.
|
|
// Both the src and dst pointers must be aligned to 16 bit.
|
|
// Shared memory uses 32 bit addressing, the pointer is passed as unsigned int.
|
|
// Generic pointers can be converted to 32 bit shared memory pointers using __cvta_generic_to_shared.
|
|
// Only the 16 bit copy is exposed because 4 and 8 bit copies did not yield performance improvements.
|
|
template <int preload>
|
|
static __device__ __forceinline__ void cp_async_cg_16(const unsigned int dst, const void * src) {
|
|
static_assert(preload == 0 || preload == 64 || preload == 128 || preload == 256, "bad preload");
|
|
#ifdef CP_ASYNC_AVAILABLE
|
|
#if CUDART_VERSION >= 11040
|
|
if (preload == 256) {
|
|
asm volatile("cp.async.cg.shared.global.L2::256B [%0], [%1], 16;"
|
|
: : "r"(dst), "l"(src));
|
|
} else if (preload == 128) {
|
|
asm volatile("cp.async.cg.shared.global.L2::128B [%0], [%1], 16;"
|
|
: : "r"(dst), "l"(src));
|
|
} else if (preload == 64) {
|
|
asm volatile("cp.async.cg.shared.global.L2::64B [%0], [%1], 16;"
|
|
: : "r"(dst), "l"(src));
|
|
} else
|
|
#endif // CUDART_VERSION >= 11040
|
|
{
|
|
asm volatile("cp.async.cg.shared.global [%0], [%1], 16;"
|
|
: : "r"(dst), "l"(src));
|
|
}
|
|
#else
|
|
GGML_UNUSED(dst);
|
|
GGML_UNUSED(src);
|
|
NO_DEVICE_CODE;
|
|
#endif // CP_ASYNC_AVAILABLE
|
|
}
|
|
|
|
// Makes each thread wait until its asynchronous data copies are done.
|
|
// This does NOT provide any additional synchronization.
|
|
// In particular, when copying data with multiple warps a call to __syncthreads will be needed.
|
|
static __device__ __forceinline__ void cp_async_wait_all() {
|
|
#ifdef CP_ASYNC_AVAILABLE
|
|
asm volatile("cp.async.wait_all;");
|
|
#else
|
|
NO_DEVICE_CODE;
|
|
#endif // CP_ASYNC_AVAILABLE
|
|
}
|