src/rocm/ds4_rocm_runtime.cuh

static const void *g_model_host_base;
static const char *g_model_device_base;
static uint64_t g_model_registered_size;
static int g_model_device_owned;
static int g_model_range_mapping_supported = 1;
static int g_model_fd = -1;
static const void *g_model_fd_host_base;
static int g_model_direct_fd = -1;
static uint64_t g_model_direct_align = 1;
static uint64_t g_model_file_size;
static int g_model_cache_full;
static int g_ssd_streaming_mode;
static cudaStream_t g_model_upload_stream;
static cudaStream_t g_stream_selected_upload_stream;
static cudaStream_t g_selected_readback_stream;
static cudaEvent_t g_selected_readback_event;
static uint64_t g_selected_readback_event_value;
static cublasHandle_t g_cublas;
static int g_cublas_ready;
#ifdef __HIP_PLATFORM_AMD__
#include "ds4_rocm_hipblaslt.cuh"
#endif
static int g_quality_mode;

enum {
    DS4_ROCM_N_EXPERT = 256u,
    DS4_ROCM_MAX_N_EXPERT = 384u,
    DS4_ROCM_N_EXPERT_USED = 8u,
    DS4_ROCM_STREAM_READ_WORKERS = DS4_ROCM_N_EXPERT_USED * 3u,
    DS4_ROCM_STREAM_READ_DEFAULT_WORKERS = 16u,
    DS4_ROCM_STREAM_READ_MAX_JOBS = DS4_ROCM_MAX_N_EXPERT * 3u,
    DS4_ROCM_STREAM_CACHE_LAYER_STATS_MAX = 128u,
    DS4_ROCM_COMPRESSOR_MAX_RATIO = 128u
};
#define DS4_ROCM_EXPERT_WEIGHT_SCALE 1.5f
#define DS4_ROCM_EXPERT_WEIGHT_SCALE_TOL 1.0e-6f

struct cuda_model_range {
    const void *host_base;
    uint64_t offset;
    uint64_t bytes;
    char *device_ptr;
    void *registered_base;
    char *registered_device_base;
    uint64_t registered_bytes;
    int host_registered;
    int arena_allocated;
};

struct cuda_model_arena {
    char *device_ptr;
    uint64_t bytes;
    uint64_t used;
};

struct cuda_model_image {
    const void *host_base;
    uint64_t size;
    char *device_ptr;
};

struct cuda_q8_f16_range {
    const void *host_base;
    uint64_t offset;
    uint64_t weight_bytes;
    uint64_t in_dim;
    uint64_t out_dim;
    __half *device_ptr;
};

struct cuda_q8_f16_transpose_range {
    const void *host_base;
    uint64_t offset;
    uint64_t weight_bytes;
    uint64_t in_dim;
    uint64_t out_dim;
    __half *device_ptr;
};

struct cuda_stream_selected_cache {
    int loaded;
    const void *model_map;
    uint32_t layer;
    uint32_t n_total_expert;
    uint32_t n_selected;
    uint64_t gate_expert_bytes;
    uint64_t down_expert_bytes;
    int32_t selected_ids[DS4_ROCM_N_EXPERT_USED];
    char *gate;
    char *up;
    char *down;
    uint64_t gate_capacity;
    uint64_t down_capacity;
    int32_t *slot_ids;
    const char **gate_ptrs;
    const char **up_ptrs;
    const char **down_ptrs;
    const char **gate_ptrs_stage;
    const char **up_ptrs_stage;
    const char **down_ptrs_stage;
    ds4_gpu_tensor slot_tensor;
};

struct cuda_stream_resident_expert {
    const void *model_map;
    uint32_t layer;
    int32_t expert;
    uint64_t gate_expert_bytes;
    uint64_t down_expert_bytes;
    uint64_t gate_offset;
    uint64_t up_offset;
    uint64_t down_offset;
    char *base;
    char *gate;
    char *up;
    char *down;
    uint64_t bytes;
    uint64_t last_used;
    int pooled;
};

/*
 * Streamed experts share one size class per model (2*gate + down bytes), so
 * the resident cache carves fixed slots out of large slabs instead of doing a
 * cudaMalloc/cudaFree round trip per cache miss.  Allocator traffic was a
 * large fraction of decode time: hundreds of 11.8 MiB device allocations and
 * frees per generated token, each costing GTT page-table work.
 */
struct cuda_stream_expert_slab {
    char *base;
    uint64_t bytes;
};

struct cuda_stream_resident_key {
    const void *model_map;
    uint32_t layer;
    int32_t expert;
    uint64_t gate_offset;
    uint64_t up_offset;
    uint64_t down_offset;
    uint64_t gate_expert_bytes;
    uint64_t down_expert_bytes;

    bool operator==(const cuda_stream_resident_key &o) const {
        return model_map == o.model_map &&
               layer == o.layer &&
               expert == o.expert &&
               gate_offset == o.gate_offset &&
               up_offset == o.up_offset &&
               down_offset == o.down_offset &&
               gate_expert_bytes == o.gate_expert_bytes &&
               down_expert_bytes == o.down_expert_bytes;
    }
};

struct cuda_stream_resident_key_hash {
    size_t operator()(const cuda_stream_resident_key &k) const {
        uint64_t h = (uint64_t)(uintptr_t)k.model_map;
        h ^= (uint64_t)k.layer + 0x9e3779b97f4a7c15ull + (h << 6) + (h >> 2);
        h ^= (uint64_t)(uint32_t)k.expert + 0x9e3779b97f4a7c15ull + (h << 6) + (h >> 2);
        h ^= k.gate_offset + 0x9e3779b97f4a7c15ull + (h << 6) + (h >> 2);
        h ^= k.up_offset + 0x9e3779b97f4a7c15ull + (h << 6) + (h >> 2);
        h ^= k.down_offset + 0x9e3779b97f4a7c15ull + (h << 6) + (h >> 2);
        h ^= k.gate_expert_bytes + 0x9e3779b97f4a7c15ull + (h << 6) + (h >> 2);
        h ^= k.down_expert_bytes + 0x9e3779b97f4a7c15ull + (h << 6) + (h >> 2);
        return (size_t)h;
    }
};

struct cuda_stream_batch_selected_cache {
    int loaded;
    const void *model_map;
    uint32_t layer;
    uint32_t n_total_expert;
    uint32_t n_selected;
    uint32_t n_tokens;
    uint32_t n_unique;
    uint64_t gate_offset;
    uint64_t up_offset;
    uint64_t down_offset;
    uint64_t gate_expert_bytes;
    uint64_t down_expert_bytes;
    int32_t *selected_ids;
    uint64_t selected_capacity;
    uint8_t *pair_missing;
    uint64_t pair_missing_capacity;
    const char **gate_ptrs;
    const char **up_ptrs;
    const char **down_ptrs;
    const char **resident_gate_ptrs;
    const char **resident_up_ptrs;
    const char **missing_gate_ptrs;
    const char **missing_up_ptrs;
    uint32_t ptr_capacity;
    int32_t *selected_stage;
    uint64_t selected_stage_capacity;
    uint8_t *pair_missing_stage;
    uint64_t pair_missing_stage_capacity;
    const char **gate_ptrs_stage;
    const char **up_ptrs_stage;
    const char **down_ptrs_stage;
    const char **resident_gate_ptrs_stage;
    const char **resident_up_ptrs_stage;
    const char **missing_gate_ptrs_stage;
    const char **missing_up_ptrs_stage;
    uint32_t ptr_stage_capacity;
    ds4_gpu_tensor selected_tensor;
};

struct cuda_stream_layer_expert_cache {
    int active;
    const void *model_map;
    uint32_t layer;
    uint32_t n_total_expert;
    uint64_t gate_offset;
    uint64_t up_offset;
    uint64_t down_offset;
    uint64_t gate_expert_bytes;
    uint64_t down_expert_bytes;
    uint64_t bytes;
    uint64_t capacity;
    char *base;
    char *gate;
    char *up;
    char *down;
};

struct cuda_stream_cache_stats {
    uint64_t selected_calls;
    uint64_t selected_slots;
    uint64_t selected_hits;
    uint64_t selected_misses;
    uint64_t batch_calls;
    uint64_t batch_unique;
    uint64_t batch_hits;
    uint64_t batch_misses;
    uint64_t seed_calls;
    uint64_t seed_unique;
    uint64_t layer_loads;
    uint64_t layer_load_bytes;
    uint64_t layer_resident_flushes;
    uint64_t allocs;
    uint64_t alloc_bytes;
    uint64_t evictions;
    uint64_t evict_bytes;
    uint64_t max_resident_count;
    uint64_t max_resident_bytes;
};

struct cuda_stream_cache_layer_stats {
    uint64_t selected_calls;
    uint64_t selected_slots;
    uint64_t selected_hits;
    uint64_t selected_misses;
    uint64_t batch_calls;
    uint64_t batch_unique;
    uint64_t batch_hits;
    uint64_t batch_misses;
};

static std::vector<cuda_model_range> g_model_ranges;
static std::vector<cuda_model_arena> g_model_arenas;
static std::vector<cuda_model_image> g_model_images;
static std::unordered_map<uint64_t, size_t> g_model_range_by_offset;
static std::vector<cuda_q8_f16_range> g_q8_f16_ranges;
static std::unordered_map<uint64_t, size_t> g_q8_f16_by_offset;
static std::vector<cuda_q8_f16_transpose_range> g_q8_f16_transpose_ranges;
static std::unordered_map<uint64_t, size_t> g_q8_f16_transpose_by_offset;
static uint64_t g_model_range_bytes;
static uint64_t g_q8_f16_bytes;
static int g_q8_f16_disabled_after_oom;
static int g_q8_f16_disabled_for_multi_model;
static int g_q8_f16_budget_notice_printed;
static uint64_t g_model_load_progress_next;
static double g_model_load_progress_last;
static int g_model_load_progress_started;
static int g_model_load_progress_tty;
static void *g_cuda_tmp;
static uint64_t g_cuda_tmp_bytes;
static void *g_model_stage_raw[4];
static void *g_model_stage[4];
static cudaEvent_t g_model_stage_event[4];
static uint64_t g_model_stage_bytes;
static uint32_t g_stream_expert_cache_budget;
static cuda_stream_selected_cache g_stream_selected_cache;
static cuda_stream_batch_selected_cache g_stream_batch_selected_cache;
static cuda_stream_layer_expert_cache g_stream_layer_expert_cache[2];
static std::vector<cuda_stream_resident_expert> g_stream_resident_experts;
static std::unordered_map<cuda_stream_resident_key,
                          size_t,
                          cuda_stream_resident_key_hash> g_stream_resident_index;
static uint64_t g_stream_resident_bytes;
static uint64_t g_stream_resident_clock;
static std::vector<cuda_stream_expert_slab> g_stream_expert_slabs;
static std::vector<char *> g_stream_expert_free_slots;
static uint64_t g_stream_expert_slot_bytes;
static uint32_t g_stream_expert_slot_count;
static cuda_stream_cache_stats g_stream_cache_stats;
static cuda_stream_cache_layer_stats
    g_stream_cache_layer_stats[DS4_ROCM_STREAM_CACHE_LAYER_STATS_MAX];
static int g_stream_cache_stats_enabled = -1;
static int g_stream_cache_layer_stats_enabled = -1;
static int g_stream_evict_past_layers_first_enabled = -1;
static int32_t g_routed_moe_selected_override[DS4_ROCM_N_EXPERT_USED];
static uint32_t g_routed_moe_selected_override_n;
static cudaEvent_t g_stream_selected_reuse_event;
static int g_stream_selected_reuse_event_pending;
static cudaEvent_t g_stream_selected_upload_ready_event;
static int g_stream_selected_upload_event_pending;
static cudaEvent_t g_stream_batch_selected_reuse_event;
static int g_stream_batch_selected_reuse_event_pending;
static cudaEvent_t g_stream_batch_selected_upload_ready_event;
static int g_stream_batch_selected_upload_event_pending;
static void *g_stream_read_stage_raw[DS4_ROCM_STREAM_READ_WORKERS];
static uint64_t g_stream_read_stage_bytes[DS4_ROCM_STREAM_READ_WORKERS];
static cudaStream_t g_stream_read_upload_streams[DS4_ROCM_STREAM_READ_WORKERS];
static pthread_t g_stream_read_threads[DS4_ROCM_STREAM_READ_WORKERS];
static uint32_t g_stream_read_thread_ids[DS4_ROCM_STREAM_READ_WORKERS];
static pthread_mutex_t g_stream_read_mutex = PTHREAD_MUTEX_INITIALIZER;
static pthread_cond_t g_stream_read_work_cond = PTHREAD_COND_INITIALIZER;
static pthread_cond_t g_stream_read_done_cond = PTHREAD_COND_INITIALIZER;
static int g_stream_read_pool_started;
static uint32_t g_stream_read_pool_workers;
static int g_stream_read_pool_stop;
static struct cuda_stream_read_job *g_stream_read_active_jobs;
static uint32_t g_stream_read_active_count;
static uint32_t g_stream_read_active_next;
static uint32_t g_stream_read_active_done;
static int g_stream_read_active_ok;
static pthread_t g_stream_read_active_owner;
static int g_stream_read_active_owner_set;
static int g_stream_read_profile_enabled = -1;
static int g_stream_read_direct_disabled = -1;
static int g_stream_read_profile_registered;
static uint64_t g_stream_read_profile_jobs;
static uint64_t g_stream_read_profile_bytes;
static uint64_t g_stream_read_profile_read_us;
static uint64_t g_stream_read_profile_upload_us;
static uint64_t g_stream_read_profile_wait_calls;
static uint64_t g_stream_read_profile_wait_us;
static uint64_t g_stream_read_profile_start_calls;
static uint64_t g_stream_read_profile_exact_groups;
static uint64_t g_stream_read_profile_gap1m_groups;
static uint64_t g_stream_read_profile_gap4m_groups;
static uint64_t g_stream_read_profile_gap1m_extra;
static uint64_t g_stream_read_profile_gap4m_extra;

static int cuda_ok(cudaError_t err, const char *what);
static double cuda_wall_sec(void);
static uint64_t cuda_model_copy_chunk_bytes(void);
static void cuda_model_drop_file_pages(uint64_t offset, uint64_t bytes);
static uint64_t cuda_round_down(uint64_t v, uint64_t align);
static uint64_t cuda_round_up(uint64_t v, uint64_t align);
static int cuda_model_stage_pool_alloc(uint64_t bytes);
static int cuda_pread_full(int fd, void *buf, uint64_t bytes, uint64_t offset);
static int cuda_model_stage_read(void *stage, uint64_t stage_bytes,
                                 uint64_t offset, uint64_t bytes,
                                 const char **payload);
static int cuda_stream_selected_reuse_wait(const char *what);
static int cuda_stream_selected_upload_wait_host(const char *what);
static void cuda_stream_selected_stage_release(void);
static int cuda_stream_batch_selected_reuse_wait(const char *what);
static int cuda_stream_batch_selected_upload_wait_host(const char *what);
static void cuda_stream_batch_selected_stage_release(void);
static void cuda_stream_read_pool_shutdown(void);

static int cuda_u64_mul_checked(uint64_t a, uint64_t b, uint64_t *out) {
    if (!out) return 0;
    if (a != 0u && b > UINT64_MAX / a) return 0;
    *out = a * b;
    return 1;
}

static int cuda_u64_add_checked(uint64_t a, uint64_t b, uint64_t *out) {
    if (!out || a > UINT64_MAX - b) return 0;
    *out = a + b;
    return 1;
}

static int cuda_u64_mul3_checked(uint64_t a, uint64_t b, uint64_t c, uint64_t *out) {
    uint64_t tmp = 0;
    return cuda_u64_mul_checked(a, b, &tmp) && cuda_u64_mul_checked(tmp, c, out);
}

static int cuda_stream_cache_layer_stats_on(void) {
    if (g_stream_cache_layer_stats_enabled < 0) {
        g_stream_cache_layer_stats_enabled =
            getenv("DS4_ROCM_STREAM_CACHE_LAYER_STATS") != NULL ? 1 : 0;
    }
    return g_stream_cache_layer_stats_enabled;
}

static int cuda_stream_cache_stats_on(void) {
    if (g_stream_cache_stats_enabled < 0) {
        g_stream_cache_stats_enabled =
            (getenv("DS4_ROCM_STREAM_CACHE_STATS") != NULL ||
             cuda_stream_cache_layer_stats_on()) ? 1 : 0;
    }
    return g_stream_cache_stats_enabled;
}

static int cuda_stream_evict_past_layers_first(void) {
    if (g_stream_evict_past_layers_first_enabled < 0) {
        const char *env = getenv("DS4_ROCM_STREAM_EVICT_PAST_LAYERS_FIRST");
        g_stream_evict_past_layers_first_enabled =
            (env != NULL && env[0] != '\0' && strcmp(env, "0") != 0) ? 1 : 0;
    }
    return g_stream_evict_past_layers_first_enabled;
}

static void cuda_stream_cache_stats_note_resident(void) {
    if (!cuda_stream_cache_stats_on()) return;
    const uint64_t count = (uint64_t)g_stream_resident_experts.size();
    if (count > g_stream_cache_stats.max_resident_count) {
        g_stream_cache_stats.max_resident_count = count;
    }
    if (g_stream_resident_bytes > g_stream_cache_stats.max_resident_bytes) {
        g_stream_cache_stats.max_resident_bytes = g_stream_resident_bytes;
    }
}

static void cuda_stream_cache_stats_print(const char *label) {
    if (!cuda_stream_cache_stats_on()) return;
    fprintf(stderr,
            DS4_GPU_LOG_PREFIX "stream cache stats %s: "
            "selected calls=%llu slots=%llu hits=%llu misses=%llu; "
            "batch calls=%llu unique=%llu hits=%llu misses=%llu; "
            "seed calls=%llu unique=%llu; "
            "full-layer loads=%llu bytes=%.2f GiB resident-flushes=%llu; "
            "allocs=%llu alloc=%.2f GiB evictions=%llu evicted=%.2f GiB; "
            "resident current=%zu/%.2f GiB max=%llu/%.2f GiB budget=%u\n",
            label ? label : "",
            (unsigned long long)g_stream_cache_stats.selected_calls,
            (unsigned long long)g_stream_cache_stats.selected_slots,
            (unsigned long long)g_stream_cache_stats.selected_hits,
            (unsigned long long)g_stream_cache_stats.selected_misses,
            (unsigned long long)g_stream_cache_stats.batch_calls,
            (unsigned long long)g_stream_cache_stats.batch_unique,
            (unsigned long long)g_stream_cache_stats.batch_hits,
            (unsigned long long)g_stream_cache_stats.batch_misses,
            (unsigned long long)g_stream_cache_stats.seed_calls,
            (unsigned long long)g_stream_cache_stats.seed_unique,
            (unsigned long long)g_stream_cache_stats.layer_loads,
            (double)g_stream_cache_stats.layer_load_bytes / 1073741824.0,
            (unsigned long long)g_stream_cache_stats.layer_resident_flushes,
            (unsigned long long)g_stream_cache_stats.allocs,
            (double)g_stream_cache_stats.alloc_bytes / 1073741824.0,
            (unsigned long long)g_stream_cache_stats.evictions,
            (double)g_stream_cache_stats.evict_bytes / 1073741824.0,
            g_stream_resident_experts.size(),
            (double)g_stream_resident_bytes / 1073741824.0,
            (unsigned long long)g_stream_cache_stats.max_resident_count,
            (double)g_stream_cache_stats.max_resident_bytes / 1073741824.0,
            g_stream_expert_cache_budget);
    if (!cuda_stream_cache_layer_stats_on()) return;
    for (uint32_t layer = 0;
         layer < DS4_ROCM_STREAM_CACHE_LAYER_STATS_MAX;
         layer++) {
        const cuda_stream_cache_layer_stats *s =
            &g_stream_cache_layer_stats[layer];
        if (s->selected_calls == 0 && s->batch_calls == 0) continue;
        const double selected_hit_pct =
            s->selected_slots ?
                100.0 * (double)s->selected_hits /
                    (double)s->selected_slots : 0.0;
        const double batch_hit_pct =
            s->batch_unique ?
                100.0 * (double)s->batch_hits /
                    (double)s->batch_unique : 0.0;
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "stream cache layer %u %s: "
                "selected calls=%llu slots=%llu hits=%llu misses=%llu hit=%.1f%%; "
                "batch calls=%llu unique=%llu hits=%llu misses=%llu hit=%.1f%%\n",
                layer,
                label ? label : "",
                (unsigned long long)s->selected_calls,
                (unsigned long long)s->selected_slots,
                (unsigned long long)s->selected_hits,
                (unsigned long long)s->selected_misses,
                selected_hit_pct,
                (unsigned long long)s->batch_calls,
                (unsigned long long)s->batch_unique,
                (unsigned long long)s->batch_hits,
                (unsigned long long)s->batch_misses,
                batch_hit_pct);
    }
}

static int cuda_model_range_fits(uint64_t model_size, uint64_t offset, uint64_t bytes) {
    return offset <= model_size && bytes <= model_size - offset;
}

static int cuda_tensor_has_bytes(const ds4_gpu_tensor *t, uint64_t bytes) {
    return t && t->ptr && t->bytes >= bytes;
}

static int cuda_tensor_has_elems(const ds4_gpu_tensor *t, uint64_t elems, uint64_t elem_size) {
    uint64_t bytes = 0;
    return cuda_u64_mul_checked(elems, elem_size, &bytes) && cuda_tensor_has_bytes(t, bytes);
}

static int cuda_tensor_has_elems2(const ds4_gpu_tensor *t, uint64_t a, uint64_t b, uint64_t elem_size) {
    uint64_t bytes = 0;
    return cuda_u64_mul3_checked(a, b, elem_size, &bytes) && cuda_tensor_has_bytes(t, bytes);
}

static int cuda_tensor_has_elems3(const ds4_gpu_tensor *t, uint64_t a, uint64_t b, uint64_t c, uint64_t elem_size) {
    uint64_t ab = 0, elems = 0, bytes = 0;
    return cuda_u64_mul_checked(a, b, &ab) &&
           cuda_u64_mul_checked(ab, c, &elems) &&
           cuda_u64_mul_checked(elems, elem_size, &bytes) &&
           cuda_tensor_has_bytes(t, bytes);
}

static int cuda_tensor_has_f32(const ds4_gpu_tensor *t, uint64_t elems) {
    return cuda_tensor_has_elems(t, elems, sizeof(float));
}

static int cuda_tensor_has_i32(const ds4_gpu_tensor *t, uint64_t elems) {
    return cuda_tensor_has_elems(t, elems, sizeof(int32_t));
}

static int cuda_tensor_has_f16(const ds4_gpu_tensor *t, uint64_t elems) {
    return cuda_tensor_has_elems(t, elems, sizeof(__half));
}

static int cuda_tensor_has_u16(const ds4_gpu_tensor *t, uint64_t elems) {
    return cuda_tensor_has_elems(t, elems, sizeof(uint16_t));
}

static const char *cuda_model_range_ptr_from_fd(
        const void *model_map,
        uint64_t offset,
        uint64_t bytes,
        const char *what);
static const char *cuda_model_range_ptr(
        const void *model_map,
        uint64_t offset,
        uint64_t bytes,
        const char *what);
static int cuda_model_range_is_cached(
        const void *model_map,
        uint64_t offset,
        uint64_t bytes);
__global__ static void dequant_q8_0_to_f16_kernel(
        __half *out,
        const unsigned char *w,
        uint64_t in_dim,
        uint64_t out_dim,
        uint64_t blocks);
__global__ static void dequant_q8_0_to_f32_kernel(
        float *out,
        const unsigned char *w,
        uint64_t in_dim,
        uint64_t out_dim,
        uint64_t blocks);
__global__ static void dequant_q8_0_to_f16_transpose_kernel(
        __half *out,
        const unsigned char *w,
        uint64_t in_dim,
        uint64_t out_dim,
        uint64_t blocks);

static void cuda_shared_gate_up_async_cleanup(void);

static void *cuda_tmp_alloc(uint64_t bytes, const char *what) {
    if (bytes == 0) return NULL;
    if (g_cuda_tmp_bytes >= bytes) return g_cuda_tmp;
    if (g_cuda_tmp) {
        (void)cudaFree(g_cuda_tmp);
        g_cuda_tmp = NULL;
        g_cuda_tmp_bytes = 0;
    }
    void *ptr = NULL;
    cudaError_t err = cudaMalloc(&ptr, (size_t)bytes);
    if (err != cudaSuccess) {
        fprintf(stderr, DS4_GPU_LOG_PREFIX "temp alloc failed for %s (%.2f MiB): %s\n",
                what ? what : "scratch", (double)bytes / 1048576.0, cudaGetErrorString(err));
        (void)cudaGetLastError();
        return NULL;
    }
    g_cuda_tmp = ptr;
    g_cuda_tmp_bytes = bytes;
    return g_cuda_tmp;
}

static int cuda_attention_score_buffer_fits(uint32_t n_comp) {
    return n_comp <= DS4_ROCM_ATTENTION_SCORE_CAP - DS4_ROCM_ATTENTION_RAW_SCORE_CAP;
}

static int cuda_model_image_find(const void *model_map) {
    if (!model_map) return -1;
    for (size_t i = 0; i < g_model_images.size(); i++) {
        if (g_model_images[i].host_base == model_map) return (int)i;
    }
    return -1;
}

static const char *cuda_model_image_ptr(const void *model_map, uint64_t offset) {
    const int idx = cuda_model_image_find(model_map);
    if (idx < 0) return NULL;
    const cuda_model_image &img = g_model_images[(size_t)idx];
    if (offset > img.size) return NULL;
    return img.device_ptr + offset;
}

static int cuda_model_image_owned(const void *model_map) {
    return cuda_model_image_find(model_map) >= 0;
}

static uint64_t cuda_model_image_bytes(void) {
    uint64_t bytes = 0;
    for (const cuda_model_image &img : g_model_images) bytes += img.size;
    return bytes;
}

static void cuda_model_image_release_all(void) {
    for (const cuda_model_image &img : g_model_images) {
        if (img.device_ptr) (void)cudaFree(img.device_ptr);
    }
    g_model_images.clear();
}

static int cuda_stream_resident_reclaim_wait(const char *what) {
    if (!cuda_stream_selected_reuse_wait(what)) return 0;
    if (!cuda_stream_batch_selected_reuse_wait(what)) return 0;
    return 1;
}

static void cuda_stream_resident_cache_release(void) {
    if (!cuda_stream_resident_reclaim_wait(
                "streaming resident expert cache release")) {
        return;
    }
    for (cuda_stream_resident_expert &e : g_stream_resident_experts) {
        if (e.base && !e.pooled) (void)cudaFree(e.base);
    }
    g_stream_resident_experts.clear();
    g_stream_resident_index.clear();
    g_stream_resident_bytes = 0;
    g_stream_resident_clock = 0;
    for (cuda_stream_expert_slab &slab : g_stream_expert_slabs) {
        if (slab.base) (void)cudaFree(slab.base);
    }
    g_stream_expert_slabs.clear();
    g_stream_expert_free_slots.clear();
    g_stream_expert_slot_bytes = 0;
    g_stream_expert_slot_count = 0;
}

static void cuda_stream_layer_expert_cache_release(void) {
    bool any_active = false;
    for (uint32_t i = 0; i < 2u; i++) {
        if (g_stream_layer_expert_cache[i].base) {
            any_active = true;
            break;
        }
    }
    if (any_active) {
        cudaError_t sync_err = cudaDeviceSynchronize();
        if (sync_err != cudaSuccess) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming full-layer expert cache "
                    "release sync failed: %s\n",
                    cudaGetErrorString(sync_err));
            (void)cudaGetLastError();
        }
    }
    for (uint32_t i = 0; i < 2u; i++) {
        cuda_stream_layer_expert_cache &c = g_stream_layer_expert_cache[i];
        if (c.base) (void)cudaFree(c.base);
        memset(&c, 0, sizeof(c));
    }
}

static void cuda_stream_read_stage_release(void) {
    cuda_stream_read_pool_shutdown();
    for (uint32_t i = 0; i < DS4_ROCM_STREAM_READ_WORKERS; i++) {
        if (g_stream_read_stage_raw[i]) {
            (void)cudaFreeHost(g_stream_read_stage_raw[i]);
            g_stream_read_stage_raw[i] = NULL;
            g_stream_read_stage_bytes[i] = 0;
        }
    }
    for (uint32_t i = 0; i < DS4_ROCM_STREAM_READ_WORKERS; i++) {
        if (g_stream_read_upload_streams[i]) {
            (void)cudaStreamDestroy(g_stream_read_upload_streams[i]);
            g_stream_read_upload_streams[i] = NULL;
        }
    }
}

static void cuda_stream_batch_selected_cache_release(void) {
    (void)cuda_stream_batch_selected_reuse_wait(
            "streaming batch selected cache release");
    (void)cuda_stream_batch_selected_upload_wait_host(
            "streaming batch selected cache release upload");
    if (g_stream_batch_selected_cache.selected_ids) {
        (void)cudaFree(g_stream_batch_selected_cache.selected_ids);
    }
    if (g_stream_batch_selected_cache.pair_missing) {
        (void)cudaFree(g_stream_batch_selected_cache.pair_missing);
    }
    if (g_stream_batch_selected_cache.gate_ptrs) {
        (void)cudaFree(g_stream_batch_selected_cache.gate_ptrs);
    }
    if (g_stream_batch_selected_cache.up_ptrs) {
        (void)cudaFree(g_stream_batch_selected_cache.up_ptrs);
    }
    if (g_stream_batch_selected_cache.down_ptrs) {
        (void)cudaFree(g_stream_batch_selected_cache.down_ptrs);
    }
    if (g_stream_batch_selected_cache.resident_gate_ptrs) {
        (void)cudaFree(g_stream_batch_selected_cache.resident_gate_ptrs);
    }
    if (g_stream_batch_selected_cache.resident_up_ptrs) {
        (void)cudaFree(g_stream_batch_selected_cache.resident_up_ptrs);
    }
    if (g_stream_batch_selected_cache.missing_gate_ptrs) {
        (void)cudaFree(g_stream_batch_selected_cache.missing_gate_ptrs);
    }
    if (g_stream_batch_selected_cache.missing_up_ptrs) {
        (void)cudaFree(g_stream_batch_selected_cache.missing_up_ptrs);
    }
    if (g_stream_batch_selected_reuse_event) {
        (void)cudaEventDestroy(g_stream_batch_selected_reuse_event);
        g_stream_batch_selected_reuse_event = NULL;
    }
    if (g_stream_batch_selected_upload_ready_event) {
        (void)cudaEventDestroy(g_stream_batch_selected_upload_ready_event);
        g_stream_batch_selected_upload_ready_event = NULL;
    }
    g_stream_batch_selected_reuse_event_pending = 0;
    g_stream_batch_selected_upload_event_pending = 0;
    cuda_stream_batch_selected_stage_release();
    memset(&g_stream_batch_selected_cache, 0, sizeof(g_stream_batch_selected_cache));
}

static void cuda_stream_selected_cache_release(void) {
    (void)cuda_stream_selected_reuse_wait("streaming selected cache release");
    (void)cuda_stream_selected_upload_wait_host(
            "streaming selected cache release upload");
    if (g_stream_selected_cache.gate) (void)cudaFree(g_stream_selected_cache.gate);
    if (g_stream_selected_cache.up) (void)cudaFree(g_stream_selected_cache.up);
    if (g_stream_selected_cache.down) (void)cudaFree(g_stream_selected_cache.down);
    if (g_stream_selected_cache.slot_ids) (void)cudaFree(g_stream_selected_cache.slot_ids);
    if (g_stream_selected_cache.gate_ptrs) (void)cudaFree(g_stream_selected_cache.gate_ptrs);
    if (g_stream_selected_cache.up_ptrs) (void)cudaFree(g_stream_selected_cache.up_ptrs);
    if (g_stream_selected_cache.down_ptrs) (void)cudaFree(g_stream_selected_cache.down_ptrs);
    if (g_stream_selected_reuse_event) {
        (void)cudaEventDestroy(g_stream_selected_reuse_event);
        g_stream_selected_reuse_event = NULL;
    }
    if (g_stream_selected_upload_ready_event) {
        (void)cudaEventDestroy(g_stream_selected_upload_ready_event);
        g_stream_selected_upload_ready_event = NULL;
    }
    g_stream_selected_reuse_event_pending = 0;
    g_stream_selected_upload_event_pending = 0;
    cuda_stream_selected_stage_release();
    memset(&g_stream_selected_cache, 0, sizeof(g_stream_selected_cache));
    cuda_stream_batch_selected_cache_release();
    cuda_stream_resident_cache_release();
    cuda_stream_layer_expert_cache_release();
    cuda_stream_read_stage_release();
    g_routed_moe_selected_override_n = 0;
}

static int cuda_stream_selected_ensure_stream(void) {
    if (g_stream_selected_upload_stream) return 1;
    cudaError_t err = cudaStreamCreateWithFlags(&g_stream_selected_upload_stream, cudaStreamNonBlocking);
    if (err != cudaSuccess) {
        fprintf(stderr, DS4_GPU_LOG_PREFIX "streaming selected upload stream creation failed: %s\n",
                cudaGetErrorString(err));
        (void)cudaGetLastError();
        return 0;
    }
    return 1;
}

static int cuda_stream_selected_reuse_ensure_event(void) {
    if (g_stream_selected_reuse_event) return 1;
    cudaError_t err =
        cudaEventCreateWithFlags(&g_stream_selected_reuse_event,
                                 cudaEventDisableTiming);
    if (err != cudaSuccess) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "streaming selected reuse event creation failed: %s\n",
                cudaGetErrorString(err));
        (void)cudaGetLastError();
        return 0;
    }
    return 1;
}

static int cuda_stream_selected_reuse_wait(const char *what) {
    if (!g_stream_selected_reuse_event_pending) return 1;
    cudaError_t err = cudaEventSynchronize(g_stream_selected_reuse_event);
    if (err != cudaSuccess) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "%s wait failed: %s\n",
                what ? what : "streaming selected cache reuse",
                cudaGetErrorString(err));
        (void)cudaGetLastError();
        return 0;
    }
    g_stream_selected_reuse_event_pending = 0;
    return 1;
}

static int cuda_stream_selected_upload_ensure_event(void) {
    if (g_stream_selected_upload_ready_event) return 1;
    cudaError_t err =
        cudaEventCreateWithFlags(&g_stream_selected_upload_ready_event,
                                 cudaEventDisableTiming);
    if (err != cudaSuccess) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "streaming selected upload event creation failed: %s\n",
                cudaGetErrorString(err));
        (void)cudaGetLastError();
        return 0;
    }
    return 1;
}

static int cuda_stream_selected_upload_wait_host(const char *what) {
    if (!g_stream_selected_upload_event_pending) return 1;
    cudaError_t err = cudaEventSynchronize(g_stream_selected_upload_ready_event);
    if (err != cudaSuccess) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "%s wait failed: %s\n",
                what ? what : "streaming selected upload",
                cudaGetErrorString(err));
        (void)cudaGetLastError();
        return 0;
    }
    g_stream_selected_upload_event_pending = 0;
    return 1;
}

