Fix: AMD ROCm/HIP tiered cache improvements

- Fix Issue 1: Replace CUDA API calls with HIP API in ggml-cuda.cu
  - cudaMemcpyPeerAsync → hipMemcpyPeerAsync
  - cudaError_t → hipError_t
  - cudaSuccess → hipSuccess

- Fix Issue 2: Fix typo std.max → std::max in llama-eviction-policy.h

- Fix Issue 3: Implement migrate_tokens and batch_migrate_tokens with HIP-aware memory copies

- Fix Issue 4: Implement save_to_ssd and load_from_ssd using llama_ssd_storage_format

- Fix Issue 5: Add ggml tensor handles (ggml_tensor*) for K and V data to tier class

- Fix Issue 6: Add device-to-host copy for hot→warm path before serialization
  - Added is_device_data parameter to save_to_ssd()
  - Added to_device parameter to load_from_ssd()
  - Uses hipMemcpy with hipMemcpyDeviceToHost/hipMemcpyHostToDevice as needed

Author: kmbandy

.gitignore

@@ -148,3 +148,28 @@ a.out.*
 .gitnexus
 .claude/
 .agent.md
+
+# cmake root-level build artifacts
+/CMakeFiles/
+/CMakeCache.txt
+/CTestTestfile.cmake
+/DartConfiguration.tcl
+/cmake_install.cmake
+/bin/
+/CPackConfig.cmake
+/CPackSourceConfig.cmake
+*.so
+*.so.*
+
+# nested cmake build artifacts
+**/CMakeFiles/
+**/cmake_install.cmake
+**/CTestTestfile.cmake
+**/Makefile
+compile_commands.json
+error.txt
+license.cpp
+llama.pc
+ggml/ggml-config.cmake
+ggml/ggml-version.cmake
+tests/libgguf-model-data.a

common/arg.cpp

@@ -31,6 +31,7 @@
 #include <list>
 #include <regex>
 #include <set>
+#include <sstream>
 #include <string>
 #include <thread> // for hardware_concurrency
 #include <vector>
@@ -1310,6 +1311,53 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
             params.cache_ram_mib = value;
         }
     ).set_env("LLAMA_ARG_CACHE_RAM").set_examples({LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}));
