feat(hip): AMD ROCm/HIP port, INT8 KV cache, and runtime VRAM stability

CMakeLists.txt

@@ -1,4 +1,4 @@
-# Strata Engine - top-level build (P1.S1).
+# Strata Engine - top-level build (P1.S1).
 #
 # Targets (phase-1 §P1.S1): `strata_artifact` (no CUDA dependency), `tests`.  `strata_plan` and the C++
 # `strata-pack` are added when their sources exist - a target with no sources would make the build graph
@@ -1630,7 +1630,18 @@ endif()
 #
 # A second engine beside `strata`.  It shares the quantized matvec kernels and the GGUF reader, and takes the
 # CPU experts' dot products from ggml-cpu.
-if(STRATA_ENABLE_CUDA AND STRATA_NATIVE_EXPERTS AND NOT STRATA_ENABLE_HIP AND NOT STRATA_HIP_GFX906)
+if((STRATA_ENABLE_CUDA OR STRATA_ENABLE_HIP) AND STRATA_NATIVE_EXPERTS AND NOT STRATA_HIP_GFX906)
+  if(NOT _strata_gg)
+    if(STRATA_GGML_DIR)
+      set(_strata_gg "${STRATA_GGML_DIR}/ggml")
+    else()
+      set(_strata_gg "${strata_llamacpp_SOURCE_DIR}/ggml")
+    endif()
+  endif()
+  if(NOT _strata_mmq)
+    set(_strata_mmq "${_strata_gg}/src/ggml-cuda")
+  endif()
+
   # The prompt path's quantized products: llama.cpp's MMQ kernels (ggml-cuda, MIT) for the types the model's
   # GGUFs keep their matrices in (the list in src/poly/mmq.cu), from the checkout strata_mmq is compiled from and
   # with the same host symbols.  A library of its own, so that what `strata` multiplies with MMQ stays as it is.
@@ -1639,31 +1650,52 @@ if(STRATA_ENABLE_CUDA AND STRATA_NATIVE_EXPERTS AND NOT STRATA_ENABLE_HIP AND NO
              iq4_xs)
     list(APPEND _strata_poly_mmq_instances ${_strata_mmq}/template-instances/mmq-instance-${_t}.cu)
   endforeach()
+  if(STRATA_ENABLE_HIP)
+    set_source_files_properties(${_strata_mmq}/quantize.cu ${_strata_poly_mmq_instances}
+                                PROPERTIES LANGUAGE HIP)
+  endif()
   add_library(strata_poly_mmq STATIC src/poly/mmq.cu src/prefill/ggml_cuda_host.cu ${_strata_mmq}/quantize.cu
               ${_strata_poly_mmq_instances})
   target_include_directories(strata_poly_mmq PRIVATE ${_strata_gg}/include ${_strata_gg}/src ${_strata_mmq}
                                                     ${CMAKE_CURRENT_SOURCE_DIR}/include)
+  if(STRATA_ENABLE_HIP)
+    target_compile_definitions(strata_poly_mmq PRIVATE GGML_USE_HIP=1)
+  endif()
   target_compile_options(strata_poly_mmq PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:--extended-lambda -use_fast_math>)
   if(MSVC)
     target_compile_options(strata_poly_mmq PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=/bigobj>)
   endif()
-  target_link_libraries(strata_poly_mmq PUBLIC ggml-base ${_strata_gpu_runtime_target})
+  if(STRATA_ENABLE_HIP)
+    target_link_libraries(strata_poly_mmq PUBLIC ggml-base hip::host roc::hipblas)
+  else()
+    target_link_libraries(strata_poly_mmq PUBLIC ggml-base ${_strata_gpu_runtime_target})
+  endif()
 
   add_executable(strata-poly src/poly/main.cpp src/poly/engine.cpp src/poly/glm5next.cpp src/poly/qwen35.cpp
                             src/poly/gemma4.cpp src/poly/cpu_experts.cpp src/poly/kernels.cu
                             src/core/conversation_memory.cpp)
   target_include_directories(strata-poly PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/include ${_strata_gpu_include_directories})
   # the driver's virtual memory calls: the attention caches grow with the conversation
   target_link_libraries(strata-poly PRIVATE strata_poly_mmq strata_kernels strata_spec ggml-cpu ggml-base
-                                           ${_strata_gpu_blas_target} ${_strata_gpu_runtime_target} CUDA::cuda_driver
+                                           ${_strata_gpu_blas_target} ${_strata_gpu_runtime_target}
                                            Threads::Threads)