static int cuda_stream_selected_upload_record_ready(void) {
    if (!cuda_stream_selected_upload_ensure_event()) return 0;
    cudaError_t err = cudaEventRecord(g_stream_selected_upload_ready_event,
                                      g_stream_selected_upload_stream);
    if (err != cudaSuccess) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "streaming selected upload event record failed: %s\n",
                cudaGetErrorString(err));
        (void)cudaGetLastError();
        return 0;
    }
    g_stream_selected_upload_event_pending = 1;
    return 1;
}

static int cuda_stream_selected_wait_upload_ready(void) {
    if (!g_stream_selected_upload_event_pending) return 1;
#ifdef __HIP_PLATFORM_AMD__
    cudaError_t err =
        hipStreamWaitEvent(0, g_stream_selected_upload_ready_event, 0);
#else
    cudaError_t err =
        cudaStreamWaitEvent(0, g_stream_selected_upload_ready_event, 0);
#endif
    if (err != cudaSuccess) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "streaming selected upload stream wait failed: %s\n",
                cudaGetErrorString(err));
        (void)cudaGetLastError();
        return 0;
    }
    return 1;
}

static int cuda_stream_selected_mark_inflight(void) {
    if (!g_ssd_streaming_mode) return 1;
    if (!cuda_stream_selected_reuse_ensure_event()) return 0;
    cudaError_t err = cudaEventRecord(g_stream_selected_reuse_event, 0);
    if (err != cudaSuccess) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "streaming selected reuse event record failed: %s\n",
                cudaGetErrorString(err));
        (void)cudaGetLastError();
        return 0;
    }
    g_stream_selected_reuse_event_pending = 1;
    g_stream_selected_upload_event_pending = 0;
    return 1;
}

static int cuda_stream_batch_selected_reuse_ensure_event(void) {
    if (g_stream_batch_selected_reuse_event) return 1;
    cudaError_t err =
        cudaEventCreateWithFlags(&g_stream_batch_selected_reuse_event,
                                 cudaEventDisableTiming);
    if (err != cudaSuccess) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "streaming batch selected reuse event creation failed: %s\n",
                cudaGetErrorString(err));
        (void)cudaGetLastError();
        return 0;
    }
    return 1;
}

static int cuda_stream_batch_selected_upload_ensure_event(void) {
    if (g_stream_batch_selected_upload_ready_event) return 1;
    cudaError_t err =
        cudaEventCreateWithFlags(&g_stream_batch_selected_upload_ready_event,
                                 cudaEventDisableTiming);
    if (err != cudaSuccess) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "streaming batch selected upload event creation failed: %s\n",
                cudaGetErrorString(err));
        (void)cudaGetLastError();
        return 0;
    }
    return 1;
}

static int cuda_stream_batch_selected_upload_wait_host(const char *what) {
    if (!g_stream_batch_selected_upload_event_pending) return 1;
    cudaError_t err =
        cudaEventSynchronize(g_stream_batch_selected_upload_ready_event);
    if (err != cudaSuccess) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "%s wait failed: %s\n",
                what ? what : "streaming batch selected upload",
                cudaGetErrorString(err));
        (void)cudaGetLastError();
        return 0;
    }
    g_stream_batch_selected_upload_event_pending = 0;
    return 1;
}

static int cuda_stream_batch_selected_upload_record_ready(void) {
    if (!cuda_stream_batch_selected_upload_ensure_event()) return 0;
    cudaError_t err = cudaEventRecord(g_stream_batch_selected_upload_ready_event,
                                      g_stream_selected_upload_stream);
    if (err != cudaSuccess) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "streaming batch selected upload event record failed: %s\n",
                cudaGetErrorString(err));
        (void)cudaGetLastError();
        return 0;
    }
    g_stream_batch_selected_upload_event_pending = 1;
    return 1;
}

static int cuda_stream_batch_selected_wait_upload_ready(void) {
    if (!g_stream_batch_selected_upload_event_pending) return 1;
#ifdef __HIP_PLATFORM_AMD__
    cudaError_t err =
        hipStreamWaitEvent(0, g_stream_batch_selected_upload_ready_event, 0);
#else
    cudaError_t err =
        cudaStreamWaitEvent(0, g_stream_batch_selected_upload_ready_event, 0);
#endif
    if (err != cudaSuccess) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "streaming batch selected upload stream wait failed: %s\n",
                cudaGetErrorString(err));
        (void)cudaGetLastError();
        return 0;
    }
    return 1;
}

static int cuda_stream_batch_selected_reuse_wait(const char *what) {
    if (!g_stream_batch_selected_reuse_event_pending) return 1;
    cudaError_t err = cudaEventSynchronize(g_stream_batch_selected_reuse_event);
    if (err != cudaSuccess) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "%s wait failed: %s\n",
                what ? what : "streaming batch selected cache reuse",
                cudaGetErrorString(err));
        (void)cudaGetLastError();
        return 0;
    }
    g_stream_batch_selected_reuse_event_pending = 0;
    return 1;
}

static int cuda_stream_batch_selected_mark_inflight(void) {
    if (!g_ssd_streaming_mode) return 1;
    if (!cuda_stream_batch_selected_reuse_ensure_event()) return 0;
    cudaError_t err = cudaEventRecord(g_stream_batch_selected_reuse_event, 0);
    if (err != cudaSuccess) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "streaming batch selected reuse event record failed: %s\n",
                cudaGetErrorString(err));
        (void)cudaGetLastError();
        return 0;
    }
    g_stream_batch_selected_reuse_event_pending = 1;
    g_stream_batch_selected_upload_event_pending = 0;
    return 1;
}

static void cuda_stream_selected_stage_release(void) {
    if (g_stream_selected_cache.gate_ptrs_stage) {
        (void)cudaFreeHost((void *)g_stream_selected_cache.gate_ptrs_stage);
    }
    if (g_stream_selected_cache.up_ptrs_stage) {
        (void)cudaFreeHost((void *)g_stream_selected_cache.up_ptrs_stage);
    }
    if (g_stream_selected_cache.down_ptrs_stage) {
        (void)cudaFreeHost((void *)g_stream_selected_cache.down_ptrs_stage);
    }
    g_stream_selected_cache.gate_ptrs_stage = NULL;
    g_stream_selected_cache.up_ptrs_stage = NULL;
    g_stream_selected_cache.down_ptrs_stage = NULL;
}

static int cuda_stream_selected_ensure_buffers(uint64_t gate_bytes, uint64_t down_bytes) {
    if (gate_bytes == 0 || down_bytes == 0) return 0;
    if (!cuda_stream_selected_upload_wait_host(
                "streaming selected cache upload")) {
        return 0;
    }
    cudaError_t err = cudaSuccess;
    if (g_stream_selected_cache.gate_capacity < gate_bytes) {
        if (g_stream_selected_cache.gate) (void)cudaFree(g_stream_selected_cache.gate);
        if (g_stream_selected_cache.up) (void)cudaFree(g_stream_selected_cache.up);
        g_stream_selected_cache.gate = NULL;
        g_stream_selected_cache.up = NULL;
        g_stream_selected_cache.gate_capacity = 0;
        err = cudaMalloc((void **)&g_stream_selected_cache.gate, (size_t)gate_bytes);
        if (err == cudaSuccess) {
            err = cudaMalloc((void **)&g_stream_selected_cache.up, (size_t)gate_bytes);
        }
        if (err != cudaSuccess) {
            fprintf(stderr, DS4_GPU_LOG_PREFIX "streaming selected gate/up alloc failed (%.2f MiB): %s\n",
                    (double)gate_bytes / 1048576.0,
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            return 0;
        }
        g_stream_selected_cache.gate_capacity = gate_bytes;
    }
    if (g_stream_selected_cache.down_capacity < down_bytes) {
        if (g_stream_selected_cache.down) (void)cudaFree(g_stream_selected_cache.down);
        g_stream_selected_cache.down = NULL;
        g_stream_selected_cache.down_capacity = 0;
        err = cudaMalloc((void **)&g_stream_selected_cache.down, (size_t)down_bytes);
        if (err != cudaSuccess) {
            fprintf(stderr, DS4_GPU_LOG_PREFIX "streaming selected down alloc failed (%.2f MiB): %s\n",
                    (double)down_bytes / 1048576.0,
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            return 0;
        }
        g_stream_selected_cache.down_capacity = down_bytes;
    }
    if (!g_stream_selected_cache.slot_ids) {
        err = cudaMalloc((void **)&g_stream_selected_cache.slot_ids,
                         DS4_ROCM_N_EXPERT_USED * sizeof(int32_t));
        if (err != cudaSuccess) {
            fprintf(stderr, DS4_GPU_LOG_PREFIX "streaming selected slot-id alloc failed: %s\n",
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            return 0;
        }
        int32_t slots[DS4_ROCM_N_EXPERT_USED];
        for (uint32_t i = 0; i < DS4_ROCM_N_EXPERT_USED; i++) slots[i] = (int32_t)i;
        err = cudaMemcpy(g_stream_selected_cache.slot_ids,
                         slots,
                         sizeof(slots),
                         cudaMemcpyHostToDevice);
        if (err != cudaSuccess) {
            fprintf(stderr, DS4_GPU_LOG_PREFIX "streaming selected slot-id upload failed: %s\n",
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            return 0;
        }
        g_stream_selected_cache.slot_tensor.ptr = g_stream_selected_cache.slot_ids;
        g_stream_selected_cache.slot_tensor.bytes =
            DS4_ROCM_N_EXPERT_USED * sizeof(int32_t);
        g_stream_selected_cache.slot_tensor.owner = 0;
    }
    if (!g_stream_selected_cache.gate_ptrs) {
        err = cudaMalloc((void **)&g_stream_selected_cache.gate_ptrs,
                         DS4_ROCM_N_EXPERT_USED * sizeof(char *));
        if (err == cudaSuccess) {
            err = cudaMalloc((void **)&g_stream_selected_cache.up_ptrs,
                             DS4_ROCM_N_EXPERT_USED * sizeof(char *));
        }
        if (err == cudaSuccess) {
            err = cudaMalloc((void **)&g_stream_selected_cache.down_ptrs,
                             DS4_ROCM_N_EXPERT_USED * sizeof(char *));
        }
        if (err != cudaSuccess) {
            fprintf(stderr, DS4_GPU_LOG_PREFIX "streaming selected pointer table alloc failed: %s\n",
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            return 0;
        }
    }
    if (!g_stream_selected_cache.gate_ptrs_stage) {
        void *stage = NULL;
        err = cudaMallocHost(&stage, DS4_ROCM_N_EXPERT_USED * sizeof(char *));
        g_stream_selected_cache.gate_ptrs_stage =
            err == cudaSuccess ? (const char **)stage : NULL;
        stage = NULL;
        if (err == cudaSuccess) {
            err = cudaMallocHost(&stage, DS4_ROCM_N_EXPERT_USED * sizeof(char *));
            g_stream_selected_cache.up_ptrs_stage =
                err == cudaSuccess ? (const char **)stage : NULL;
            stage = NULL;
        }
        if (err == cudaSuccess) {
            err = cudaMallocHost(&stage, DS4_ROCM_N_EXPERT_USED * sizeof(char *));
            g_stream_selected_cache.down_ptrs_stage =
                err == cudaSuccess ? (const char **)stage : NULL;
        }
        if (err != cudaSuccess) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming selected pointer staging alloc failed: %s\n",
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            cuda_stream_selected_stage_release();
            return 0;
        }
    }
    return 1;
}

static void cuda_stream_batch_selected_stage_release(void) {
    if (g_stream_batch_selected_cache.selected_stage) {
        (void)cudaFreeHost(g_stream_batch_selected_cache.selected_stage);
    }
    if (g_stream_batch_selected_cache.pair_missing_stage) {
        (void)cudaFreeHost(g_stream_batch_selected_cache.pair_missing_stage);
    }
    if (g_stream_batch_selected_cache.gate_ptrs_stage) {
        (void)cudaFreeHost((void *)g_stream_batch_selected_cache.gate_ptrs_stage);
    }
    if (g_stream_batch_selected_cache.up_ptrs_stage) {
        (void)cudaFreeHost((void *)g_stream_batch_selected_cache.up_ptrs_stage);
    }
    if (g_stream_batch_selected_cache.down_ptrs_stage) {
        (void)cudaFreeHost((void *)g_stream_batch_selected_cache.down_ptrs_stage);
    }
    if (g_stream_batch_selected_cache.resident_gate_ptrs_stage) {
        (void)cudaFreeHost((void *)g_stream_batch_selected_cache.resident_gate_ptrs_stage);
    }
    if (g_stream_batch_selected_cache.resident_up_ptrs_stage) {
        (void)cudaFreeHost((void *)g_stream_batch_selected_cache.resident_up_ptrs_stage);
    }
    if (g_stream_batch_selected_cache.missing_gate_ptrs_stage) {
        (void)cudaFreeHost((void *)g_stream_batch_selected_cache.missing_gate_ptrs_stage);
    }
    if (g_stream_batch_selected_cache.missing_up_ptrs_stage) {
        (void)cudaFreeHost((void *)g_stream_batch_selected_cache.missing_up_ptrs_stage);
    }
    g_stream_batch_selected_cache.selected_stage = NULL;
    g_stream_batch_selected_cache.selected_stage_capacity = 0;
    g_stream_batch_selected_cache.pair_missing_stage = NULL;
    g_stream_batch_selected_cache.pair_missing_stage_capacity = 0;
    g_stream_batch_selected_cache.gate_ptrs_stage = NULL;
    g_stream_batch_selected_cache.up_ptrs_stage = NULL;
    g_stream_batch_selected_cache.down_ptrs_stage = NULL;
    g_stream_batch_selected_cache.resident_gate_ptrs_stage = NULL;
    g_stream_batch_selected_cache.resident_up_ptrs_stage = NULL;
    g_stream_batch_selected_cache.missing_gate_ptrs_stage = NULL;
    g_stream_batch_selected_cache.missing_up_ptrs_stage = NULL;
    g_stream_batch_selected_cache.ptr_stage_capacity = 0;
}

static int cuda_stream_batch_selected_ensure_buffers(
        uint64_t n_ids,
        uint32_t n_unique) {
    if (n_ids == 0 || n_unique == 0) return 0;
    if (!cuda_stream_batch_selected_reuse_wait(
                "streaming batch selected cache reuse")) {
        return 0;
    }
    if (!cuda_stream_batch_selected_upload_wait_host(
                "streaming batch selected cache upload")) {
        return 0;
    }
    cudaError_t err = cudaSuccess;
    const uint64_t selected_bytes = n_ids * sizeof(int32_t);
    if (g_stream_batch_selected_cache.selected_capacity < selected_bytes) {
        if (g_stream_batch_selected_cache.selected_ids) {
            (void)cudaFree(g_stream_batch_selected_cache.selected_ids);
            g_stream_batch_selected_cache.selected_ids = NULL;
            g_stream_batch_selected_cache.selected_capacity = 0;
        }
        err = cudaMalloc((void **)&g_stream_batch_selected_cache.selected_ids,
                         (size_t)selected_bytes);
        if (err != cudaSuccess) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming batch selected-id alloc failed "
                    "(%.2f MiB): %s\n",
                    (double)selected_bytes / 1048576.0,
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            return 0;
        }
        g_stream_batch_selected_cache.selected_capacity = selected_bytes;
    }
    if (g_stream_batch_selected_cache.pair_missing_capacity < n_ids) {
        if (g_stream_batch_selected_cache.pair_missing) {
            (void)cudaFree(g_stream_batch_selected_cache.pair_missing);
            g_stream_batch_selected_cache.pair_missing = NULL;
            g_stream_batch_selected_cache.pair_missing_capacity = 0;
        }
        err = cudaMalloc((void **)&g_stream_batch_selected_cache.pair_missing,
                         (size_t)n_ids);
        if (err != cudaSuccess) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming batch selected split-flag alloc failed "
                    "(%.2f MiB): %s\n",
                    (double)n_ids / 1048576.0,
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            return 0;
        }
        g_stream_batch_selected_cache.pair_missing_capacity = n_ids;
    }
    if (g_stream_batch_selected_cache.ptr_capacity < n_unique) {
        if (g_stream_batch_selected_cache.gate_ptrs) {
            (void)cudaFree(g_stream_batch_selected_cache.gate_ptrs);
            (void)cudaFree(g_stream_batch_selected_cache.up_ptrs);
            (void)cudaFree(g_stream_batch_selected_cache.down_ptrs);
            (void)cudaFree(g_stream_batch_selected_cache.resident_gate_ptrs);
            (void)cudaFree(g_stream_batch_selected_cache.resident_up_ptrs);
            (void)cudaFree(g_stream_batch_selected_cache.missing_gate_ptrs);
            (void)cudaFree(g_stream_batch_selected_cache.missing_up_ptrs);
            g_stream_batch_selected_cache.gate_ptrs = NULL;
            g_stream_batch_selected_cache.up_ptrs = NULL;
            g_stream_batch_selected_cache.down_ptrs = NULL;
            g_stream_batch_selected_cache.resident_gate_ptrs = NULL;
            g_stream_batch_selected_cache.resident_up_ptrs = NULL;
            g_stream_batch_selected_cache.missing_gate_ptrs = NULL;
            g_stream_batch_selected_cache.missing_up_ptrs = NULL;
            g_stream_batch_selected_cache.ptr_capacity = 0;
        }
        err = cudaMalloc((void **)&g_stream_batch_selected_cache.gate_ptrs,
                         (size_t)n_unique * sizeof(char *));
        if (err == cudaSuccess) {
            err = cudaMalloc((void **)&g_stream_batch_selected_cache.up_ptrs,
                             (size_t)n_unique * sizeof(char *));
        }
        if (err == cudaSuccess) {
            err = cudaMalloc((void **)&g_stream_batch_selected_cache.down_ptrs,
                             (size_t)n_unique * sizeof(char *));
        }
        if (err == cudaSuccess) {
            err = cudaMalloc((void **)&g_stream_batch_selected_cache.resident_gate_ptrs,
                             (size_t)n_unique * sizeof(char *));
        }
        if (err == cudaSuccess) {
            err = cudaMalloc((void **)&g_stream_batch_selected_cache.resident_up_ptrs,
                             (size_t)n_unique * sizeof(char *));
        }
        if (err == cudaSuccess) {
            err = cudaMalloc((void **)&g_stream_batch_selected_cache.missing_gate_ptrs,
                             (size_t)n_unique * sizeof(char *));
        }
        if (err == cudaSuccess) {
            err = cudaMalloc((void **)&g_stream_batch_selected_cache.missing_up_ptrs,
                             (size_t)n_unique * sizeof(char *));
        }
        if (err != cudaSuccess) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming batch pointer-table alloc failed "
                    "(unique=%u): %s\n",
                    n_unique,
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            return 0;
        }
        g_stream_batch_selected_cache.ptr_capacity = n_unique;
    }
    if (g_stream_batch_selected_cache.selected_stage_capacity < selected_bytes ||
        g_stream_batch_selected_cache.pair_missing_stage_capacity < n_ids ||
        g_stream_batch_selected_cache.ptr_stage_capacity < n_unique) {
        cuda_stream_batch_selected_stage_release();
        err = cudaMallocHost((void **)&g_stream_batch_selected_cache.selected_stage,
                             (size_t)selected_bytes);
        if (err == cudaSuccess) {
            err = cudaMallocHost((void **)&g_stream_batch_selected_cache.pair_missing_stage,
                                 (size_t)n_ids);
        }
        void *stage = NULL;
        if (err == cudaSuccess) {
            err = cudaMallocHost(&stage, (size_t)n_unique * sizeof(char *));
            g_stream_batch_selected_cache.gate_ptrs_stage =
                err == cudaSuccess ? (const char **)stage : NULL;
            stage = NULL;
        }
        if (err == cudaSuccess) {
            err = cudaMallocHost(&stage, (size_t)n_unique * sizeof(char *));
            g_stream_batch_selected_cache.up_ptrs_stage =
                err == cudaSuccess ? (const char **)stage : NULL;
            stage = NULL;
        }
        if (err == cudaSuccess) {
            err = cudaMallocHost(&stage, (size_t)n_unique * sizeof(char *));
            g_stream_batch_selected_cache.down_ptrs_stage =
                err == cudaSuccess ? (const char **)stage : NULL;
            stage = NULL;
        }
        if (err == cudaSuccess) {
            err = cudaMallocHost(&stage, (size_t)n_unique * sizeof(char *));
            g_stream_batch_selected_cache.resident_gate_ptrs_stage =
                err == cudaSuccess ? (const char **)stage : NULL;
            stage = NULL;
        }
        if (err == cudaSuccess) {
            err = cudaMallocHost(&stage, (size_t)n_unique * sizeof(char *));
            g_stream_batch_selected_cache.resident_up_ptrs_stage =
                err == cudaSuccess ? (const char **)stage : NULL;
            stage = NULL;
        }
        if (err == cudaSuccess) {
            err = cudaMallocHost(&stage, (size_t)n_unique * sizeof(char *));
            g_stream_batch_selected_cache.missing_gate_ptrs_stage =
                err == cudaSuccess ? (const char **)stage : NULL;
            stage = NULL;
        }
        if (err == cudaSuccess) {
            err = cudaMallocHost(&stage, (size_t)n_unique * sizeof(char *));
            g_stream_batch_selected_cache.missing_up_ptrs_stage =
                err == cudaSuccess ? (const char **)stage : NULL;
        }
        if (err != cudaSuccess) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming batch selected staging alloc failed "
                    "(ids=%.2f MiB unique=%u): %s\n",
                    (double)selected_bytes / 1048576.0,
                    n_unique,
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            cuda_stream_batch_selected_stage_release();
            return 0;
        }
        g_stream_batch_selected_cache.selected_stage_capacity = selected_bytes;
        g_stream_batch_selected_cache.pair_missing_stage_capacity = n_ids;
        g_stream_batch_selected_cache.ptr_stage_capacity = n_unique;
    }
    g_stream_batch_selected_cache.selected_tensor.ptr =
        g_stream_batch_selected_cache.selected_ids;
    g_stream_batch_selected_cache.selected_tensor.bytes = selected_bytes;
    g_stream_batch_selected_cache.selected_tensor.owner = 0;
    return 1;
}

static int cuda_stream_selected_is_current(
        const cuda_stream_resident_expert &e,
        uint32_t layer,
        const int32_t *selected_ids,
        uint32_t n_selected) {
    if (!selected_ids || e.layer != layer) return 0;
    for (uint32_t i = 0; i < n_selected; i++) {
        if (e.expert == selected_ids[i]) return 1;
    }
    return 0;
}

static cuda_stream_resident_key cuda_stream_resident_make_key(
        const void *model_map,
        uint32_t layer,
        int32_t expert,
        uint64_t gate_offset,
        uint64_t up_offset,
        uint64_t down_offset,
        uint64_t gate_expert_bytes,
        uint64_t down_expert_bytes) {
    cuda_stream_resident_key k;
    k.model_map = model_map;
    k.layer = layer;
    k.expert = expert;
    k.gate_offset = gate_offset;
    k.up_offset = up_offset;
    k.down_offset = down_offset;
    k.gate_expert_bytes = gate_expert_bytes;
    k.down_expert_bytes = down_expert_bytes;
    return k;
}

static cuda_stream_resident_key cuda_stream_resident_entry_key(
        const cuda_stream_resident_expert &e) {
    return cuda_stream_resident_make_key(e.model_map,
                                         e.layer,
                                         e.expert,
                                         e.gate_offset,
                                         e.up_offset,
                                         e.down_offset,
                                         e.gate_expert_bytes,
                                         e.down_expert_bytes);
}

static int cuda_stream_resident_find(
        const void *model_map,
        uint32_t layer,
        int32_t expert,
        uint64_t gate_offset,
        uint64_t up_offset,
        uint64_t down_offset,
        uint64_t gate_expert_bytes,
        uint64_t down_expert_bytes) {
    const cuda_stream_resident_key key =
        cuda_stream_resident_make_key(model_map,
                                      layer,
                                      expert,
                                      gate_offset,
                                      up_offset,
                                      down_offset,
                                      gate_expert_bytes,
                                      down_expert_bytes);
    const auto it = g_stream_resident_index.find(key);
    if (it != g_stream_resident_index.end() &&
        it->second < g_stream_resident_experts.size()) {
        return (int)it->second;
    }
    return -1;
}

static int cuda_stream_resident_evict_at(size_t idx) {
    if (idx >= g_stream_resident_experts.size()) return 0;
    if (!cuda_stream_resident_reclaim_wait(
                "streaming resident expert cache eviction")) {
        return 0;
    }
    cuda_stream_resident_expert &e = g_stream_resident_experts[idx];
    const cuda_stream_resident_key evicted_key =
        cuda_stream_resident_entry_key(e);
    if (cuda_stream_cache_stats_on()) {
        g_stream_cache_stats.evictions++;
        g_stream_cache_stats.evict_bytes += e.bytes;
    }
    if (e.base) {
        if (e.pooled) g_stream_expert_free_slots.push_back(e.base);
        else (void)cudaFree(e.base);
    }
    if (g_stream_resident_bytes >= e.bytes) {
        g_stream_resident_bytes -= e.bytes;
    } else {
        g_stream_resident_bytes = 0;
    }
    g_stream_resident_index.erase(evicted_key);
    const size_t last = g_stream_resident_experts.size() - 1u;
    if (idx != last) {
        g_stream_resident_experts[idx] = g_stream_resident_experts[last];
        g_stream_resident_index[cuda_stream_resident_entry_key(
                g_stream_resident_experts[idx])] = idx;
    }
    g_stream_resident_experts.pop_back();
    return 1;
}

static int cuda_stream_resident_evict_one(
        uint32_t layer,
        const int32_t *selected_ids,
        uint32_t n_selected) {
    size_t victim = (size_t)-1;
    uint64_t oldest = UINT64_MAX;
    if (cuda_stream_evict_past_layers_first()) {
        for (size_t i = 0; i < g_stream_resident_experts.size(); i++) {
            const cuda_stream_resident_expert &e =
                g_stream_resident_experts[i];
            if (e.layer > layer ||
                cuda_stream_selected_is_current(e,
                                                layer,
                                                selected_ids,
                                                n_selected)) {
                continue;
            }
            if (e.last_used < oldest) {
                oldest = e.last_used;
                victim = i;
            }
        }
        if (victim != (size_t)-1) {
            return cuda_stream_resident_evict_at(victim);
        }
    }
    oldest = UINT64_MAX;
    for (size_t i = 0; i < g_stream_resident_experts.size(); i++) {
        const cuda_stream_resident_expert &e = g_stream_resident_experts[i];
        if (cuda_stream_selected_is_current(e, layer, selected_ids, n_selected)) {
            continue;
        }
        if (e.last_used < oldest) {
            oldest = e.last_used;
            victim = i;
        }
    }
    if (victim == (size_t)-1) return 0;
    return cuda_stream_resident_evict_at(victim);
}

static uint64_t cuda_stream_resident_free_reserve_bytes(void) {
    /*
     * Headroom kept free on the (unified-memory) device while growing the
     * expert cache.  It must cover decode scratch and transient graph buffers
     * but every reserved GiB is a GiB not spent caching experts, so keep it
     * tight.  Override with DS4_ROCM_STREAM_FREE_RESERVE_GB.
     */
    static int64_t cached = -1;
    if (cached < 0) {
        const char *env = getenv("DS4_ROCM_STREAM_FREE_RESERVE_GB");
        uint64_t gib = 16;
        if (env && env[0]) {
            char *end = NULL;
            errno = 0;
            unsigned long v = strtoul(env, &end, 10);
            if (end != env && *end == '\0' && errno == 0 && v >= 2 && v <= 64) {
                gib = (uint64_t)v;
            }
        }
        cached = (int64_t)(gib * 1024ull * 1024ull * 1024ull);
    }
    return (uint64_t)cached;
}

static int cuda_stream_resident_make_room(
        uint64_t bytes,
        uint32_t layer,
        const int32_t *selected_ids,
        uint32_t n_selected) {
    while (g_stream_resident_experts.size() >= g_stream_expert_cache_budget) {
        if (!cuda_stream_resident_evict_one(layer, selected_ids, n_selected)) {
            break;
        }
    }

    size_t free_b = 0;
    size_t total_b = 0;
    const uint64_t reserve = cuda_stream_resident_free_reserve_bytes();
    while (cudaMemGetInfo(&free_b, &total_b) == cudaSuccess) {
        (void)total_b;
        if ((uint64_t)free_b >= reserve &&
            bytes <= (uint64_t)free_b - reserve) {
            return 1;
        }
        if (!cuda_stream_resident_evict_one(layer, selected_ids, n_selected)) {
            return 0;
        }
    }
    (void)cudaGetLastError();
    return 1;
}

static int cuda_stream_expert_slab_grow(uint64_t slot_bytes) {
    const uint64_t slab_target_bytes = 1024ull * 1024ull * 1024ull;
    if (g_stream_expert_slot_count >= g_stream_expert_cache_budget) return 0;
    uint32_t slab_slots =
        slot_bytes >= slab_target_bytes ?
            1u : (uint32_t)(slab_target_bytes / slot_bytes);
    const uint32_t want =
        g_stream_expert_cache_budget - g_stream_expert_slot_count;
    if (slab_slots > want) slab_slots = want;
    while (slab_slots != 0) {
        uint64_t slab_bytes = 0;
        if (!cuda_u64_mul_checked(slab_slots, slot_bytes, &slab_bytes)) {
            slab_slots >>= 1u;
            continue;
        }
        size_t free_b = 0;
        size_t total_b = 0;
        if (cudaMemGetInfo(&free_b, &total_b) == cudaSuccess) {
            (void)total_b;
            const uint64_t reserve = cuda_stream_resident_free_reserve_bytes();
            if ((uint64_t)free_b < reserve ||
                slab_bytes > (uint64_t)free_b - reserve) {
                slab_slots >>= 1u;
                continue;
            }
        } else {
            (void)cudaGetLastError();
        }
        void *base = NULL;
        cudaError_t err = cudaMalloc(&base, (size_t)slab_bytes);
        if (err != cudaSuccess) {
            (void)cudaGetLastError();
            slab_slots >>= 1u;
            continue;
        }
        g_stream_expert_slabs.push_back({(char *)base, slab_bytes});
        g_stream_expert_free_slots.reserve(
                g_stream_expert_free_slots.size() + slab_slots);
        for (uint32_t i = 0; i < slab_slots; i++) {
            g_stream_expert_free_slots.push_back(
                    (char *)base + (uint64_t)i * slot_bytes);
        }
        g_stream_expert_slot_count += slab_slots;
        return 1;
    }
    return 0;
}

static char *cuda_stream_expert_slot_acquire(
        uint64_t bytes,
        uint32_t layer,
        const int32_t *selected_ids,
        uint32_t n_selected) {
    if (g_stream_expert_slot_bytes == 0) g_stream_expert_slot_bytes = bytes;
    if (bytes != g_stream_expert_slot_bytes) return NULL;
    for (;;) {
        if (!g_stream_expert_free_slots.empty()) {
            char *slot = g_stream_expert_free_slots.back();
            g_stream_expert_free_slots.pop_back();
            return slot;
        }
        if (g_stream_expert_slot_count < g_stream_expert_cache_budget &&
            cuda_stream_expert_slab_grow(bytes)) {
            continue;
        }
        if (!cuda_stream_resident_evict_one(layer, selected_ids, n_selected)) {
            return NULL;
        }
    }
}

static int cuda_stream_resident_alloc(
        const void *model_map,
        uint32_t layer,
        int32_t expert,
        const int32_t *selected_ids,
        uint32_t n_selected,
        uint64_t gate_offset,
        uint64_t up_offset,
        uint64_t down_offset,
        uint64_t gate_expert_bytes,
        uint64_t down_expert_bytes) {
    if (g_stream_expert_cache_budget == 0) return -1;
    if (!cuda_stream_resident_reclaim_wait(
                "streaming resident expert cache allocation")) {
        return -1;
    }
    uint64_t bytes = 0;
    uint64_t gate_pair = 0;
    if (!cuda_u64_mul_checked(2u, gate_expert_bytes, &gate_pair) ||
        gate_pair > UINT64_MAX - down_expert_bytes) {
        return -1;
    }
    bytes = gate_pair + down_expert_bytes;

    void *base = cuda_stream_expert_slot_acquire(bytes,
                                                 layer,
                                                 selected_ids,
                                                 n_selected);
    const int pooled = base != NULL;
    if (!pooled && bytes == g_stream_expert_slot_bytes) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "streaming expert cache cannot reserve a "
                "%.2f MiB slot for layer=%u expert=%d\n",
                (double)bytes / 1048576.0,
                layer,
                expert);
        return -1;
    }
    if (!pooled) {
        /*
         * Mixed-precision layers whose experts do not match the slab size
         * class fall back to dedicated allocations.
         */
        if (!cuda_stream_resident_make_room(bytes, layer, selected_ids, n_selected)) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming expert cache cannot keep %.2f MiB "
                    "for layer=%u expert=%d while preserving %.2f GiB free\n",
                    (double)bytes / 1048576.0,
                    layer,
                    expert,
                    (double)cuda_stream_resident_free_reserve_bytes() / 1073741824.0);
            return -1;
        }

        cudaError_t err = cudaMalloc(&base, (size_t)bytes);
        while (err != cudaSuccess && cuda_stream_resident_evict_one(layer, selected_ids, n_selected)) {
            (void)cudaGetLastError();
            err = cudaMalloc(&base, (size_t)bytes);
        }
        if (err != cudaSuccess) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming expert cache allocation failed "
                    "for layer=%u expert=%d (%.2f MiB): %s\n",
                    layer,
                    expert,
                    (double)bytes / 1048576.0,
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            return -1;
        }
    }

    cuda_stream_resident_expert e;
    memset(&e, 0, sizeof(e));
    e.model_map = model_map;
    e.layer = layer;
    e.expert = expert;
    e.gate_expert_bytes = gate_expert_bytes;
    e.down_expert_bytes = down_expert_bytes;
    e.gate_offset = gate_offset;
    e.up_offset = up_offset;
    e.down_offset = down_offset;
    e.base = (char *)base;
    e.gate = e.base;
    e.up = e.base + gate_expert_bytes;
    e.down = e.base + 2u * gate_expert_bytes;
    e.bytes = bytes;
    e.last_used = ++g_stream_resident_clock;
    e.pooled = pooled;
    g_stream_resident_experts.push_back(e);
    g_stream_resident_index[cuda_stream_resident_entry_key(e)] =
        g_stream_resident_experts.size() - 1u;
    g_stream_resident_bytes += bytes;
    if (cuda_stream_cache_stats_on()) {
        g_stream_cache_stats.allocs++;
        g_stream_cache_stats.alloc_bytes += bytes;
    }
    cuda_stream_cache_stats_note_resident();
    return (int)g_stream_resident_experts.size() - 1;
}

typedef struct cuda_stream_read_job {
    char *dst;
    uint64_t offset;
    uint64_t bytes;
    void *host_raw;
    void *host_buf;
    int ok;
    int uploaded;
    int errnum;
    int direct;
} cuda_stream_read_job;