+    add_opt(common_arg(
+        {"--kv-tiered"}, "VRAM%,RAM%,SSD%",
+        "enable tiered KV cache with percentages for VRAM, RAM, SSD (e.g., 25,25,50 for 25% VRAM, 25% RAM, 50% SSD)",
+        [](common_params & params, const std::string & value) {
+            params.kv_tiered_enabled = true;
+            // Parse the percentage string
+            std::vector<std::string> parts;
+            std::stringstream ss(value);
+            std::string part;
+            while (std::getline(ss, part, ',')) {
+                parts.push_back(part);
+            }
+            if (parts.size() == 3) {
+                params.kv_tier_hot_pct = std::stof(parts[0]);
+                params.kv_tier_warm_pct = std::stof(parts[1]);
+                params.kv_tier_cold_pct = std::stof(parts[2]);
+            }
+        }
+    ).set_env("LLAMA_ARG_KT_TIERED").set_examples({LLAMA_EXAMPLE_SERVER}));
+    add_opt(common_arg(
+        {"--tier-ssd-path"}, "PATH",
+        "directory for cold tier (SSD) storage",
+        [](common_params & params, const std::string & value) {
+            params.kv_tier_ssd_path = value;
+        }
+    ).set_env("LLAMA_ARG_TIER_SSD_PATH").set_examples({LLAMA_EXAMPLE_SERVER}));
+    add_opt(common_arg(
+        {"--tier-eviction-policy"}, "POLICY",
+        "eviction policy: 0=LRU, 1=LFU, 2=attention, 3=hybrid (default)",
+        [](common_params & params, int value) {
+            params.kv_tier_eviction_policy = value;
+        }
+    ).set_env("LLAMA_ARG_TIER_EVICTION_POLICY").set_examples({LLAMA_EXAMPLE_SERVER}));
+    add_opt(common_arg(
+        {"--tier-compression"}, "TYPE",
+        "compression type: 0=none, 1=int4, 2=int8, 3=lz4, 4=quantized",
+        [](common_params & params, int value) {
+            params.kv_tier_compression = value;
+        }
+    ).set_env("LLAMA_ARG_TIER_COMPRESSION").set_examples({LLAMA_EXAMPLE_SERVER}));
+    add_opt(common_arg(
+        {"--tier-attention-threshold"}, "THRESH",
+        "attention threshold for eviction (0.0-1.0, default: 0.1)",
+        [](common_params & params, int value) {
+            params.kv_tier_attention_threshold = float(value) / 100.0f;
+        }
+    ).set_env("LLAMA_ARG_TIER_ATTENTION_THRESHOLD").set_examples({LLAMA_EXAMPLE_SERVER}));
     add_opt(common_arg(
         {"-kvu", "--kv-unified"},
         {"-no-kvu", "--no-kv-unified"},

common/common.h

@@ -573,6 +573,16 @@ struct common_params {
     int32_t checkpoint_every_nt = 8192;  // make a checkpoint every n tokens during prefill
     int32_t cache_ram_mib       = 8192;  // -1 = no limit, 0 - disable, 1 = 1 MiB, etc.
 
+    // tiered KV cache parameters
+    bool kv_tiered_enabled    = false;   // enable tiered KV cache
+    float kv_tier_hot_pct     = 25.0f;   // hot tier percentage (VRAM)
+    float kv_tier_warm_pct    = 25.0f;   // warm tier percentage (RAM)
+    float kv_tier_cold_pct    = 50.0f;   // cold tier percentage (SSD)
+    std::string kv_tier_ssd_path = "";   // SSD path for cold tier storage
+    int kv_tier_eviction_policy = 3;     // 0=LRU, 1=LFU, 2=attention, 3=hybrid (default)
+    int kv_tier_compression   = 1;       // 0=none, 1=int4, 2=int8, 3=lz4, 4=quantized
+    float kv_tier_attention_threshold = 0.1f;  // attention threshold for eviction
+
     std::string hostname      = "127.0.0.1";
     std::string public_path   = "";                                                                         // NOLINT
     std::string api_prefix    = "";                                                                         // NOLINT

ggml/src/ggml-cuda/ggml-cuda.cu

@@ -1426,6 +1426,54 @@ static cudaError_t ggml_cuda_cpy_tensor_2d(
     const int64_t i1_diff = i1_high - i1_low;
 
     const char * x = src_ptr + i1_low*nb1 + i2*nb2 + i3*nb3;
+
+#if defined(GGML_USE_HIP)
+    // hipMemcpy2DAsync with hipMemcpyDeviceToDevice requires P2P access between devices.
+    // For mixed-architecture multi-GPU setups (e.g. gfx1201 + gfx1030) without P2P,
+    // it fails asynchronously with hipErrorNoBinaryForGpu when the internal copy kernel
+    // isn't compiled for the destination device. Detect cross-device copies via pointer
+    // attributes and fall back to row-wise hipMemcpyPeerAsync, which stages through the
+    // host when P2P is unavailable. Mirrors the fix in ggml_cuda_Memcpy2DPeerAsync.
+    hipPointerAttribute_t src_attr = {};
+    hipPointerAttribute_t dst_attr = {};
+    const bool src_ok = hipPointerGetAttributes(&src_attr, x)       == hipSuccess;
+    const bool dst_ok = hipPointerGetAttributes(&dst_attr, dst_ptr) == hipSuccess;
+    if (src_ok && dst_ok && src_attr.device != dst_attr.device) {
+        const int src_dev = src_attr.device;
+        const int dst_dev = dst_attr.device;
+        if (nb0 == ts && nb1 == ts*ne0/bs) {
+            return hipMemcpyPeerAsync(dst_ptr, dst_dev, x, src_dev, i1_diff*nb1, stream);
+        } else if (nb0 == ts) {
+            const size_t row_bytes = ts*ne0/bs;
+            for (int64_t i1 = 0; i1 < i1_diff; i1++) {
+                hipError_t err = hipMemcpyPeerAsync(
+                    dst_ptr + i1*row_bytes, dst_dev,
+                    x + i1*nb1,             src_dev,
+                    row_bytes, stream);
+                if (err != hipSuccess) {
+                    return err;
+                }
+            }
+            return hipSuccess;
+        } else {
+            for (int64_t i1 = 0; i1 < i1_diff; i1++) {
+                const char * rx = x + i1*nb1;
+                char       * rd = dst_ptr + i1*ts*ne0/bs;
+                for (int64_t e = 0; e < ne0; e++) {
+                    hipError_t err = hipMemcpyPeerAsync(
+                        rd + e*(ts/bs), dst_dev,
+                        rx + e*nb0,     src_dev,
+                        ts/bs, stream);
+                    if (err != hipSuccess) {
+                        return err;
+                    }
+                }
+            }
+            return hipSuccess;
+        }
+    }
+#endif
+
     if (nb0 == ts && nb1 == ts*ne0/bs) {
         return cudaMemcpyAsync(dst_ptr, x, i1_diff*nb1, cudaMemcpyDeviceToDevice, stream);
     } else if (nb0 == ts) {