+  if(TARGET CUDA::cuda_driver)
+    target_link_libraries(strata-poly PRIVATE CUDA::cuda_driver)
+  endif()
   # needs the model file, so it is run by hand: poly_expert_parity MODEL.gguf
   add_executable(poly_expert_parity src/poly/expert_parity.cpp src/poly/cpu_experts.cpp src/poly/kernels.cu)
   target_include_directories(poly_expert_parity PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/include
                                                        ${_strata_gpu_include_directories})
   target_link_libraries(poly_expert_parity PRIVATE strata_poly_mmq strata_kernels ggml-cpu ggml-base
                                                   ${_strata_gpu_blas_target} ${_strata_gpu_runtime_target}
                                                   Threads::Threads)
+  # INT8 KV cache attention parity test vs FP16 ground truth
+  add_executable(poly_attn_parity src/poly/attn_int8_test.cpp src/poly/kernels.cu)
+  target_include_directories(poly_attn_parity PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/include
+                                                      ${_strata_gpu_include_directories})
+  target_link_libraries(poly_attn_parity PRIVATE strata_kernels
+                                                ${_strata_gpu_blas_target} ${_strata_gpu_runtime_target}
+                                                Threads::Threads)
   # the pool of the CPU experts, without a model
   add_executable(poly_workers_test src/poly/workers_test.cpp src/poly/cpu_experts.cpp)
   target_include_directories(poly_workers_test PRIVATE ${_strata_gg}/include)

include/strata/hip_compat/cublas_v2.h