struct cuda_stream_batch_selected_pending {
    int active;
    const void *model_map;
    uint32_t layer;
    uint32_t n_total_expert;
    uint32_t n_selected;
    uint32_t n_tokens;
    uint32_t n_unique;
    uint32_t resident_count;
    uint32_t missing_count;
    uint64_t gate_offset;
    uint64_t up_offset;
    uint64_t down_offset;
    uint64_t gate_expert_bytes;
    uint64_t down_expert_bytes;
    cuda_stream_read_job read_jobs[DS4_ROCM_STREAM_READ_MAX_JOBS];
    uint32_t read_job_count;
};

static cuda_stream_batch_selected_pending g_stream_batch_selected_pending;

struct cuda_stream_selected_pending {
    int active;
    const void *model_map;
    uint32_t layer;
    uint32_t n_total_expert;
    uint32_t n_selected;
    uint64_t gate_offset;
    uint64_t up_offset;
    uint64_t down_offset;
    uint64_t gate_expert_bytes;
    uint64_t down_expert_bytes;
    uint32_t resident_mask;
    uint32_t missing_mask;
    int32_t selected_ids[DS4_ROCM_N_EXPERT_USED];
    cuda_stream_read_job read_jobs[DS4_ROCM_N_EXPERT_USED * 3u];
    uint32_t read_job_count;
};

static cuda_stream_selected_pending g_stream_selected_pending;

typedef struct cuda_stream_read_profile_span {
    uint64_t offset;
    uint64_t end;
} cuda_stream_read_profile_span;

static uint64_t cuda_stream_read_profile_us(double seconds) {
    return seconds <= 0.0 ? 0u : (uint64_t)(seconds * 1000000.0 + 0.5);
}

static void cuda_stream_read_profile_coalesce(
        cuda_stream_read_profile_span *spans,
        uint32_t count,
        uint64_t max_gap,
        uint64_t *groups_out,
        uint64_t *extra_out) {
    uint64_t groups = 0;
    uint64_t extra = 0;
    uint64_t cur_end = 0;
    for (uint32_t i = 0; i < count; i++) {
        if (spans[i].end <= spans[i].offset) continue;
        if (groups == 0) {
            groups = 1;
            cur_end = spans[i].end;
            continue;
        }
        if (spans[i].offset <= cur_end) {
            if (spans[i].end > cur_end) cur_end = spans[i].end;
            continue;
        }
        const uint64_t gap = spans[i].offset - cur_end;
        if (gap <= max_gap) {
            extra += gap;
            if (spans[i].end > cur_end) cur_end = spans[i].end;
            continue;
        }
        groups++;
        cur_end = spans[i].end;
    }
    if (groups_out) *groups_out = groups;
    if (extra_out) *extra_out = extra;
}

static void cuda_stream_read_profile_note_jobs(
        const cuda_stream_read_job *jobs,
        uint32_t count) {
    if (g_stream_read_profile_enabled != 1 || !jobs || count == 0) return;
    if (count > DS4_ROCM_STREAM_READ_MAX_JOBS) return;

    cuda_stream_read_profile_span spans[DS4_ROCM_STREAM_READ_MAX_JOBS];
    uint32_t n = 0;
    for (uint32_t i = 0; i < count; i++) {
        if (jobs[i].bytes == 0) continue;
        spans[n].offset = jobs[i].offset;
        spans[n].end = jobs[i].offset > UINT64_MAX - jobs[i].bytes ?
            UINT64_MAX : jobs[i].offset + jobs[i].bytes;
        n++;
    }
    for (uint32_t i = 1; i < n; i++) {
        const cuda_stream_read_profile_span x = spans[i];
        uint32_t j = i;
        while (j != 0 && spans[j - 1].offset > x.offset) {
            spans[j] = spans[j - 1];
            j--;
        }
        spans[j] = x;
    }

    uint64_t exact_groups = 0;
    uint64_t gap1m_groups = 0;
    uint64_t gap4m_groups = 0;
    uint64_t exact_extra = 0;
    uint64_t gap1m_extra = 0;
    uint64_t gap4m_extra = 0;
    cuda_stream_read_profile_coalesce(spans, n, 0,
                                      &exact_groups, &exact_extra);
    cuda_stream_read_profile_coalesce(spans, n, 1ull * 1048576ull,
                                      &gap1m_groups, &gap1m_extra);
    cuda_stream_read_profile_coalesce(spans, n, 4ull * 1048576ull,
                                      &gap4m_groups, &gap4m_extra);

    (void)exact_extra;
    g_stream_read_profile_start_calls++;
    g_stream_read_profile_exact_groups += exact_groups;
    g_stream_read_profile_gap1m_groups += gap1m_groups;
    g_stream_read_profile_gap4m_groups += gap4m_groups;
    g_stream_read_profile_gap1m_extra += gap1m_extra;
    g_stream_read_profile_gap4m_extra += gap4m_extra;
}

static void cuda_stream_read_profile_print(void) {
    if (g_stream_read_profile_jobs == 0u &&
        g_stream_read_profile_wait_calls == 0u) {
        return;
    }
    const double jobs = g_stream_read_profile_jobs ?
        (double)g_stream_read_profile_jobs : 1.0;
    const double waits = g_stream_read_profile_wait_calls ?
        (double)g_stream_read_profile_wait_calls : 1.0;
    fprintf(stderr,
            DS4_GPU_LOG_PREFIX "stream read profile jobs=%llu bytes=%.2f GiB "
            "read=%.3f ms upload=%.3f ms wait=%.3f ms "
            "avg_job=%.3f ms avg_wait=%.3f ms\n",
            (unsigned long long)g_stream_read_profile_jobs,
            (double)g_stream_read_profile_bytes / 1073741824.0,
            (double)g_stream_read_profile_read_us / 1000.0,
            (double)g_stream_read_profile_upload_us / 1000.0,
            (double)g_stream_read_profile_wait_us / 1000.0,
            (double)(g_stream_read_profile_read_us +
                     g_stream_read_profile_upload_us) / 1000.0 / jobs,
            (double)g_stream_read_profile_wait_us / 1000.0 / waits);
    fprintf(stderr,
            DS4_GPU_LOG_PREFIX "stream read locality batches=%llu "
            "exact_jobs=%llu gap1m_jobs=%llu gap1m_extra=%.2f GiB "
            "gap4m_jobs=%llu gap4m_extra=%.2f GiB\n",
            (unsigned long long)g_stream_read_profile_start_calls,
            (unsigned long long)g_stream_read_profile_exact_groups,
            (unsigned long long)g_stream_read_profile_gap1m_groups,
            (double)g_stream_read_profile_gap1m_extra / 1073741824.0,
            (unsigned long long)g_stream_read_profile_gap4m_groups,
            (double)g_stream_read_profile_gap4m_extra / 1073741824.0);
}

static int cuda_stream_read_profile_enabled(void) {
    if (g_stream_read_profile_enabled < 0) {
        const char *env = getenv("DS4_ROCM_STREAM_READ_PROFILE");
        g_stream_read_profile_enabled =
            (env != NULL && env[0] != '\0' && strcmp(env, "0") != 0) ? 1 : 0;
        if (g_stream_read_profile_enabled &&
            !g_stream_read_profile_registered) {
            atexit(cuda_stream_read_profile_print);
            g_stream_read_profile_registered = 1;
        }
    }
    return g_stream_read_profile_enabled;
}

static int cuda_stream_read_direct_disabled(void) {
    if (g_stream_read_direct_disabled < 0) {
        const char *env = getenv("DS4_ROCM_STREAM_NO_DIRECT");
        g_stream_read_direct_disabled =
            (env != NULL && env[0] != '\0' && strcmp(env, "0") != 0) ? 1 : 0;
    }
    return g_stream_read_direct_disabled;
}

static void cuda_stream_read_job_run(cuda_stream_read_job *job,
                                     void *stage,
                                     uint64_t stage_bytes) {
    if (!job) return;
    job->ok = 0;
    job->uploaded = 0;
    job->errnum = 0;
    job->direct = 0;
    if (!stage || job->bytes == 0 || g_model_fd < 0) {
        job->errnum = EINVAL;
        return;
    }
    job->host_raw = stage;
    job->host_buf = stage;
#if defined(__linux__) && defined(O_DIRECT)
    /*
     * Direct reads skip the page cache: no page allocation, no extra copy and
     * no POSIX_FADV_DONTNEED churn per streamed expert.  Alignment slack is
     * pre-reserved in the staging buffers.  On any direct-read failure fall
     * back to the buffered fd for this job only; the shared direct fd is left
     * alone so concurrent workers are unaffected.
     */
    if (!cuda_stream_read_direct_disabled() &&
        g_model_direct_fd >= 0 &&
        g_model_direct_align > 1 &&
        g_model_file_size != 0) {
        const uint64_t aligned_off =
            cuda_round_down(job->offset, g_model_direct_align);
        const uint64_t delta = job->offset - aligned_off;
        const uint64_t read_size =
            cuda_round_up(delta + job->bytes, g_model_direct_align);
        if (read_size <= stage_bytes &&
            aligned_off <= g_model_file_size &&
            read_size <= g_model_file_size - aligned_off &&
            cuda_pread_full(g_model_direct_fd, stage, read_size, aligned_off)) {
            job->host_buf = (char *)stage + delta;
            job->direct = 1;
            job->ok = 1;
            return;
        }
    }
#endif
    if (cuda_pread_full(g_model_fd, job->host_buf, job->bytes, job->offset)) {
        job->ok = 1;
    } else {
        job->errnum = errno ? errno : EIO;
    }
}

static int cuda_stream_read_job_upload(
        cuda_stream_read_job *job,
        cudaStream_t stream) {
    if (!job || !job->ok || !job->dst || !job->host_buf || !stream) {
        if (job) job->errnum = EINVAL;
        return 0;
    }
    cudaError_t err = cudaMemcpyAsync(job->dst,
                                      job->host_buf,
                                      (size_t)job->bytes,
                                      cudaMemcpyHostToDevice,
                                      stream);
    if (err == cudaSuccess) err = cudaStreamSynchronize(stream);
    if (err != cudaSuccess) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "streaming read-worker upload failed: %s\n",
                cudaGetErrorString(err));
        (void)cudaGetLastError();
        job->ok = 0;
        job->errnum = EIO;
        return 0;
    }
    job->uploaded = 1;
    if (!job->direct) cuda_model_drop_file_pages(job->offset, job->bytes);
    return 1;
}

static void *cuda_stream_read_worker(void *arg) {
    const uint32_t worker_id = arg ? *(const uint32_t *)arg : 0u;
    (void)cudaSetDevice(0);
    for (;;) {
        pthread_mutex_lock(&g_stream_read_mutex);
        while (!g_stream_read_pool_stop &&
               (!g_stream_read_active_jobs ||
                g_stream_read_active_next >= g_stream_read_active_count)) {
            pthread_cond_wait(&g_stream_read_work_cond, &g_stream_read_mutex);
        }
        if (g_stream_read_pool_stop) {
            pthread_mutex_unlock(&g_stream_read_mutex);
            break;
        }
        const uint32_t idx = g_stream_read_active_next++;
        cuda_stream_read_job *job = &g_stream_read_active_jobs[idx];
        void *stage = NULL;
        uint64_t stage_bytes = 0;
        if (worker_id < DS4_ROCM_STREAM_READ_WORKERS) {
            stage = g_stream_read_stage_raw[worker_id];
            stage_bytes = g_stream_read_stage_bytes[worker_id];
        }
        pthread_mutex_unlock(&g_stream_read_mutex);

        const int profile = g_stream_read_profile_enabled == 1;
        const double read_t0 = profile ? cuda_wall_sec() : 0.0;
        cuda_stream_read_job_run(job, stage, stage_bytes);
        uint64_t read_us = 0;
        uint64_t upload_us = 0;
        if (profile) {
            read_us = cuda_stream_read_profile_us(cuda_wall_sec() - read_t0);
        }
        if (job->ok) {
            const double upload_t0 = profile ? cuda_wall_sec() : 0.0;
            (void)cuda_stream_read_job_upload(
                    job,
                    worker_id < DS4_ROCM_STREAM_READ_WORKERS ?
                        g_stream_read_upload_streams[worker_id] : NULL);
            if (profile) {
                upload_us =
                    cuda_stream_read_profile_us(cuda_wall_sec() - upload_t0);
            }
        }

        pthread_mutex_lock(&g_stream_read_mutex);
        if (profile) {
            g_stream_read_profile_jobs++;
            g_stream_read_profile_bytes += job->bytes;
            g_stream_read_profile_read_us += read_us;
            g_stream_read_profile_upload_us += upload_us;
        }
        if (!job->ok) g_stream_read_active_ok = 0;
        g_stream_read_active_done++;
        if (g_stream_read_active_done >= g_stream_read_active_count) {
            pthread_cond_signal(&g_stream_read_done_cond);
        }
        pthread_mutex_unlock(&g_stream_read_mutex);
    }
    return NULL;
}

static uint32_t cuda_stream_read_worker_count(void) {
    const char *env = getenv("DS4_ROCM_STREAM_READ_WORKERS");
    if (env && env[0]) {
        char *end = NULL;
        errno = 0;
        unsigned long v = strtoul(env, &end, 10);
        if (end != env && errno == 0) {
            if (v == 0) return 1u;
            if (v > DS4_ROCM_STREAM_READ_WORKERS) {
                return DS4_ROCM_STREAM_READ_WORKERS;
            }
            return (uint32_t)v;
        }
    }
    return DS4_ROCM_STREAM_READ_DEFAULT_WORKERS;
}

static void cuda_stream_read_upload_streams_destroy(void) {
    for (uint32_t i = 0; i < DS4_ROCM_STREAM_READ_WORKERS; i++) {
        if (g_stream_read_upload_streams[i]) {
            (void)cudaStreamDestroy(g_stream_read_upload_streams[i]);
            g_stream_read_upload_streams[i] = NULL;
        }
    }
}

static int cuda_stream_read_upload_streams_ensure(void) {
    const uint32_t workers = cuda_stream_read_worker_count();
    for (uint32_t i = 0; i < workers; i++) {
        if (g_stream_read_upload_streams[i]) continue;
        cudaError_t err = cudaStreamCreateWithFlags(
                &g_stream_read_upload_streams[i],
                cudaStreamNonBlocking);
        if (err != cudaSuccess) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming read upload stream creation failed: %s\n",
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            cuda_stream_read_upload_streams_destroy();
            return 0;
        }
    }
    return 1;
}

static int cuda_stream_read_pool_ensure(void) {
    if (g_stream_read_pool_started) return 1;
    pthread_mutex_lock(&g_stream_read_mutex);
    if (g_stream_read_pool_started) {
        pthread_mutex_unlock(&g_stream_read_mutex);
        return 1;
    }
    g_stream_read_pool_stop = 0;
    g_stream_read_active_jobs = NULL;
    g_stream_read_active_count = 0;
    g_stream_read_active_next = 0;
    g_stream_read_active_done = 0;
    g_stream_read_active_ok = 1;
    g_stream_read_active_owner_set = 0;
    g_stream_read_pool_workers = cuda_stream_read_worker_count();
    (void)cuda_stream_read_profile_enabled();
    if (!cuda_stream_read_upload_streams_ensure()) {
        pthread_mutex_unlock(&g_stream_read_mutex);
        return 0;
    }
    for (uint32_t i = 0; i < g_stream_read_pool_workers; i++) {
        g_stream_read_thread_ids[i] = i;
        const int rc = pthread_create(&g_stream_read_threads[i],
                                      NULL,
                                      cuda_stream_read_worker,
                                      &g_stream_read_thread_ids[i]);
        if (rc != 0) {
            g_stream_read_pool_stop = 1;
            pthread_cond_broadcast(&g_stream_read_work_cond);
            pthread_mutex_unlock(&g_stream_read_mutex);
            for (uint32_t j = 0; j < i; j++) {
                (void)pthread_join(g_stream_read_threads[j], NULL);
            }
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming read worker creation failed: %s\n",
                    strerror(rc));
            cuda_stream_read_upload_streams_destroy();
            return 0;
        }
    }
    g_stream_read_pool_started = 1;
    pthread_mutex_unlock(&g_stream_read_mutex);
    return 1;
}

static void cuda_stream_read_pool_shutdown(void) {
    if (!g_stream_read_pool_started) return;
    pthread_mutex_lock(&g_stream_read_mutex);
    g_stream_read_pool_stop = 1;
    pthread_cond_broadcast(&g_stream_read_work_cond);
    pthread_mutex_unlock(&g_stream_read_mutex);
    for (uint32_t i = 0; i < g_stream_read_pool_workers; i++) {
        (void)pthread_join(g_stream_read_threads[i], NULL);
    }
    pthread_mutex_lock(&g_stream_read_mutex);
    g_stream_read_pool_started = 0;
    g_stream_read_pool_workers = 0;
    g_stream_read_pool_stop = 0;
    g_stream_read_active_jobs = NULL;
    g_stream_read_active_count = 0;
    g_stream_read_active_next = 0;
    g_stream_read_active_done = 0;
    g_stream_read_active_ok = 1;
    g_stream_read_active_owner_set = 0;
    pthread_mutex_unlock(&g_stream_read_mutex);
}

static int cuda_stream_read_jobs_prepare(cuda_stream_read_job *jobs, uint32_t count) {
    if (!jobs || count == 0) return 1;
    if (count > DS4_ROCM_STREAM_READ_MAX_JOBS) return 0;

    uint64_t max_bytes = 0;
    for (uint32_t i = 0; i < count; i++) {
        jobs[i].ok = 0;
        jobs[i].uploaded = 0;
        jobs[i].errnum = 0;
        jobs[i].direct = 0;
        jobs[i].host_raw = NULL;
        jobs[i].host_buf = NULL;
        if (jobs[i].bytes > max_bytes) max_bytes = jobs[i].bytes;
    }
    /* Slack so direct reads can align their offset and length. */
    if (g_model_direct_align > 1u &&
        max_bytes <= UINT64_MAX - 2u * g_model_direct_align) {
        max_bytes += 2u * g_model_direct_align;
    }

    const uint32_t workers = g_stream_read_pool_started ?
        g_stream_read_pool_workers : cuda_stream_read_worker_count();
    for (uint32_t i = 0; i < workers; i++) {
        if (g_stream_read_stage_bytes[i] < max_bytes) {
            if (g_stream_read_stage_raw[i]) {
                (void)cudaFreeHost(g_stream_read_stage_raw[i]);
                g_stream_read_stage_raw[i] = NULL;
                g_stream_read_stage_bytes[i] = 0;
            }
            cudaError_t err = cudaMallocHost(&g_stream_read_stage_raw[i],
                                             (size_t)max_bytes);
            if (err != cudaSuccess) {
                fprintf(stderr,
                        DS4_GPU_LOG_PREFIX "streaming read pinned allocation failed "
                        "(%.2f MiB): %s\n",
                        (double)max_bytes / 1048576.0,
                        cudaGetErrorString(err));
                (void)cudaGetLastError();
                return 0;
            }
            g_stream_read_stage_bytes[i] = max_bytes;
        }
    }
    return 1;
}

static int cuda_stream_read_jobs_start(cuda_stream_read_job *jobs, uint32_t count) {
    if (!jobs || count == 0) return 1;
    if (!cuda_stream_read_pool_ensure()) return 0;
    const pthread_t self = pthread_self();
    pthread_mutex_lock(&g_stream_read_mutex);
    while (g_stream_read_active_jobs != NULL) {
        if (g_stream_read_active_owner_set &&
            pthread_equal(g_stream_read_active_owner, self)) {
            pthread_mutex_unlock(&g_stream_read_mutex);
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming read pool already has active work for this thread\n");
            return 0;
        }
        pthread_cond_wait(&g_stream_read_done_cond, &g_stream_read_mutex);
    }
    if (!cuda_stream_read_jobs_prepare(jobs, count)) {
        pthread_mutex_unlock(&g_stream_read_mutex);
        return 0;
    }
    cuda_stream_read_profile_note_jobs(jobs, count);
    g_stream_read_active_jobs = jobs;
    g_stream_read_active_count = count;
    g_stream_read_active_next = 0;
    g_stream_read_active_done = 0;
    g_stream_read_active_ok = 1;
    g_stream_read_active_owner = self;
    g_stream_read_active_owner_set = 1;
    pthread_cond_broadcast(&g_stream_read_work_cond);
    pthread_mutex_unlock(&g_stream_read_mutex);
    return 1;
}

static int cuda_stream_read_jobs_wait(cuda_stream_read_job *jobs, uint32_t count) {
    if (!jobs || count == 0) return 1;
    const int profile = g_stream_read_profile_enabled == 1;
    const double wait_t0 = profile ? cuda_wall_sec() : 0.0;
    pthread_mutex_lock(&g_stream_read_mutex);
    if (g_stream_read_active_jobs != jobs) {
        pthread_mutex_unlock(&g_stream_read_mutex);
        fprintf(stderr, DS4_GPU_LOG_PREFIX "streaming read wait received inactive job set\n");
        return 0;
    }
    while (g_stream_read_active_done < g_stream_read_active_count) {
        pthread_cond_wait(&g_stream_read_done_cond, &g_stream_read_mutex);
    }
    const int pool_ok = g_stream_read_active_ok;
    g_stream_read_active_jobs = NULL;
    g_stream_read_active_count = 0;
    g_stream_read_active_next = 0;
    g_stream_read_active_done = 0;
    g_stream_read_active_ok = 1;
    g_stream_read_active_owner_set = 0;
    if (profile) {
        g_stream_read_profile_wait_calls++;
        g_stream_read_profile_wait_us +=
            cuda_stream_read_profile_us(cuda_wall_sec() - wait_t0);
    }
    pthread_cond_broadcast(&g_stream_read_done_cond);
    pthread_mutex_unlock(&g_stream_read_mutex);

    int ok = 1;
    if (!pool_ok) ok = 0;
    for (uint32_t i = 0; i < count; i++) {
        if (!jobs[i].ok) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming read failed at offset %.2f GiB "
                    "size %.2f MiB: %s\n",
                    (double)jobs[i].offset / 1073741824.0,
                    (double)jobs[i].bytes / 1048576.0,
                    strerror(jobs[i].errnum ? jobs[i].errnum : EIO));
            ok = 0;
        }
    }
    return ok;
}

static int cuda_stream_read_jobs_parallel(cuda_stream_read_job *jobs, uint32_t count) {
    if (!jobs || count == 0) return 1;
    return cuda_stream_read_jobs_start(jobs, count) &&
           cuda_stream_read_jobs_wait(jobs, count);
}

static void cuda_stream_read_jobs_free(cuda_stream_read_job *jobs, uint32_t count) {
    if (!jobs) return;
    for (uint32_t i = 0; i < count; i++) {
        jobs[i].host_raw = NULL;
        jobs[i].host_buf = NULL;
        jobs[i].uploaded = 0;
        jobs[i].direct = 0;
    }
}

static int cuda_stream_selected_upload_read_jobs(
        cuda_stream_read_job *jobs,
        uint32_t count);

static int cuda_stream_batch_selected_pending_matches(
        const void *model_map,
        uint32_t layer,
        uint32_t n_total_expert,
        uint32_t n_selected,
        uint32_t n_tokens,
        uint64_t gate_offset,
        uint64_t up_offset,
        uint64_t down_offset,
        uint64_t gate_expert_bytes,
        uint64_t down_expert_bytes) {
    return g_stream_batch_selected_pending.active &&
           g_stream_batch_selected_pending.model_map == model_map &&
           g_stream_batch_selected_pending.layer == layer &&
           g_stream_batch_selected_pending.n_total_expert == n_total_expert &&
           g_stream_batch_selected_pending.n_selected == n_selected &&
           g_stream_batch_selected_pending.n_tokens == n_tokens &&
           g_stream_batch_selected_pending.gate_offset == gate_offset &&
           g_stream_batch_selected_pending.up_offset == up_offset &&
           g_stream_batch_selected_pending.down_offset == down_offset &&
           g_stream_batch_selected_pending.gate_expert_bytes == gate_expert_bytes &&
           g_stream_batch_selected_pending.down_expert_bytes == down_expert_bytes;
}

static int cuda_stream_batch_selected_apply_split(
        const void *model_map,
        uint32_t layer,
        uint32_t n_total_expert,
        uint32_t n_selected,
        uint32_t n_tokens,
        uint64_t gate_offset,
        uint64_t up_offset,
        uint64_t down_offset,
        uint64_t gate_expert_bytes,
        uint64_t down_expert_bytes,
        const ds4_gpu_tensor **selected_exec,
        const char ***resident_gate_ptrs,
        const char ***resident_up_ptrs,
        const char ***missing_gate_ptrs,
        const char ***missing_up_ptrs,
        const char ***down_ptrs,
        const uint8_t **pair_missing,
        uint32_t *resident_count,
        uint32_t *missing_count,
        uint32_t *unique_out) {
    if (!selected_exec || !resident_gate_ptrs || !resident_up_ptrs ||
        !missing_gate_ptrs || !missing_up_ptrs || !down_ptrs ||
        !pair_missing ||
        !resident_count || !missing_count || !unique_out ||
        !cuda_stream_batch_selected_pending_matches(model_map,
                                                    layer,
                                                    n_total_expert,
                                                    n_selected,
                                                    n_tokens,
                                                    gate_offset,
                                                    up_offset,
                                                    down_offset,
                                                    gate_expert_bytes,
                                                    down_expert_bytes) ||
        !g_stream_batch_selected_cache.selected_ids ||
        !g_stream_batch_selected_cache.resident_gate_ptrs ||
        !g_stream_batch_selected_cache.resident_up_ptrs ||
        !g_stream_batch_selected_cache.missing_gate_ptrs ||
        !g_stream_batch_selected_cache.missing_up_ptrs ||
        !g_stream_batch_selected_cache.down_ptrs ||
        !g_stream_batch_selected_cache.pair_missing) {
        return 0;
    }
    *selected_exec = &g_stream_batch_selected_cache.selected_tensor;
    *resident_gate_ptrs = g_stream_batch_selected_cache.resident_gate_ptrs;
    *resident_up_ptrs = g_stream_batch_selected_cache.resident_up_ptrs;
    *missing_gate_ptrs = g_stream_batch_selected_cache.missing_gate_ptrs;
    *missing_up_ptrs = g_stream_batch_selected_cache.missing_up_ptrs;
    *down_ptrs = g_stream_batch_selected_cache.down_ptrs;
    *pair_missing = g_stream_batch_selected_cache.pair_missing;
    *resident_count = g_stream_batch_selected_pending.resident_count;
    *missing_count = g_stream_batch_selected_pending.missing_count;
    *unique_out = g_stream_batch_selected_pending.n_unique;
    return 1;
}

static int cuda_stream_batch_selected_finish_pending_missing(void) {
    if (!g_stream_batch_selected_pending.active) return 1;
    const uint32_t read_job_count =
        g_stream_batch_selected_pending.read_job_count;
    if (!cuda_stream_read_jobs_wait(g_stream_batch_selected_pending.read_jobs,
                                    read_job_count) ||
        !cuda_stream_selected_upload_read_jobs(
                g_stream_batch_selected_pending.read_jobs,
                read_job_count)) {
        cuda_stream_read_jobs_free(g_stream_batch_selected_pending.read_jobs,
                                   read_job_count);
        memset(&g_stream_batch_selected_pending, 0,
               sizeof(g_stream_batch_selected_pending));
        cuda_stream_resident_cache_release();
        return 0;
    }
    cuda_stream_read_jobs_free(g_stream_batch_selected_pending.read_jobs,
                               read_job_count);
    g_stream_batch_selected_cache.loaded = 0;
    memset(&g_stream_batch_selected_pending, 0,
           sizeof(g_stream_batch_selected_pending));
    return 1;
}

static void cuda_stream_batch_selected_abort_pending(void) {
    if (!g_stream_batch_selected_pending.active) return;
    const uint32_t read_job_count =
        g_stream_batch_selected_pending.read_job_count;
    (void)cuda_stream_read_jobs_wait(g_stream_batch_selected_pending.read_jobs,
                                     read_job_count);
    cuda_stream_read_jobs_free(g_stream_batch_selected_pending.read_jobs,
                               read_job_count);
    memset(&g_stream_batch_selected_pending, 0,
           sizeof(g_stream_batch_selected_pending));
}

static void cuda_stream_selected_abort_pending(void) {
    if (!g_stream_selected_pending.active) return;
    const uint32_t read_job_count = g_stream_selected_pending.read_job_count;
    if (read_job_count != 0) {
        (void)cuda_stream_read_jobs_wait(g_stream_selected_pending.read_jobs,
                                         read_job_count);
        cuda_stream_read_jobs_free(g_stream_selected_pending.read_jobs,
                                   read_job_count);
    }
    memset(&g_stream_selected_pending, 0, sizeof(g_stream_selected_pending));
}

static int cuda_stream_selected_upload_read_jobs(
        cuda_stream_read_job *jobs,
        uint32_t count) {
    if (!jobs || count == 0) return 1;
    int need_upload = 0;
    for (uint32_t i = 0; i < count; i++) {
        if (!jobs[i].uploaded) {
            need_upload = 1;
            break;
        }
    }
    if (!need_upload) return 1;
    if (!cuda_stream_selected_ensure_stream()) return 0;
    for (uint32_t i = 0; i < count; i++) {
        if (jobs[i].uploaded) continue;
        cudaError_t err = cudaMemcpyAsync(jobs[i].dst,
                                          jobs[i].host_buf,
                                          (size_t)jobs[i].bytes,
                                          cudaMemcpyHostToDevice,
                                          g_stream_selected_upload_stream);
        if (err != cudaSuccess) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming selected cached upload failed: %s\n",
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            return 0;
        }
        if (!jobs[i].direct) {
            cuda_model_drop_file_pages(jobs[i].offset, jobs[i].bytes);
        }
    }
    cudaError_t err = cudaStreamSynchronize(g_stream_selected_upload_stream);
    if (err != cudaSuccess) {
        fprintf(stderr, DS4_GPU_LOG_PREFIX "streaming selected upload sync failed: %s\n",
                cudaGetErrorString(err));
        (void)cudaGetLastError();
        return 0;
    }
    return 1;
}

static int cuda_stream_flush_read_jobs(
        cuda_stream_read_job *jobs,
        uint32_t *count) {
    if (!count || *count == 0) return 1;
    if (!cuda_stream_read_jobs_parallel(jobs, *count) ||
        !cuda_stream_selected_upload_read_jobs(jobs, *count)) {
        cuda_stream_read_jobs_free(jobs, *count);
        *count = 0;
        return 0;
    }
    cuda_stream_read_jobs_free(jobs, *count);
    *count = 0;
    return 1;
}

static int cuda_stream_layer_expert_cache_apply(
        const void *model_map,
        uint32_t layer,
        uint32_t n_total_expert,
        uint64_t gate_offset,
        uint64_t up_offset,
        uint64_t down_offset,
        uint64_t gate_expert_bytes,
        uint64_t down_expert_bytes,
        const char **gate_w,
        const char **up_w,
        const char **down_w) {
    if (!g_ssd_streaming_mode || !gate_w || !up_w || !down_w) return 0;
    for (uint32_t i = 0; i < 2u; i++) {
        const cuda_stream_layer_expert_cache &c = g_stream_layer_expert_cache[i];
        if (c.active &&
            c.model_map == model_map &&
            c.layer == layer &&
            c.n_total_expert == n_total_expert &&
            c.gate_offset == gate_offset &&
            c.up_offset == up_offset &&
            c.down_offset == down_offset &&
            c.gate_expert_bytes == gate_expert_bytes &&
            c.down_expert_bytes == down_expert_bytes &&
            c.gate && c.up && c.down) {
            *gate_w = c.gate;
            *up_w = c.up;
            *down_w = c.down;
            return 1;
        }
    }
    return 0;
}

static int cuda_stream_layer_expert_cache_load(
        const void *model_map,
        uint64_t model_size,
        uint32_t layer,
        uint32_t n_total_expert,
        uint64_t gate_offset,
        uint64_t up_offset,
        uint64_t down_offset,
        uint64_t gate_expert_bytes,
        uint64_t down_expert_bytes) {
    if (!g_ssd_streaming_mode ||
        !model_map ||
        model_size == 0 ||
        n_total_expert == 0 ||
        n_total_expert > DS4_ROCM_STREAM_READ_MAX_JOBS / 3u ||
        gate_expert_bytes == 0 ||
        down_expert_bytes == 0 ||
        g_model_fd < 0 ||
        (g_model_fd_host_base != NULL && model_map != g_model_fd_host_base)) {
        return 0;
    }

    uint64_t gate_bytes = 0;
    uint64_t down_bytes = 0;
    uint64_t gate_pair_bytes = 0;
    uint64_t total_bytes = 0;
    if (!cuda_u64_mul_checked(n_total_expert, gate_expert_bytes, &gate_bytes) ||
        !cuda_u64_mul_checked(n_total_expert, down_expert_bytes, &down_bytes) ||
        !cuda_u64_mul_checked(2u, gate_bytes, &gate_pair_bytes) ||
        gate_pair_bytes > UINT64_MAX - down_bytes) {
        return 0;
    }
    total_bytes = gate_pair_bytes + down_bytes;
    if (cuda_stream_cache_stats_on()) {
        g_stream_cache_stats.layer_loads++;
        g_stream_cache_stats.layer_load_bytes += total_bytes;
    }

    if (gate_offset > model_size ||
        up_offset > model_size ||
        down_offset > model_size ||
        gate_bytes > model_size - gate_offset ||
        gate_bytes > model_size - up_offset ||
        down_bytes > model_size - down_offset) {
        return 0;
    }

    cuda_stream_layer_expert_cache &slot =
        g_stream_layer_expert_cache[layer & 1u];
    slot.active = 0;
    if (slot.capacity < total_bytes) {
        if (slot.base) {
            (void)cudaFree(slot.base);
            memset(&slot, 0, sizeof(slot));
        }
        if (cuda_stream_cache_stats_on() &&
            !g_stream_resident_experts.empty()) {
            g_stream_cache_stats.layer_resident_flushes++;
        }
        cuda_stream_resident_cache_release();
        void *base = NULL;
        cudaError_t err = cudaMalloc(&base, (size_t)total_bytes);
        if (err != cudaSuccess) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming full-layer expert cache allocation "
                    "failed for layer=%u (%.2f GiB): %s\n",
                    layer,
                    (double)total_bytes / 1073741824.0,
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            return 0;
        }
        slot.base = (char *)base;
        slot.capacity = total_bytes;
    }

    slot.bytes = total_bytes;
    slot.gate = slot.base;
    slot.up = slot.base + gate_bytes;
    slot.down = slot.base + gate_pair_bytes;

    const uint64_t read_chunk = 32ull * 1048576ull;
    const uint64_t gate_chunks =
        (gate_bytes + read_chunk - 1u) / read_chunk;
    const uint64_t down_chunks =
        (down_bytes + read_chunk - 1u) / read_chunk;
    const uint64_t read_job_count64 = gate_chunks * 2u + down_chunks;
    if (read_job_count64 == 0 ||
        read_job_count64 > DS4_ROCM_STREAM_READ_MAX_JOBS ||
        read_job_count64 > UINT32_MAX) {
        return 0;
    }
    const uint32_t read_job_count = (uint32_t)read_job_count64;
    cuda_stream_read_job *jobs =
        (cuda_stream_read_job *)calloc((size_t)read_job_count, sizeof(jobs[0]));
    if (!jobs) return 0;

    int ok = 1;
    uint32_t j = 0;
    for (uint64_t off = 0; off < gate_bytes; off += read_chunk) {
        const uint64_t n = gate_bytes - off < read_chunk ? gate_bytes - off : read_chunk;
        jobs[j++] = {slot.gate + off, gate_offset + off, n, NULL, NULL, 0, 0};
    }
    for (uint64_t off = 0; off < gate_bytes; off += read_chunk) {
        const uint64_t n = gate_bytes - off < read_chunk ? gate_bytes - off : read_chunk;
        jobs[j++] = {slot.up + off, up_offset + off, n, NULL, NULL, 0, 0};
    }
    for (uint64_t off = 0; off < down_bytes; off += read_chunk) {
        const uint64_t n = down_bytes - off < read_chunk ? down_bytes - off : read_chunk;
        jobs[j++] = {slot.down + off, down_offset + off, n, NULL, NULL, 0, 0};
    }
    if (j != read_job_count ||
        !cuda_stream_read_jobs_parallel(jobs, read_job_count)) {
        ok = 0;
    }
    cuda_stream_read_jobs_free(jobs, read_job_count);
    free(jobs);
    if (!ok) return 0;

    slot.active = 1;
    slot.model_map = model_map;
    slot.layer = layer;
    slot.n_total_expert = n_total_expert;
    slot.gate_offset = gate_offset;
    slot.up_offset = up_offset;
    slot.down_offset = down_offset;
    slot.gate_expert_bytes = gate_expert_bytes;
    slot.down_expert_bytes = down_expert_bytes;
    return 1;
}

static int cuda_stream_resident_seed_experts(
        const void *model_map,
        uint64_t model_size,
        uint32_t layer,
        const int32_t *expert_ids,
        const uint32_t *expert_priorities,
        uint32_t n_experts,
        uint32_t n_total_expert,
        uint64_t gate_offset,
        uint64_t up_offset,
        uint64_t down_offset,
        uint64_t gate_expert_bytes,
        uint64_t down_expert_bytes) {
    if (!g_ssd_streaming_mode) return 1;
    if (!model_map || !expert_ids || n_experts == 0 ||
        n_total_expert == 0 ||
        n_total_expert > DS4_ROCM_MAX_N_EXPERT ||
        gate_expert_bytes == 0 ||
        down_expert_bytes == 0) {
        return 0;
    }
    if (g_stream_expert_cache_budget == 0) return 1;

    uint64_t gate_bytes = 0;
    uint64_t down_bytes = 0;
    if (!cuda_u64_mul_checked(n_total_expert, gate_expert_bytes, &gate_bytes) ||
        !cuda_u64_mul_checked(n_total_expert, down_expert_bytes, &down_bytes) ||
        gate_offset > model_size ||
        up_offset > model_size ||
        down_offset > model_size ||
        gate_bytes > model_size - gate_offset ||
        gate_bytes > model_size - up_offset ||
        down_bytes > model_size - down_offset) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "streaming hotlist seed expert range outside model map\n");
        return 0;
    }

    uint32_t seed_cap = n_experts < g_stream_expert_cache_budget ?
        n_experts : g_stream_expert_cache_budget;
    if (seed_cap > DS4_ROCM_MAX_N_EXPERT) seed_cap = DS4_ROCM_MAX_N_EXPERT;
    if (seed_cap == 0) return 1;

    bool seen[DS4_ROCM_MAX_N_EXPERT] = {0};
    uint32_t best_index[DS4_ROCM_MAX_N_EXPERT];
    uint32_t best_priority[DS4_ROCM_MAX_N_EXPERT];
    for (uint32_t i = 0; i < n_experts; i++) {
        const int32_t expert_i = expert_ids[i];
        if (expert_i < 0 || (uint32_t)expert_i >= n_total_expert) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming hotlist seed expert id %d outside 0..%u "
                    "(layer=%u)\n",
                    expert_i,
                    n_total_expert,
                    layer);
            return 0;
        }
        const uint32_t expert = (uint32_t)expert_i;
        const uint32_t priority =
            expert_priorities ? expert_priorities[i] : (n_experts - i);
        if (!seen[expert] || priority > best_priority[expert]) {
            seen[expert] = true;
            best_index[expert] = i;
            best_priority[expert] = priority;
        }
    }

    uint32_t chosen_indices[DS4_ROCM_MAX_N_EXPERT];
    uint32_t chosen_priorities[DS4_ROCM_MAX_N_EXPERT];
    uint32_t chosen_count = 0;
    for (uint32_t expert = 0; expert < n_total_expert; expert++) {
        if (!seen[expert]) continue;
        const uint32_t priority = best_priority[expert];
        uint32_t pos = 0;
        while (pos < chosen_count && priority <= chosen_priorities[pos]) {
            pos++;
        }
        if (chosen_count < seed_cap) {
            for (uint32_t j = chosen_count; j > pos; j--) {
                chosen_indices[j] = chosen_indices[j - 1u];
                chosen_priorities[j] = chosen_priorities[j - 1u];
            }
            chosen_indices[pos] = best_index[expert];
            chosen_priorities[pos] = priority;
            chosen_count++;
        } else if (pos < chosen_count) {
            for (uint32_t j = chosen_count - 1u; j > pos; j--) {
                chosen_indices[j] = chosen_indices[j - 1u];
                chosen_priorities[j] = chosen_priorities[j - 1u];
            }
            chosen_indices[pos] = best_index[expert];
            chosen_priorities[pos] = priority;
        }
    }
    if (chosen_count == 0) return 1;

    int32_t protected_ids[DS4_ROCM_MAX_N_EXPERT];
    for (uint32_t i = 0; i < chosen_count; i++) {
        protected_ids[i] = expert_ids[chosen_indices[i]];
    }
    if (cuda_stream_cache_stats_on()) {
        g_stream_cache_stats.seed_calls++;
        g_stream_cache_stats.seed_unique += chosen_count;
    }

    const int use_fd =
        g_model_fd >= 0 &&
        (g_model_fd_host_base == NULL || model_map == g_model_fd_host_base);
    if (!use_fd && !cuda_stream_selected_ensure_stream()) return 1;

    cuda_stream_read_job read_jobs[DS4_ROCM_STREAM_READ_MAX_JOBS];
    memset(read_jobs, 0, sizeof(read_jobs));
    uint32_t read_job_count = 0;
    int ok = 1;
    uint32_t loaded = 0;

    for (uint32_t ri = 0; ok && ri < chosen_count; ri++) {
        const uint32_t chosen_i = chosen_indices[chosen_count - 1u - ri];
        const int32_t expert_i = expert_ids[chosen_i];
        int idx = cuda_stream_resident_find(model_map,
                                            layer,
                                            expert_i,
                                            gate_offset,
                                            up_offset,
                                            down_offset,
                                            gate_expert_bytes,
                                            down_expert_bytes);
        if (idx >= 0) {
            g_stream_resident_experts[(size_t)idx].last_used =
                ++g_stream_resident_clock;
            loaded++;
            continue;
        }

        idx = cuda_stream_resident_alloc(model_map,
                                         layer,
                                         expert_i,
                                         protected_ids,
                                         chosen_count,
                                         gate_offset,
                                         up_offset,
                                         down_offset,
                                         gate_expert_bytes,
                                         down_expert_bytes);
        if (idx < 0) {
            ok = 0;
            break;
        }

        const uint64_t expert = (uint64_t)(uint32_t)expert_i;
        uint64_t gate_rel = 0;
        uint64_t down_rel = 0;
        if (!cuda_u64_mul_checked(expert, gate_expert_bytes, &gate_rel) ||
            !cuda_u64_mul_checked(expert, down_expert_bytes, &down_rel)) {
            ok = 0;
            break;
        }

        cuda_stream_resident_expert &entry =
            g_stream_resident_experts[(size_t)idx];
        if (use_fd) {
            if (read_job_count + 3u > DS4_ROCM_STREAM_READ_MAX_JOBS) {
                ok = 0;
                break;
            }
            read_jobs[read_job_count++] =
                {entry.gate, gate_offset + gate_rel, gate_expert_bytes,
                 NULL, NULL, 0, 0};
            read_jobs[read_job_count++] =
                {entry.up, up_offset + gate_rel, gate_expert_bytes,
                 NULL, NULL, 0, 0};
            read_jobs[read_job_count++] =
                {entry.down, down_offset + down_rel, down_expert_bytes,
                 NULL, NULL, 0, 0};
        } else {
            cudaError_t err = cudaMemcpyAsync(entry.gate,
                                              (const char *)model_map + gate_offset + gate_rel,
                                              (size_t)gate_expert_bytes,
                                              cudaMemcpyHostToDevice,
                                              g_stream_selected_upload_stream);
            if (err == cudaSuccess) {
                err = cudaMemcpyAsync(entry.up,
                                      (const char *)model_map + up_offset + gate_rel,
                                      (size_t)gate_expert_bytes,
                                      cudaMemcpyHostToDevice,
                                      g_stream_selected_upload_stream);
            }
            if (err == cudaSuccess) {
                err = cudaMemcpyAsync(entry.down,
                                      (const char *)model_map + down_offset + down_rel,
                                      (size_t)down_expert_bytes,
                                      cudaMemcpyHostToDevice,
                                      g_stream_selected_upload_stream);
            }
            if (err != cudaSuccess) {
                fprintf(stderr,
                        DS4_GPU_LOG_PREFIX "streaming hotlist seed upload failed: %s\n",
                        cudaGetErrorString(err));
                (void)cudaGetLastError();
                ok = 0;
                break;
            }
        }
        loaded++;
    }

    if (ok && use_fd && read_job_count != 0) {
        ok = cuda_stream_read_jobs_parallel(read_jobs, read_job_count) &&
             cuda_stream_selected_upload_read_jobs(read_jobs, read_job_count);
        cuda_stream_read_jobs_free(read_jobs, read_job_count);
    } else if (ok && !use_fd) {
        cudaError_t err = cudaStreamSynchronize(g_stream_selected_upload_stream);
        if (err != cudaSuccess) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming hotlist seed upload sync failed: %s\n",
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            ok = 0;
        }
    }

    if (!ok) {
        cuda_stream_read_jobs_free(read_jobs, read_job_count);
        cuda_stream_resident_cache_release();
        if (getenv("DS4_ROCM_STREAMING_EXPERT_CACHE_VERBOSE") != NULL) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming hotlist seed skipped after partial load "
                    "(layer=%u requested=%u loaded=%u)\n",
                    layer,
                    n_experts,
                    loaded);
        }
        return 1;
    }

    if (getenv("DS4_ROCM_STREAMING_EXPERT_CACHE_VERBOSE") != NULL) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "streaming hotlist seeded layer=%u requested=%u cached=%u budget=%u\n",
                layer,
                n_experts,
                chosen_count,
                g_stream_expert_cache_budget);
    }
    return 1;
}

static void cuda_stream_selected_cache_header(
        const void *model_map,
        uint32_t layer,
        uint32_t n_total_expert,
        uint32_t n_selected,
        const int32_t *selected_ids,
        uint64_t gate_expert_bytes,
        uint64_t down_expert_bytes) {
    g_stream_selected_cache.model_map = model_map;
    g_stream_selected_cache.layer = layer;
    g_stream_selected_cache.n_total_expert = n_total_expert;
    g_stream_selected_cache.n_selected = n_selected;
    g_stream_selected_cache.gate_expert_bytes = gate_expert_bytes;
    g_stream_selected_cache.down_expert_bytes = down_expert_bytes;
    for (uint32_t i = 0; i < n_selected; i++) {
        g_stream_selected_cache.selected_ids[i] = selected_ids[i];
    }
}

static int cuda_stream_selected_compact_mask(
        const void *model_map,
        uint32_t layer,
        const int32_t *selected_ids,
        uint32_t n_total_expert,
        uint32_t n_selected,
        uint64_t gate_offset,
        uint64_t up_offset,
        uint64_t down_offset,
        uint64_t gate_expert_bytes,
        uint64_t down_expert_bytes,
        uint32_t mask) {
    (void)n_total_expert;
    if (mask == 0) return 1;
    if (!selected_ids || !cuda_stream_selected_ensure_stream()) return 0;
    cudaError_t err = cudaSuccess;
    for (uint32_t i = 0; i < n_selected; i++) {
        if ((mask & (1u << i)) == 0) continue;
        int idx = cuda_stream_resident_find(model_map,
                                            layer,
                                            selected_ids[i],
                                            gate_offset,
                                            up_offset,
                                            down_offset,
                                            gate_expert_bytes,
                                            down_expert_bytes);
        if (idx < 0) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming selected resident expert missing during compact\n");
            return 0;
        }
        cuda_stream_resident_expert &entry =
            g_stream_resident_experts[(size_t)idx];
        entry.last_used = ++g_stream_resident_clock;
        err = cudaMemcpyAsync(g_stream_selected_cache.gate +
                                  (uint64_t)i * gate_expert_bytes,
                              entry.gate,
                              (size_t)gate_expert_bytes,
                              cudaMemcpyDeviceToDevice,
                              g_stream_selected_upload_stream);
        if (err == cudaSuccess) {
            err = cudaMemcpyAsync(g_stream_selected_cache.up +
                                      (uint64_t)i * gate_expert_bytes,
                                  entry.up,
                                  (size_t)gate_expert_bytes,
                                  cudaMemcpyDeviceToDevice,
                                  g_stream_selected_upload_stream);
        }
        if (err == cudaSuccess) {
            err = cudaMemcpyAsync(g_stream_selected_cache.down +
                                      (uint64_t)i * down_expert_bytes,
                                  entry.down,
                                  (size_t)down_expert_bytes,
                                  cudaMemcpyDeviceToDevice,
                                  g_stream_selected_upload_stream);
        }
        if (err != cudaSuccess) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming selected compact copy failed: %s\n",
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            return 0;
        }
    }
    if (!cuda_stream_selected_upload_record_ready()) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "streaming selected compact event record failed\n");
        return 0;
    }
    return 1;
}

static int cuda_stream_selected_prepare_ptrs(
        const void *model_map,
        uint32_t layer,
        const int32_t *selected_ids,
        uint32_t n_selected,
        uint64_t gate_offset,
        uint64_t up_offset,
        uint64_t down_offset,
        uint64_t gate_expert_bytes,
        uint64_t down_expert_bytes) {
    if (!selected_ids ||
        n_selected == 0 ||
        n_selected > DS4_ROCM_N_EXPERT_USED ||
        !g_stream_selected_cache.gate_ptrs ||
        !g_stream_selected_cache.up_ptrs ||
        !g_stream_selected_cache.down_ptrs ||
        !cuda_stream_selected_ensure_stream()) {
        return 0;
    }
    const char *gate_ptrs[DS4_ROCM_N_EXPERT_USED] = {0};
    const char *up_ptrs[DS4_ROCM_N_EXPERT_USED] = {0};
    const char *down_ptrs[DS4_ROCM_N_EXPERT_USED] = {0};
    for (uint32_t i = 0; i < n_selected; i++) {
        int idx = cuda_stream_resident_find(model_map,
                                            layer,
                                            selected_ids[i],
                                            gate_offset,
                                            up_offset,
                                            down_offset,
                                            gate_expert_bytes,
                                            down_expert_bytes);
        if (idx < 0) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming selected pointer expert missing\n");
            return 0;
        }
        cuda_stream_resident_expert &entry =
            g_stream_resident_experts[(size_t)idx];
        entry.last_used = ++g_stream_resident_clock;
        gate_ptrs[i] = entry.gate;
        up_ptrs[i] = entry.up;
        down_ptrs[i] = entry.down;
    }
    if (!g_stream_selected_cache.gate_ptrs_stage ||
        !g_stream_selected_cache.up_ptrs_stage ||
        !g_stream_selected_cache.down_ptrs_stage) {
        return 0;
    }
    const size_t ptr_bytes = n_selected * sizeof(gate_ptrs[0]);
    memcpy((void *)g_stream_selected_cache.gate_ptrs_stage,
           gate_ptrs,
           ptr_bytes);
    memcpy((void *)g_stream_selected_cache.up_ptrs_stage,
           up_ptrs,
           ptr_bytes);
    memcpy((void *)g_stream_selected_cache.down_ptrs_stage,
           down_ptrs,
           ptr_bytes);
    cudaError_t err = cudaMemcpyAsync(g_stream_selected_cache.gate_ptrs,
                                      g_stream_selected_cache.gate_ptrs_stage,
                                      ptr_bytes,
                                      cudaMemcpyHostToDevice,
                                      g_stream_selected_upload_stream);
    if (err == cudaSuccess) {
        err = cudaMemcpyAsync(g_stream_selected_cache.up_ptrs,
                              g_stream_selected_cache.up_ptrs_stage,
                              ptr_bytes,
                              cudaMemcpyHostToDevice,
                              g_stream_selected_upload_stream);
    }
    if (err == cudaSuccess) {
        err = cudaMemcpyAsync(g_stream_selected_cache.down_ptrs,
                              g_stream_selected_cache.down_ptrs_stage,
                              ptr_bytes,
                              cudaMemcpyHostToDevice,
                              g_stream_selected_upload_stream);
    }
    int upload_record_ok = 1;
    if (err == cudaSuccess) {
        upload_record_ok = cuda_stream_selected_upload_record_ready();
    }
    if (err != cudaSuccess || !upload_record_ok) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "streaming selected pointer upload failed%s%s\n",
                err != cudaSuccess ? ": " : "",
                err != cudaSuccess ? cudaGetErrorString(err) : "");
        if (err != cudaSuccess) (void)cudaGetLastError();
        return 0;
    }
    return 1;
}

static int cuda_stream_batch_selected_prepare_from_host(
        const void *model_map,
        uint64_t model_size,
        uint32_t layer,
        const int32_t *ids,
        uint32_t n_tokens,
        uint32_t n_total_expert,
        uint32_t n_selected,
        uint64_t gate_offset,
        uint64_t up_offset,
        uint64_t down_offset,
        uint64_t gate_expert_bytes,
        uint64_t down_expert_bytes,
        const ds4_gpu_tensor **selected_exec,
        const char ***gate_ptrs,
        const char ***up_ptrs,
        const char ***down_ptrs,
        uint32_t *unique_out,
        int begin_pending) {
    if (!g_ssd_streaming_mode ||
        !model_map ||
        !ids ||
        !selected_exec ||
        !gate_ptrs ||
        !up_ptrs ||
        !down_ptrs ||
        !unique_out ||
        n_tokens <= 1 ||
        n_total_expert == 0 ||
        n_total_expert > DS4_ROCM_MAX_N_EXPERT ||
        n_selected == 0 ||
        n_selected > DS4_ROCM_N_EXPERT_USED ||
        gate_expert_bytes == 0 ||
        down_expert_bytes == 0) {
        return 0;
    }
    if (g_stream_batch_selected_pending.active) {
        cuda_stream_batch_selected_abort_pending();
    }
    if (g_stream_selected_pending.active) {
        cuda_stream_selected_abort_pending();
    }
    g_stream_batch_selected_cache.loaded = 0;

    uint64_t n_ids64 = 0;
    if (!cuda_u64_mul_checked(n_tokens, n_selected, &n_ids64) ||
        n_ids64 > SIZE_MAX / sizeof(int32_t)) {
        return 0;
    }
    int32_t *compact_ids = (int32_t *)malloc((size_t)n_ids64 * sizeof(compact_ids[0]));
    if (!compact_ids) {
        free(compact_ids);
        return 0;
    }
    uint8_t *pair_missing = (uint8_t *)malloc((size_t)n_ids64);
    if (!pair_missing) {
        free(compact_ids);
        return 0;
    }

    int ok = 1;
    int32_t map[DS4_ROCM_MAX_N_EXPERT];
    int32_t unique_ids[DS4_ROCM_MAX_N_EXPERT];
    uint8_t unique_missing[DS4_ROCM_MAX_N_EXPERT] = {0};
    for (uint32_t i = 0; i < DS4_ROCM_MAX_N_EXPERT; i++) map[i] = -1;
    uint32_t unique_count = 0;
    if (ok) {
        for (uint64_t i = 0; i < n_ids64; i++) {
            const int32_t expert = ids[i];
            if (expert < 0 || (uint32_t)expert >= n_total_expert) {
                fprintf(stderr,
                        DS4_GPU_LOG_PREFIX "streaming batch selected expert id %d outside 0..%u "
                        "(layer=%u)\n",
                        expert,
                        n_total_expert,
                        layer);
                ok = 0;
                break;
            }
            int32_t slot = map[(uint32_t)expert];
            if (slot < 0) {
                if (unique_count >= DS4_ROCM_MAX_N_EXPERT) {
                    ok = 0;
                    break;
                }
                slot = (int32_t)unique_count;
                map[(uint32_t)expert] = slot;
                unique_ids[unique_count++] = expert;
            }
            compact_ids[i] = slot;
        }
    }
    if (ok && unique_count == 0) ok = 0;
    if (ok && !cuda_stream_batch_selected_ensure_buffers(n_ids64, unique_count)) {
        ok = 0;
    }
    if (ok && !cuda_stream_selected_ensure_stream()) ok = 0;
    const int stats_on = cuda_stream_cache_stats_on();
    const int layer_stats_on =
        cuda_stream_cache_layer_stats_on() &&
        layer < DS4_ROCM_STREAM_CACHE_LAYER_STATS_MAX;
    if (ok && stats_on) {
        g_stream_cache_stats.batch_calls++;
        g_stream_cache_stats.batch_unique += unique_count;
    }
    if (ok && layer_stats_on) {
        cuda_stream_cache_layer_stats *ls = &g_stream_cache_layer_stats[layer];
        ls->batch_calls++;
        ls->batch_unique += unique_count;
    }

    cuda_stream_read_job read_jobs[DS4_ROCM_STREAM_READ_MAX_JOBS];
    memset(read_jobs, 0, sizeof(read_jobs));
    uint32_t read_job_count = 0;
    const int use_fd =
        g_model_fd >= 0 &&
        (g_model_fd_host_base == NULL || model_map == g_model_fd_host_base);

    const char *gate_host[DS4_ROCM_MAX_N_EXPERT] = {0};
    const char *up_host[DS4_ROCM_MAX_N_EXPERT] = {0};
    const char *down_host[DS4_ROCM_MAX_N_EXPERT] = {0};
    const char *resident_gate_host[DS4_ROCM_MAX_N_EXPERT] = {0};
    const char *resident_up_host[DS4_ROCM_MAX_N_EXPERT] = {0};
    const char *missing_gate_host[DS4_ROCM_MAX_N_EXPERT] = {0};
    const char *missing_up_host[DS4_ROCM_MAX_N_EXPERT] = {0};
    uint32_t resident_count = 0;
    uint32_t missing_count = 0;

    for (uint32_t u = 0; ok && u < unique_count; u++) {
        const int32_t expert_i = unique_ids[u];
        const uint64_t expert = (uint64_t)(uint32_t)expert_i;
        uint64_t gate_rel = 0;
        uint64_t down_rel = 0;
        if (!cuda_u64_mul_checked(expert, gate_expert_bytes, &gate_rel) ||
            !cuda_u64_mul_checked(expert, down_expert_bytes, &down_rel) ||
            gate_rel > model_size ||
            down_rel > model_size ||
            gate_offset > model_size ||
            up_offset > model_size ||
            down_offset > model_size ||
            gate_rel > model_size - gate_offset ||
            gate_rel > model_size - up_offset ||
            down_rel > model_size - down_offset ||
            gate_expert_bytes > model_size - gate_offset - gate_rel ||
            gate_expert_bytes > model_size - up_offset - gate_rel ||
            down_expert_bytes > model_size - down_offset - down_rel) {
            fprintf(stderr, DS4_GPU_LOG_PREFIX "streaming batch selected expert offset overflow\n");
            ok = 0;
            break;
        }

        int idx = cuda_stream_resident_find(model_map,
                                            layer,
                                            expert_i,
                                            gate_offset,
                                            up_offset,
                                            down_offset,
                                            gate_expert_bytes,
                                            down_expert_bytes);
        const int was_resident = idx >= 0;
        if (stats_on) {
            if (was_resident) {
                g_stream_cache_stats.batch_hits++;
            } else {
                g_stream_cache_stats.batch_misses++;
            }
        }
        if (layer_stats_on) {
            cuda_stream_cache_layer_stats *ls = &g_stream_cache_layer_stats[layer];
            if (was_resident) {
                ls->batch_hits++;
            } else {
                ls->batch_misses++;
            }
        }
        if (idx < 0) {
            idx = cuda_stream_resident_alloc(model_map,
                                             layer,
                                             expert_i,
                                             unique_ids,
                                             unique_count,
                                             gate_offset,
                                             up_offset,
                                             down_offset,
                                             gate_expert_bytes,
                                             down_expert_bytes);
            if (idx < 0) {
                ok = 0;
                break;
            }
            cuda_stream_resident_expert &entry =
                g_stream_resident_experts[(size_t)idx];
            if (use_fd) {
                if (read_job_count + 3u > DS4_ROCM_STREAM_READ_MAX_JOBS) {
                    if (!cuda_stream_flush_read_jobs(read_jobs, &read_job_count)) {
                        ok = 0;
                        break;
                    }
                }
                read_jobs[read_job_count++] =
                    {entry.gate, gate_offset + gate_rel, gate_expert_bytes,
                     NULL, NULL, 0, 0};
                read_jobs[read_job_count++] =
                    {entry.up, up_offset + gate_rel, gate_expert_bytes,
                     NULL, NULL, 0, 0};
                read_jobs[read_job_count++] =
                    {entry.down, down_offset + down_rel, down_expert_bytes,
                     NULL, NULL, 0, 0};
            } else {
                cudaError_t err = cudaMemcpyAsync(entry.gate,
                                                  (const char *)model_map + gate_offset + gate_rel,
                                                  (size_t)gate_expert_bytes,
                                                  cudaMemcpyHostToDevice,
                                                  g_stream_selected_upload_stream);
                if (err == cudaSuccess) {
                    err = cudaMemcpyAsync(entry.up,
                                          (const char *)model_map + up_offset + gate_rel,
                                          (size_t)gate_expert_bytes,
                                          cudaMemcpyHostToDevice,
                                          g_stream_selected_upload_stream);
                }
                if (err == cudaSuccess) {
                    err = cudaMemcpyAsync(entry.down,
                                          (const char *)model_map + down_offset + down_rel,
                                          (size_t)down_expert_bytes,
                                          cudaMemcpyHostToDevice,
                                          g_stream_selected_upload_stream);
                }
                if (err != cudaSuccess) {
                    fprintf(stderr,
                            DS4_GPU_LOG_PREFIX "streaming batch selected cached copy failed: %s\n",
                            cudaGetErrorString(err));
                    (void)cudaGetLastError();
                    ok = 0;
                    break;
                }
            }
        }
        if (idx >= 0) {
            cuda_stream_resident_expert &entry =
                g_stream_resident_experts[(size_t)idx];
            entry.last_used = ++g_stream_resident_clock;
            gate_host[u] = entry.gate;
            up_host[u] = entry.up;
            down_host[u] = entry.down;
            if (was_resident) {
                resident_gate_host[u] = entry.gate;
                resident_up_host[u] = entry.up;
                resident_count++;
            } else {
                unique_missing[u] = 1;
                missing_gate_host[u] = entry.gate;
                missing_up_host[u] = entry.up;
                missing_count++;
            }
        }
    }

    if (ok) {
        for (uint64_t i = 0; i < n_ids64; i++) {
            const int32_t slot = compact_ids[i];
            if (slot < 0 || (uint32_t)slot >= unique_count) {
                ok = 0;
                break;
            }
            pair_missing[i] = unique_missing[(uint32_t)slot];
        }
    }

    if (ok && begin_pending) {
        memset(&g_stream_batch_selected_pending, 0,
               sizeof(g_stream_batch_selected_pending));
        g_stream_batch_selected_pending.active = 1;
        g_stream_batch_selected_pending.model_map = model_map;
        g_stream_batch_selected_pending.layer = layer;
        g_stream_batch_selected_pending.n_total_expert = n_total_expert;
        g_stream_batch_selected_pending.n_selected = n_selected;
        g_stream_batch_selected_pending.n_tokens = n_tokens;
        g_stream_batch_selected_pending.n_unique = unique_count;
        g_stream_batch_selected_pending.resident_count = resident_count;
        g_stream_batch_selected_pending.missing_count = missing_count;
        g_stream_batch_selected_pending.gate_offset = gate_offset;
        g_stream_batch_selected_pending.up_offset = up_offset;
        g_stream_batch_selected_pending.down_offset = down_offset;
        g_stream_batch_selected_pending.gate_expert_bytes = gate_expert_bytes;
        g_stream_batch_selected_pending.down_expert_bytes = down_expert_bytes;
        g_stream_batch_selected_pending.read_job_count = read_job_count;
        memcpy(g_stream_batch_selected_pending.read_jobs,
               read_jobs,
               (size_t)read_job_count * sizeof(read_jobs[0]));
        if (use_fd && read_job_count != 0 &&
            !cuda_stream_read_jobs_start(
                    g_stream_batch_selected_pending.read_jobs,
                    read_job_count)) {
            memset(&g_stream_batch_selected_pending, 0,
                   sizeof(g_stream_batch_selected_pending));
            ok = 0;
        } else if (use_fd && read_job_count != 0) {
            read_job_count = 0;
        }
    }
    if (ok && !cuda_stream_flush_read_jobs(read_jobs, &read_job_count)) {
        ok = 0;
    }
    if (ok && !use_fd) {
        cudaError_t err = cudaStreamSynchronize(g_stream_selected_upload_stream);
        if (err != cudaSuccess) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming batch selected upload sync failed: %s\n",
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            ok = 0;
        }
    }
    if (ok) {
        if (!g_stream_batch_selected_cache.selected_stage ||
            !g_stream_batch_selected_cache.pair_missing_stage ||
            !g_stream_batch_selected_cache.gate_ptrs_stage ||
            !g_stream_batch_selected_cache.up_ptrs_stage ||
            !g_stream_batch_selected_cache.down_ptrs_stage ||
            !g_stream_batch_selected_cache.resident_gate_ptrs_stage ||
            !g_stream_batch_selected_cache.resident_up_ptrs_stage ||
            !g_stream_batch_selected_cache.missing_gate_ptrs_stage ||
            !g_stream_batch_selected_cache.missing_up_ptrs_stage) {
            ok = 0;
        }
    }
    if (ok) {
        const size_t selected_bytes =
            (size_t)n_ids64 * sizeof(compact_ids[0]);
        const size_t pair_missing_bytes = (size_t)n_ids64;
        const size_t ptr_bytes = (size_t)unique_count * sizeof(gate_host[0]);
        memcpy(g_stream_batch_selected_cache.selected_stage,
               compact_ids,
               selected_bytes);
        memcpy(g_stream_batch_selected_cache.pair_missing_stage,
               pair_missing,
               pair_missing_bytes);
        memcpy((void *)g_stream_batch_selected_cache.gate_ptrs_stage,
               gate_host,
               ptr_bytes);
        memcpy((void *)g_stream_batch_selected_cache.up_ptrs_stage,
               up_host,
               ptr_bytes);
        memcpy((void *)g_stream_batch_selected_cache.down_ptrs_stage,
               down_host,
               ptr_bytes);
        memcpy((void *)g_stream_batch_selected_cache.resident_gate_ptrs_stage,
               resident_gate_host,
               ptr_bytes);
        memcpy((void *)g_stream_batch_selected_cache.resident_up_ptrs_stage,
               resident_up_host,
               ptr_bytes);
        memcpy((void *)g_stream_batch_selected_cache.missing_gate_ptrs_stage,
               missing_gate_host,
               ptr_bytes);
        memcpy((void *)g_stream_batch_selected_cache.missing_up_ptrs_stage,
               missing_up_host,
               ptr_bytes);

        cudaError_t err = cudaMemcpyAsync(g_stream_batch_selected_cache.selected_ids,
                                          g_stream_batch_selected_cache.selected_stage,
                                          selected_bytes,
                                          cudaMemcpyHostToDevice,
                                          g_stream_selected_upload_stream);
        if (err == cudaSuccess) {
            err = cudaMemcpyAsync(g_stream_batch_selected_cache.pair_missing,
                                  g_stream_batch_selected_cache.pair_missing_stage,
                                  pair_missing_bytes,
                                  cudaMemcpyHostToDevice,
                                  g_stream_selected_upload_stream);
        }
        if (err == cudaSuccess) {
            err = cudaMemcpyAsync(g_stream_batch_selected_cache.gate_ptrs,
                                  g_stream_batch_selected_cache.gate_ptrs_stage,
                                  ptr_bytes,
                                  cudaMemcpyHostToDevice,
                                  g_stream_selected_upload_stream);
        }
        if (err == cudaSuccess) {
            err = cudaMemcpyAsync(g_stream_batch_selected_cache.up_ptrs,
                                  g_stream_batch_selected_cache.up_ptrs_stage,
                                  ptr_bytes,
                                  cudaMemcpyHostToDevice,
                                  g_stream_selected_upload_stream);
        }
        if (err == cudaSuccess) {
            err = cudaMemcpyAsync(g_stream_batch_selected_cache.down_ptrs,
                                  g_stream_batch_selected_cache.down_ptrs_stage,
                                  ptr_bytes,
                                  cudaMemcpyHostToDevice,
                                  g_stream_selected_upload_stream);
        }
        if (err == cudaSuccess) {
            err = cudaMemcpyAsync(g_stream_batch_selected_cache.resident_gate_ptrs,
                                  g_stream_batch_selected_cache.resident_gate_ptrs_stage,
                                  ptr_bytes,
                                  cudaMemcpyHostToDevice,
                                  g_stream_selected_upload_stream);
        }
        if (err == cudaSuccess) {
            err = cudaMemcpyAsync(g_stream_batch_selected_cache.resident_up_ptrs,
                                  g_stream_batch_selected_cache.resident_up_ptrs_stage,
                                  ptr_bytes,
                                  cudaMemcpyHostToDevice,
                                  g_stream_selected_upload_stream);
        }
        if (err == cudaSuccess) {
            err = cudaMemcpyAsync(g_stream_batch_selected_cache.missing_gate_ptrs,
                                  g_stream_batch_selected_cache.missing_gate_ptrs_stage,
                                  ptr_bytes,
                                  cudaMemcpyHostToDevice,
                                  g_stream_selected_upload_stream);
        }
        if (err == cudaSuccess) {
            err = cudaMemcpyAsync(g_stream_batch_selected_cache.missing_up_ptrs,
                                  g_stream_batch_selected_cache.missing_up_ptrs_stage,
                                  ptr_bytes,
                                  cudaMemcpyHostToDevice,
                                  g_stream_selected_upload_stream);
        }
        int upload_record_ok = 1;
        if (err == cudaSuccess) {
            upload_record_ok =
                cuda_stream_batch_selected_upload_record_ready();
        }
        if (err != cudaSuccess || !upload_record_ok) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming batch selected table upload failed%s%s\n",
                    err != cudaSuccess ? ": " : "",
                    err != cudaSuccess ? cudaGetErrorString(err) : "");
            if (err != cudaSuccess) (void)cudaGetLastError();
            if (g_stream_batch_selected_pending.active) {
                cuda_stream_batch_selected_abort_pending();
            }
            ok = 0;
        }
    }

    if (ok) {
        g_stream_batch_selected_cache.loaded = 0;
        g_stream_batch_selected_cache.model_map = model_map;
        g_stream_batch_selected_cache.layer = layer;
        g_stream_batch_selected_cache.n_total_expert = n_total_expert;
        g_stream_batch_selected_cache.n_selected = n_selected;
        g_stream_batch_selected_cache.n_tokens = n_tokens;
        g_stream_batch_selected_cache.n_unique = unique_count;
        g_stream_batch_selected_cache.gate_offset = gate_offset;
        g_stream_batch_selected_cache.up_offset = up_offset;
        g_stream_batch_selected_cache.down_offset = down_offset;
        g_stream_batch_selected_cache.gate_expert_bytes = gate_expert_bytes;
        g_stream_batch_selected_cache.down_expert_bytes = down_expert_bytes;
        *selected_exec = &g_stream_batch_selected_cache.selected_tensor;
        *gate_ptrs = g_stream_batch_selected_cache.gate_ptrs;
        *up_ptrs = g_stream_batch_selected_cache.up_ptrs;
        *down_ptrs = g_stream_batch_selected_cache.down_ptrs;
        *unique_out = unique_count;
    } else {
        g_stream_batch_selected_cache.loaded = 0;
        if (g_stream_batch_selected_pending.active) {
            cuda_stream_batch_selected_abort_pending();
        }
        if (read_job_count != 0) cuda_stream_read_jobs_free(read_jobs, read_job_count);
    }

    free(compact_ids);
    free(pair_missing);
    return ok;
}