@@ -43,6 +43,7 @@ inline rocblas_status cublasGemmEx(rocblas_handle h, rocblas_operation trans_a,
 #define CUBLAS_COMPUTE_32F HIPBLAS_COMPUTE_32F
 #define CUBLAS_DEFAULT_MATH HIPBLAS_DEFAULT_MATH
 #define CUBLAS_GEMM_DEFAULT HIPBLAS_GEMM_DEFAULT
+#define CUBLAS_GEMM_DEFAULT_TENSOR_OP HIPBLAS_GEMM_DEFAULT
 #define CUBLAS_OP_N HIPBLAS_OP_N
 #define CUBLAS_OP_T HIPBLAS_OP_T
 #define CUBLAS_STATUS_SUCCESS HIPBLAS_STATUS_SUCCESS
@@ -52,9 +53,12 @@ inline rocblas_status cublasGemmEx(rocblas_handle h, rocblas_operation trans_a,
 #define cublasCreate hipblasCreate
 #define cublasDestroy hipblasDestroy
 #define cublasGemmEx hipblasGemmEx
+#define cublasGemmStridedBatchedEx hipblasGemmStridedBatchedEx
 #define cublasHandle_t hipblasHandle_t
 #define cublasSetMathMode hipblasSetMathMode
 #define cublasSetStream hipblasSetStream
 #define cublasSetWorkspace hipblasSetWorkspace
+#define cublasSgemm hipblasSgemm
+#define cublasSgemmStridedBatched hipblasSgemmStridedBatched
 #define cublasStatus_t hipblasStatus_t
 #endif

include/strata/hip_compat/cuda_runtime.h

@@ -94,6 +94,7 @@ hipError_t mem_get_info(size_t* free_bytes, size_t* total_bytes);
 #define cudaMemGetInfo hipMemGetInfo
 #endif
 #define cudaMemcpy hipMemcpy
+#define cudaMemcpy2D hipMemcpy2D
 #define cudaMemcpy2DAsync hipMemcpy2DAsync
 #define cudaMemcpyAsync hipMemcpyAsync
 #define cudaMemcpyPeerAsync hipMemcpyPeerAsync
@@ -140,6 +141,7 @@ inline hipError_t cudaGraphInstantiate(hipGraphExec_t* exec, hipGraph_t graph,
                                      hipGraphNode_t* error, char* log, size_t size) {
     return hipGraphInstantiate(exec, graph, error, log, size);
 }
+#define cudaGraphInstantiateWithFlags hipGraphInstantiateWithFlags
 #define __trap() __builtin_trap()   // the compiled-out sm_80 paths (never selected on AMD)
 #define cudaMemcpyToSymbol(symbol, ...) hipMemcpyToSymbol(HIP_SYMBOL(symbol), __VA_ARGS__)
 

include/strata/poly/kernels.hpp

@@ -62,6 +62,12 @@ void kda_step(float* qkv, float* g1, const float* a, const float* beta, float* s
 void store_f16(const float* x, uint16_t* cache, int n, int64_t pos, int rows, void* stream);
 /// The same from the row the card holds at `at` when the launch runs; with `ring`, a row is its number % ring.
 void store_f16_at(const float* x, uint16_t* cache, int n, const long long* at, int rows, int ring, void* stream);
+/// Stores the tokens' keys or values as int8 codes with fp16 group scales (KV_Q8_GROUP = 64) from row `pos`.
+void store_int8(const float* x, int8_t* cache_q, uint16_t* cache_scale, int kv_heads, int head_dim, int64_t pos,
+                int rows, void* stream);
+/// The same from the row the card holds at `at` when the launch runs; with `ring`, a row is its number % ring.
+void store_int8_at(const float* x, int8_t* cache_q, uint16_t* cache_scale, int kv_heads, int head_dim,
+                   const long long* at, int rows, int ring, void* stream);
 /// Attention over cached latents for the tokens from position `pos`.  q: n_head * rank absorbed queries a token.
 /// The first dense_rows tokens attend to every cache row up to their own.  The others each have a list in sel
 /// (sel_stride apart): top_rows cache rows, then the rows of the pool their own position is still filling.
@@ -173,6 +179,17 @@ void attend(const float* q, const uint16_t* kcache, const uint16_t* vcache, int6
 void attend_batch(void* blas, const float* q, const uint16_t* kcache, const uint16_t* vcache, int64_t pos, int rows,
                   float* score, size_t score_floats, float* work, float* out, int dim, int n_head, int kv_heads,
                   float scale, void* stream, int window = 0, int ring = 0, bool whole = false);
+/// attend() with int8 KV cache and group scales (KV_Q8_GROUP = 64).
+void attend_int8(const float* q, const int8_t* kcache_q, const uint16_t* kcache_scale,
+                 const int8_t* vcache_q, const uint16_t* vcache_scale, int64_t pos, int rows, float* score,
+                 int chunk, float* out, int dim, int n_head, int kv_heads, float scale, void* stream, int window = 0,
+                 int ring = 0, bool whole = false, const long long* at = nullptr, int room = 0);
+/// attend_batch() with int8 KV cache and group scales.
+void attend_batch_int8(void* blas, const float* q, const int8_t* kcache_q, const uint16_t* kcache_scale,
+                       const int8_t* vcache_q, const uint16_t* vcache_scale, int64_t pos, int rows,
+                       float* score, size_t score_floats, float* work, uint16_t* kdequant, uint16_t* vdequant,
+                       float* out, int dim, int n_head, int kv_heads, float scale, void* stream, int window = 0,
+                       int ring = 0, bool whole = false);
 /// x = cap * tanh(x / cap).
 void softcap(float* x, size_t n, float cap, void* stream);
 /// The first of the largest of n logits and the softmax's probability of it, summed without the logits more than

setup.py

@@ -5244,9 +5244,7 @@ def poly_setup(a, family, gpu, hip, cuda_tk, multi, ram, data, elsewhere, port,
     config and the start script."""
     fam = FAMILIES[family]
     title = fam["title"]
-    if hip:
-        fail(f"{title} runs on NVIDIA cards only: its engine is a CUDA program",
-             "choose another model for the AMD card")
+    # HIP backend support enabled for AMD ROCm
     if cuda_tk == 12:
         fail(f"{title}'s engine is built for RTX 20 cards and newer: there is none for Pascal / Volta cards or "
              "--cuda 12",