static int cuda_stream_batch_selected_prepare(
        const void *model_map,
        uint64_t model_size,
        uint32_t layer,
        const ds4_gpu_tensor *selected,
        uint32_t n_tokens,
        uint32_t n_total_expert,
        uint32_t n_selected,
        uint64_t gate_offset,
        uint64_t up_offset,
        uint64_t down_offset,
        uint64_t gate_expert_bytes,
        uint64_t down_expert_bytes,
        const ds4_gpu_tensor **selected_exec,
        const char ***gate_ptrs,
        const char ***up_ptrs,
        const char ***down_ptrs,
        uint32_t *unique_out) {
    if (!selected ||
        !cuda_tensor_has_elems2(selected, n_tokens, n_selected, sizeof(int32_t))) {
        return 0;
    }
    uint64_t n_ids64 = 0;
    if (!cuda_u64_mul_checked(n_tokens, n_selected, &n_ids64) ||
        n_ids64 > SIZE_MAX / sizeof(int32_t)) {
        return 0;
    }
    int32_t *ids = (int32_t *)malloc((size_t)n_ids64 * sizeof(ids[0]));
    if (!ids) return 0;

    const int copy_ok = cuda_ok(cudaMemcpy(ids,
                                           selected->ptr,
                                           (size_t)n_ids64 * sizeof(ids[0]),
                                           cudaMemcpyDeviceToHost),
                                "streaming batch selected ids copy");
    const int ok = copy_ok &&
        cuda_stream_batch_selected_prepare_from_host(model_map,
                                                     model_size,
                                                     layer,
                                                     ids,
                                                     n_tokens,
                                                     n_total_expert,
                                                     n_selected,
                                                     gate_offset,
                                                     up_offset,
                                                     down_offset,
                                                     gate_expert_bytes,
                                                     down_expert_bytes,
                                                     selected_exec,
                                                     gate_ptrs,
                                                     up_ptrs,
                                                     down_ptrs,
                                                     unique_out,
                                                     0);
    free(ids);
    return ok;
}

static int cuda_stream_layer_expert_cache_prepare_batch(
        const void *model_map,
        uint32_t layer,
        const ds4_gpu_tensor *selected,
        uint32_t n_tokens,
        uint32_t n_total_expert,
        uint32_t n_selected,
        uint64_t gate_offset,
        uint64_t up_offset,
        uint64_t down_offset,
        uint64_t gate_expert_bytes,
        uint64_t down_expert_bytes,
        const ds4_gpu_tensor **selected_exec,
        const char ***gate_ptrs,
        const char ***up_ptrs,
        const char ***down_ptrs,
        uint32_t *unique_out) {
    if (!selected ||
        !selected_exec ||
        !gate_ptrs ||
        !up_ptrs ||
        !down_ptrs ||
        !unique_out ||
        !cuda_tensor_has_elems2(selected, n_tokens, n_selected, sizeof(int32_t)) ||
        n_tokens <= 1 ||
        n_total_expert == 0 ||
        n_total_expert > DS4_ROCM_MAX_N_EXPERT ||
        n_selected == 0 ||
        n_selected > DS4_ROCM_N_EXPERT_USED ||
        gate_expert_bytes == 0 ||
        down_expert_bytes == 0) {
        return 0;
    }
    const char *layer_gate = NULL;
    const char *layer_up = NULL;
    const char *layer_down = NULL;
    if (!cuda_stream_layer_expert_cache_apply(model_map,
                                              layer,
                                              n_total_expert,
                                              gate_offset,
                                              up_offset,
                                              down_offset,
                                              gate_expert_bytes,
                                              down_expert_bytes,
                                              &layer_gate,
                                              &layer_up,
                                              &layer_down)) {
        return 0;
    }

    uint64_t n_ids64 = 0;
    if (!cuda_u64_mul_checked(n_tokens, n_selected, &n_ids64) ||
        n_ids64 > SIZE_MAX / sizeof(int32_t)) {
        return 0;
    }
    int32_t *ids = (int32_t *)malloc((size_t)n_ids64 * sizeof(ids[0]));
    int32_t *compact_ids =
        (int32_t *)malloc((size_t)n_ids64 * sizeof(compact_ids[0]));
    if (!ids || !compact_ids) {
        free(ids);
        free(compact_ids);
        return 0;
    }

    int ok = cuda_ok(cudaMemcpy(ids,
                                selected->ptr,
                                (size_t)n_ids64 * sizeof(ids[0]),
                                cudaMemcpyDeviceToHost),
                     "streaming full-layer selected ids copy");

    int32_t map[DS4_ROCM_MAX_N_EXPERT];
    int32_t unique_ids[DS4_ROCM_MAX_N_EXPERT];
    for (uint32_t i = 0; i < DS4_ROCM_MAX_N_EXPERT; i++) map[i] = -1;
    uint32_t unique_count = 0;
    for (uint64_t i = 0; ok && i < n_ids64; i++) {
        const int32_t expert = ids[i];
        if (expert < 0 || (uint32_t)expert >= n_total_expert) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming full-layer selected expert id %d "
                    "outside 0..%u (layer=%u)\n",
                    expert,
                    n_total_expert,
                    layer);
            ok = 0;
            break;
        }
        int32_t slot = map[(uint32_t)expert];
        if (slot < 0) {
            if (unique_count >= DS4_ROCM_MAX_N_EXPERT) {
                ok = 0;
                break;
            }
            slot = (int32_t)unique_count;
            map[(uint32_t)expert] = slot;
            unique_ids[unique_count++] = expert;
        }
        compact_ids[i] = slot;
    }
    if (ok && unique_count == 0) ok = 0;
    if (ok && !cuda_stream_batch_selected_ensure_buffers(n_ids64, unique_count)) {
        ok = 0;
    }
    if (ok && !cuda_stream_selected_ensure_stream()) ok = 0;

    const char *gate_host[DS4_ROCM_MAX_N_EXPERT] = {0};
    const char *up_host[DS4_ROCM_MAX_N_EXPERT] = {0};
    const char *down_host[DS4_ROCM_MAX_N_EXPERT] = {0};
    for (uint32_t u = 0; ok && u < unique_count; u++) {
        const uint64_t expert = (uint64_t)(uint32_t)unique_ids[u];
        uint64_t gate_rel = 0;
        uint64_t down_rel = 0;
        if (!cuda_u64_mul_checked(expert, gate_expert_bytes, &gate_rel) ||
            !cuda_u64_mul_checked(expert, down_expert_bytes, &down_rel)) {
            ok = 0;
            break;
        }
        gate_host[u] = layer_gate + gate_rel;
        up_host[u] = layer_up + gate_rel;
        down_host[u] = layer_down + down_rel;
    }

    if (ok) {
        cudaError_t err = cudaMemcpyAsync(g_stream_batch_selected_cache.selected_ids,
                                          compact_ids,
                                          (size_t)n_ids64 * sizeof(compact_ids[0]),
                                          cudaMemcpyHostToDevice,
                                          g_stream_selected_upload_stream);
        if (err == cudaSuccess) {
            err = cudaMemcpyAsync(g_stream_batch_selected_cache.gate_ptrs,
                                  gate_host,
                                  unique_count * sizeof(gate_host[0]),
                                  cudaMemcpyHostToDevice,
                                  g_stream_selected_upload_stream);
        }
        if (err == cudaSuccess) {
            err = cudaMemcpyAsync(g_stream_batch_selected_cache.up_ptrs,
                                  up_host,
                                  unique_count * sizeof(up_host[0]),
                                  cudaMemcpyHostToDevice,
                                  g_stream_selected_upload_stream);
        }
        if (err == cudaSuccess) {
            err = cudaMemcpyAsync(g_stream_batch_selected_cache.down_ptrs,
                                  down_host,
                                  unique_count * sizeof(down_host[0]),
                                  cudaMemcpyHostToDevice,
                                  g_stream_selected_upload_stream);
        }
        if (err == cudaSuccess) err = cudaStreamSynchronize(g_stream_selected_upload_stream);
        if (err != cudaSuccess) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming full-layer selected table upload failed: %s\n",
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            ok = 0;
        }
    }

    if (ok) {
        g_stream_batch_selected_cache.loaded = 0;
        g_stream_batch_selected_cache.model_map = model_map;
        g_stream_batch_selected_cache.layer = layer;
        g_stream_batch_selected_cache.n_total_expert = n_total_expert;
        g_stream_batch_selected_cache.n_selected = n_selected;
        g_stream_batch_selected_cache.n_tokens = n_tokens;
        g_stream_batch_selected_cache.n_unique = unique_count;
        g_stream_batch_selected_cache.gate_offset = gate_offset;
        g_stream_batch_selected_cache.up_offset = up_offset;
        g_stream_batch_selected_cache.down_offset = down_offset;
        g_stream_batch_selected_cache.gate_expert_bytes = gate_expert_bytes;
        g_stream_batch_selected_cache.down_expert_bytes = down_expert_bytes;
        *selected_exec = &g_stream_batch_selected_cache.selected_tensor;
        *gate_ptrs = g_stream_batch_selected_cache.gate_ptrs;
        *up_ptrs = g_stream_batch_selected_cache.up_ptrs;
        *down_ptrs = g_stream_batch_selected_cache.down_ptrs;
        *unique_out = unique_count;
    } else {
        g_stream_batch_selected_cache.loaded = 0;
    }

    free(ids);
    free(compact_ids);
    return ok;
}

static int cuda_stream_layer_expert_cache_seed_selected(
        const void *model_map,
        uint64_t model_size,
        uint32_t layer,
        const ds4_gpu_tensor *selected,
        uint32_t n_tokens,
        uint32_t n_seed_tokens,
        uint32_t n_total_expert,
        uint32_t n_selected,
        uint64_t gate_offset,
        uint64_t up_offset,
        uint64_t down_offset,
        uint64_t gate_expert_bytes,
        uint64_t down_expert_bytes) {
    if (!g_ssd_streaming_mode ||
        !model_map ||
        !selected ||
        n_tokens == 0 ||
        n_seed_tokens == 0 ||
        n_total_expert == 0 ||
        n_total_expert > DS4_ROCM_MAX_N_EXPERT ||
        n_selected == 0 ||
        n_selected > DS4_ROCM_N_EXPERT_USED ||
        gate_expert_bytes == 0 ||
        down_expert_bytes == 0 ||
        !cuda_tensor_has_elems2(selected, n_tokens, n_selected, sizeof(int32_t))) {
        return 0;
    }

    uint64_t gate_bytes = 0;
    uint64_t down_bytes = 0;
    if (!cuda_u64_mul_checked(n_total_expert, gate_expert_bytes, &gate_bytes) ||
        !cuda_u64_mul_checked(n_total_expert, down_expert_bytes, &down_bytes) ||
        gate_offset > model_size ||
        up_offset > model_size ||
        down_offset > model_size ||
        gate_bytes > model_size - gate_offset ||
        gate_bytes > model_size - up_offset ||
        down_bytes > model_size - down_offset) {
        return 0;
    }

    const char *layer_gate = NULL;
    const char *layer_up = NULL;
    const char *layer_down = NULL;
    if (!cuda_stream_layer_expert_cache_apply(model_map,
                                              layer,
                                              n_total_expert,
                                              gate_offset,
                                              up_offset,
                                              down_offset,
                                              gate_expert_bytes,
                                              down_expert_bytes,
                                              &layer_gate,
                                              &layer_up,
                                              &layer_down)) {
        if (!cuda_model_range_is_cached(model_map, gate_offset, gate_bytes) ||
            !cuda_model_range_is_cached(model_map, up_offset, gate_bytes) ||
            !cuda_model_range_is_cached(model_map, down_offset, down_bytes)) {
            return 0;
        }
        layer_gate = cuda_model_range_ptr(model_map,
                                          gate_offset,
                                          gate_bytes,
                                          "streaming full-layer seed gate");
        layer_up = cuda_model_range_ptr(model_map,
                                        up_offset,
                                        gate_bytes,
                                        "streaming full-layer seed up");
        layer_down = cuda_model_range_ptr(model_map,
                                          down_offset,
                                          down_bytes,
                                          "streaming full-layer seed down");
        if (!layer_gate || !layer_up || !layer_down) return 0;
    }

    if (n_seed_tokens > n_tokens) n_seed_tokens = n_tokens;
    const uint64_t n_ids64 = (uint64_t)n_seed_tokens * n_selected;
    if (n_ids64 == 0 || n_ids64 > SIZE_MAX / sizeof(int32_t)) return 0;
    int32_t ids_stack[DS4_ROCM_N_EXPERT_USED * 16u];
    int32_t *ids_heap = NULL;
    int32_t *ids = ids_stack;
    if (n_ids64 > sizeof(ids_stack) / sizeof(ids_stack[0])) {
        ids_heap = (int32_t *)malloc((size_t)n_ids64 * sizeof(ids_heap[0]));
        if (!ids_heap) return 0;
        ids = ids_heap;
    }

    const uint64_t src_off =
        (uint64_t)(n_tokens - n_seed_tokens) * n_selected * sizeof(int32_t);
    int ok = cuda_ok(cudaMemcpy(ids,
                                (const char *)selected->ptr + src_off,
                                (size_t)n_ids64 * sizeof(ids[0]),
                                cudaMemcpyDeviceToHost),
                     "streaming full-layer seed selected ids copy");

    int32_t unique_stack[DS4_ROCM_N_EXPERT_USED * 16u];
    int32_t *unique_heap = NULL;
    int32_t *unique = unique_stack;
    if (n_ids64 > sizeof(unique_stack) / sizeof(unique_stack[0])) {
        unique_heap = (int32_t *)malloc((size_t)n_ids64 * sizeof(unique_heap[0]));
        if (!unique_heap) ok = 0;
        unique = unique_heap;
    }
    uint32_t unique_count = 0;
    bool seen[DS4_ROCM_MAX_N_EXPERT] = {0};
    for (uint64_t i = 0; ok && i < n_ids64; i++) {
        const int32_t expert = ids[i];
        if (expert < 0 || (uint32_t)expert >= n_total_expert) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming full-layer seed expert id %d "
                    "outside 0..%u (layer=%u)\n",
                    expert,
                    n_total_expert,
                    layer);
            ok = 0;
            break;
        }
        if (seen[(uint32_t)expert]) continue;
        seen[(uint32_t)expert] = true;
        unique[unique_count++] = expert;
    }
    if (ok && cuda_stream_cache_stats_on()) {
        g_stream_cache_stats.seed_calls++;
        g_stream_cache_stats.seed_unique += unique_count;
    }

    if (ok && unique_count != 0 && !cuda_stream_selected_ensure_stream()) {
        ok = 0;
    }

    for (uint32_t u = 0; ok && u < unique_count; u++) {
        const int32_t expert_i32 = unique[u];
        int idx = cuda_stream_resident_find(model_map,
                                            layer,
                                            expert_i32,
                                            gate_offset,
                                            up_offset,
                                            down_offset,
                                            gate_expert_bytes,
                                            down_expert_bytes);
        if (idx >= 0) {
            g_stream_resident_experts[(size_t)idx].last_used =
                ++g_stream_resident_clock;
            continue;
        }

        idx = cuda_stream_resident_alloc(model_map,
                                         layer,
                                         expert_i32,
                                         unique,
                                         unique_count,
                                         gate_offset,
                                         up_offset,
                                         down_offset,
                                         gate_expert_bytes,
                                         down_expert_bytes);
        if (idx < 0) {
            ok = 0;
            break;
        }

        const uint64_t expert = (uint64_t)(uint32_t)expert_i32;
        uint64_t gate_rel = 0;
        uint64_t down_rel = 0;
        if (!cuda_u64_mul_checked(expert, gate_expert_bytes, &gate_rel) ||
            !cuda_u64_mul_checked(expert, down_expert_bytes, &down_rel)) {
            ok = 0;
            break;
        }

        cuda_stream_resident_expert &entry =
            g_stream_resident_experts[(size_t)idx];
        cudaError_t err = cudaMemcpyAsync(entry.gate,
                                          layer_gate + gate_rel,
                                          (size_t)gate_expert_bytes,
                                          cudaMemcpyDeviceToDevice,
                                          g_stream_selected_upload_stream);
        if (err == cudaSuccess) {
            err = cudaMemcpyAsync(entry.up,
                                  layer_up + gate_rel,
                                  (size_t)gate_expert_bytes,
                                  cudaMemcpyDeviceToDevice,
                                  g_stream_selected_upload_stream);
        }
        if (err == cudaSuccess) {
            err = cudaMemcpyAsync(entry.down,
                                  layer_down + down_rel,
                                  (size_t)down_expert_bytes,
                                  cudaMemcpyDeviceToDevice,
                                  g_stream_selected_upload_stream);
        }
        if (err != cudaSuccess) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming full-layer seed D2D copy failed: %s\n",
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            ok = 0;
            break;
        }
    }

    if (ok && unique_count != 0) {
        cudaError_t err = cudaStreamSynchronize(g_stream_selected_upload_stream);
        if (err != cudaSuccess) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming full-layer seed sync failed: %s\n",
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            ok = 0;
        }
    }

    if (!ok) cuda_stream_resident_cache_release();
    free(ids_heap);
    free(unique_heap);
    return ok;
}

static int cuda_stream_selected_load(
        const void *model_map,
        uint64_t model_size,
        uint32_t layer,
        const int32_t *selected_ids,
        uint32_t n_total_expert,
        uint32_t n_selected,
        uint64_t gate_offset,
        uint64_t up_offset,
        uint64_t down_offset,
        uint64_t gate_expert_bytes,
        uint64_t down_expert_bytes) {
    g_stream_selected_cache.loaded = 0;
    if (g_stream_selected_pending.active) {
        cuda_stream_selected_abort_pending();
    }
    if (g_stream_batch_selected_pending.active) {
        cuda_stream_batch_selected_abort_pending();
    }
    if (!g_ssd_streaming_mode) return 1;
    if (!model_map || !selected_ids ||
        n_total_expert == 0 ||
        n_total_expert > DS4_ROCM_MAX_N_EXPERT ||
        n_selected == 0 ||
        n_selected > DS4_ROCM_N_EXPERT_USED ||
        gate_expert_bytes == 0 ||
        down_expert_bytes == 0) {
        return 0;
    }
    uint64_t gate_bytes = 0;
    uint64_t down_bytes = 0;
    if (!cuda_u64_mul_checked(n_selected, gate_expert_bytes, &gate_bytes) ||
        !cuda_u64_mul_checked(n_selected, down_expert_bytes, &down_bytes)) {
        return 0;
    }
    if (!cuda_stream_selected_reuse_wait("streaming selected cache reuse")) {
        return 0;
    }
    if (!cuda_stream_selected_ensure_buffers(gate_bytes, down_bytes)) return 0;
    if (!cuda_stream_selected_ensure_stream()) return 0;
    cuda_stream_selected_cache_header(model_map,
                                      layer,
                                      n_total_expert,
                                      n_selected,
                                      selected_ids,
                                      gate_expert_bytes,
                                      down_expert_bytes);
    const int stats_on = cuda_stream_cache_stats_on();
    const int layer_stats_on =
        cuda_stream_cache_layer_stats_on() &&
        layer < DS4_ROCM_STREAM_CACHE_LAYER_STATS_MAX;
    if (stats_on) {
        g_stream_cache_stats.selected_calls++;
        g_stream_cache_stats.selected_slots += n_selected;
    }
    if (layer_stats_on) {
        cuda_stream_cache_layer_stats *ls = &g_stream_cache_layer_stats[layer];
        ls->selected_calls++;
        ls->selected_slots += n_selected;
    }

    cuda_stream_read_job read_jobs[DS4_ROCM_N_EXPERT_USED * 3u];
    memset(read_jobs, 0, sizeof(read_jobs));
    uint32_t read_job_count = 0;
    uint32_t resident_mask = 0;
    uint32_t missing_mask = 0;
    const int use_fd =
        g_model_fd >= 0 &&
        (g_model_fd_host_base == NULL || model_map == g_model_fd_host_base);

    for (uint32_t i = 0; i < n_selected; i++) {
        if (selected_ids[i] < 0 || (uint32_t)selected_ids[i] >= n_total_expert) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming selected expert id %d outside 0..%u "
                    "(layer=%u slot=%u selected=[%d,%d,%d,%d,%d,%d,%d,%d])\n",
                    selected_ids[i],
                    n_total_expert,
                    layer,
                    i,
                    n_selected > 0 ? selected_ids[0] : -1,
                    n_selected > 1 ? selected_ids[1] : -1,
                    n_selected > 2 ? selected_ids[2] : -1,
                    n_selected > 3 ? selected_ids[3] : -1,
                    n_selected > 4 ? selected_ids[4] : -1,
                    n_selected > 5 ? selected_ids[5] : -1,
                    n_selected > 6 ? selected_ids[6] : -1,
                    n_selected > 7 ? selected_ids[7] : -1);
            return 0;
        }
        const uint64_t expert = (uint64_t)(uint32_t)selected_ids[i];
        uint64_t gate_rel = 0;
        uint64_t down_rel = 0;
        if (!cuda_u64_mul_checked(expert, gate_expert_bytes, &gate_rel) ||
            !cuda_u64_mul_checked(expert, down_expert_bytes, &down_rel) ||
            gate_rel > model_size ||
            down_rel > model_size ||
            gate_offset > model_size ||
            up_offset > model_size ||
            down_offset > model_size ||
            gate_rel > model_size - gate_offset ||
            gate_rel > model_size - up_offset ||
            down_rel > model_size - down_offset ||
            gate_expert_bytes > model_size - gate_offset - gate_rel ||
            gate_expert_bytes > model_size - up_offset - gate_rel ||
            down_expert_bytes > model_size - down_offset - down_rel) {
            fprintf(stderr, DS4_GPU_LOG_PREFIX "streaming selected expert offset overflow\n");
            return 0;
        }

        int idx = cuda_stream_resident_find(model_map,
                                            layer,
                                            selected_ids[i],
                                            gate_offset,
                                            up_offset,
                                            down_offset,
                                            gate_expert_bytes,
                                            down_expert_bytes);
        if (stats_on) {
            if (idx >= 0) {
                g_stream_cache_stats.selected_hits++;
            } else {
                g_stream_cache_stats.selected_misses++;
            }
        }
        if (layer_stats_on) {
            cuda_stream_cache_layer_stats *ls = &g_stream_cache_layer_stats[layer];
            if (idx >= 0) {
                ls->selected_hits++;
            } else {
                ls->selected_misses++;
            }
        }
        if (idx >= 0) {
            g_stream_resident_experts[(size_t)idx].last_used =
                ++g_stream_resident_clock;
            resident_mask |= 1u << i;
            continue;
        }

        idx = cuda_stream_resident_alloc(model_map,
                                         layer,
                                         selected_ids[i],
                                         selected_ids,
                                         n_selected,
                                         gate_offset,
                                         up_offset,
                                         down_offset,
                                         gate_expert_bytes,
                                         down_expert_bytes);
        if (idx < 0) return 0;
        missing_mask |= 1u << i;
        cuda_stream_resident_expert &entry =
            g_stream_resident_experts[(size_t)idx];

        if (use_fd) {
            if (read_job_count + 3u > DS4_ROCM_N_EXPERT_USED * 3u) return 0;
            read_jobs[read_job_count++] =
                {entry.gate, gate_offset + gate_rel, gate_expert_bytes,
                 NULL, NULL, 0, 0};
            read_jobs[read_job_count++] =
                {entry.up, up_offset + gate_rel, gate_expert_bytes,
                 NULL, NULL, 0, 0};
            read_jobs[read_job_count++] =
                {entry.down, down_offset + down_rel, down_expert_bytes,
                 NULL, NULL, 0, 0};
        } else {
            cudaError_t err = cudaMemcpyAsync(entry.gate,
                                              (const char *)model_map + gate_offset + gate_rel,
                                              (size_t)gate_expert_bytes,
                                              cudaMemcpyHostToDevice,
                                              g_stream_selected_upload_stream);
            if (err == cudaSuccess) {
                err = cudaMemcpyAsync(entry.up,
                                      (const char *)model_map + up_offset + gate_rel,
                                      (size_t)gate_expert_bytes,
                                      cudaMemcpyHostToDevice,
                                      g_stream_selected_upload_stream);
            }
            if (err == cudaSuccess) {
                err = cudaMemcpyAsync(entry.down,
                                      (const char *)model_map + down_offset + down_rel,
                                      (size_t)down_expert_bytes,
                                      cudaMemcpyHostToDevice,
                                      g_stream_selected_upload_stream);
            }
            if (err != cudaSuccess) {
                fprintf(stderr,
                        DS4_GPU_LOG_PREFIX "streaming selected cached copy failed: %s\n",
                        cudaGetErrorString(err));
                (void)cudaGetLastError();
                cuda_stream_resident_cache_release();
                return 0;
            }
        }
    }

    if (resident_mask != 0 && missing_mask == 0) {
        g_stream_selected_pending.active = 1;
        g_stream_selected_pending.model_map = model_map;
        g_stream_selected_pending.layer = layer;
        g_stream_selected_pending.n_total_expert = n_total_expert;
        g_stream_selected_pending.n_selected = n_selected;
        g_stream_selected_pending.gate_offset = gate_offset;
        g_stream_selected_pending.up_offset = up_offset;
        g_stream_selected_pending.down_offset = down_offset;
        g_stream_selected_pending.gate_expert_bytes = gate_expert_bytes;
        g_stream_selected_pending.down_expert_bytes = down_expert_bytes;
        g_stream_selected_pending.resident_mask = resident_mask;
        g_stream_selected_pending.missing_mask = 0;
        g_stream_selected_pending.read_job_count = 0;
        for (uint32_t i = 0; i < n_selected; i++) {
            g_stream_selected_pending.selected_ids[i] = selected_ids[i];
        }
        if (!cuda_stream_selected_prepare_ptrs(model_map,
                                               layer,
                                               selected_ids,
                                               n_selected,
                                               gate_offset,
                                               up_offset,
                                               down_offset,
                                               gate_expert_bytes,
                                               down_expert_bytes)) {
            memset(&g_stream_selected_pending, 0,
                   sizeof(g_stream_selected_pending));
            cuda_stream_resident_cache_release();
            return 0;
        }
        return 1;
    }

    if (use_fd && read_job_count != 0) {
        g_stream_selected_pending.active = 1;
        g_stream_selected_pending.model_map = model_map;
        g_stream_selected_pending.layer = layer;
        g_stream_selected_pending.n_total_expert = n_total_expert;
        g_stream_selected_pending.n_selected = n_selected;
        g_stream_selected_pending.gate_offset = gate_offset;
        g_stream_selected_pending.up_offset = up_offset;
        g_stream_selected_pending.down_offset = down_offset;
        g_stream_selected_pending.gate_expert_bytes = gate_expert_bytes;
        g_stream_selected_pending.down_expert_bytes = down_expert_bytes;
        g_stream_selected_pending.resident_mask = resident_mask;
        g_stream_selected_pending.missing_mask = missing_mask;
        g_stream_selected_pending.read_job_count = read_job_count;
        for (uint32_t i = 0; i < n_selected; i++) {
            g_stream_selected_pending.selected_ids[i] = selected_ids[i];
        }
        memcpy(g_stream_selected_pending.read_jobs,
               read_jobs,
               (size_t)read_job_count * sizeof(read_jobs[0]));
        if (!cuda_stream_read_jobs_start(g_stream_selected_pending.read_jobs,
                                         read_job_count)) {
            memset(&g_stream_selected_pending, 0, sizeof(g_stream_selected_pending));
            cuda_stream_read_jobs_free(read_jobs, read_job_count);
            cuda_stream_resident_cache_release();
            return 0;
        }
        if (!cuda_stream_selected_prepare_ptrs(model_map,
                                               layer,
                                               selected_ids,
                                               n_selected,
                                               gate_offset,
                                               up_offset,
                                               down_offset,
                                               gate_expert_bytes,
                                               down_expert_bytes)) {
            (void)cuda_stream_read_jobs_wait(g_stream_selected_pending.read_jobs,
                                             read_job_count);
            cuda_stream_read_jobs_free(g_stream_selected_pending.read_jobs,
                                       read_job_count);
            memset(&g_stream_selected_pending, 0, sizeof(g_stream_selected_pending));
            cuda_stream_read_jobs_free(read_jobs, read_job_count);
            cuda_stream_resident_cache_release();
            return 0;
        }
        return 1;
    }

    if (resident_mask != 0 &&
        !cuda_stream_selected_compact_mask(model_map,
                                           layer,
                                           selected_ids,
                                           n_total_expert,
                                           n_selected,
                                           gate_offset,
                                           up_offset,
                                           down_offset,
                                           gate_expert_bytes,
                                           down_expert_bytes,
                                           resident_mask)) {
        cuda_stream_read_jobs_free(read_jobs, read_job_count);
        cuda_stream_resident_cache_release();
        return 0;
    }

    if (read_job_count != 0) {
        if (!cuda_stream_read_jobs_parallel(read_jobs, read_job_count) ||
            !cuda_stream_selected_upload_read_jobs(read_jobs, read_job_count)) {
            cuda_stream_read_jobs_free(read_jobs, read_job_count);
            cuda_stream_resident_cache_release();
            return 0;
        }
    } else {
        cudaError_t err = cudaStreamSynchronize(g_stream_selected_upload_stream);
        if (err != cudaSuccess) {
            fprintf(stderr, DS4_GPU_LOG_PREFIX "streaming selected upload sync failed: %s\n",
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            cuda_stream_read_jobs_free(read_jobs, read_job_count);
            cuda_stream_resident_cache_release();
            return 0;
        }
    }
    cuda_stream_read_jobs_free(read_jobs, read_job_count);

    {
        const uint32_t all_mask =
            n_selected >= 32u ? 0xffffffffu : ((1u << n_selected) - 1u);
        const uint32_t compact_mask = resident_mask != 0 ? missing_mask : all_mask;
        if (!cuda_stream_selected_compact_mask(model_map,
                                               layer,
                                               selected_ids,
                                               n_total_expert,
                                               n_selected,
                                               gate_offset,
                                               up_offset,
                                               down_offset,
                                               gate_expert_bytes,
                                               down_expert_bytes,
                                               compact_mask)) {
            cuda_stream_resident_cache_release();
            return 0;
        }
    }

    g_stream_selected_cache.loaded = 1;
    return 1;
}

static int cuda_stream_selected_pending_matches(
        const void *model_map,
        uint32_t layer,
        uint32_t n_total_expert,
        uint32_t n_selected,
        uint64_t gate_expert_bytes,
        uint64_t down_expert_bytes) {
    if (!g_stream_selected_pending.active ||
        g_routed_moe_selected_override_n != n_selected ||
        g_stream_selected_pending.model_map != model_map ||
        g_stream_selected_pending.layer != layer ||
        g_stream_selected_pending.n_total_expert != n_total_expert ||
        g_stream_selected_pending.n_selected != n_selected ||
        g_stream_selected_pending.gate_expert_bytes != gate_expert_bytes ||
        g_stream_selected_pending.down_expert_bytes != down_expert_bytes) {
        return 0;
    }
    for (uint32_t i = 0; i < n_selected; i++) {
        if (g_stream_selected_pending.selected_ids[i] !=
            g_routed_moe_selected_override[i]) {
            return 0;
        }
    }
    return 1;
}

static int cuda_stream_selected_finish_pending_missing(uint32_t compact_mask);

static int cuda_stream_selected_apply_split(
        const void *model_map,
        uint32_t layer,
        uint32_t n_total_expert,
        uint32_t n_selected,
        uint64_t gate_expert_bytes,
        uint64_t down_expert_bytes,
        const ds4_gpu_tensor **selected_exec,
        const char **gate_w,
        const char **up_w,
        const char **down_w,
        const char ***gate_ptrs,
        const char ***up_ptrs,
        const char ***down_ptrs,
        uint32_t *resident_mask,
        uint32_t *missing_mask) {
    if (!g_ssd_streaming_mode ||
        !selected_exec ||
        !gate_w ||
        !up_w ||
        !down_w ||
        !gate_ptrs ||
        !up_ptrs ||
        !down_ptrs ||
        !resident_mask ||
        !missing_mask ||
        !cuda_stream_selected_pending_matches(model_map,
                                              layer,
                                              n_total_expert,
                                              n_selected,
                                              gate_expert_bytes,
                                              down_expert_bytes)) {
        return 0;
    }
    if ((g_stream_selected_pending.resident_mask |
         g_stream_selected_pending.missing_mask) == 0 ||
        !g_stream_selected_cache.gate_ptrs ||
        !g_stream_selected_cache.up_ptrs ||
        !g_stream_selected_cache.down_ptrs) {
        return 0;
    }
    *selected_exec = &g_stream_selected_cache.slot_tensor;
    *gate_w = g_stream_selected_cache.gate;
    *up_w = g_stream_selected_cache.up;
    *down_w = g_stream_selected_cache.down;
    *gate_ptrs = g_stream_selected_cache.gate_ptrs;
    *up_ptrs = g_stream_selected_cache.up_ptrs;
    *down_ptrs = g_stream_selected_cache.down_ptrs;
    *resident_mask = g_stream_selected_pending.resident_mask;
    *missing_mask = g_stream_selected_pending.missing_mask;
    g_routed_moe_selected_override_n = 0;
    return 1;
}

static int cuda_stream_selected_finish_pending_missing(uint32_t compact_mask) {
    if (!g_stream_selected_pending.active) return 1;
    const uint32_t read_job_count = g_stream_selected_pending.read_job_count;
    if (!cuda_stream_read_jobs_wait(g_stream_selected_pending.read_jobs,
                                    read_job_count) ||
        !cuda_stream_selected_upload_read_jobs(g_stream_selected_pending.read_jobs,
                                              read_job_count)) {
        cuda_stream_read_jobs_free(g_stream_selected_pending.read_jobs,
                                   read_job_count);
        memset(&g_stream_selected_pending, 0,
               sizeof(g_stream_selected_pending));
        cuda_stream_resident_cache_release();
        return 0;
    }
    cuda_stream_read_jobs_free(g_stream_selected_pending.read_jobs,
                               read_job_count);
    if (compact_mask != 0 &&
        !cuda_stream_selected_compact_mask(
                g_stream_selected_pending.model_map,
                g_stream_selected_pending.layer,
                g_stream_selected_pending.selected_ids,
                g_stream_selected_pending.n_total_expert,
                g_stream_selected_pending.n_selected,
                g_stream_selected_pending.gate_offset,
                g_stream_selected_pending.up_offset,
                g_stream_selected_pending.down_offset,
                g_stream_selected_pending.gate_expert_bytes,
                g_stream_selected_pending.down_expert_bytes,
                compact_mask)) {
        memset(&g_stream_selected_pending, 0, sizeof(g_stream_selected_pending));
        cuda_stream_resident_cache_release();
        return 0;
    }
    g_stream_selected_cache.loaded = compact_mask != 0 ? 1 : 0;
    memset(&g_stream_selected_pending, 0, sizeof(g_stream_selected_pending));
    return 1;
}

static int cuda_stream_selected_apply(
        const void *model_map,
        uint32_t layer,
        uint32_t n_total_expert,
        uint32_t n_selected,
        uint64_t gate_expert_bytes,
        uint64_t down_expert_bytes,
        const ds4_gpu_tensor **selected_exec,
        const char **gate_w,
        const char **up_w,
        const char **down_w) {
    if (g_ssd_streaming_mode &&
        !g_stream_selected_cache.loaded &&
        getenv("DS4_ROCM_DISABLE_STREAMING_SPLIT_SELECTED") != NULL &&
        cuda_stream_selected_pending_matches(model_map,
                                             layer,
                                             n_total_expert,
                                             n_selected,
                                             gate_expert_bytes,
                                             down_expert_bytes)) {
        const uint32_t compact_mask =
            g_stream_selected_pending.resident_mask |
            g_stream_selected_pending.missing_mask;
        if (!cuda_stream_selected_finish_pending_missing(compact_mask)) {
            return 0;
        }
    }
    if (!g_ssd_streaming_mode ||
        !g_stream_selected_cache.loaded ||
        !selected_exec ||
        !gate_w ||
        !up_w ||
        !down_w ||
        g_routed_moe_selected_override_n != n_selected ||
        g_stream_selected_cache.model_map != model_map ||
        g_stream_selected_cache.layer != layer ||
        g_stream_selected_cache.n_total_expert != n_total_expert ||
        g_stream_selected_cache.n_selected != n_selected ||
        g_stream_selected_cache.gate_expert_bytes != gate_expert_bytes ||
        g_stream_selected_cache.down_expert_bytes != down_expert_bytes) {
        return 0;
    }
    for (uint32_t i = 0; i < n_selected; i++) {
        if (g_stream_selected_cache.selected_ids[i] !=
            g_routed_moe_selected_override[i]) {
            return 0;
        }
    }
    *selected_exec = &g_stream_selected_cache.slot_tensor;
    *gate_w = g_stream_selected_cache.gate;
    *up_w = g_stream_selected_cache.up;
    *down_w = g_stream_selected_cache.down;
    g_routed_moe_selected_override_n = 0;
    return 1;
}

static const char *cuda_model_ptr(const void *model_map, uint64_t offset) {
    const char *owned = cuda_model_image_ptr(model_map, offset);
    if (owned) return owned;
    if (model_map == g_model_host_base && g_model_device_base) return g_model_device_base + offset;
    return (const char *)model_map + offset;
}

static const char *cuda_model_range_copy_uncached(
        const void *model_map,
        uint64_t offset,
        uint64_t bytes,
        const char *what) {
    void *dev = NULL;
    cudaError_t err = cudaMalloc(&dev, (size_t)bytes);
    if (err != cudaSuccess) {
        fprintf(stderr, DS4_GPU_LOG_PREFIX "model range alloc failed for %s (%.2f MiB): %s\n",
                what ? what : "weights", (double)bytes / 1048576.0, cudaGetErrorString(err));
        (void)cudaGetLastError();
        return NULL;
    }
    const char *src = (const char *)model_map + offset;
    err = cudaMemcpy(dev, src, (size_t)bytes, cudaMemcpyHostToDevice);
    if (err != cudaSuccess) {
        fprintf(stderr, DS4_GPU_LOG_PREFIX "model range copy failed for %s: %s\n",
                what ? what : "weights", cudaGetErrorString(err));
        (void)cudaFree(dev);
        (void)cudaGetLastError();
        return NULL;
    }
    g_model_ranges.push_back({model_map, offset, bytes, (char *)dev, NULL, NULL, 0, 0, 0});
    g_model_range_bytes += bytes;
    return (const char *)dev;
}

static const char *cuda_model_range_ptr(const void *model_map, uint64_t offset, uint64_t bytes, const char *what) {
    if (bytes == 0) return cuda_model_ptr(model_map, offset);
    if (cuda_model_image_owned(model_map)) return cuda_model_ptr(model_map, offset);

    const uint64_t end = offset + bytes;
    auto exact = g_model_range_by_offset.find(offset);
    if (exact != g_model_range_by_offset.end()) {
        const cuda_model_range &r = g_model_ranges[exact->second];
        if (r.host_base == model_map && end >= offset && bytes <= r.bytes) return r.device_ptr;
    }
    for (const cuda_model_range &r : g_model_ranges) {
        if (r.host_base == model_map && offset >= r.offset && end >= offset && end <= r.offset + r.bytes) {
            return r.device_ptr + (offset - r.offset);
        }
        if (r.host_base == model_map && r.host_registered && r.registered_base && r.registered_device_base) {
            const uintptr_t h0 = (uintptr_t)((const char *)model_map + offset);
            const uintptr_t h1 = h0 + bytes;
            const uintptr_t r0 = (uintptr_t)r.registered_base;
            const uintptr_t r1 = r0 + r.registered_bytes;
            if (h1 >= h0 && h0 >= r0 && h1 <= r1) return r.registered_device_base + (h0 - r0);
        }
    }

    const char *fd_ptr = cuda_model_range_ptr_from_fd(model_map, offset, bytes, what);
    if (fd_ptr) return fd_ptr;
    if (g_ssd_streaming_mode && model_map == g_model_host_base) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "streaming model range %s at offset %.2f GiB "
                "is not device-cached; refusing host-pointer fallback\n",
                what ? what : "weights",
                (double)offset / 1073741824.0);
        return NULL;
    }

    if (model_map != g_model_host_base) {
        return cuda_model_range_copy_uncached(model_map, offset, bytes, what);
    }

    cudaError_t err = cudaSuccess;
    if (g_model_range_mapping_supported && model_map == g_model_host_base) {
        const long page_sz_l = sysconf(_SC_PAGESIZE);
        const uint64_t page_sz = page_sz_l > 0 ? (uint64_t)page_sz_l : 4096u;
        const uintptr_t host_addr = (uintptr_t)((const char *)model_map + offset);
        const uintptr_t reg_addr = host_addr & ~(uintptr_t)(page_sz - 1u);
        const uint64_t reg_delta = (uint64_t)(host_addr - reg_addr);
        const uint64_t reg_bytes = (reg_delta + bytes + page_sz - 1u) & ~(page_sz - 1u);
        void *reg_dev = NULL;
        err = cudaHostRegister((void *)reg_addr,
                               (size_t)reg_bytes,
                               cudaHostRegisterMapped | cudaHostRegisterReadOnly);
        if (err == cudaSuccess) {
            err = cudaHostGetDevicePointer(&reg_dev, (void *)reg_addr, 0);
            if (err == cudaSuccess && reg_dev) {
                char *dev_ptr = (char *)reg_dev + reg_delta;
                g_model_ranges.push_back({model_map, offset, bytes, dev_ptr, (void *)reg_addr, (char *)reg_dev, reg_bytes, 1, 0});
                g_model_range_by_offset[offset] = g_model_ranges.size() - 1u;
                return dev_ptr;
            }
            fprintf(stderr, DS4_GPU_LOG_PREFIX "model range map pointer failed for %s: %s\n",
                    what ? what : "weights", cudaGetErrorString(err));
            (void)cudaHostUnregister((void *)reg_addr);
            (void)cudaGetLastError();
        } else {
            if (err == cudaErrorNotSupported || err == cudaErrorInvalidValue) g_model_range_mapping_supported = 0;
            (void)cudaGetLastError();
        }
    }

    void *dev = NULL;
    err = cudaMalloc(&dev, (size_t)bytes);
    if (err != cudaSuccess) {
        (void)cudaGetLastError();
        fprintf(stderr, DS4_GPU_LOG_PREFIX "model range alloc failed for %s (%.2f MiB): %s\n",
                what ? what : "weights", (double)bytes / 1048576.0, cudaGetErrorString(err));
        return NULL;
    }

    const char *src = (const char *)model_map + offset;
    const uint64_t chunk = 64ull * 1024ull * 1024ull;
    for (uint64_t done = 0; done < bytes; done += chunk) {
        uint64_t n = bytes - done < chunk ? bytes - done : chunk;
        err = cudaMemcpy((char *)dev + done, src + done, (size_t)n, cudaMemcpyHostToDevice);
        if (err != cudaSuccess) {
            fprintf(stderr, DS4_GPU_LOG_PREFIX "model range copy failed for %s at %.2f/%.2f MiB: %s\n",
                    what ? what : "weights",
                    (double)done / 1048576.0,
                    (double)bytes / 1048576.0,
                    cudaGetErrorString(err));
            (void)cudaFree(dev);
            (void)cudaGetLastError();
            return NULL;
        }
    }
    g_model_ranges.push_back({model_map, offset, bytes, (char *)dev, NULL, NULL, 0, 0, 0});
    g_model_range_by_offset[offset] = g_model_ranges.size() - 1u;
    g_model_range_bytes += bytes;
    return (const char *)dev;
}

static int cuda_model_range_is_cached(const void *model_map, uint64_t offset, uint64_t bytes) {
    if (bytes == 0) return 1;
    if (cuda_model_image_owned(model_map)) return 1;

    const uint64_t end = offset + bytes;
    if (end < offset) return 0;
    for (const cuda_model_range &r : g_model_ranges) {
        if (r.host_base == model_map &&
            offset >= r.offset &&
            end <= r.offset + r.bytes) {
            return 1;
        }
        if (r.host_base == model_map &&
            r.host_registered &&
            r.registered_base &&
            r.registered_device_base) {
            const uintptr_t h0 = (uintptr_t)((const char *)model_map + offset);
            const uintptr_t h1 = h0 + bytes;
            const uintptr_t r0 = (uintptr_t)r.registered_base;
            const uintptr_t r1 = r0 + r.registered_bytes;
            if (h1 >= h0 && h0 >= r0 && h1 <= r1) return 1;
        }
    }
    return 0;
}

static void cuda_q8_f16_cache_release_all(void) {
    for (const cuda_q8_f16_transpose_range &r : g_q8_f16_transpose_ranges) {
        (void)cudaFree(r.device_ptr);
    }
    for (const cuda_q8_f16_range &r : g_q8_f16_ranges) {
        (void)cudaFree(r.device_ptr);
    }
    g_q8_f16_transpose_ranges.clear();
    g_q8_f16_transpose_by_offset.clear();
    g_q8_f16_ranges.clear();
    g_q8_f16_by_offset.clear();
    g_q8_f16_bytes = 0;
}

static int cuda_env_present(const char *env) {
    if (env != NULL) return env[0] != '\0' && strcmp(env, "0") != 0;
    return 0;
}

static uint32_t cuda_rows_per_block_or_default(uint32_t v, uint32_t def) {
    return (v == 1u || v == 2u || v == 4u || v == 8u || v == 16u || v == 32u) ? v : def;
}

static uint32_t cuda_rows_per_block_env_or_default(const char *name, uint32_t def) {
    const char *env = name ? getenv(name) : NULL;
    if (!env || !env[0]) return def;
    char *end = NULL;
    errno = 0;
    unsigned long v = strtoul(env, &end, 10);
    if (end == env || errno != 0) return def;
    return cuda_rows_per_block_or_default((uint32_t)v, def);
}

struct ds4_rocm_runtime_config {
    int initialized;
    int disable_splitk_attn_out_low;
    int disable_shared_gate_up_fused_w32;
    int attention_output_cublas_all;
    int shared_down_cublas;
    int graph_dump;
    uint32_t moe_decode_rpb;
    uint32_t moe_decode_gate_rpb;
    uint32_t moe_decode_down_rpb;
    int oldhip_attention_decode;
};

static ds4_rocm_runtime_config g_rocm_cfg;

static const ds4_rocm_runtime_config *cuda_runtime_config(void) {
    if (!g_rocm_cfg.initialized) {
        g_rocm_cfg.disable_splitk_attn_out_low = !g_quality_mode;
        g_rocm_cfg.disable_shared_gate_up_fused_w32 = !g_quality_mode;
        g_rocm_cfg.attention_output_cublas_all = !g_quality_mode;
        g_rocm_cfg.shared_down_cublas = !g_quality_mode;
        g_rocm_cfg.graph_dump =
            cuda_env_present(getenv("DS4_ROCM_GRAPH_DUMP_PREFIX")) ||
            cuda_env_present(getenv("DS4_METAL_GRAPH_DUMP_PREFIX"));
        const char *moe_decode_rpb_env = getenv("DS4_ROCM_MOE_DECODE_RPB");
        const int moe_decode_rpb_env_present =
            moe_decode_rpb_env != NULL && moe_decode_rpb_env[0] != '\0';
        g_rocm_cfg.moe_decode_rpb =
            cuda_rows_per_block_env_or_default("DS4_ROCM_MOE_DECODE_RPB",
                                               g_quality_mode ? 8u :
                                               (g_ssd_streaming_mode ? 2u : 1u));
        g_rocm_cfg.moe_decode_gate_rpb =
            cuda_rows_per_block_env_or_default("DS4_ROCM_MOE_DECODE_GATE_RPB",
                                               (!g_quality_mode &&
                                                g_ssd_streaming_mode &&
                                                !moe_decode_rpb_env_present) ?
                                                    1u : g_rocm_cfg.moe_decode_rpb);
        g_rocm_cfg.moe_decode_down_rpb =
            cuda_rows_per_block_env_or_default("DS4_ROCM_MOE_DECODE_DOWN_RPB",
                                               (!g_quality_mode &&
                                                g_ssd_streaming_mode &&
                                                !moe_decode_rpb_env_present) ?
                                                    2u : g_rocm_cfg.moe_decode_rpb);
        g_rocm_cfg.oldhip_attention_decode = !g_quality_mode;
        g_rocm_cfg.initialized = 1;
    }
    return &g_rocm_cfg;
}

static uint64_t cuda_q8_f16_cache_limit_bytes(void) {
    if (!g_ssd_streaming_mode) return UINT64_MAX;
    const char *env = getenv("DS4_ROCM_STREAM_Q8_F16_CACHE_GB");
    if (env && env[0]) {
        char *end = NULL;
        errno = 0;
        unsigned long long gib = strtoull(env, &end, 10);
        if (end != env && *end == '\0' && errno == 0 &&
            gib <= UINT64_MAX / 1073741824ull) {
            return (uint64_t)gib * 1073741824ull;
        }
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "invalid DS4_ROCM_STREAM_Q8_F16_CACHE_GB=%s; "
                "using automatic q8 fp16 streaming cache limit\n",
                env);
    }

    size_t free_b = 0;
    size_t total_b = 0;
    if (cudaMemGetInfo(&free_b, &total_b) != cudaSuccess || total_b == 0) {
        (void)cudaGetLastError();
        return 8ull * 1073741824ull;
    }
    (void)free_b;
    uint64_t limit = (uint64_t)total_b / 8ull;
    const uint64_t min_limit = 2ull * 1073741824ull;
    const uint64_t max_limit = 16ull * 1073741824ull;
    if (limit < min_limit) limit = min_limit;
    if (limit > max_limit) limit = max_limit;
    return limit;
}

static uint64_t cuda_q8_f16_cache_reserve_bytes(uint64_t total_bytes) {
    if (g_ssd_streaming_mode) {
        return cuda_stream_resident_free_reserve_bytes();
    }
    if (total_bytes >= 112ull * 1024ull * 1024ull * 1024ull) {
        return 512ull * 1048576ull;
    }

    /* The expanded Q8->F16 cache is only an acceleration path.  Keep enough
     * device memory free for cuBLAS workspaces, transient graph buffers, and
     * driver bookkeeping instead of letting optional cached weights consume the
     * last few GiB on 96 GiB cards. */
    const uint64_t min_reserve = 4096ull * 1048576ull;
    const uint64_t pct_reserve = total_bytes / 20u; /* 5% */
    return pct_reserve > min_reserve ? pct_reserve : min_reserve;
}

static void cuda_q8_f16_cache_budget_notice(
        const char *reason,
        uint64_t request_bytes,
        uint64_t free_bytes,
        uint64_t total_bytes,
        uint64_t reserve_bytes,
        uint64_t limit_bytes) {
    if (g_q8_f16_budget_notice_printed) return;
    g_q8_f16_budget_notice_printed = 1;
    if (limit_bytes != UINT64_MAX && free_bytes == 0 && total_bytes == 0 && reserve_bytes == 0) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "q8 fp16 cache %s; using q8 kernels "
                "(request=%.2f MiB cached=%.2f GiB limit=%.2f GiB)\n",
                reason,
                (double)request_bytes / 1048576.0,
                (double)g_q8_f16_bytes / 1073741824.0,
                (double)limit_bytes / 1073741824.0);
    } else if (limit_bytes == UINT64_MAX) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "q8 fp16 cache %s; using q8 kernels "
                "(request=%.2f MiB cached=%.2f GiB free=%.2f GiB reserve=%.2f GiB total=%.2f GiB)\n",
                reason,
                (double)request_bytes / 1048576.0,
                (double)g_q8_f16_bytes / 1073741824.0,
                (double)free_bytes / 1073741824.0,
                (double)reserve_bytes / 1073741824.0,
                (double)total_bytes / 1073741824.0);
    } else {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "q8 fp16 cache %s; using q8 kernels "
                "(request=%.2f MiB cached=%.2f GiB limit=%.2f GiB free=%.2f GiB reserve=%.2f GiB total=%.2f GiB)\n",
                reason,
                (double)request_bytes / 1048576.0,
                (double)g_q8_f16_bytes / 1073741824.0,
                (double)limit_bytes / 1073741824.0,
                (double)free_bytes / 1073741824.0,
                (double)reserve_bytes / 1073741824.0,
                (double)total_bytes / 1073741824.0);
    }
}

static int cuda_q8_f16_cache_has_budget(uint64_t request_bytes, const char *label) {
    (void)label;
    const uint64_t limit = cuda_q8_f16_cache_limit_bytes();
    if (limit == 0) return 0;
    if (g_q8_f16_bytes > limit || request_bytes > limit - g_q8_f16_bytes) {
        cuda_q8_f16_cache_budget_notice("limit reached", request_bytes, 0, 0, 0, limit);
        return 0;
    }

    size_t free_b = 0;
    size_t total_b = 0;
    cudaError_t err = cudaMemGetInfo(&free_b, &total_b);
    if (err != cudaSuccess) {
        fprintf(stderr, DS4_GPU_LOG_PREFIX "q8 fp16 cache memory query failed: %s; using q8 kernels\n",
                cudaGetErrorString(err));
        (void)cudaGetLastError();
        return 0;
    }

    const uint64_t free_bytes = (uint64_t)free_b;
    const uint64_t total_bytes = (uint64_t)total_b;
    const uint64_t reserve_bytes = cuda_q8_f16_cache_reserve_bytes(total_bytes);
    if (request_bytes > free_bytes ||
        free_bytes - request_bytes < reserve_bytes) {
        cuda_q8_f16_cache_budget_notice("budget exhausted", request_bytes,
                                        free_bytes, total_bytes,
                                        reserve_bytes, limit);
        return 0;
    }
    return 1;
}

static void cuda_q8_f16_cache_disable_after_failure(const char *what, uint64_t request_bytes) {
    if (!g_q8_f16_disabled_after_oom) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "q8 fp16 cache disabled after %s "
                "(request=%.2f MiB cached=%.2f GiB); using q8 kernels\n",
                what ? what : "allocation failure",
                (double)request_bytes / 1048576.0,
                (double)g_q8_f16_bytes / 1073741824.0);
    }
    g_q8_f16_disabled_after_oom = 1;
    if (!g_q8_f16_ranges.empty() || !g_q8_f16_transpose_ranges.empty()) {
        (void)cudaDeviceSynchronize();
        cuda_q8_f16_cache_release_all();
    }
    (void)cudaGetLastError();
}

static int cuda_q8_f16_cache_allowed(const char *label, uint64_t in_dim, uint64_t out_dim) {
    if (g_quality_mode) return 0;
    if (g_q8_f16_disabled_after_oom) return 0;
    if (g_q8_f16_disabled_for_multi_model) return 0;
    if (getenv("DS4_CUDA_NO_Q8_F16_CACHE") != NULL) return 0;
    if (!label) return 0;
    if (strstr(label, "attn_output_a") != NULL ||
        strstr(label, "attn_output_b") != NULL ||
        strstr(label, "attention_output_a") != NULL ||
        strstr(label, "attention_output_b") != NULL) {
        return 1;
    }
    if (strstr(label, "attn_q_b") != NULL) {
        return 1;
    }
    if (strstr(label, "ffn_gate_shexp") != NULL ||
        strstr(label, "ffn_up_shexp") != NULL ||
        strstr(label, "ffn_down_shexp") != NULL) {
        return 1;
    }
    return (in_dim == 4096u && out_dim == 2048u) ||
           (in_dim == 2048u && out_dim == 4096u) ||
           (in_dim == 4096u && out_dim == 1024u) ||
           (in_dim == 4096u && out_dim == 512u) ||
           (in_dim == 1024u && out_dim == 32768u);
}

static int cuda_q8_label_is_attention_output(const char *label) {
    return label &&
           (strstr(label, "attn_output_a") != NULL ||
            strstr(label, "attn_output_b") != NULL ||
            strstr(label, "attention_output_a") != NULL ||
            strstr(label, "attention_output_b") != NULL);
}

static int cuda_q8_f16_preload_allowed(const char *label, uint64_t in_dim, uint64_t out_dim) {
    if (cuda_q8_label_is_attention_output(label) &&
        !cuda_runtime_config()->attention_output_cublas_all) {
        return 0;
    }
    return cuda_q8_f16_cache_allowed(label, in_dim, out_dim);
}

static const __half *cuda_q8_f16_ptr(
        const void *model_map,
        uint64_t offset,
        uint64_t weight_bytes,
        uint64_t in_dim,
        uint64_t out_dim,
        const char *label) {
    auto exact = g_q8_f16_by_offset.find(offset);
    if (exact != g_q8_f16_by_offset.end()) {
        const cuda_q8_f16_range &r = g_q8_f16_ranges[exact->second];
        if (r.host_base == model_map && r.weight_bytes == weight_bytes &&
            r.in_dim == in_dim && r.out_dim == out_dim) {
            return r.device_ptr;
        }
    }
    for (const cuda_q8_f16_range &r : g_q8_f16_ranges) {
        if (r.host_base == model_map && r.offset == offset &&
            r.weight_bytes == weight_bytes &&
            r.in_dim == in_dim && r.out_dim == out_dim) {
            return r.device_ptr;
        }
    }
    if (!cuda_q8_f16_cache_allowed(label, in_dim, out_dim)) return NULL;

    uint64_t out_bytes = 0;
    if (in_dim == 0u || out_dim == 0u ||
        !cuda_u64_mul3_checked(in_dim, out_dim, sizeof(__half), &out_bytes)) return NULL;
    if (!cuda_q8_f16_cache_has_budget(out_bytes, label)) return NULL;

    const char *q8 = cuda_model_range_ptr(model_map, offset, weight_bytes, "q8_0");
    if (!q8) return NULL;

    __half *dev = NULL;
    cudaError_t err = cudaMalloc(&dev, (size_t)out_bytes);
    if (err != cudaSuccess) {
        fprintf(stderr, DS4_GPU_LOG_PREFIX "q8 fp16 cache alloc failed (%.2f MiB): %s\n",
                (double)out_bytes / 1048576.0, cudaGetErrorString(err));
        cuda_q8_f16_cache_disable_after_failure("allocation failure", out_bytes);
        return NULL;
    }
    const uint64_t blocks = (in_dim + 31) / 32;
    const uint64_t n = in_dim * out_dim;
    dequant_q8_0_to_f16_kernel<<<(n + 255) / 256, 256>>>(dev,
                                                          (const unsigned char *)q8,
                                                          in_dim,
                                                          out_dim,
                                                          blocks);
    if (!cuda_ok(cudaGetLastError(), "q8 fp16 dequant launch")) {
        (void)cudaFree(dev);
        cuda_q8_f16_cache_disable_after_failure("dequant launch failure", out_bytes);
        return NULL;
    }
    g_q8_f16_ranges.push_back({model_map, offset, weight_bytes, in_dim, out_dim, dev});
    g_q8_f16_by_offset[offset] = g_q8_f16_ranges.size() - 1u;
    g_q8_f16_bytes += out_bytes;
    return dev;
}

static const __half *cuda_q8_f16_transpose_ptr(
        const void *model_map,
        uint64_t offset,
        uint64_t weight_bytes,
        uint64_t in_dim,
        uint64_t out_dim,
        const char *label) {
    auto exact = g_q8_f16_transpose_by_offset.find(offset);
    if (exact != g_q8_f16_transpose_by_offset.end()) {
        const cuda_q8_f16_transpose_range &r = g_q8_f16_transpose_ranges[exact->second];
        if (r.host_base == model_map && r.weight_bytes == weight_bytes &&
            r.in_dim == in_dim && r.out_dim == out_dim) {
            return r.device_ptr;
        }
    }
    for (const cuda_q8_f16_transpose_range &r : g_q8_f16_transpose_ranges) {
        if (r.host_base == model_map && r.offset == offset &&
            r.weight_bytes == weight_bytes &&
            r.in_dim == in_dim && r.out_dim == out_dim) {
            return r.device_ptr;
        }
    }
    if (!cuda_q8_f16_cache_allowed(label, in_dim, out_dim)) return NULL;
    uint64_t out_bytes = 0;
    if (in_dim == 0u || out_dim == 0u ||
        !cuda_u64_mul3_checked(in_dim, out_dim, sizeof(__half), &out_bytes)) return NULL;
    if (!cuda_q8_f16_cache_has_budget(out_bytes, label)) return NULL;
    const char *q8 = cuda_model_range_ptr(model_map, offset, weight_bytes, "q8_0");
    if (!q8) return NULL;
    __half *dev = NULL;
    cudaError_t err = cudaMalloc(&dev, (size_t)out_bytes);
    if (err != cudaSuccess) {
        fprintf(stderr, DS4_GPU_LOG_PREFIX "q8 fp16 transpose cache alloc failed (%.2f MiB): %s\n",
                (double)out_bytes / 1048576.0, cudaGetErrorString(err));
        cuda_q8_f16_cache_disable_after_failure("transpose allocation failure", out_bytes);
        return NULL;
    }
    const uint64_t blocks = (in_dim + 31u) / 32u;
    const uint64_t n = in_dim * out_dim;
    dequant_q8_0_to_f16_transpose_kernel<<<(n + 255u) / 256u, 256>>>(dev,
                                                                     (const unsigned char *)q8,
                                                                     in_dim,
                                                                     out_dim,
                                                                     blocks);
    if (!cuda_ok(cudaGetLastError(), "q8 fp16 transpose dequant launch")) {
        (void)cudaFree(dev);
        cuda_q8_f16_cache_disable_after_failure("transpose launch failure", out_bytes);
        return NULL;
    }
    g_q8_f16_transpose_ranges.push_back({model_map, offset, weight_bytes, in_dim, out_dim, dev});
    g_q8_f16_transpose_by_offset[offset] = g_q8_f16_transpose_ranges.size() - 1u;
    g_q8_f16_bytes += out_bytes;
    return dev;
}

static uint32_t cuda_prefill_warmup_tokens(void) {
    uint32_t n_tok = 2048u;
    const char *chunk_env = getenv("DS4_METAL_PREFILL_CHUNK");
    if (chunk_env && chunk_env[0]) {
        char *end = NULL;
        unsigned long long v = strtoull(chunk_env, &end, 10);
        if (end != chunk_env && *end == '\0' && v > 0 && v <= 4096u) n_tok = (uint32_t)v;
    }
    return n_tok;
}

static void cuda_q8_f16_warmup_attention_output_a_gemm(const __half *out_a_f16,
                                                       uint64_t group_dim,
                                                       uint64_t rank,
                                                       uint32_t n_groups) {
    static int warmed = 0;
    if (warmed || !g_cublas_ready || !out_a_f16 || group_dim == 0 || rank == 0 || n_groups == 0) return;
    const ds4_rocm_runtime_config *cfg = cuda_runtime_config();
    if (!cfg->attention_output_cublas_all) return;
    warmed = 1;
    const uint32_t n_tok = cuda_prefill_warmup_tokens();
    const uint64_t heads_h_count = (uint64_t)n_groups * n_tok * group_dim;
    const uint64_t low_h_count = (uint64_t)n_groups * n_tok * rank;
    const uint64_t heads_h_bytes = heads_h_count * sizeof(__half);
    const uint64_t low_h_off = (heads_h_bytes + 255ull) & ~255ull;
    if (low_h_count > (UINT64_MAX - low_h_off) / sizeof(__half)) return;
    void *tmp = cuda_tmp_alloc(low_h_off + low_h_count * sizeof(__half), "attention output a warmup");
    if (!tmp) return;
    __half *heads_h = (__half *)tmp;
    __half *low_h = (__half *)((char *)tmp + low_h_off);
    if (cudaMemset(heads_h, 0, (size_t)heads_h_bytes) != cudaSuccess) return;
    const float alpha = 1.0f;
    const float beta = 0.0f;
    cublasStatus_t st = cublasGemmStridedBatchedEx(g_cublas,
                                                   CUBLAS_OP_T,
                                                   CUBLAS_OP_N,
                                                   (int)rank,
                                                   (int)n_tok,
                                                   (int)group_dim,
                                                   &alpha,
                                                   out_a_f16,
                                                   CUDA_R_16F,
                                                   (int)group_dim,
                                                   (long long)rank * (long long)group_dim,
                                                   heads_h,
                                                   CUDA_R_16F,
                                                   (int)group_dim,
                                                   (long long)n_tok * (long long)group_dim,
                                                   &beta,
                                                   low_h,
                                                   CUDA_R_16F,
                                                   (int)(n_groups * rank),
                                                   (long long)rank,
                                                   (int)n_groups,
                                                   CUBLAS_COMPUTE_32F,
                                                   CUBLAS_GEMM_DEFAULT);
    if (st == CUBLAS_STATUS_SUCCESS) (void)cudaDeviceSynchronize();
}

static void cuda_q8_f16_warmup_attention_output_b_gemm(const __half *out_b_f16_t,
                                                       uint64_t low_dim,
                                                       uint64_t out_dim) {
    static int warmed = 0;
    if (warmed || !g_cublas_ready || !out_b_f16_t || low_dim == 0 || out_dim == 0) return;
    if (!cuda_runtime_config()->attention_output_cublas_all) return;
    warmed = 1;
    const uint32_t n_tok = cuda_prefill_warmup_tokens();
    const uint64_t low_h_count = (uint64_t)n_tok * low_dim;
    const uint64_t out_count = (uint64_t)n_tok * out_dim;
    const uint64_t low_h_bytes = low_h_count * sizeof(__half);
    const uint64_t out_off = (low_h_bytes + 255ull) & ~255ull;
    if (out_count > (UINT64_MAX - out_off) / sizeof(float)) return;
    void *tmp = cuda_tmp_alloc(out_off + out_count * sizeof(float), "attention output b warmup");
    if (!tmp) return;
    __half *low_h = (__half *)tmp;
    float *out = (float *)((char *)tmp + out_off);
    if (cudaMemset(low_h, 0, (size_t)low_h_bytes) != cudaSuccess) return;
    const float alpha = 1.0f;
    const float beta = 0.0f;
    cublasStatus_t st = cublasGemmEx(g_cublas,
                                     CUBLAS_OP_N,
                                     CUBLAS_OP_N,
                                     (int)out_dim,
                                     (int)n_tok,
                                     (int)low_dim,
                                     &alpha,
                                     out_b_f16_t,
                                     CUDA_R_16F,
                                     (int)out_dim,
                                     low_h,
                                     CUDA_R_16F,
                                     (int)low_dim,
                                     &beta,
                                     out,
                                     CUDA_R_32F,
                                     (int)out_dim,
                                     CUBLAS_COMPUTE_32F,
                                     CUBLAS_GEMM_DEFAULT);
    if (st == CUBLAS_STATUS_SUCCESS) (void)cudaDeviceSynchronize();
}

static int cuda_ok(cudaError_t err, const char *what) {
    if (err == cudaSuccess) return 1;
    fprintf(stderr, DS4_GPU_LOG_PREFIX "%s failed: %s\n", what, cudaGetErrorString(err));
    return 0;
}

static double cuda_wall_sec(void) {
    struct timespec ts;
    clock_gettime(CLOCK_MONOTONIC, &ts);
    return (double)ts.tv_sec + (double)ts.tv_nsec * 1.0e-9;
}

static int cuda_model_load_progress_enabled(void) {
    return 1;
}

static void cuda_model_load_progress_reset(void) {
    g_model_load_progress_next = 0;
    g_model_load_progress_last = 0.0;
    g_model_load_progress_started = 0;
    g_model_load_progress_tty = 0;
}

static void cuda_model_load_progress_note(uint64_t cached_bytes) {
    if (!cuda_model_load_progress_enabled()) return;

    const double now = cuda_wall_sec();
    if (!g_model_load_progress_started) {
        g_model_load_progress_started = 1;
        g_model_load_progress_tty = isatty(STDERR_FILENO) != 0;
        g_model_load_progress_next = (g_model_load_progress_tty ? 2ull : 16ull) *
                                     1024ull * 1024ull * 1024ull;
        g_model_load_progress_last = now;
        if (g_model_load_progress_tty) {
            fprintf(stderr, DS4_GPU_LOG_PREFIX "loading model tensors into device cache: 0.00 GiB");
        } else {
            fprintf(stderr, DS4_GPU_LOG_PREFIX "loading model tensors into device cache\n");
        }
    }

    if (cached_bytes < g_model_load_progress_next &&
        now - g_model_load_progress_last < (g_model_load_progress_tty ? 2.0 : 10.0)) {
        return;
    }

    if (g_model_load_progress_tty) {
        fprintf(stderr, "\r" DS4_GPU_LOG_PREFIX "loading model tensors into device cache: %.2f GiB",
                (double)cached_bytes / 1073741824.0);
    } else {
        fprintf(stderr, DS4_GPU_LOG_PREFIX "loading model tensors %.2f GiB cached\n",
                (double)cached_bytes / 1073741824.0);
    }
    fflush(stderr);
    g_model_load_progress_last = now;
    const uint64_t step = (g_model_load_progress_tty ? 2ull : 16ull) *
                          1024ull * 1024ull * 1024ull;
    while (g_model_load_progress_next <= cached_bytes) {
        g_model_load_progress_next += step;
    }
}

static uint64_t cuda_model_copy_chunk_bytes(void) {
    return 64ull * 1048576ull;
}

static void cuda_model_discard_source_pages(const void *model_map, uint64_t model_size, uint64_t offset, uint64_t bytes) {
#if defined(POSIX_MADV_DONTNEED)
    if (!model_map || bytes == 0 || offset > model_size) return;
    if (bytes > model_size - offset) bytes = model_size - offset;
    const long page_sz_l = sysconf(_SC_PAGESIZE);
    const uint64_t page_sz = page_sz_l > 0 ? (uint64_t)page_sz_l : 4096u;
    const uintptr_t h0 = (uintptr_t)((const char *)model_map + offset);
    const uintptr_t h1 = h0 + bytes;
    const uintptr_t p0 = h0 & ~(uintptr_t)(page_sz - 1u);
    const uintptr_t p1 = (h1 + page_sz - 1u) & ~(uintptr_t)(page_sz - 1u);
    if (p1 > p0) (void)posix_madvise((void *)p0, (size_t)(p1 - p0), POSIX_MADV_DONTNEED);
#else
    (void)model_map;
    (void)model_size;
    (void)offset;
    (void)bytes;
#endif
}

static void cuda_model_drop_file_pages(uint64_t offset, uint64_t bytes) {
#if defined(POSIX_FADV_DONTNEED)
    if (g_model_fd < 0 || bytes == 0) return;
    (void)posix_fadvise(g_model_fd, (off_t)offset, (off_t)bytes, POSIX_FADV_DONTNEED);
#else
    (void)offset;
    (void)bytes;
#endif
}

static uint64_t cuda_round_down(uint64_t v, uint64_t align) {
    if (align <= 1) return v;
    return (v / align) * align;
}

static uint64_t cuda_round_up(uint64_t v, uint64_t align) {
    if (align <= 1) return v;
    const uint64_t rem = v % align;
    return rem == 0 ? v : v + (align - rem);
}

static void *cuda_align_ptr(void *ptr, uint64_t align) {
    if (align <= 1) return ptr;
    uintptr_t p = (uintptr_t)ptr;
    uintptr_t a = (uintptr_t)align;
    return (void *)(((p + a - 1u) / a) * a);
}

static int cuda_model_stage_pool_alloc(uint64_t bytes) {
    if (g_model_stage_bytes >= bytes) return 1;
    for (size_t i = 0; i < 4; i++) {
        if (g_model_stage_event[i]) {
            (void)cudaEventDestroy(g_model_stage_event[i]);
            g_model_stage_event[i] = NULL;
        }
        if (g_model_stage_raw[i]) {
            (void)cudaFreeHost(g_model_stage_raw[i]);
            g_model_stage_raw[i] = NULL;
            g_model_stage[i] = NULL;
        }
    }
    g_model_stage_bytes = 0;
    if (!g_model_upload_stream) {
        cudaError_t err = cudaStreamCreateWithFlags(&g_model_upload_stream, cudaStreamNonBlocking);
        if (err != cudaSuccess) {
            fprintf(stderr, DS4_GPU_LOG_PREFIX "model upload stream creation failed: %s\n", cudaGetErrorString(err));
            (void)cudaGetLastError();
            return 0;
        }
    }
    uint64_t alloc_bytes = bytes;
    if (g_model_direct_align > 1u) {
        const uint64_t pad = g_model_direct_align - 1u;
        if (alloc_bytes > UINT64_MAX - pad) return 0;
        alloc_bytes += pad;
    }
    if (alloc_bytes > (uint64_t)SIZE_MAX) return 0;
    for (size_t i = 0; i < 4; i++) {
        cudaError_t err = cudaMallocHost(&g_model_stage_raw[i], (size_t)alloc_bytes);
        if (err != cudaSuccess) {
            fprintf(stderr, DS4_GPU_LOG_PREFIX "pinned model staging allocation failed: %s\n", cudaGetErrorString(err));
            (void)cudaGetLastError();
            return 0;
        }
        g_model_stage[i] = cuda_align_ptr(g_model_stage_raw[i], g_model_direct_align);
        err = cudaEventCreateWithFlags(&g_model_stage_event[i], cudaEventDisableTiming);
        if (err != cudaSuccess) {
            fprintf(stderr, DS4_GPU_LOG_PREFIX "model staging event creation failed: %s\n", cudaGetErrorString(err));
            (void)cudaGetLastError();
            return 0;
        }
    }
    g_model_stage_bytes = bytes;
    return 1;
}

static int cuda_pread_full(int fd, void *buf, uint64_t bytes, uint64_t offset) {
    uint64_t done = 0;
    while (done < bytes) {
        const size_t n_req = (bytes - done > (uint64_t)SSIZE_MAX) ? (size_t)SSIZE_MAX : (size_t)(bytes - done);
        ssize_t n = pread(fd, (char *)buf + done, n_req, (off_t)(offset + done));
        if (n < 0) {
            if (errno == EINTR) continue;
            return 0;
        }
        if (n == 0) return 0;
        done += (uint64_t)n;
    }
    return 1;
}

static int cuda_model_stage_read(void *stage, uint64_t stage_bytes,
                                 uint64_t offset, uint64_t bytes,
                                 const char **payload) {
    *payload = (const char *)stage;
#if defined(__linux__) && defined(O_DIRECT)
    if (g_model_direct_fd >= 0 && g_model_direct_align > 1 && g_model_file_size != 0) {
        const uint64_t aligned_off = cuda_round_down(offset, g_model_direct_align);
        const uint64_t delta = offset - aligned_off;
        uint64_t read_size = cuda_round_up(delta + bytes, g_model_direct_align);
        if (aligned_off <= g_model_file_size &&
            read_size <= stage_bytes &&
            read_size <= g_model_file_size - aligned_off) {
            const int saved_errno = errno;
            errno = 0;
            if (cuda_pread_full(g_model_direct_fd, stage, read_size, aligned_off)) {
                *payload = (const char *)stage + delta;
                errno = saved_errno;
                return 1;
            }
            const int direct_errno = errno;
            if (direct_errno == EINVAL || direct_errno == EFAULT || direct_errno == ENOTSUP || direct_errno == EOPNOTSUPP) {
                (void)close(g_model_direct_fd);
                g_model_direct_fd = -1;
                g_model_direct_align = 1;
            }
            errno = direct_errno;
        }
    }
#else
    (void)stage_bytes;
#endif
    return cuda_pread_full(g_model_fd, stage, bytes, offset);
}

static uint64_t cuda_model_cache_limit_bytes(void) {
    if (!g_ssd_streaming_mode) return UINT64_MAX;
    const char *env = getenv("DS4_ROCM_STREAM_MODEL_CACHE_GB");
    if (env && env[0]) {
        char *end = NULL;
        errno = 0;
        unsigned long long gib = strtoull(env, &end, 10);
        if (end != env && *end == '\0' && errno == 0 && gib != 0 &&
            gib <= UINT64_MAX / 1073741824ull) {
            return (uint64_t)gib * 1073741824ull;
        }
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "invalid DS4_ROCM_STREAM_MODEL_CACHE_GB=%s; "
                "using automatic streaming model span cache limit\n",
                env);
    }

    size_t free_b = 0;
    size_t total_b = 0;
    const uint64_t fallback = 32ull * 1073741824ull;
    if (cudaMemGetInfo(&free_b, &total_b) != cudaSuccess || total_b == 0) {
        (void)cudaGetLastError();
        return fallback;
    }
    (void)free_b;
    uint64_t limit = (uint64_t)total_b / 3ull;
    const uint64_t min_limit = 8ull * 1073741824ull;
    const uint64_t max_limit = 48ull * 1073741824ull;
    if (limit < min_limit) limit = min_limit;
    if (limit > max_limit) limit = max_limit;
    return limit;
}

static int cuda_stream_model_cache_prepare_memory(
        uint64_t request_bytes,
        const char *what) {
    if (!g_ssd_streaming_mode || request_bytes == 0 || g_q8_f16_bytes == 0) {
        return 1;
    }

    size_t free_b = 0;
    size_t total_b = 0;
    if (cudaMemGetInfo(&free_b, &total_b) != cudaSuccess) {
        (void)cudaGetLastError();
        return 1;
    }
    (void)total_b;
    const uint64_t reserve = cuda_stream_resident_free_reserve_bytes();
    const uint64_t free_bytes = (uint64_t)free_b;
    if (free_bytes >= reserve && request_bytes <= free_bytes - reserve) {
        return 1;
    }

    if (!cuda_ok(cudaDeviceSynchronize(),
                 "streaming model cache q8 fp16 release sync")) {
        return 0;
    }
    fprintf(stderr,
            DS4_GPU_LOG_PREFIX "releasing %.2f GiB q8 fp16 cache for %s "
            "(request=%.2f GiB free=%.2f GiB reserve=%.2f GiB)\n",
            (double)g_q8_f16_bytes / 1073741824.0,
            what ? what : "streaming model spans",
            (double)request_bytes / 1073741824.0,
            (double)free_bytes / 1073741824.0,
            (double)reserve / 1073741824.0);
    cuda_q8_f16_cache_release_all();
    g_q8_f16_budget_notice_printed = 0;
    return 1;
}

static uint64_t cuda_model_arena_chunk_bytes(uint64_t need) {
    uint64_t bytes = 1792ull * 1048576ull;
    if (bytes < need) {
        const uint64_t align = 256ull * 1048576ull;
        bytes = (need + align - 1u) & ~(align - 1u);
    }
    return bytes;
}

static char *cuda_model_arena_alloc(uint64_t bytes, const char *what) {
    if (bytes == 0) return NULL;
    if (g_model_cache_full) return NULL;
    const uint64_t align = 256u;
    const uint64_t aligned = (bytes + align - 1u) & ~(align - 1u);

    for (cuda_model_arena &a : g_model_arenas) {
        const uint64_t used = (a.used + align - 1u) & ~(align - 1u);
        if (used <= a.bytes && aligned <= a.bytes - used) {
            char *ptr = a.device_ptr + used;
            a.used = used + aligned;
            return ptr;
        }
    }

    const uint64_t limit = cuda_model_cache_limit_bytes();
    if (g_model_range_bytes > limit || aligned > limit - g_model_range_bytes) return NULL;

    const uint64_t chunk = cuda_model_arena_chunk_bytes(aligned);
    void *dev = NULL;
    cudaError_t err = cudaMalloc(&dev, (size_t)chunk);
    if (err != cudaSuccess) {
        fprintf(stderr, DS4_GPU_LOG_PREFIX "model arena alloc failed for %s (%.2f MiB chunk): %s\n",
                what ? what : "weights",
                (double)chunk / 1048576.0,
                cudaGetErrorString(err));
        (void)cudaGetLastError();
        g_model_cache_full = 1;
        return NULL;
    }
    g_model_arenas.push_back({(char *)dev, chunk, aligned});
    return (char *)dev;
}

static const char *cuda_model_range_ptr_from_fd(
        const void *model_map,
        uint64_t offset,
        uint64_t bytes,
        const char *what) {
    if (g_model_fd < 0 || bytes == 0) return NULL;
    if (g_model_fd_host_base != NULL && model_map != g_model_fd_host_base) return NULL;
    const uint64_t limit = cuda_model_cache_limit_bytes();
    if (g_model_range_bytes > limit || bytes > limit - g_model_range_bytes) {
        if (g_ssd_streaming_mode) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming model cache limit prevents "
                    "loading %s range %.2f MiB; refusing host-pointer fallback\n",
                    what ? what : "weights",
                    (double)bytes / 1048576.0);
            return NULL;
        }
        return cuda_model_ptr(model_map, offset);
    }

    const uint64_t chunk = cuda_model_copy_chunk_bytes();
    const uint64_t stage_bytes =
        chunk + (g_model_direct_align > 1 ? g_model_direct_align : 1);
    if (!cuda_model_stage_pool_alloc(stage_bytes)) return NULL;

    char *dev = cuda_model_arena_alloc(bytes, what);
    if (!dev) {
        if (g_ssd_streaming_mode) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming model cache allocation failed "
                    "for %s range %.2f MiB; refusing host-pointer fallback\n",
                    what ? what : "weights",
                    (double)bytes / 1048576.0);
            return NULL;
        }
        return cuda_model_ptr(model_map, offset);
    }
    cudaError_t err = cudaSuccess;

    uint64_t copied = 0;
    uint64_t chunk_idx = 0;
    while (copied < bytes) {
        const uint64_t n = (bytes - copied < chunk) ? (bytes - copied) : chunk;
        const uint64_t bi = chunk_idx % 4u;
        if (chunk_idx >= 4u) {
            err = cudaEventSynchronize(g_model_stage_event[bi]);
            if (err != cudaSuccess) {
                fprintf(stderr, DS4_GPU_LOG_PREFIX "model staging wait failed for %s: %s\n",
                        what ? what : "weights", cudaGetErrorString(err));
                (void)cudaGetLastError();
                return NULL;
            }
        }
        const char *payload = NULL;
        if (!cuda_model_stage_read(g_model_stage[bi], g_model_stage_bytes,
                                   offset + copied, n, &payload)) {
            fprintf(stderr, DS4_GPU_LOG_PREFIX "model range read failed for %s at %.2f MiB: %s\n",
                    what ? what : "weights",
                    (double)copied / 1048576.0,
                    strerror(errno));
            return NULL;
        }
        err = cudaMemcpyAsync(dev + copied, payload, (size_t)n,
                              cudaMemcpyHostToDevice, g_model_upload_stream);
        if (err != cudaSuccess) {
            fprintf(stderr, DS4_GPU_LOG_PREFIX "model range copy failed for %s at %.2f MiB: %s\n",
                    what ? what : "weights",
                    (double)copied / 1048576.0,
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            return NULL;
        }
        err = cudaEventRecord(g_model_stage_event[bi], g_model_upload_stream);
        if (err != cudaSuccess) {
            fprintf(stderr, DS4_GPU_LOG_PREFIX "model staging record failed for %s: %s\n",
                    what ? what : "weights", cudaGetErrorString(err));
            (void)cudaGetLastError();
            return NULL;
        }
        cuda_model_drop_file_pages(offset + copied, n);
        cuda_model_discard_source_pages(model_map, g_model_registered_size, offset + copied, n);
        copied += n;
        cuda_model_load_progress_note(g_model_range_bytes + copied);
        chunk_idx++;
    }
    err = cudaStreamSynchronize(g_model_upload_stream);
    if (err != cudaSuccess) {
        fprintf(stderr, DS4_GPU_LOG_PREFIX "model range upload sync failed for %s: %s\n",
                what ? what : "weights", cudaGetErrorString(err));
        (void)cudaGetLastError();
        return NULL;
    }

    g_model_ranges.push_back({model_map, offset, bytes, dev, NULL, NULL, 0, 0, 1});
    g_model_range_by_offset[offset] = g_model_ranges.size() - 1u;
    g_model_range_bytes += bytes;
    cuda_model_load_progress_note(g_model_range_bytes);
    return (const char *)dev;
}

static int cuda_model_copy_chunked(const void *model_map, uint64_t model_size, uint64_t map_offset, uint64_t map_size) {
    if (!model_map || model_size == 0 || map_offset > model_size || map_size > model_size - map_offset) return 0;
    if (cuda_model_image_owned(model_map)) {
        g_model_host_base = model_map;
        g_model_device_base = cuda_model_image_ptr(model_map, 0);
        g_model_registered_size = model_size;
        g_model_device_owned = 1;
        return 1;
    }

    void *dev = NULL;
    const double t0 = cuda_wall_sec();
    cudaError_t err = cudaMalloc(&dev, (size_t)model_size);
    if (err != cudaSuccess) {
        fprintf(stderr, DS4_GPU_LOG_PREFIX "model allocation skipped: %s\n", cudaGetErrorString(err));
        (void)cudaGetLastError();
        return 0;
    }

    fprintf(stderr, DS4_GPU_LOG_PREFIX "chunk-copying %.2f GiB model image\n",
            (double)model_size / 1073741824.0);

    const uint64_t chunk = cuda_model_copy_chunk_bytes();
    const uint64_t stage_bytes = chunk + (g_model_direct_align > 1 ? g_model_direct_align : 1);
    if (!cuda_model_stage_pool_alloc(stage_bytes)) {
        (void)cudaFree(dev);
        return 0;
    }

    uint64_t copied = 0;
    uint64_t chunk_idx = 0;
    while (copied < model_size) {
        const uint64_t n = (model_size - copied < chunk) ? (model_size - copied) : chunk;
        const uint64_t bi = chunk_idx % 4u;
        if (chunk_idx >= 4u) {
            err = cudaEventSynchronize(g_model_stage_event[bi]);
            if (err != cudaSuccess) {
                fprintf(stderr, DS4_GPU_LOG_PREFIX "model staging wait failed: %s\n", cudaGetErrorString(err));
                (void)cudaFree(dev);
                (void)cudaGetLastError();
                return 0;
            }
        }
        const char *payload = NULL;
        if (!cuda_model_stage_read(g_model_stage[bi], g_model_stage_bytes,
                                   copied, n, &payload)) {
            fprintf(stderr, DS4_GPU_LOG_PREFIX "model staged read failed at %.2f GiB: %s\n",
                    (double)copied / 1073741824.0, strerror(errno));
            (void)cudaFree(dev);
            return 0;
        }
        err = cudaMemcpyAsync((char *)dev + copied, payload, (size_t)n,
                              cudaMemcpyHostToDevice, g_model_upload_stream);
        if (err != cudaSuccess) {
            fprintf(stderr, DS4_GPU_LOG_PREFIX "model chunk copy failed at %.2f GiB: %s\n",
                    (double)copied / 1073741824.0, cudaGetErrorString(err));
            (void)cudaFree(dev);
            (void)cudaGetLastError();
            return 0;
        }
        err = cudaEventRecord(g_model_stage_event[bi], g_model_upload_stream);
        if (err != cudaSuccess) {
            fprintf(stderr, DS4_GPU_LOG_PREFIX "model staging record failed: %s\n", cudaGetErrorString(err));
            (void)cudaFree(dev);
            (void)cudaGetLastError();
            return 0;
        }
        cuda_model_drop_file_pages(copied, n);
        cuda_model_discard_source_pages(model_map, model_size, copied, n);
        copied += n;
        chunk_idx++;
        cuda_model_load_progress_note(copied > map_offset ? copied - map_offset : 0);
    }
    err = cudaStreamSynchronize(g_model_upload_stream);
    if (err != cudaSuccess) {
        fprintf(stderr, DS4_GPU_LOG_PREFIX "model upload sync failed: %s\n", cudaGetErrorString(err));
        (void)cudaFree(dev);
        (void)cudaGetLastError();
        return 0;
    }
    g_model_images.push_back({model_map, model_size, (char *)dev});
    g_model_host_base = model_map;
    g_model_device_base = (const char *)dev;
    g_model_registered_size = model_size;
    g_model_device_owned = 1;
    const double t1 = cuda_wall_sec();
    fprintf(stderr,
            DS4_GPU_LOG_PREFIX "model chunk copy complete in %.3fs (%.2f GiB tensors)\n",
            t1 - t0,
            (double)map_size / 1073741824.0);
    return 1;
}

static void cuda_model_range_release_ranges_only(void) {
    for (const cuda_model_range &r : g_model_ranges) {
        if (r.host_registered && r.registered_base) {
            (void)cudaHostUnregister(r.registered_base);
        } else if (r.device_ptr && !r.arena_allocated) {
            (void)cudaFree(r.device_ptr);
        }
    }
    for (const cuda_model_arena &a : g_model_arenas) {
        if (a.device_ptr) (void)cudaFree(a.device_ptr);
    }
    g_model_arenas.clear();
    g_model_ranges.clear();
    g_model_range_by_offset.clear();
    g_model_range_bytes = 0;
    g_model_cache_full = 0;
}

static void cuda_model_range_release_all(void) {
    cuda_model_range_release_ranges_only();
    g_stream_selected_cache.loaded = 0;
    cuda_stream_resident_cache_release();
    cuda_model_load_progress_reset();
}

static int cublas_ok(cublasStatus_t st, const char *what) {
    if (st == CUBLAS_STATUS_SUCCESS) return 1;
    fprintf(stderr, "ds4: " DS4_GPU_BLAS_NAME " %s failed: status %d\n", what, (int)st);
    return 0;
}


/* =========================================================================
 * Strix Halo / ROCm runtime diagnostics.
 * =========================================================================
 *
 * These helpers print a runtime profile and surface common Strix Halo memory
 * and driver misconfigurations. They are ROCm-only and live here so the rest
 * of the backend can stay CUDA-portable through the hip-mapped compatibility
 * layer.
 */

#include <stdarg.h>
#include <unistd.h>

/* Count of WARNING-level diagnostics emitted this process. Used by the smoke
 * test's strict mode (ROCM_SMOKE_STRICT) to turn configuration warnings into
 * test failures. */
static int g_ds4_rocm_warnings = 0;

int ds4_rocm_warning_count(void) { return g_ds4_rocm_warnings; }

size_t ds4_gpu_hip_free_bytes(void) {
#ifdef __HIP_PLATFORM_AMD__
    size_t free_b = 0, total_b = 0;
    if (hipMemGetInfo(&free_b, &total_b) == hipSuccess) return free_b;
#endif
    return 0;
}

/* Per-model residency estimate, populated by ds4_gpu_set_model_map. The type
 * itself is defined in ds4_gpu.h (ds4_rocm_model_load_estimate) so the public
 * API and the runtime agree on the layout. */
static ds4_rocm_model_load_estimate g_ds4_rocm_last_model_estimate = {0};

const ds4_rocm_model_load_estimate *
ds4_rocm_last_model_load_estimate(void) {
    return &g_ds4_rocm_last_model_estimate;
}

/* Read total system RAM in bytes from /proc/meminfo. Returns 0 on failure. */
static uint64_t ds4_rocm_total_system_memory(void) {
    FILE *f = fopen("/proc/meminfo", "r");
    if (!f) return 0;
    uint64_t kb = 0;
    char line[256];
    while (fgets(line, sizeof(line), f)) {
        if (sscanf(line, "MemTotal: %llu kB", (unsigned long long *)&kb) == 1) break;
    }
    fclose(f);
    return kb * 1024ull;
}

/* Read the TTM/GTT pages_limit from the relevant module parameter and return
 * the mapped byte limit, or 0 if it cannot be determined. Module names differ
 * across ROCm/KFD versions: ttm, amdttm, and amd_ttm. pages_limit is expressed
 * in 4 KiB pages. */
static uint64_t ds4_rocm_ttm_pages_limit_bytes(void) {
    static const char *const paths[] = {
        "/sys/module/ttm/parameters/pages_limit",
        "/sys/module/amdttm/parameters/pages_limit",
        "/sys/module/amd_ttm/parameters/pages_limit",
        NULL
    };
    for (int i = 0; paths[i]; i++) {
        FILE *f = fopen(paths[i], "r");
        if (!f) continue;
        long long pages = -1;
        if (fscanf(f, "%lld", &pages) == 1 && pages > 0) {
            fclose(f);
            return (uint64_t)pages * 4096ull;
        }
        fclose(f);
    }
    return 0;
}

/* Resolve the TTM/GTT byte limit. DS4_ROCM_TTM_PAGES overrides the kernel
 * value (4 KiB pages). Returns 0 if neither is available. */
static uint64_t ds4_rocm_effective_ttm_bytes(void) {
    const char *env = getenv("DS4_ROCM_TTM_PAGES");
    if (env && env[0] != '\0') {
        char *end = NULL;
        long long pages = strtoll(env, &end, 10);
        if (end != env && pages > 0) return (uint64_t)pages * 4096ull;
    }
    return ds4_rocm_ttm_pages_limit_bytes();
}

/* Attempt to raise the TTM/GTT limit via amd-ttm when the effective limit is
 * below the requested bytes and DS4_ROCM_TTM_AUTORAISE is set. Honors
 * DS4_ROCM_AUTO_RAISE_ONCE: when set (non-0), at most one amd-ttm call is
 * made per process (subsequent calls are no-ops, and the current limit is
 * returned as-is). Returns the new limit in bytes (re-read after the call),
 * or 0 if it could not be raised / was not attempted. */
static uint64_t ds4_rocm_try_autoraise_ttm(uint64_t want_bytes) {
    const char *autoenv = getenv("DS4_ROCM_TTM_AUTORAISE");
    if (autoenv == NULL || autoenv[0] == '\0' || strcmp(autoenv, "0") == 0)
        return 0;
    const char *onceenv = getenv("DS4_ROCM_AUTO_RAISE_ONCE");
    static int s_already_raised = 0;
    if (onceenv && onceenv[0] != '\0' && strcmp(onceenv, "0") != 0) {
        if (s_already_raised) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "DS4_ROCM_AUTO_RAISE_ONCE: skipping "
                    "another amd-ttm call (already raised this process).\n");
            return 0;
        }
    }
    const char *amdttm = "/usr/bin/amd-ttm";
    struct stat st;
    if (stat(amdttm, &st) != 0) return 0;
    const uint64_t pages = (want_bytes + 4095ull) / 4096ull;
    char cmd[512];
    snprintf(cmd, sizeof(cmd), "%s --set-pages %llu",
             amdttm, (unsigned long long)pages);
    fprintf(stderr,
            DS4_GPU_LOG_PREFIX "attempting to raise TTM/GTT limit via amd-ttm "
            "(%llu pages)...\n", (unsigned long long)pages);
    const int rc = system(cmd);
    if (rc != 0) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "amd-ttm --set-pages failed (rc=%d). Re-run as "
                "root or raise the limit manually.\n", rc);
        return 0;
    }
    s_already_raised = 1;
    return ds4_rocm_ttm_pages_limit_bytes();
}

/* Format a HIP version packed as (major*100 + minor)*100 + patch into a string
 * buffer. Returns the same buffer for convenience. */
static const char *ds4_rocm_format_version(int version, char *buf, size_t buflen) {
    if (!buf || buflen == 0) return "(invalid)";
    int major = version / 10000;
    int minor = (version % 10000) / 100;
    int patch = version % 100;
    snprintf(buf, buflen, "%d.%d.%d", major, minor, patch);
    return buf;
}

/* Warn when the machine exposes less than 75 percent of RAM through TTM/GTT, a
 * common BIOS VRAM carveout misconfiguration on Strix Halo. Also print the
 * exact `amd-ttm --set-pages` value that would reach 90% of RAM so the user can
 * copy-paste the fix. */
static void ds4_rocm_check_ttm_limit(uint64_t total_ram, uint64_t ttm_bytes) {
    if (total_ram == 0 || ttm_bytes == 0) return;
    const double frac = (double)ttm_bytes / (double)total_ram;
    if (frac < 0.75) {
        const uint64_t target = (uint64_t)(0.90 * (double)total_ram);
        const uint64_t target_pages = (target + 4095ull) / 4096ull;
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "WARNING: TTM/GTT mapping limit is %.1f%% of "
                "system RAM (%.1f/%.1f GiB). Increase the GTT limit via "
                "amd-ttm or the ttm pages_limit; a large BIOS VRAM carveout is "
                "likely starving the OS of memory.\n",
                100.0 * frac,
                (double)ttm_bytes / 1073741824.0,
                (double)total_ram / 1073741824.0);
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "SUGGESTED: sudo amd-ttm --set-pages %llu "
                "(raises GTT to ~90%% of %.1f GiB RAM)\n",
                (unsigned long long)target_pages,
                (double)total_ram / 1073741824.0);
        g_ds4_rocm_warnings++;
    }
}

/* Warn when a model is sized so close to the TTM/GTT limit that less than
 * roughly 8 GiB remains for runtime state. */
static void ds4_rocm_check_model_fits_ttm(uint64_t model_bytes, uint64_t ttm_bytes) {
    if (model_bytes == 0 || ttm_bytes == 0) return;
    const uint64_t reserved = 8ull * 1073741824ull;
    if (model_bytes > ttm_bytes || ttm_bytes - model_bytes < reserved) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "WARNING: model range (%.1f GiB) leaves only "
                "%.1f GiB of the %.1f GiB TTM/GTT mapping limit for runtime "
                "state. Lower the BIOS VRAM carveout or raise the GTT limit.\n",
                (double)model_bytes / 1073741824.0,
                (double)(ttm_bytes > model_bytes ? ttm_bytes - model_bytes : 0) / 1073741824.0,
                (double)ttm_bytes / 1073741824.0);
        g_ds4_rocm_warnings++;
    }
}

/* Optional machine-readable diagnostics. When DS4_ROCM_DIAG is set to a path,
 * the runtime profile is written there. The file is truncated (not appended) on
 * each run, and a timestamp marks the block. Set DS4_ROCM_DIAG_JSON=1 to emit
 * JSON instead of key=value lines. DS4_ROCM_DIAG_FIELDS controls verbosity:
 *   basic = device, arch, versions, TTM limit
 *   full  = + memory flags, HIP-visible memory, system RAM (default)
 *   all   = + gguf magic, cmdline gttsize, model-fit estimate
 */
static int ds4_rocm_diag_json(void) {
    const char *env = getenv("DS4_ROCM_DIAG_JSON");
    return env && env[0] != '\0' && strcmp(env, "0") != 0;
}

static int ds4_rocm_diag_level(void) {
    const char *env = getenv("DS4_ROCM_DIAG_FIELDS");
    if (env == NULL || env[0] == '\0' || strcmp(env, "full") == 0) return 1; /* full */
    if (strcmp(env, "basic") == 0) return 0;
    if (strcmp(env, "all") == 0) return 2;
    return 1; /* unknown -> default to full */
}

static int ds4_rocm_diag_include(int min_level) {
    return ds4_rocm_diag_level() >= min_level;
}

static FILE *ds4_rocm_diag_open(void) {
    const char *path = getenv("DS4_ROCM_DIAG");
    if (path == NULL || path[0] == '\0') return NULL;
    FILE *f = fopen(path, "w");
    if (!f) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "could not open DS4_ROCM_DIAG=%s: %s\n",
                path, strerror(errno));
        return NULL;
    }
    if (ds4_rocm_diag_json()) {
        fprintf(f, "{\n");
    } else {
        time_t now = time(NULL);
        char buf[64];
        struct tm tm_now;
        if (localtime_r(&now, &tm_now) &&
            strftime(buf, sizeof(buf), "%Y-%m-%dT%H:%M:%S%z", &tm_now)) {
            fprintf(f, "# ds4 rocm diag %s\n", buf);
        }
    }
    return f;
}

/* Emit a diag key/value. min_level: 0=basic, 1=full (default), 2=all. */
static void ds4_rocm_diag_kv_lvl(FILE *f, int min_level, const char *key, const char *fmt, ...) {
    if (!f) return;
    if (!ds4_rocm_diag_include(min_level)) return;
    va_list ap;
    va_start(ap, fmt);
    if (ds4_rocm_diag_json()) {
        static int s_first = 1;
        fprintf(f, "%s  \"%s\": \"", s_first ? "\n" : ",\n", key);
        vfprintf(f, fmt, ap);
        fprintf(f, "\"");
        s_first = 0;
    } else {
        fprintf(f, "%s=", key);
        vfprintf(f, fmt, ap);
        fprintf(f, "\n");
    }
    va_end(ap);
}

static void ds4_rocm_diag_kv(FILE *f, const char *key, const char *fmt, ...) {
    if (!f) return;
    va_list ap;
    va_start(ap, fmt);
    if (ds4_rocm_diag_json()) {
        static int s_first = 1;
        fprintf(f, "%s  \"%s\": \"", s_first ? "\n" : ",\n", key);
        vfprintf(f, fmt, ap);
        fprintf(f, "\"");
        s_first = 0;
    } else {
        fprintf(f, "%s=", key);
        vfprintf(f, fmt, ap);
        fprintf(f, "\n");
    }
    va_end(ap);
}

/* Close the diag file, terminating JSON if needed. */
static void ds4_rocm_diag_close(FILE *f) {
    if (!f) return;
    if (ds4_rocm_diag_json()) fprintf(f, "\n}\n");
    fclose(f);
}

/* Probe amd-smi / rocm-smi availability. On headless Strix Halo boxes these may
 * be missing or return limited data; note it so users don't blame the engine. */
static void ds4_rocm_probe_smi(void) {
    struct stat st;
    int have_amd_smi = (stat("/usr/bin/amd-smi", &st) == 0);
    int have_rocm_smi = (stat("/usr/bin/rocm-smi", &st) == 0);
    if (!have_amd_smi && !have_rocm_smi) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "note: amd-smi/rocm-smi not found; GPU "
                "telemetry is unavailable (engine is unaffected). Install the "
                "rocm-smi package for utilization/power readings.\n");
    }
}

/* Read the amdgpu.gttsize boot parameter (MiB) from /proc/cmdline, or 0 if not
 * set / unreadable. Used to detect when a configured GTT size did not take
 * effect against the live pages_limit. */
static uint64_t ds4_rocm_cmdline_gttsize_mib(void) {
    FILE *f = fopen("/proc/cmdline", "r");
    if (!f) return 0;
    char line[2048];
    uint64_t mib = 0;
    if (fgets(line, sizeof(line), f)) {
        const char *p = strstr(line, "amdgpu.gttsize=");
        if (p) {
            const char *v = p + strlen("amdgpu.gttsize=");
            char *end = NULL;
            long long val = strtoll(v, &end, 10);
            if (end != v && val > 0) mib = (uint64_t)val;
        }
    }
    fclose(f);
    return mib;
}

/* Detect the "amdgpu.gttsize=0" / "no GTT special" case: the driver default
 * (no boot parameter override) is usually a small fraction of RAM. We treat
 * the *absence* of amdgpu.gttsize as the driver-default case, and a literal
 * "0" as "explicitly disabled". */
static int ds4_rocm_cmdline_has_gttsize_param(void) {
    FILE *f = fopen("/proc/cmdline", "r");
    if (!f) return 0;
    char line[2048];
    int found = 0;
    if (fgets(line, sizeof(line), f)) found = strstr(line, "amdgpu.gttsize=") != NULL;
    fclose(f);
    return found;
}

/* Minimal GGUF header peek. Reads magic + tensor_count (GGUF v3 layout) and
 * returns 0 on success, -1 on a non-GGUF or truncated file. */
static int ds4_rocm_peek_gguf_header(int fd, char magic[8], uint64_t *tensor_count) {
    magic[0] = '\0';
    *tensor_count = 0;
    if (lseek(fd, 0, SEEK_SET) < 0) return -1;
    unsigned char hdr[24] = {0};
    ssize_t got = read(fd, hdr, sizeof(hdr));
    if (got < 8) return -1;
    if (memcmp(hdr, "GGUF", 4) != 0) return -1;
    /* magic + version(uint32) + tensor_count(uint64) + kv_count(uint64) */
    memcpy(magic, hdr, 4);
    magic[4] = '\0';
    /* Read tensor_count (offset 8, uint64 LE). */
    if (got >= 16) {
        uint64_t tc = 0;
        for (int i = 0; i < 8; i++) tc |= ((uint64_t)hdr[8 + i]) << (i * 8);
        *tensor_count = tc;
    }
    return 0;
}

/* Estimate whether `model_bytes` will fit in the TTM/GTT limit, populating
 * `g_ds4_rocm_last_model_estimate` as a side effect. */
static void ds4_rocm_record_model_estimate(const void *model_map, uint64_t model_bytes) {
    g_ds4_rocm_last_model_estimate.model_bytes = model_bytes;
    g_ds4_rocm_last_model_estimate.gguf_magic[0] = '\0';
    g_ds4_rocm_last_model_estimate.gguf_tensor_count = 0;

    /* Try to peek a GGUF header. We don't need to be authoritative: the file
     * size is the real residency signal, but a magic + tensor count is useful
     * context in the diagnostic. */
    const char *path = NULL;
    /* The model map may be an mmap of a file we don't know the path of, so we
     * best-effort: try /proc/self/maps to find the backing path. */
    char buf[4096];
    buf[0] = '\0';
    FILE *maps = fopen("/proc/self/maps", "r");
    if (maps) {
        uintptr_t want = (uintptr_t)model_map;
        while (fgets(buf, sizeof(buf), maps)) {
            uintptr_t lo = 0, hi = 0;
            char perms[8] = {0};
            char path_buf[1024] = {0};
            if (sscanf(buf, "%lx-%lx %7s %*x %*s %*u %1023s",
                       (unsigned long *)&lo, (unsigned long *)&hi, perms, path_buf) >= 3) {
                if (want >= lo && want < hi && path_buf[0] == '/') {
                    path = strdup(path_buf);
                    break;
                }
            }
        }
        fclose(maps);
    }
    if (path) {
        int fd = open(path, O_RDONLY);
        if (fd >= 0) {
            (void)ds4_rocm_peek_gguf_header(fd, g_ds4_rocm_last_model_estimate.gguf_magic,
                                            &g_ds4_rocm_last_model_estimate.gguf_tensor_count);
            close(fd);
        }
        free((void *)path);
    }

    const uint64_t ttm = ds4_rocm_effective_ttm_bytes();
    g_ds4_rocm_last_model_estimate.ttm_limit_bytes = ttm;
    if (ttm == 0) {
        g_ds4_rocm_last_model_estimate.headroom_bytes = 0;
        g_ds4_rocm_last_model_estimate.would_have_oomed = 0;
        g_ds4_rocm_last_model_estimate.headroom_below_8g = 0;
        return;
    }
    if (model_bytes > ttm) {
        g_ds4_rocm_last_model_estimate.headroom_bytes = 0;
        g_ds4_rocm_last_model_estimate.would_have_oomed = 1;
        g_ds4_rocm_last_model_estimate.headroom_below_8g = 1;
    } else {
        const uint64_t reserved = 8ull * 1073741824ull;
        const uint64_t headroom = ttm - model_bytes;
        g_ds4_rocm_last_model_estimate.headroom_bytes = headroom;
        g_ds4_rocm_last_model_estimate.would_have_oomed = 0;
        g_ds4_rocm_last_model_estimate.headroom_below_8g = headroom < reserved;
    }
}

static void ds4_rocm_print_profile(void) {
#ifdef __HIP_PLATFORM_AMD__
    int dev = 0;
    if (!cuda_ok(cudaSetDevice(dev), "set device")) return;
    cudaDeviceProp prop;
    if (cudaGetDeviceProperties(&prop, dev) != cudaSuccess) return;

    char build_ver[32], runtime_ver[32], driver_ver[32];
    int runtime_version = 0;
    int driver_version = 0;
    (void)hipRuntimeGetVersion(&runtime_version);
    (void)hipDriverGetVersion(&driver_version);
    fprintf(stderr,
            DS4_GPU_LOG_PREFIX "Strix Halo profile: device=%s gcnArchName=%s\n",
            prop.name, prop.gcnArchName);
    fprintf(stderr,
            DS4_GPU_LOG_PREFIX "ROCm build=%s runtime=%s driver=%s "
            "(tested baseline ROCm 7.2.3)\n",
            ds4_rocm_format_version(HIP_VERSION, build_ver, sizeof(build_ver)),
            ds4_rocm_format_version(runtime_version, runtime_ver, sizeof(runtime_ver)),
            ds4_rocm_format_version(driver_version, driver_ver, sizeof(driver_ver)));

    if (strcmp(prop.gcnArchName, "gfx1151") != 0 &&
        strstr(prop.gcnArchName, "gfx1151") == NULL) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "WARNING: device architecture is %s, not "
                "gfx1151 (Strix Halo). Kernels are compiled for gfx1151 and may "
                "not run correctly.\n", prop.gcnArchName);
        g_ds4_rocm_warnings++;
    }

    const int hip_major = HIP_VERSION_MAJOR;
    const int hip_minor = HIP_VERSION_MINOR;
    if (hip_major < 7 || (hip_major == 7 && hip_minor < 2)) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "WARNING: HIP %s is older than the tested "
                "ROCm 7.2.x baseline. Some kernels or the ROCm backend may "
                "behave unexpectedly.\n",
                ds4_rocm_format_version(HIP_VERSION, build_ver, sizeof(build_ver)));
        g_ds4_rocm_warnings++;
    }

    int managed = 0, concurrent_managed = 0, pageable = 0;
    (void)hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory, dev);
    (void)hipDeviceGetAttribute(&concurrent_managed,
                                hipDeviceAttributeConcurrentManagedAccess, dev);
    (void)hipDeviceGetAttribute(&pageable, hipDeviceAttributePageableMemoryAccess, dev);
    fprintf(stderr,
            DS4_GPU_LOG_PREFIX "managed-memory=%d concurrent-managed=%d "
            "pageable-access=%d\n", managed, concurrent_managed, pageable);
    if (pageable == 0) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "NOTE: pageable-memory access is disabled. "
                "The managed-KV-cache demand-paging path assumes pageable access "
                "for good throughput; without it, large KV caches may be slower "
                "or use more device memory.\n");
    }

    ds4_rocm_probe_smi();

    size_t free_b = 0, total_b = 0;
    if (cudaMemGetInfo(&free_b, &total_b) == cudaSuccess) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "HIP-visible memory: %.1f GiB free / %.1f GiB total\n",
                (double)free_b / 1073741824.0, (double)total_b / 1073741824.0);
    }

    const uint64_t total_ram = ds4_rocm_total_system_memory();
    const uint64_t ttm_env = ds4_rocm_effective_ttm_bytes();
    const uint64_t ttm_kernel = ds4_rocm_ttm_pages_limit_bytes();
    if (total_ram > 0) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "system memory: %.1f GiB total\n",
                (double)total_ram / 1073741824.0);
    }
    if (ttm_env > 0) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "TTM/GTT mapping limit: %.1f GiB%s\n",
                (double)ttm_env / 1073741824.0,
                (ttm_kernel > 0 && ttm_env != ttm_kernel) ? " (DS4_ROCM_TTM_PAGES override)" : "");
        ds4_rocm_check_ttm_limit(total_ram, ttm_env);
        const uint64_t gtt_mib = ds4_rocm_cmdline_gttsize_mib();
        const int has_gttsize = ds4_rocm_cmdline_has_gttsize_param();
        if (!has_gttsize) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "NOTE: no amdgpu.gttsize boot parameter "
                    "is set; the driver is choosing a default GTT limit. If "
                    "the default is too small, set amdgpu.gttsize=<MiB> in "
                    "GRUB, or raise the limit at runtime with amd-ttm.\n");
        } else if (gtt_mib > 0) {
            const uint64_t gtt_bytes = gtt_mib * 1048576ull;
            if (ttm_env + 1073741824ull < gtt_bytes) {
                fprintf(stderr,
                        DS4_GPU_LOG_PREFIX "NOTE: amdgpu.gttsize=%llu MiB "
                        "(%.1f GiB) is set but the live TTM/GTT limit is only "
                        "%.1f GiB. The boot parameter may not have taken effect "
                        "(re-check GRUB/kernel cmdline) or a larger BIOS VRAM "
                        "carveout is reducing the available GTT.\n",
                        (unsigned long long)gtt_mib, (double)gtt_bytes / 1073741824.0,
                        (double)ttm_env / 1073741824.0);
            }
        }
    } else {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "TTM/GTT mapping limit: unknown (could not "
                "read ttm/amd_ttm pages_limit). A low GTT limit can OOM large "
                "models; check amd-ttm.\n");
    }

    FILE *diag = ds4_rocm_diag_open();
    if (diag) {
        ds4_rocm_diag_kv_lvl(diag, 0, "ds4_rocm_device", "%s", prop.name);
        ds4_rocm_diag_kv_lvl(diag, 0, "ds4_rocm_gcn_arch", "%s", prop.gcnArchName);
        ds4_rocm_diag_kv_lvl(diag, 0, "ds4_rocm_build", "%s",
                             ds4_rocm_format_version(HIP_VERSION, build_ver, sizeof(build_ver)));
        ds4_rocm_diag_kv_lvl(diag, 0, "ds4_rocm_runtime", "%s",
                             ds4_rocm_format_version(runtime_version, runtime_ver, sizeof(runtime_ver)));
        ds4_rocm_diag_kv_lvl(diag, 0, "ds4_rocm_driver", "%s",
                             ds4_rocm_format_version(driver_version, driver_ver, sizeof(driver_ver)));
        ds4_rocm_diag_kv_lvl(diag, 1, "ds4_rocm_managed_memory", "%d", managed);
        ds4_rocm_diag_kv_lvl(diag, 1, "ds4_rocm_concurrent_managed", "%d", concurrent_managed);
        ds4_rocm_diag_kv_lvl(diag, 1, "ds4_rocm_pageable_access", "%d", pageable);
        if (total_b)
            ds4_rocm_diag_kv_lvl(diag, 1, "ds4_rocm_hip_visible_bytes", "%llu", (unsigned long long)total_b);
        if (total_ram)
            ds4_rocm_diag_kv_lvl(diag, 1, "ds4_rocm_system_ram_bytes", "%llu", (unsigned long long)total_ram);
        if (ttm_env)
            ds4_rocm_diag_kv_lvl(diag, 0, "ds4_rocm_ttm_limit_bytes", "%llu", (unsigned long long)ttm_env);
        const uint64_t gtt_mib = ds4_rocm_cmdline_gttsize_mib();
        if (gtt_mib)
            ds4_rocm_diag_kv_lvl(diag, 2, "ds4_rocm_cmdline_gttsize_mib", "%llu", (unsigned long long)gtt_mib);
        if (g_ds4_rocm_last_model_estimate.model_bytes)
            ds4_rocm_diag_kv_lvl(diag, 2, "ds4_rocm_model_bytes", "%llu",
                                 (unsigned long long)g_ds4_rocm_last_model_estimate.model_bytes);
        if (g_ds4_rocm_last_model_estimate.gguf_magic[0])
            ds4_rocm_diag_kv_lvl(diag, 2, "ds4_rocm_gguf_magic", "%s",
                                 g_ds4_rocm_last_model_estimate.gguf_magic);
        if (g_ds4_rocm_last_model_estimate.gguf_tensor_count)
            ds4_rocm_diag_kv_lvl(diag, 2, "ds4_rocm_gguf_tensor_count", "%llu",
                                 (unsigned long long)g_ds4_rocm_last_model_estimate.gguf_tensor_count);
        ds4_rocm_diag_close(diag);
    }
#else
    (void)ds4_rocm_check_model_fits_ttm;
    (void)ds4_rocm_try_autoraise_ttm;
#endif
}

extern "C" int ds4_gpu_init(void) {
    int dev = 0;
    if (!cuda_ok(cudaSetDevice(dev), "set device")) return 0;
    cudaDeviceProp prop;
    if (cudaGetDeviceProperties(&prop, dev) == cudaSuccess) {
        fprintf(stderr, DS4_GPU_LOG_PREFIX "backend initialized on %s (sm_%d%d)\n",
                prop.name, prop.major, prop.minor);
    }
#ifdef __HIP_PLATFORM_AMD__
    ds4_rocm_print_profile();
#endif
    if (!g_cublas_ready) {
        if (!cublas_ok(cublasCreate(&g_cublas), "create handle")) return 0;
        const cublasMath_t math_mode = g_quality_mode ? CUBLAS_DEFAULT_MATH : CUBLAS_TF32_TENSOR_OP_MATH;
        (void)cublasSetMathMode(g_cublas, math_mode);
        g_cublas_ready = 1;
    }
#ifdef __HIP_PLATFORM_AMD__
    if (!g_hipblaslt_ready) {
        if (hipblaslt_ok(hipblasLtCreate(&g_hipblaslt), "create handle")) {
            g_hipblaslt_ready = 1;
        }
    }
#endif
    return 1;
}

extern "C" void ds4_gpu_cleanup(void) {
    (void)cudaDeviceSynchronize();
    cuda_stream_cache_stats_print("cleanup");
    cuda_shared_gate_up_async_cleanup();
#ifdef __HIP_PLATFORM_AMD__
    hipblaslt_gemm_plan_clear();
#endif
    if (g_cublas_ready) {
        (void)cublasDestroy(g_cublas);
        g_cublas_ready = 0;
        g_cublas = NULL;
    }
#ifdef __HIP_PLATFORM_AMD__
    if (g_hipblaslt_ready) {
        (void)hipblasLtDestroy(g_hipblaslt);
        g_hipblaslt_ready = 0;
        g_hipblaslt = NULL;
    }
#endif
    cuda_model_range_release_all();
    cuda_q8_f16_cache_release_all();
    cuda_stream_selected_cache_release();
    g_q8_f16_disabled_after_oom = 0;
    g_q8_f16_disabled_for_multi_model = 0;
    g_q8_f16_budget_notice_printed = 0;
    if (g_cuda_tmp) {
        (void)cudaFree(g_cuda_tmp);
        g_cuda_tmp = NULL;
        g_cuda_tmp_bytes = 0;
    }
    for (size_t i = 0; i < 4; i++) {
        if (g_model_stage_event[i]) {
            (void)cudaEventDestroy(g_model_stage_event[i]);
            g_model_stage_event[i] = NULL;
        }
        if (g_model_stage_raw[i]) {
            (void)cudaFreeHost(g_model_stage_raw[i]);
            g_model_stage_raw[i] = NULL;
            g_model_stage[i] = NULL;
        }
    }
    g_model_stage_bytes = 0;
    if (g_model_upload_stream) {
        (void)cudaStreamDestroy(g_model_upload_stream);
        g_model_upload_stream = NULL;
    }
    if (g_stream_selected_upload_stream) {
        (void)cudaStreamDestroy(g_stream_selected_upload_stream);
        g_stream_selected_upload_stream = NULL;
    }
    if (g_selected_readback_stream) {
        (void)cudaStreamDestroy(g_selected_readback_stream);
        g_selected_readback_stream = NULL;
    }
    if (g_selected_readback_event) {
        (void)cudaEventDestroy(g_selected_readback_event);
        g_selected_readback_event = NULL;
    }
    g_selected_readback_event_value = 0;
    cuda_model_image_release_all();
    g_model_host_base = NULL;
    g_model_device_base = NULL;
    g_model_registered_size = 0;
    g_model_device_owned = 0;
    g_model_range_mapping_supported = 1;
    g_model_fd = -1;
    if (g_model_direct_fd >= 0) {
        (void)close(g_model_direct_fd);
        g_model_direct_fd = -1;
    }
    g_model_direct_align = 1;
    g_model_file_size = 0;
    g_model_cache_full = 0;
}

__global__ static void fill_f32_kernel(float *x, uint64_t n, float v);

extern "C" ds4_gpu_tensor *ds4_gpu_tensor_alloc(uint64_t bytes) {
    if (bytes == 0) bytes = 1;
    ds4_gpu_tensor *t = (ds4_gpu_tensor *)calloc(1, sizeof(*t));
    if (!t) return NULL;
    if (!cuda_ok(cudaMalloc(&t->ptr, (size_t)bytes), "tensor alloc")) {
        free(t);
        return NULL;
    }
    t->bytes = bytes;
    t->owner = 1;
    return t;
}

extern "C" ds4_gpu_tensor *ds4_gpu_tensor_alloc_managed(uint64_t bytes) {
    if (bytes == 0) bytes = 1;
    ds4_gpu_tensor *t = (ds4_gpu_tensor *)calloc(1, sizeof(*t));
    if (!t) return NULL;
    if (!cuda_ok(cudaMallocManaged(&t->ptr, (size_t)bytes), "managed tensor alloc")) {
        free(t);
        return NULL;
    }
    t->bytes = bytes;
    t->owner = 1;
    return t;
}

static uint64_t cuda_managed_kv_reserve_bytes(uint64_t total_bytes) {
    const uint64_t min_reserve = 8ull * 1073741824ull;
    const uint64_t max_reserve = 40ull * 1073741824ull;
    uint64_t reserve = total_bytes / 4u;
    if (reserve < min_reserve) reserve = min_reserve;
    if (reserve > max_reserve) reserve = max_reserve;
    return reserve;
}

extern "C" int ds4_gpu_should_use_managed_kv_cache(uint64_t kv_cache_bytes, uint64_t context_bytes) {
    if (kv_cache_bytes == 0) return 0;

    /* Very large KV caches are where device-only cudaMalloc() can make a
     * unified-memory machine unresponsive.  Managed memory restores the old
     * demand-paged behavior for this one long-lived allocation class only. */
    const uint64_t huge_kv = 8ull * 1073741824ull;
    if (kv_cache_bytes >= huge_kv) return 1;

    const uint64_t large_context = 8ull * 1073741824ull;
    if (context_bytes < large_context) return 0;

    size_t free_b = 0;
    size_t total_b = 0;
    cudaError_t err = cudaMemGetInfo(&free_b, &total_b);
    if (err != cudaSuccess) {
        (void)cudaGetLastError();
        return 0;
    }

    const uint64_t free_bytes = (uint64_t)free_b;
    const uint64_t total_bytes = (uint64_t)total_b;
    const uint64_t reserve_bytes = cuda_managed_kv_reserve_bytes(total_bytes);
    if (context_bytes > free_bytes) return 1;
    return free_bytes - context_bytes < reserve_bytes;
}

extern "C" ds4_gpu_tensor *ds4_gpu_tensor_view(const ds4_gpu_tensor *base, uint64_t offset, uint64_t bytes) {
    if (!base || offset > base->bytes || bytes > base->bytes - offset) return NULL;
    ds4_gpu_tensor *t = (ds4_gpu_tensor *)calloc(1, sizeof(*t));
    if (!t) return NULL;
    t->ptr = (char *)base->ptr + offset;
    t->bytes = bytes;
    t->owner = 0;
    return t;
}

extern "C" void ds4_gpu_tensor_free(ds4_gpu_tensor *tensor) {
    if (!tensor) return;
    if (tensor->owner && tensor->ptr) (void)cudaFree(tensor->ptr);
    free(tensor);
}

extern "C" uint64_t ds4_gpu_tensor_bytes(const ds4_gpu_tensor *tensor) {
    return tensor ? tensor->bytes : 0;
}

extern "C" void *ds4_gpu_tensor_contents(ds4_gpu_tensor *tensor) {
    if (!tensor) return NULL;
    (void)cudaDeviceSynchronize();
    return tensor->ptr;
}

extern "C" int ds4_gpu_tensor_fill_f32(ds4_gpu_tensor *tensor, float value, uint64_t count) {
    if (!tensor || count > tensor->bytes / sizeof(float)) return 0;
    if (count == 0) return 1;
    fill_f32_kernel<<<(count + 255u) / 256u, 256>>>((float *)tensor->ptr, count, value);
    return cuda_ok(cudaGetLastError(), "tensor fill f32 launch");
}

extern "C" int ds4_gpu_tensor_write(ds4_gpu_tensor *tensor, uint64_t offset, const void *data, uint64_t bytes) {
    if (!tensor || !data || offset > tensor->bytes || bytes > tensor->bytes - offset) return 0;
    return cuda_ok(cudaMemcpy((char *)tensor->ptr + offset, data, (size_t)bytes, cudaMemcpyHostToDevice), "tensor write");
}

extern "C" int ds4_gpu_tensor_read(const ds4_gpu_tensor *tensor, uint64_t offset, void *data, uint64_t bytes) {
    if (!tensor || !data || offset > tensor->bytes || bytes > tensor->bytes - offset) return 0;
    return cuda_ok(cudaMemcpy(data, (const char *)tensor->ptr + offset, (size_t)bytes, cudaMemcpyDeviceToHost), "tensor read");
}

extern "C" int ds4_gpu_tensor_copy(ds4_gpu_tensor *dst, uint64_t dst_offset,
                                     const ds4_gpu_tensor *src, uint64_t src_offset,
                                     uint64_t bytes) {
    if (!dst || !src || dst_offset > dst->bytes || src_offset > src->bytes ||
        bytes > dst->bytes - dst_offset || bytes > src->bytes - src_offset) {
        return 0;
    }
    if (bytes == 0) return 1;
    return cuda_ok(cudaMemcpyAsync((char *)dst->ptr + dst_offset,
                                   (const char *)src->ptr + src_offset,
                                   (size_t)bytes,
                                   cudaMemcpyDeviceToDevice,
                                   0),
                   "tensor copy enqueue");
}

extern "C" int ds4_gpu_begin_commands(void) { return 1; }
extern "C" int ds4_gpu_flush_commands(void) { return cuda_ok(cudaDeviceSynchronize(), "flush"); }
extern "C" int ds4_gpu_flush_encoder(void) { return ds4_gpu_flush_commands(); }
extern "C" int ds4_gpu_commands_active(void) { return 0; }
extern "C" int ds4_gpu_signal_selected_readback_ready(uint64_t *event_value) {
    if (!event_value) return 0;
    *event_value = 0;
    if (!g_selected_readback_event) {
        cudaError_t err =
            cudaEventCreateWithFlags(&g_selected_readback_event,
                                     cudaEventDisableTiming);
        if (err != cudaSuccess) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "selected readback event creation failed: %s\n",
                    cudaGetErrorString(err));
            (void)cudaGetLastError();
            return 0;
        }
    }
    cudaError_t err = cudaEventRecord(g_selected_readback_event, 0);
    if (err != cudaSuccess) {
        fprintf(stderr,
                DS4_GPU_LOG_PREFIX "selected readback event record failed: %s\n",
                cudaGetErrorString(err));
        (void)cudaGetLastError();
        return 0;
    }
    *event_value = ++g_selected_readback_event_value;
    return 1;
}
extern "C" int ds4_gpu_commit_and_wait_selected_readback(uint64_t event_value, const char *label) {
    if (event_value == 0 || !g_selected_readback_event) return 0;
    return cuda_ok(cudaEventSynchronize(g_selected_readback_event),
                   label ? label : "selected readback");
}
extern "C" int ds4_gpu_wait_selected_readback_ready(uint64_t event_value, const char *label) {
    if (event_value == 0 || !g_selected_readback_event) return 0;
    return cuda_ok(cudaEventSynchronize(g_selected_readback_event),
                   label ? label : "selected readback wait");
}
extern "C" int ds4_gpu_end_commands(void) {
    return cuda_ok(cudaDeviceSynchronize(), "end commands");
}
extern "C" int ds4_gpu_synchronize(void) { return cuda_ok(cudaDeviceSynchronize(), "synchronize"); }

extern "C" int ds4_gpu_set_model_map(const void *model_map, uint64_t model_size) {
    if (!model_map || model_size == 0) return 0;
    if (g_model_host_base == model_map && g_model_registered_size == model_size) return 1;
    const int multi_model =
        g_model_host_base != NULL &&
        (g_model_host_base != model_map || g_model_registered_size != model_size);
    cuda_model_range_release_all();
    cuda_q8_f16_cache_release_all();
    g_q8_f16_disabled_after_oom = 0;
    g_q8_f16_budget_notice_printed = 0;
    if (multi_model) {
        /*
         * MTP loads a second GGUF mapping.  Its weights are small, but on UMA
         * ROCm systems the optional expanded Q8->F16 cache can consume the
         * memory margin needed for session/context tensors once both model
         * mappings are resident.  The cache is only a speed path; the normal
         * Q8 kernels remain available and keep MTP startup reliable.
         */
        g_q8_f16_disabled_for_multi_model = 1;
    }
    g_model_host_base = model_map;
    g_model_device_base = cuda_model_image_owned(model_map) ?
                          cuda_model_image_ptr(model_map, 0) :
                          (const char *)model_map;
    g_model_registered_size = model_size;
    g_model_device_owned = cuda_model_image_owned(model_map);
    g_model_range_mapping_supported = 1;
    g_model_cache_full = 0;
    if (g_model_fd >= 0 && g_model_fd_host_base == NULL) {
        g_model_fd_host_base = model_map;
    }

    /* Strix Halo uses the staged full-copy path in ds4_gpu_set_model_map_range().
     * Avoid host-registering the mmap here: that would make the staged copier
     * believe the model is already device-resident. */
#ifdef __HIP_PLATFORM_AMD__
    uint64_t ttm_bytes = ds4_rocm_effective_ttm_bytes();
    const uint64_t reserved = 8ull * 1073741824ull;
    if (ttm_bytes > 0 && model_size + reserved > ttm_bytes) {
        const uint64_t want = model_size + reserved;
        const uint64_t raised = ds4_rocm_try_autoraise_ttm(want);
        if (raised > ttm_bytes) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "raised TTM/GTT limit to %.1f GiB\n",
                    (double)raised / 1073741824.0);
            ttm_bytes = raised;
        }
    }
    if (ttm_bytes > 0) ds4_rocm_check_model_fits_ttm(model_size, ttm_bytes);
    /* Record the per-model residency verdict for tooling (smoke, doctor). */
    ds4_rocm_record_model_estimate(model_map, model_size);
#else
    (void)ds4_rocm_check_model_fits_ttm;
#endif
    return 1;
}

extern "C" int ds4_gpu_set_model_map_range(const void *model_map, uint64_t model_size, uint64_t map_offset, uint64_t map_size, uint64_t max_tensor_bytes) {
    (void)max_tensor_bytes;
    if (!model_map || model_size == 0 ||
        map_offset > model_size ||
        map_size > model_size - map_offset) {
        return 0;
    }
    if (!ds4_gpu_set_model_map(model_map, model_size)) return 0;
    if (g_ssd_streaming_mode) {
        const uint64_t limit = cuda_model_cache_limit_bytes();
        if (map_size > limit) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming model range %.2f GiB exceeds "
                    "cache limit %.2f GiB; increase DS4_ROCM_STREAM_MODEL_CACHE_GB\n",
                    (double)map_size / 1073741824.0,
                    (double)limit / 1073741824.0);
            return 0;
        }
        if (g_model_range_bytes > limit ||
            map_size > limit - g_model_range_bytes) {
            if (!cuda_ok(cudaDeviceSynchronize(),
                         "streaming model range cache eviction sync")) {
                return 0;
            }
            cuda_model_range_release_ranges_only();
        }
        if (!cuda_stream_model_cache_prepare_memory(map_size,
                                                    "streaming model range")) {
            return 0;
        }
        if (!cuda_model_range_ptr(model_map, map_offset, map_size, "stream_range")) return 0;
        return cuda_model_range_is_cached(model_map, map_offset, map_size);
    }
    /*
     * Do not eagerly copy a contiguous model image here.  On Strix Halo the
     * caller immediately follows with accelerator_cache_model_tensors(), which
     * prepares the exact tensor spans selected by --layers.  Copying here would
     * either allocate the whole GGUF image or, for sparse span sets, an oversized
     * envelope before the precise tensor-span cache gets a chance to run.
     */
    return 1;
}

extern "C" int ds4_gpu_set_model_map_spans(
        const void *model_map,
        uint64_t model_size,
        const uint64_t *offsets,
        const uint64_t *sizes,
        uint32_t count,
        uint64_t max_tensor_bytes) {
    (void)max_tensor_bytes;
    if (!model_map || model_size == 0 || !offsets || !sizes || count == 0) return 0;
    for (uint32_t i = 0; i < count; i++) {
        if (offsets[i] > model_size ||
            sizes[i] == 0 ||
            sizes[i] > model_size - offsets[i]) {
            return 0;
        }
    }
    if (!ds4_gpu_set_model_map(model_map, model_size)) return 0;
    if (g_ssd_streaming_mode) {
        uint64_t request_bytes = 0;
        for (uint32_t i = 0; i < count; i++) {
            if (!cuda_u64_add_checked(request_bytes, sizes[i], &request_bytes)) {
                fprintf(stderr,
                        DS4_GPU_LOG_PREFIX "streaming model span byte count overflow\n");
                return 0;
            }
        }
        const uint64_t limit = cuda_model_cache_limit_bytes();
        if (request_bytes > limit) {
            fprintf(stderr,
                    DS4_GPU_LOG_PREFIX "streaming model spans %.2f GiB exceed "
                    "cache limit %.2f GiB; increase DS4_ROCM_STREAM_MODEL_CACHE_GB\n",
                    (double)request_bytes / 1073741824.0,
                    (double)limit / 1073741824.0);
            return 0;
        }
        if (g_model_range_bytes > limit ||
            request_bytes > limit - g_model_range_bytes) {
            if (!cuda_ok(cudaDeviceSynchronize(),
                         "streaming model span cache eviction sync")) {
                return 0;
            }
            cuda_model_range_release_ranges_only();
        }
        if (!cuda_stream_model_cache_prepare_memory(request_bytes,
                                                    "streaming model spans")) {
            return 0;
        }
        for (uint32_t i = 0; i < count; i++) {
            if (!cuda_model_range_ptr(model_map, offsets[i], sizes[i], "stream_span")) return 0;
            if (!cuda_model_range_is_cached(model_map, offsets[i], sizes[i])) return 0;
        }
        return 1;
    }
    /*
     * The spans can be sparse distributed layer slices.  Materializing their
     * min..max envelope can be much larger than the actual selected tensors.
     * Leave the precise per-tensor preparation to accelerator_cache_model_tensors().
     */
    return 1;
}

extern "C" int ds4_gpu_set_model_fd(int fd) {
    g_model_fd = fd;
    g_model_fd_host_base = g_model_host_base;
    g_model_file_size = 0;
    if (g_model_direct_fd >= 0) {
        (void)close(g_model_direct_fd);
        g_model_direct_fd = -1;
    }
    g_model_direct_align = 1;
    if (fd >= 0) {
        struct stat st;
        if (fstat(fd, &st) == 0 && st.st_size > 0) {
            g_model_file_size = (uint64_t)st.st_size;
            if (st.st_blksize > 1) g_model_direct_align = (uint64_t)st.st_blksize;
        }
#if defined(__linux__) && defined(O_DIRECT)
        {
            char proc_path[64];
            snprintf(proc_path, sizeof(proc_path), "/proc/self/fd/%d", fd);
            int direct_fd = open(proc_path, O_RDONLY | O_DIRECT);
            if (direct_fd >= 0) {
                g_model_direct_fd = direct_fd;
                if (g_model_direct_align < 512) g_model_direct_align = 512;
            }
        }
#endif
    }
    return 1;
}

extern "C" int ds4_gpu_cache_model_range(const void *model_map, uint64_t model_size, uint64_t offset, uint64_t bytes, const char *label) {
    if (!model_map || bytes == 0) return 1;
    if (offset > model_size || bytes > model_size - offset) return 0;
    if (!cuda_model_range_ptr(model_map, offset, bytes, label ? label : "model_tensor")) return 0;
    return cuda_model_range_is_cached(model_map, offset, bytes);
}

extern "C" int ds4_gpu_cache_q8_f16_range(const void *model_map, uint64_t model_size, uint64_t offset, uint64_t bytes, uint64_t in_dim, uint64_t out_dim, const char *label) {
    if (!model_map || bytes == 0) return 1;
    if (offset > model_size || bytes > model_size - offset) return 0;
    static int optional_q8_preload_disabled = 0;
    if (optional_q8_preload_disabled) return 1;
    const char *cache_label = label ? label : "q8_0";
    if (!cuda_q8_f16_preload_allowed(cache_label, in_dim, out_dim)) return 1;
    const int preload_transposed_b = !g_quality_mode &&
                                     strstr(cache_label, "attn_output_b") != NULL;
    if (preload_transposed_b) {
        const __half *f16_t = cuda_q8_f16_transpose_ptr(model_map, offset, bytes, in_dim, out_dim, cache_label);
        if (f16_t) {
            if (strstr(cache_label, "attn_output_b") != NULL && in_dim == 8192u && out_dim == 4096u) {
                cuda_q8_f16_warmup_attention_output_b_gemm(f16_t, in_dim, out_dim);
            }
            return 1;
        }
    } else {
        const __half *f16 = cuda_q8_f16_ptr(model_map, offset, bytes, in_dim, out_dim, cache_label);
        if (f16) {
            if (strstr(cache_label, "attn_output_a") != NULL && in_dim == 4096u && out_dim == 8192u) {
                cuda_q8_f16_warmup_attention_output_a_gemm(f16, in_dim, 1024u, 8u);
            }
            return 1;
        }
    }
    optional_q8_preload_disabled = 1;
    return 1;
}

extern "C" void ds4_gpu_release_q8_f16_cache(void) {
    cuda_q8_f16_cache_release_all();
    g_q8_f16_disabled_after_oom = 0;
    g_q8_f16_budget_notice_printed = 0;
}

extern "C" void ds4_gpu_print_memory_report(const char *label) {
    size_t free_b = 0, total_b = 0;
    cudaError_t err = cudaMemGetInfo(&free_b, &total_b);
    if (err != cudaSuccess) {
        fprintf(stderr, DS4_GPU_LOG_PREFIX "memory %s: query failed: %s\n",
                label ? label : "", cudaGetErrorString(err));
        (void)cudaGetLastError();
        return;
    }
    const uint64_t used_b = (uint64_t)total_b - (uint64_t)free_b;
    const char *placement = cuda_model_image_bytes() ? "device_copy" : "mapped/range_cache";
    fprintf(stderr,
            DS4_GPU_LOG_PREFIX "memory %s: used=%.2f GiB free=%.2f GiB total=%.2f GiB "
            "placement=%s model_image=%.2f GiB range_cache=%.2f GiB "
            "q8_f16_cache=%.2f GiB scratch=%.2f GiB",
            label ? label : "",
            (double)used_b / 1073741824.0,
            (double)free_b / 1073741824.0,
            (double)total_b / 1073741824.0,
            placement,
            (double)cuda_model_image_bytes() / 1073741824.0,
            (double)g_model_range_bytes / 1073741824.0,
            (double)g_q8_f16_bytes / 1073741824.0,
            (double)g_cuda_tmp_bytes / 1073741824.0);
    fprintf(stderr, "\n");
}

extern "C" void ds4_gpu_set_quality(bool quality) {
    const int new_quality_mode = quality ? 1 : 0;
    if (g_quality_mode != new_quality_mode) {
        g_rocm_cfg.initialized = 0;
    }
    g_quality_mode = new_quality_mode;
    if (g_cublas_ready) {
        const cublasMath_t math_mode = g_quality_mode ? CUBLAS_DEFAULT_MATH : CUBLAS_TF32_TENSOR_OP_MATH;
        (void)cublasSetMathMode(g_cublas, math_mode);
    }
}