diff --git a/3rd/bmengine/CMakeLists.txt b/3rd/bmengine/CMakeLists.txt
index db3a15e..82d30b4 100644
--- a/3rd/bmengine/CMakeLists.txt
+++ b/3rd/bmengine/CMakeLists.txt
@@ -3,7 +3,11 @@ project(main VERSION 0.1)
 
 enable_language(C)
 enable_language(CXX)
-enable_language(CUDA)
+if(USE_HIP)
+  enable_language(HIP)
+else()
+  enable_language(CUDA)
+endif()
 
 set(CMAKE_CXX_STANDARD 17)
 set(CMAKE_CXX_STANDARD_REQUIRED True)
diff --git a/3rd/bmengine/bmengine/CMakeLists.txt b/3rd/bmengine/bmengine/CMakeLists.txt
index 6511a93..8cea27b 100644
--- a/3rd/bmengine/bmengine/CMakeLists.txt
+++ b/3rd/bmengine/bmengine/CMakeLists.txt
@@ -3,11 +3,17 @@ project(bmengine VERSION 0.1)
 
 enable_language(C)
 enable_language(CXX)
-if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES)
-  set(CMAKE_CUDA_ARCHITECTURES "80;89;90a")
-endif()
-if(NOT APPLE)
-  enable_language(CUDA)
+if(USE_HIP)
+  # enable_language(HIP) honors -DCMAKE_HIP_ARCHITECTURES, otherwise auto-detects the
+  # host GPU(s) and errors if none is found; do not pin a default arch here.
+  enable_language(HIP)
+else()
+  if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES)
+    set(CMAKE_CUDA_ARCHITECTURES "80;89;90a")
+  endif()
+  if(NOT APPLE)
+    enable_language(CUDA)
+  endif()
 endif()
 
 set(CMAKE_CXX_STANDARD 17)
@@ -15,16 +21,34 @@ set(CMAKE_CXX_STANDARD_REQUIRED True)
 set(CMAKE_MODULE_PATH ${PROJECT_SOURCE_DIR}/cmake/)
 
 
-find_library(CUDART_LIBRARY cudart ${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES})
-find_library(CUBLAS_LIBRARY cublas ${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES})
-find_library(CUBLASLT_LIBRARY cublasLt ${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES})
-find_library(CURAND_LIBRARY curand ${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES})
-find_library(CULIBOS_LIBRARY culibos ${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES})
+if(USE_HIP)
+  # cuBLAS/cuBLASLt/cuRAND -> hipBLAS/hipBLASLt/hipRAND (NCCL -> RCCL below).
+  # culibos has no ROCm analogue (its symbols live in the HIP runtime).
+  find_library(HIPBLAS_LIBRARY hipblas HINTS /opt/rocm/lib)
+  find_library(HIPBLASLT_LIBRARY hipblaslt HINTS /opt/rocm/lib)
+  find_library(HIPRAND_LIBRARY hiprand HINTS /opt/rocm/lib)
+  set(CUDART_LIBRARY "")
+  set(CUBLAS_LIBRARY ${HIPBLAS_LIBRARY})
+  set(CUBLASLT_LIBRARY ${HIPBLASLT_LIBRARY})
+  set(CURAND_LIBRARY ${HIPRAND_LIBRARY})
+  set(CULIBOS_LIBRARY "")
+else()
+  find_library(CUDART_LIBRARY cudart ${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES})
+  find_library(CUBLAS_LIBRARY cublas ${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES})
+  find_library(CUBLASLT_LIBRARY cublasLt ${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES})
+  find_library(CURAND_LIBRARY curand ${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES})
+  find_library(CULIBOS_LIBRARY culibos ${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES})
+endif()
 
 # ENABLE_NCCL_TP=on by default.
 if(NOT DEFINED ENV{ENABLE_NCCL_TP} OR "$ENV{ENABLE_NCCL_TP}" STREQUAL "on")
     add_definitions(-DENABLE_NCCL_TP=1)
-    find_package(NCCL REQUIRED)
+    if(USE_HIP)
+        find_library(NCCL_LIBRARIES rccl HINTS /opt/rocm/lib REQUIRED)
+        set(NCCL_INCLUDE_DIRS /opt/rocm/include/rccl /opt/rocm/include)
+    else()
+        find_package(NCCL REQUIRED)
+    endif()
     include_directories(SYSTEM ${NCCL_INCLUDE_DIRS})
 else()
     message(STATUS "Not compiling with NCCL support.")
@@ -44,6 +68,12 @@ set_property(TARGET bmengine PROPERTY POSITION_INDEPENDENT_CODE ON)
 set_property(TARGET bmengine PROPERTY CMAKE_CXX_VISIBILITY_PRESET hidden)
 set_property(TARGET bmengine PROPERTY CMAKE_CUDA_VISIBILITY_PRESET hidden)
 
+if(USE_HIP)
+    set_source_files_properties(${FILES_BMENGINE_CUDA} PROPERTIES LANGUAGE HIP)
+    set_target_properties(bmengine PROPERTIES HIP_ARCHITECTURES "${CMAKE_HIP_ARCHITECTURES}")
+    target_compile_options(bmengine PRIVATE $<$<COMPILE_LANGUAGE:HIP>:--offload-compress>)
+endif()
+
 target_include_directories(bmengine 
     PUBLIC "include"
     PUBLIC "include/private/3rd/"
diff --git a/3rd/bmengine/bmengine/core/engine.cpp b/3rd/bmengine/bmengine/core/engine.cpp
index 1077ec2..f34097a 100644
--- a/3rd/bmengine/bmengine/core/engine.cpp
+++ b/3rd/bmengine/bmengine/core/engine.cpp
@@ -70,13 +70,18 @@ DeviceHandles::DeviceHandles(int dev_id, ncclUniqueId uniqueID, int tp_rank, int
         std::cout << "CC:" << compute_capability
             << ", mp_count:" << mp_count
             << ", L2 Cache:" << (l2_cache_size / 1024 / 1024) << "MB"
+#if !defined(USE_HIP)
             << ", Max Persistent L2:" << (dev_prop.persistingL2CacheMaxSize / 1024) << "KB"
+#endif
             << ", max_smem:" << (max_shared_memory / 1024) << "KB\n";
+#if !defined(USE_HIP)
+    // The persisting-L2-cache device limit is not programmable on CDNA/gfx90a.
     int max_persist_l2 = get_int_env("MAX_PERSIST_L2", 0);
     if (max_persist_l2)  {
         BM_CUDART_ASSERT(
             cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, max_persist_l2 * 1024 * 1024));
     }
+#endif
 }
 DeviceHandles::~DeviceHandles() {
     DeviceGuard guard(dev_id);
diff --git a/3rd/bmengine/bmengine/core/exception.cpp b/3rd/bmengine/bmengine/core/exception.cpp
index 1aa702a..e982c27 100644
--- a/3rd/bmengine/bmengine/core/exception.cpp
+++ b/3rd/bmengine/bmengine/core/exception.cpp
@@ -14,11 +14,14 @@ const char* cublasGetErrorString(cublasStatus_t status) {
         case CUBLAS_STATUS_ALLOC_FAILED: return "CUBLAS_STATUS_ALLOC_FAILED";
         case CUBLAS_STATUS_INVALID_VALUE: return "CUBLAS_STATUS_INVALID_VALUE";
         case CUBLAS_STATUS_ARCH_MISMATCH: return "CUBLAS_STATUS_ARCH_MISMATCH";
-        case CUBLAS_STATUS_MAPPING_ERROR: return "CUBLAS_STATUS_MAPPING_ERROR";
         case CUBLAS_STATUS_EXECUTION_FAILED: return "CUBLAS_STATUS_EXECUTION_FAILED";
-        case CUBLAS_STATUS_INTERNAL_ERROR: return "CUBLAS_STATUS_INTERNAL_ERROR";
         case CUBLAS_STATUS_NOT_SUPPORTED: return "CUBLAS_STATUS_NOT_SUPPORTED";
+#if !defined(USE_HIP)
+        // hipBLAS has no MAPPING/INTERNAL/LICENSE status codes.
+        case CUBLAS_STATUS_MAPPING_ERROR: return "CUBLAS_STATUS_MAPPING_ERROR";
+        case CUBLAS_STATUS_INTERNAL_ERROR: return "CUBLAS_STATUS_INTERNAL_ERROR";
         case CUBLAS_STATUS_LICENSE_ERROR: return "CUBLAS_STATUS_LICENSE_ERROR";
+#endif
         default: return "Unknown CUBLAS error";
     }
 }
diff --git a/3rd/bmengine/bmengine/functions/gemm.cpp b/3rd/bmengine/bmengine/functions/gemm.cpp
index 52ae39e..becf9bc 100644
--- a/3rd/bmengine/bmengine/functions/gemm.cpp
+++ b/3rd/bmengine/bmengine/functions/gemm.cpp
@@ -31,9 +31,11 @@ struct LtGemmAlgoAttr {
     int wave_count;
 };
 
+#if !defined(USE_HIP)
 LtGemmAlgoAttr get_algo_attr(const cublasLtMatmulAlgo_t* algo);
 std::string tile_to_str(int id);
 std::string stage_to_str(int id);
+#endif
 
 // clang-format off
 class Gemm::impl {
@@ -71,8 +73,16 @@ class Gemm::impl {
             compute_type = CUBLAS_COMPUTE_32F;
             scale_type = in_type = out_type = CUDA_R_32F;
         } else if (data_type == core::DataType::kHalf) {
+#if defined(USE_HIP)
+            // hipBLASLt on gfx90a rejects 16F compute/scale (INVALID_VALUE);
+            // use 32F compute + 32F scale with fp16 I/O (matches the bf16 path).
+            compute_type = CUBLAS_COMPUTE_32F;
+            scale_type = CUDA_R_32F;
+            in_type = out_type = CUDA_R_16F;
+#else
             compute_type = CUBLAS_COMPUTE_16F;
             scale_type = in_type = out_type = CUDA_R_16F;
+#endif
         } else if (data_type == core::DataType::kBFloat16) {
             // Only support 32F
             compute_type = CUBLAS_COMPUTE_32F;
@@ -173,6 +183,50 @@ class Gemm::impl {
         }
     }
 
+#if defined(USE_HIP)
+    // hipBLASLt has no cublasLtMatmulAlgoInit / AlgoConfigGetAttribute / tile+
+    // stage enums / search-by-algo-id. Correctness-first: ask hipBLASLt for the
+    // best heuristic algo (zero-workspace preference) and cache it. The cuBLASLt
+    // algo-id filtering and tile/stage debug strings are NVIDIA-only and dropped.
+    void find_algo(
+        const core::Context& ctx,
+        cublasLtMatrixLayout_t layout_A,
+        cublasLtMatrixLayout_t layout_B,
+        cublasLtMatrixLayout_t layout_C,
+        uint32_t M,
+        uint32_t K,
+        uint32_t N,
+        uint32_t batch = 0) {
+        if (algo_id != -1 && (last_m != M || last_batch != batch)) {
+            algo_found = false;
+            hipblasLtMatmulPreference_t preference;
+            BM_CUBLAS_ASSERT(hipblasLtMatmulPreferenceCreate(&preference));
+            uint64_t workspace_size = 0;
+            BM_CUBLAS_ASSERT(hipblasLtMatmulPreferenceSetAttribute(
+                preference, HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES,
+                &workspace_size, sizeof(workspace_size)));
+            cublasLtMatmulDesc_t matmul_desc = create_desc();
+            const int kRequest = 4;
+            hipblasLtMatmulHeuristicResult_t results[kRequest];
+            int nb_result = 0;
+            auto status = hipblasLtMatmulAlgoGetHeuristic(
+                ctx.current_cublas_handle(), matmul_desc,
+                layout_B, layout_A, layout_C, layout_C,
+                preference, kRequest, &results[0], &nb_result);
+            if (status == HIPBLAS_STATUS_SUCCESS && nb_result > 0) {
+                algo = results[0].algo;
+                algo_found = true;
+            } else if (last_m == 0) {
+                std::cerr << prefix << ", N=" << N << ", K=" << K << ", M=" << M
+                          << ", batch=" << batch << ". hipBLASLt heuristic found no algo.\n";
+            }
+            hipblasLtMatmulPreferenceDestroy(preference);
+            hipblasLtMatmulDescDestroy(matmul_desc);
+            last_m = M;
+            last_batch = batch;
+        }
+    }
+#else
     void find_algo(
         const core::Context& ctx,
         cublasLtMatrixLayout_t layout_A,
@@ -254,6 +308,7 @@ class Gemm::impl {
             last_batch = batch;
         }
     }
+#endif // USE_HIP
 
     core::Tensor gemm(
         const core::Context& ctx,
@@ -599,6 +654,7 @@ void Gemm::set_B_scale(const core::Tensor& B_scale) {
     pimpl->b_scale = B_scale.data();
 }
 
+#if !defined(USE_HIP)
 LtGemmAlgoAttr get_algo_attr(const cublasLtMatmulAlgo_t* algo) {
     LtGemmAlgoAttr attr;
     BM_CUBLAS_ASSERT(cublasLtMatmulAlgoConfigGetAttribute(
@@ -683,6 +739,7 @@ std::string stage_to_str(int id) {
     };
     return stage_map.at(static_cast<cublasLtMatmulStages_t>(id)).substr(23);
 }
+#endif // !USE_HIP
 // clang-format on
 
 } // namespace functions
diff --git a/3rd/bmengine/bmengine/functions/sort.cu b/3rd/bmengine/bmengine/functions/sort.cu
index 9542fbf..6e76ea2 100644
--- a/3rd/bmengine/bmengine/functions/sort.cu
+++ b/3rd/bmengine/bmengine/functions/sort.cu
@@ -11,7 +11,34 @@ namespace functions {
 using bmengine::core::Tensor;
 
 // https://nvidia.github.io/cccl/cub/api/structcub_1_1DeviceRadixSort.html
+//
+// rocPRIM's DeviceRadixSort (hipCUB backend) only accepts native radix-sortable
+// key types; half/__hip_bfloat16 keys require a custom decomposer (CUB takes
+// __half directly). The half/bf16 key sort is not on the inference fast path,
+// so on HIP it is left unsupported (runtime error if ever reached) rather than
+// shipping an order-incorrect bit-cast sort.
 template<typename KeyT>
+struct RadixSortable { static constexpr bool value = true; };
+#if defined(USE_HIP)
+template<> struct RadixSortable<half> { static constexpr bool value = false; };
+template<> struct RadixSortable<nv_bfloat16> { static constexpr bool value = false; };
+#endif
+
+template<typename KeyT,
+         typename std::enable_if<!RadixSortable<KeyT>::value, int>::type = 0>
+std::pair<core::Tensor, core::Tensor> sort_1d_temp(
+    const core::Context& ctx,
+    const core::Tensor& keys,
+    const core::Tensor& values,
+    int max_key)
+{
+    BM_EXCEPTION("sort_pair_1d with half/bfloat16 keys is not supported on ROCm/HIP "
+                 "(rocPRIM radix sort needs a decomposer for these key types).");
+    return {};
+}
+
+template<typename KeyT,
+         typename std::enable_if<RadixSortable<KeyT>::value, int>::type = 0>
 std::pair<core::Tensor, core::Tensor> sort_1d_temp(
     const core::Context& ctx,
     const core::Tensor& keys,
diff --git a/3rd/bmengine/bmengine/functions/topk.cu b/3rd/bmengine/bmengine/functions/topk.cu
index 9d02ed0..c970afb 100644
--- a/3rd/bmengine/bmengine/functions/topk.cu
+++ b/3rd/bmengine/bmengine/functions/topk.cu
@@ -5,6 +5,19 @@ namespace bmengine {
 
 namespace functions {
 
+// XOR-permute within the warp/wavefront. On HIP the no-mask form spans the full
+// (32- or 64-lane) wavefront; CUDA keeps the full-warp 32-bit mask. The bitonic
+// groups are width-N (N <= wave size), so lane^j stays inside the wavefront and
+// the permute is correct on wave32 and wave64.
+template<typename T>
+static __device__ inline T warpXorShfl(T v, int laneMask) {
+#if defined(__HIP_PLATFORM_AMD__) || defined(USE_HIP)
+    return __shfl_xor(v, laneMask);
+#else
+    return __shfl_xor_sync(0xFFFFFFFF, v, laneMask);
+#endif
+}
+
 template<typename T, int N>
 static __device__ inline void warpBitonicSort(T& v1, int& pos, bool asc) {
     int lane_id = threadIdx.x & (N - 1);
@@ -13,8 +26,8 @@ static __device__ inline void warpBitonicSort(T& v1, int& pos, bool asc) {
         bool desc = ((lane_id & k) == 0) ^ asc;
 #pragma unroll
         for (int j = k / 2; j > 0; j /= 2) {
-            T v2 = __shfl_xor_sync(0xFFFFFFFF, v1, j);
-            int pos2 = __shfl_xor_sync(0xFFFFFFFF, pos, j);
+            T v2 = warpXorShfl(v1, j);
+            int pos2 = warpXorShfl(pos, j);
             bool upper = (lane_id & j) != 0;
 
             if (desc ^ (v1 > v2 || (v1 == v2 && pos < pos2)) ^ upper) {
@@ -35,8 +48,8 @@ static __device__ inline void warpBitonicMerge(T& v1, int& pos1, T& v2, int& pos
 // resort
 #pragma unroll
     for (int j = N / 2; j > 0; j /= 2) {
-        v2 = __shfl_xor_sync(0xFFFFFFFF, v1, j);
-        int pos2 = __shfl_xor_sync(0xFFFFFFFF, pos1, j);
+        v2 = warpXorShfl(v1, j);
+        int pos2 = warpXorShfl(pos1, j);
         bool upper = (lane_id & j) != 0;
         if ((v1 < v2 || (v1 == v2 && pos1 > pos2)) ^ upper) {
             v1 = v2;
diff --git a/3rd/bmengine/bmengine/include/bmengine/core/utils.h b/3rd/bmengine/bmengine/include/bmengine/core/utils.h
index 6ede722..968e529 100644
--- a/3rd/bmengine/bmengine/include/bmengine/core/utils.h
+++ b/3rd/bmengine/bmengine/include/bmengine/core/utils.h
@@ -48,7 +48,15 @@ inline T ceil_div(T m, Tb d) {
 
 template<typename T>
 inline T round_up_thread(T m) {
+    // Round the block size up to a whole wavefront/warp. On CDNA (gfx90a) the
+    // wavefront is 64, so rounding to 32 would launch a half-filled wavefront;
+    // the full-wavefront shuffles in the warp/block reductions then read
+    // inactive lanes and give wrong amax/amin/etc. Round to 64 on HIP.
+#if defined(USE_HIP)
+    T d = 64, limit = MAX_NUM_THREADS;
+#else
     T d = 32, limit = MAX_NUM_THREADS;
+#endif
     T x = m > limit ? limit : m;
     return ((x + d - 1) / d) * d;
 }
diff --git a/3rd/bmengine/bmengine/include/bmengine/functions/reduce.cuh b/3rd/bmengine/bmengine/include/bmengine/functions/reduce.cuh
index 9afb83c..bdb7bb6 100644
--- a/3rd/bmengine/bmengine/include/bmengine/functions/reduce.cuh
+++ b/3rd/bmengine/bmengine/include/bmengine/functions/reduce.cuh
@@ -5,6 +5,97 @@ namespace bmengine{
 
 namespace functions {
 
+// Wave-size abstraction (C10_WARP_SIZE pattern). NVIDIA warps are 32 lanes;
+// AMD wavefronts are 64 on CDNA (gfx90a/gfx94x) and 32 on RDNA (gfx10xx/11xx).
+// __GFX9__ is defined only during the device compilation pass. The warp/block
+// reductions below are driven entirely by BM_WARP_SIZE, so they are correct on
+// wave32 AND wave64 with no per-arch hack. The block-reduce shared array is
+// sized to the worst case (min warp 32 -> 1024/32 = 32 partials, +1 result).
+#if defined(__HIP_PLATFORM_AMD__) || defined(USE_HIP)
+#if defined(__GFX9__)
+#define BM_WARP_SIZE 64
+#else
+#define BM_WARP_SIZE 32
+#endif
+#else
+#define BM_WARP_SIZE 32
+#endif
+#define BM_BLOCK_REDUCE_MAX_WARPS 33
+
+// Full-wavefront shuffle helpers: on HIP the no-mask __shfl_down/__shfl span the
+// whole (32- or 64-lane) wavefront; on CUDA keep the 0xFFFFFFFF full-warp mask.
+template<typename T>
+__inline__ __device__ T warpShflDown(T x, int offset) {
+#if defined(__HIP_PLATFORM_AMD__) || defined(USE_HIP)
+    return __shfl_down(x, offset);
+#else
+    return __shfl_down_sync(0xFFFFFFFF, x, offset);
+#endif
+}
+template<typename T>
+__inline__ __device__ T warpShfl(T x, int src) {
+#if defined(__HIP_PLATFORM_AMD__) || defined(USE_HIP)
+    return __shfl(x, src);
+#else
+    return __shfl_sync(0xFFFFFFFF, x, src);
+#endif
+}
+
+// Width-W logical-warp shuffle: reduces within fixed W-lane sub-groups (W must
+// divide the wave size). Kernels that tile a 32-lane NVIDIA warp (e.g. a 128-dim
+// head as lane*4) must reduce over exactly 32 lanes on BOTH wave32 and wave64,
+// so they use the Width form rather than the full-wavefront reduce above.
+template<typename T>
+__inline__ __device__ T warpShflDownW(T x, int offset, int width) {
+#if defined(__HIP_PLATFORM_AMD__) || defined(USE_HIP)
+    return __shfl_down(x, offset, width);
+#else
+    return __shfl_down_sync(0xFFFFFFFF, x, offset, width);
+#endif
+}
+template<typename T>
+__inline__ __device__ T warpShflW(T x, int src, int width) {
+#if defined(__HIP_PLATFORM_AMD__) || defined(USE_HIP)
+    return __shfl(x, src, width);
+#else
+    return __shfl_sync(0xFFFFFFFF, x, src, width);
+#endif
+}
+
+template<typename T, int WIDTH>
+__inline__ __device__ T warpReduceSumWidth(T x) {
+#pragma unroll
+    for (int offset = WIDTH / 2; offset > 0; offset /= 2)
+        x += warpShflDownW<T>(x, offset, WIDTH);
+    return x;
+}
+template<typename T, int WIDTH>
+__inline__ __device__ T warpReduceMaxWidth(T x) {
+#pragma unroll
+    for (int offset = WIDTH / 2; offset > 0; offset /= 2) {
+        T y = warpShflDownW<T>(x, offset, WIDTH);
+        x = (x > y) ? x : y;
+    }
+    return x;
+}
+// Width-W max with broadcast of the result to every lane in the W-lane subgroup.
+// Drop-in replacement for warpReduceMaxB in kernels that tile exactly W=32 lanes,
+// keeping the reduction within the intended 32-element group on wave64.
+template<typename T, int WIDTH>
+__inline__ __device__ T warpReduceMaxWidthB(T x) {
+    x = warpReduceMaxWidth<T, WIDTH>(x);
+    return warpShflW<T>(x, 0, WIDTH);  // broadcast within the W-lane subgroup
+}
+// Width-W sum with broadcast of the result to every lane in the W-lane subgroup.
+// Broadcast counterpart of warpReduceMaxWidthB; drop-in for warpReduceSumB in
+// kernels that tile exactly W=32 lanes, keeping the reduction within the
+// intended 32-element group on wave64.
+template<typename T, int WIDTH>
+__inline__ __device__ T warpReduceSumWidthB(T x) {
+    x = warpReduceSumWidth<T, WIDTH>(x);
+    return warpShflW<T>(x, 0, WIDTH);  // broadcast within the W-lane subgroup
+}
+
 template<typename T>
 __inline__ __device__ T threadMax(T a, T b) {
     return (a > b) ? a : b;
@@ -13,16 +104,16 @@ __inline__ __device__ T threadMax(T a, T b) {
 template<typename T>
 __inline__ __device__ T warpReduceMax(T x) {
     #pragma unroll
-    for (int offset = 16; offset > 0; offset /= 2) 
-        x = threadMax<T>(x, __shfl_down_sync(0xFFFFFFFF, x, offset));
+    for (int offset = BM_WARP_SIZE / 2; offset > 0; offset /= 2)
+        x = threadMax<T>(x, warpShflDown<T>(x, offset));
     return x;
 }
 template<typename T>
 __inline__ __device__ T warpReduceMaxB(T x) {
 #pragma unroll
-    for (int offset = 16; offset > 0; offset /= 2)
-        x = threadMax<T>(x, __shfl_down_sync(0xFFFFFFFF, x, offset));
-    return __shfl_sync(0xFFFFFFFF, x, 0);  // broadcast
+    for (int offset = BM_WARP_SIZE / 2; offset > 0; offset /= 2)
+        x = threadMax<T>(x, warpShflDown<T>(x, offset));
+    return warpShfl<T>(x, 0);  // broadcast
 }
 
 template<typename T>
@@ -33,77 +124,92 @@ __inline__ __device__ T threadMin(T a, T b) {
 template<typename T>
 __inline__ __device__ T warpReduceMin(T x) {
 #pragma unroll
-    for (int offset = 16; offset > 0; offset /= 2)
-        x = threadMin<T>(x, __shfl_down_sync(0xFFFFFFFF, x, offset));
+    for (int offset = BM_WARP_SIZE / 2; offset > 0; offset /= 2)
+        x = threadMin<T>(x, warpShflDown<T>(x, offset));
     return x;
 }
 
 template<typename T>
 __inline__ __device__ T warpReduceSum(T x) {
     #pragma unroll
-    for (int offset = 16; offset > 0; offset /= 2) 
-        x += __shfl_down_sync(0xFFFFFFFF, x, offset);
+    for (int offset = BM_WARP_SIZE / 2; offset > 0; offset /= 2)
+        x += warpShflDown<T>(x, offset);
     return x;
 }
 template<typename T>
 __inline__ __device__ T warpReduceSumB(T x) {
     #pragma unroll
-    for (int offset = 16; offset > 0; offset /= 2)
-        x += __shfl_down_sync(0xFFFFFFFF, x, offset);
-    return __shfl_sync(0xFFFFFFFF, x, 0);  // broadcast
+    for (int offset = BM_WARP_SIZE / 2; offset > 0; offset /= 2)
+        x += warpShflDown<T>(x, offset);
+    return warpShfl<T>(x, 0);  // broadcast
 }
 
 template<typename T>
 __inline__ __device__ T blockReduceMax(T x) {
-    static __shared__ T shared[33];
-    int lane = threadIdx.x % 32;
-    int wid = threadIdx.x / 32;
+    static __shared__ T shared[BM_BLOCK_REDUCE_MAX_WARPS];
+    int lane = threadIdx.x % BM_WARP_SIZE;
+    int wid = threadIdx.x / BM_WARP_SIZE;
     x = warpReduceMax<T>(x);
     if (lane == 0) shared[wid] = x;
     __syncthreads();
-    x = (threadIdx.x < blockDim.x / 32) ? shared[lane] : T(-INFINITY);
+#if defined(__HIP_PLATFORM_AMD__) || defined(USE_HIP)
+    int num_warps = (blockDim.x + BM_WARP_SIZE - 1) / BM_WARP_SIZE;
+    x = (lane < num_warps) ? shared[lane] : T(-INFINITY);
+#else
+    x = (threadIdx.x < blockDim.x / BM_WARP_SIZE) ? shared[lane] : T(-INFINITY);
+#endif
     if (wid == 0) {
         x = warpReduceMax<T>(x);
-        if (lane == 0) shared[32] = x;
+        if (lane == 0) shared[BM_BLOCK_REDUCE_MAX_WARPS - 1] = x;
     }
     __syncthreads();
-    return shared[32];  // avoid RAW hazard
+    return shared[BM_BLOCK_REDUCE_MAX_WARPS - 1];  // avoid RAW hazard
 }
 
 template<typename T>
 __inline__ __device__ T blockReduceMin(T x) {
-    static __shared__ T shared[33];
-    int lane = threadIdx.x % 32;
-    int wid = threadIdx.x / 32;
+    static __shared__ T shared[BM_BLOCK_REDUCE_MAX_WARPS];
+    int lane = threadIdx.x % BM_WARP_SIZE;
+    int wid = threadIdx.x / BM_WARP_SIZE;
     x = warpReduceMin<T>(x);
     if (lane == 0)
         shared[wid] = x;
     __syncthreads();
-    x = (threadIdx.x < blockDim.x / 32) ? shared[lane] : T(INFINITY);
+#if defined(__HIP_PLATFORM_AMD__) || defined(USE_HIP)
+    int num_warps = (blockDim.x + BM_WARP_SIZE - 1) / BM_WARP_SIZE;
+    x = (lane < num_warps) ? shared[lane] : T(INFINITY);
+#else
+    x = (threadIdx.x < blockDim.x / BM_WARP_SIZE) ? shared[lane] : T(INFINITY);
+#endif
     if (wid == 0) {
         x = warpReduceMin<T>(x);
         if (lane == 0)
-            shared[32] = x;
+            shared[BM_BLOCK_REDUCE_MAX_WARPS - 1] = x;
     }
     __syncthreads();
-    return shared[32]; // avoid RAW hazard
+    return shared[BM_BLOCK_REDUCE_MAX_WARPS - 1]; // avoid RAW hazard
 }
 
 template<typename T>
 __inline__ __device__ T blockReduceSum(T x) {
-    static __shared__ T shared[33];
-    int lane = threadIdx.x % 32;
-    int wid = threadIdx.x / 32;
+    static __shared__ T shared[BM_BLOCK_REDUCE_MAX_WARPS];
+    int lane = threadIdx.x % BM_WARP_SIZE;
+    int wid = threadIdx.x / BM_WARP_SIZE;
     x = warpReduceSum<T>(x);
     if (lane == 0) shared[wid] = x;
     __syncthreads();
-    x = (threadIdx.x < blockDim.x / 32) ? shared[lane] : T(0.);
+#if defined(__HIP_PLATFORM_AMD__) || defined(USE_HIP)
+    int num_warps = (blockDim.x + BM_WARP_SIZE - 1) / BM_WARP_SIZE;
+    x = (lane < num_warps) ? shared[lane] : T(0.);
+#else
+    x = (threadIdx.x < blockDim.x / BM_WARP_SIZE) ? shared[lane] : T(0.);
+#endif
     if (wid == 0) {
         x = warpReduceSum<T>(x);
-        if (lane == 0) shared[32] = x;
+        if (lane == 0) shared[BM_BLOCK_REDUCE_MAX_WARPS - 1] = x;
     }
     __syncthreads();
-    return shared[32];  // avoid RAW hazard
+    return shared[BM_BLOCK_REDUCE_MAX_WARPS - 1];  // avoid RAW hazard
 }
 
 } // namespace functions
diff --git a/3rd/bmengine/bmengine/include/bmengine/functions/utils.cuh b/3rd/bmengine/bmengine/include/bmengine/functions/utils.cuh
index a5767b2..bf29287 100644
--- a/3rd/bmengine/bmengine/include/bmengine/functions/utils.cuh
+++ b/3rd/bmengine/bmengine/include/bmengine/functions/utils.cuh
@@ -9,8 +9,15 @@ namespace functions {
 
 template<typename T> __host__ __device__ inline T Inf();
 template<> __host__ __device__ inline constexpr float Inf() { return __builtin_huge_valf(); }
+#if defined(USE_HIP)
+// clang rejects a reinterpret_cast in a constexpr; the bit-pattern build is a
+// runtime inline here (same value, identical codegen).
+template<> __host__ __device__ inline half Inf() { const short v = 0x7c00; return *(reinterpret_cast<const half *>(&(v))); }
+template<> __host__ __device__ inline nv_bfloat16 Inf() { const short v = 0x7f80; return *(reinterpret_cast<const nv_bfloat16 *>(&(v))); }
+#else
 template<> __host__ __device__ inline constexpr half Inf() { const short v = 0x7c00; return *(reinterpret_cast<const half *>(&(v))); }
 template<> __host__ __device__ inline constexpr nv_bfloat16 Inf() { const short v = 0x7f80; return *(reinterpret_cast<const nv_bfloat16 *>(&(v))); }
+#endif
 template<> __host__ __device__ inline constexpr double Inf() { return __builtin_huge_val(); }
 template<> __host__ __device__ inline constexpr int32_t Inf() { return __INT_MAX__; }
 template<> __host__ __device__ inline constexpr int64_t Inf() { return __LONG_MAX__; }
diff --git a/CMakeLists.txt b/CMakeLists.txt
index aae5c35..1ed08a5 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -5,14 +5,64 @@ project(zhilight)
 enable_language(C)
 enable_language(CXX)
 
-set(CMAKE_CUDA_ARCHITECTURES $ENV{CMAKE_CUDA_ARCHITECTURES})
-if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES)
-    set(CMAKE_CUDA_ARCHITECTURES "80;89;90a")
+# USE_HIP selects the ROCm/HIP build. Honor an env fallback (ZHILIGHT_USE_HIP)
+# so the value survives CMake's internal configure self-rerun (triggered when
+# find_package(Torch) finalizes the HIP compiler), which can drop a plain
+# command-line -DUSE_HIP. A FORCE cache write keeps it sticky across the rerun.
+if(NOT DEFINED USE_HIP)
+    if(DEFINED ENV{ZHILIGHT_USE_HIP} AND "$ENV{ZHILIGHT_USE_HIP}" STREQUAL "1")
+        set(USE_HIP ON)
+    else()
+        set(USE_HIP OFF)
+    endif()
+endif()
+if(DEFINED ENV{ZHILIGHT_USE_HIP} AND "$ENV{ZHILIGHT_USE_HIP}" STREQUAL "1")
+    set(USE_HIP ON CACHE BOOL "Build with HIP for AMD ROCm GPUs" FORCE)
+else()
+    set(USE_HIP ${USE_HIP} CACHE BOOL "Build with HIP for AMD ROCm GPUs" FORCE)
+endif()
+
+if(USE_HIP)
+    message("CMAKE_HIP_ARCHITECTURES: ${CMAKE_HIP_ARCHITECTURES}")
+    # enable_language(HIP) honors -DCMAKE_HIP_ARCHITECTURES, otherwise auto-detects
+    # the host GPU(s) and errors if none is found (no-GPU build host sets the arch).
+    enable_language(HIP)
+    add_definitions(-DUSE_HIP=1 -DUSE_ROCM=1)
+    set(ZL_HIP_COMPAT_DIR "${PROJECT_SOURCE_DIR}/hip_compat")
+    # -ffp-contract=on matches nvcc/host FMA formation; the
+    # force-include puts the compat shims (cuda*/cublas*/nccl* names) ahead of
+    # any toolkit-named include regardless of per-file order.
+    set(CMAKE_HIP_FLAGS "${CMAKE_HIP_FLAGS} -ffp-contract=on -include ${ZL_HIP_COMPAT_DIR}/zhilight_cuda_to_hip.h")
+    # Put the toolkit-named forwarding shims (cuda_runtime.h, cublasLt.h,
+    # nccl.h, ...) ahead of the real ROCm headers for every target/subdir, so
+    # host .cpp that #include <cublasLt.h> etc. resolve to the HIP mapping.
+    # On a CUDA build this dir is never added, so that path is unchanged.
+    include_directories(BEFORE "${ZL_HIP_COMPAT_DIR}")
+    # Host C++ TUs (compiled by g++, not hipcc) include <hip/hip_runtime.h> via
+    # the compat shim and so need the ROCm headers on their include path; hipcc
+    # adds these implicitly for the .cu TUs.
+    if(NOT DEFINED ROCM_PATH)
+        set(ROCM_PATH "/opt/rocm")
+    endif()
+    include_directories(SYSTEM "${ROCM_PATH}/include")
+else()
+    set(CMAKE_CUDA_ARCHITECTURES $ENV{CMAKE_CUDA_ARCHITECTURES})
+    if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES)
+        set(CMAKE_CUDA_ARCHITECTURES "80;89;90a")
+    endif()
+    message("CMAKE_CUDA_ARCHITECTURES: ${CMAKE_CUDA_ARCHITECTURES}")
+    enable_language(CUDA)
 endif()
-message("CMAKE_CUDA_ARCHITECTURES: ${CMAKE_CUDA_ARCHITECTURES}")
-enable_language(CUDA)
 
-set(CMAKE_CXX_STANDARD 17)
+if(USE_HIP)
+    # ROCm PyTorch 2.13 headers (c10 TensorImpl SetDimsTemplate) use C++20
+    # `requires` concepts, so any TU including torch must compile as C++20.
+    set(CMAKE_CXX_STANDARD 20)
+    set(CMAKE_HIP_STANDARD 20)
+    set(CMAKE_HIP_STANDARD_REQUIRED True)
+else()
+    set(CMAKE_CXX_STANDARD 17)
+endif()
 set(CMAKE_CXX_STANDARD_REQUIRED True)
 set(CMAKE_CUDA_STANDARD 17)
 set(CMAKE_CUDA_STANDARD_REQUIRED True)
@@ -34,10 +84,12 @@ find_package(Python COMPONENTS Interpreter Development REQUIRED)
 set(PYTHON_EXECUTABLE ${Python_EXECUTABLE})
 set(PYTHON_VERSION ${Python_VERSION})
 
-if(CMAKE_CUDA_COMPILER_TOOLKIT_VERSION VERSION_LESS 12.0.0)
-    error("Require CUDA 12")
-else()
-    message("Will link against CUDA 12 complied libraries")
+if(NOT USE_HIP)
+    if(CMAKE_CUDA_COMPILER_TOOLKIT_VERSION VERSION_LESS 12.0.0)
+        error("Require CUDA 12")
+    else()
+        message("Will link against CUDA 12 complied libraries")
+    endif()
 endif()
 
 option(WITH_TESTING "Compile zhilight with unit testing" ON)
@@ -78,9 +130,32 @@ file(GLOB_RECURSE BACKEND_SOURCES
 list(FILTER BACKEND_SOURCES EXCLUDE REGEX "src/py_export/.*\\.cpp$")
 list(FILTER BACKEND_SOURCES EXCLUDE REGEX "src/nn/tests/.*$")
 
+if(USE_HIP)
+    # Tensor-core / PTX quant TUs (marlin mma.sync/ldmatrix/cp.async, awq
+    # ldmatrix, fp8 cvt.e4m3x2) are reimplement-not-port; drop them from the
+    # HIP build. Their call sites are guarded with #if !defined(USE_HIP) and
+    # route to the hipBLASLt dequant+GEMM fallback. The gptq/q_gemm path
+    # (vllm lineage, has USE_ROCM shims) and int8 kernels are kept.
+    list(FILTER BACKEND_SOURCES EXCLUDE REGEX "src/nn/quant/marlin/.*\\.cu$")
+    list(FILTER BACKEND_SOURCES EXCLUDE REGEX "src/nn/quant/awq/.*\\.cu$")
+    list(FILTER BACKEND_SOURCES EXCLUDE REGEX "src/nn/quant/fp8/.*\\.cu$")
+    # gptq kernels use NVIDIA prmt.b32 PTX, __vsub4 SIMD and ldcs-cache PTX
+    # (the in-tree USE_ROCM shims are incomplete vllm-lineage stubs); defer them.
+    list(FILTER BACKEND_SOURCES EXCLUDE REGEX "src/nn/quant/gptq/.*\\.cu$")
+endif()
+
 add_library(backend STATIC ${BACKEND_SOURCES})
 set_property(TARGET backend PROPERTY POSITION_INDEPENDENT_CODE ON)
 
+if(USE_HIP)
+    set(BACKEND_CU_SOURCES ${BACKEND_SOURCES})
+    list(FILTER BACKEND_CU_SOURCES INCLUDE REGEX "\\.cu$")
+    set_source_files_properties(${BACKEND_CU_SOURCES} PROPERTIES LANGUAGE HIP)
+    set_target_properties(backend PROPERTIES HIP_ARCHITECTURES "${CMAKE_HIP_ARCHITECTURES}")
+    target_include_directories(backend BEFORE PUBLIC "${ZL_HIP_COMPAT_DIR}")
+    target_compile_options(backend PRIVATE $<$<COMPILE_LANGUAGE:HIP>:--offload-compress>)
+endif()
+
 set(BOOST_SERIAL_LIB "")
 set(DEEP_GEMM_LIB "")
 set(FLASH_MLA_LIB "")
@@ -97,13 +172,24 @@ if("$ENV{ENABLE_DIST_INFER}" STREQUAL "on")
 endif()
 
 message("CMAKE_SYSTEM_PROCESSOR ${CMAKE_SYSTEM_PROCESSOR}")
-if(CMAKE_SYSTEM_PROCESSOR STREQUAL "x86_64")
+# deep_gemm + flash_mla are prebuilt NVIDIA/Hopper (H20) static libs; the
+# DeepSeek-V3 FP8-block and MLA fast paths they back are deferred on the first
+# AMD pass (MLA falls back to the generic attention path). Keep them x86_64
+# CUDA-only.
+if(CMAKE_SYSTEM_PROCESSOR STREQUAL "x86_64" AND NOT USE_HIP)
     set(DEEP_GEMM_LIB "deep_gemm")
     set(FLASH_MLA_LIB "flash_mla")
     add_definitions(-DENABLE_DS_DEEP_GEMM=1)
     add_definitions(-DENABLE_DS_FLASH_MLA=1)
 endif()
 
+set(FLASH_ATTN_SO "")
+if(USE_HIP)
+    # The flash-attn wheel is a CUDA binary; its ABI/symbols differ on ROCm.
+    # Make discovery non-fatal: prefill uses the in-tree attention kernels
+    # (ENABLE_FLASH_MHA left undefined). ROCm flash-attn can be wired later.
+    message(STATUS "USE_HIP: flash_attn .so dependency disabled (in-tree attention path).")
+else()
 file(GLOB FLASH_ATTN_LIB "${Python_SITELIB}/flash_attn*.so")
 if(FLASH_ATTN_LIB)
     add_definitions(-DENABLE_FLASH_MHA=1)
@@ -126,6 +212,7 @@ else()
 endif()
 get_filename_component(FLASH_ATTN_SO ${FLASH_ATTN_LIB} NAME)
 message("FLASH_ATTN_SO: " ${FLASH_ATTN_SO})
+endif()
 
 target_include_directories(backend
     PUBLIC "src"
@@ -147,17 +234,39 @@ target_link_directories(backend
     PUBLIC ${Python_SITELIB}
 )
 
+set(FLASH_ATTN_LINK "")
+if(FLASH_ATTN_SO)
+    set(FLASH_ATTN_LINK ":${FLASH_ATTN_SO}")
+endif()
+set(ZL_TORCH_LIBS "")
+if(USE_HIP)
+    # The CUDA build pulled torch (c10/TORCH_CHECK) transitively through the
+    # flash-attn .so; with flash-attn dropped on HIP, link torch explicitly so
+    # the C / internals_ modules resolve c10 symbols (flash_decoding.cpp etc.).
+    find_library(ZL_TORCH_LIB     torch     HINTS ${Python_SITELIB}/torch/lib)
+    find_library(ZL_TORCH_CPU_LIB torch_cpu HINTS ${Python_SITELIB}/torch/lib)
+    find_library(ZL_C10_LIB       c10       HINTS ${Python_SITELIB}/torch/lib)
+    set(ZL_TORCH_LIBS ${ZL_TORCH_LIB} ${ZL_TORCH_CPU_LIB} ${ZL_C10_LIB})
+endif()
 target_link_libraries(backend
     PUBLIC bmengine
     ${BOOST_SERIAL_LIB}
-    ":${FLASH_ATTN_SO}"
+    ${FLASH_ATTN_LINK}
     ${DEEP_GEMM_LIB}
     ${FLASH_MLA_LIB}
+    ${ZL_TORCH_LIBS}
 )
 
 file(GLOB_RECURSE PYBIND_SOURCES "src/py_export/*.cpp")
 
-pybind11_add_module(C ${PYBIND_SOURCES})
+if(USE_HIP)
+    # NO_EXTRAS drops pybind11's -flto: the HIP link does not finalize LTO,
+    # leaving the .so as slim bitcode with no PyInit_* (ImportError). (gpuRIR.)
+    pybind11_add_module(C NO_EXTRAS ${PYBIND_SOURCES})
+    set_target_properties(C PROPERTIES INTERPROCEDURAL_OPTIMIZATION OFF)
+else()
+    pybind11_add_module(C ${PYBIND_SOURCES})
+endif()
 
 target_link_libraries(C PRIVATE
     bmengine
diff --git a/README.md b/README.md
index 0741954..f3ad20e 100755
--- a/README.md
+++ b/README.md
@@ -43,6 +43,15 @@ cd ./ZhiLight && pip install -e .
 # Start OpenAI compatible server
 python -m zhilight.server.openai.entrypoints.api_server [options]
 ```
+
+### Build for AMD GPUs (ROCm)
+ZhiLight also builds for AMD GPUs through ROCm/HIP. It requires a ROCm install and a ROCm build of PyTorch; the CUDA libraries map to their ROCm equivalents (cuBLAS/cuBLASLt -> hipBLAS/hipBLASLt, NCCL -> RCCL). Set `USE_HIP=ON` and pass the gfx target for your GPU (for example `gfx90a` for MI200, `gfx1100` for RDNA3).
+```bash
+ROCM_PATH=${ROCM_PATH:-/opt/rocm}
+HIPCXX=$ROCM_PATH/lib/llvm/bin/clang++ ZHILIGHT_USE_HIP=1 \
+  CMAKE_ARGS="-DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a -DCMAKE_HIP_COMPILER=$ROCM_PATH/lib/llvm/bin/clang++ -DCMAKE_HIP_PLATFORM=amd" \
+  CMAKE_BUILD_PARALLEL_LEVEL=32 TESTING=0 python setup.py bdist_wheel
+```
 ## ✈️ Docker Image
 ZhiLight only depends on the CUDA runtime, cuBLAS, NCCL, and a few Python packages in requirements.txt. You can use the image below for running or building it. You can also directly refer to docker/Dockerfile.
 ```bash
diff --git a/hip_compat/cub/cub.cuh b/hip_compat/cub/cub.cuh
new file mode 100644
index 0000000..ed6760f
--- /dev/null
+++ b/hip_compat/cub/cub.cuh
@@ -0,0 +1,6 @@
+#pragma once
+#include "zhilight_cuda_to_hip.h"
+#include <hipcub/hipcub.hpp>
+#ifndef cub
+#define cub hipcub
+#endif
diff --git a/hip_compat/cub/device/device_radix_sort.cuh b/hip_compat/cub/device/device_radix_sort.cuh
new file mode 100644
index 0000000..ed6760f
--- /dev/null
+++ b/hip_compat/cub/device/device_radix_sort.cuh
@@ -0,0 +1,6 @@
+#pragma once
+#include "zhilight_cuda_to_hip.h"
+#include <hipcub/hipcub.hpp>
+#ifndef cub
+#define cub hipcub
+#endif
diff --git a/hip_compat/cublasLt.h b/hip_compat/cublasLt.h
new file mode 100644
index 0000000..bcc2853
--- /dev/null
+++ b/hip_compat/cublasLt.h
@@ -0,0 +1,74 @@
+#pragma once
+// cublasLt -> hipBLASLt name mapping (ZhiLight ROCm/HIP port).
+// The algo-introspection surface (cublasLtMatmulAlgoInit /
+// cublasLtMatmulAlgoConfigGetAttribute / tile+stage enums /
+// cublasLtMatmulSearch_t) has no hipBLASLt equivalent and is handled with
+// USE_HIP guards at the call sites (gemm.cpp), not here.
+#include "zhilight_cuda_to_hip.h"
+#include <hipblas/hipblas.h>
+#include <hipblaslt/hipblaslt.h>
+
+// plain cublas handle is referenced alongside cublasLt in the core headers
+#define cublasHandle_t                hipblasHandle_t
+#define cublasCreate                  hipblasCreate
+#define cublasDestroy                 hipblasDestroy
+#define cublasSetStream               hipblasSetStream
+#define cublasHgemm                   hipblasHgemm
+
+// handles / opaque descriptors
+#define cublasLtHandle_t              hipblasLtHandle_t
+#define cublasLtMatmulDesc_t          hipblasLtMatmulDesc_t
+#define cublasLtMatrixLayout_t        hipblasLtMatrixLayout_t
+#define cublasLtMatmulPreference_t    hipblasLtMatmulPreference_t
+#define cublasLtMatmulAlgo_t          hipblasLtMatmulAlgo_t
+#define cublasLtMatmulHeuristicResult_t hipblasLtMatmulHeuristicResult_t
+#define cublasLtEpilogue_t            hipblasLtEpilogue_t
+
+// scalar enums shared with hipblas
+#define cublasComputeType_t          hipblasComputeType_t
+#define cublasOperation_t            hipblasOperation_t
+#define cublasStatus_t               hipblasStatus_t
+
+#define CUBLAS_COMPUTE_16F           HIPBLAS_COMPUTE_16F
+#define CUBLAS_COMPUTE_32F           HIPBLAS_COMPUTE_32F
+#define CUBLAS_COMPUTE_32I           HIPBLAS_COMPUTE_32I
+#define CUBLAS_OP_N                  HIPBLAS_OP_N
+#define CUBLAS_OP_T                  HIPBLAS_OP_T
+
+#define CUBLAS_STATUS_SUCCESS        HIPBLAS_STATUS_SUCCESS
+#define CUBLAS_STATUS_NOT_INITIALIZED HIPBLAS_STATUS_NOT_INITIALIZED
+#define CUBLAS_STATUS_ALLOC_FAILED   HIPBLAS_STATUS_ALLOC_FAILED
+#define CUBLAS_STATUS_INVALID_VALUE  HIPBLAS_STATUS_INVALID_VALUE
+#define CUBLAS_STATUS_ARCH_MISMATCH  HIPBLAS_STATUS_ARCH_MISMATCH
+#define CUBLAS_STATUS_EXECUTION_FAILED HIPBLAS_STATUS_EXECUTION_FAILED
+#define CUBLAS_STATUS_NOT_SUPPORTED  HIPBLAS_STATUS_NOT_SUPPORTED
+
+// matmul descriptor attributes
+#define CUBLASLT_MATMUL_DESC_TRANSA          HIPBLASLT_MATMUL_DESC_TRANSA
+#define CUBLASLT_MATMUL_DESC_TRANSB          HIPBLASLT_MATMUL_DESC_TRANSB
+#define CUBLASLT_MATMUL_DESC_EPILOGUE        HIPBLASLT_MATMUL_DESC_EPILOGUE
+#define CUBLASLT_MATMUL_DESC_BIAS_POINTER    HIPBLASLT_MATMUL_DESC_BIAS_POINTER
+#define CUBLASLT_MATMUL_DESC_A_SCALE_POINTER HIPBLASLT_MATMUL_DESC_A_SCALE_POINTER
+#define CUBLASLT_MATMUL_DESC_B_SCALE_POINTER HIPBLASLT_MATMUL_DESC_B_SCALE_POINTER
+#define CUBLASLT_EPILOGUE_BIAS               HIPBLASLT_EPILOGUE_BIAS
+#define CUBLASLT_EPILOGUE_DEFAULT            HIPBLASLT_EPILOGUE_DEFAULT
+
+// matrix layout attributes
+#define CUBLASLT_MATRIX_LAYOUT_BATCH_COUNT           HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT
+#define CUBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET  HIPBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET
+
+// preference attributes
+#define CUBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES
+
+// functions (1:1)
+#define cublasLtMatmul                       hipblasLtMatmul
+#define cublasLtMatmulDescCreate             hipblasLtMatmulDescCreate
+#define cublasLtMatmulDescDestroy            hipblasLtMatmulDescDestroy
+#define cublasLtMatmulDescSetAttribute       hipblasLtMatmulDescSetAttribute
+#define cublasLtMatrixLayoutCreate           hipblasLtMatrixLayoutCreate
+#define cublasLtMatrixLayoutDestroy          hipblasLtMatrixLayoutDestroy
+#define cublasLtMatrixLayoutSetAttribute     hipblasLtMatrixLayoutSetAttribute
+#define cublasLtMatmulPreferenceCreate       hipblasLtMatmulPreferenceCreate
+#define cublasLtMatmulPreferenceDestroy      hipblasLtMatmulPreferenceDestroy
+#define cublasLtMatmulPreferenceSetAttribute hipblasLtMatmulPreferenceSetAttribute
+#define cublasLtMatmulAlgoGetHeuristic       hipblasLtMatmulAlgoGetHeuristic
diff --git a/hip_compat/cublas_v2.h b/hip_compat/cublas_v2.h
new file mode 100644
index 0000000..c3d950d
--- /dev/null
+++ b/hip_compat/cublas_v2.h
@@ -0,0 +1,28 @@
+#pragma once
+// cublas (v2) -> hipBLAS name mapping (ZhiLight ROCm/HIP port).
+#include "zhilight_cuda_to_hip.h"
+#include <hipblas/hipblas.h>
+
+#define cublasHandle_t       hipblasHandle_t
+#define cublasStatus_t       hipblasStatus_t
+#define cublasOperation_t    hipblasOperation_t
+#define cublasComputeType_t  hipblasComputeType_t
+
+#define cublasCreate         hipblasCreate
+#define cublasDestroy        hipblasDestroy
+#define cublasSetStream      hipblasSetStream
+#define cublasHgemm          hipblasHgemm
+
+#define CUBLAS_OP_N          HIPBLAS_OP_N
+#define CUBLAS_OP_T          HIPBLAS_OP_T
+#define CUBLAS_COMPUTE_16F   HIPBLAS_COMPUTE_16F
+#define CUBLAS_COMPUTE_32F   HIPBLAS_COMPUTE_32F
+#define CUBLAS_COMPUTE_32I   HIPBLAS_COMPUTE_32I
+
+#define CUBLAS_STATUS_SUCCESS         HIPBLAS_STATUS_SUCCESS
+#define CUBLAS_STATUS_NOT_INITIALIZED HIPBLAS_STATUS_NOT_INITIALIZED
+#define CUBLAS_STATUS_ALLOC_FAILED    HIPBLAS_STATUS_ALLOC_FAILED
+#define CUBLAS_STATUS_INVALID_VALUE   HIPBLAS_STATUS_INVALID_VALUE
+#define CUBLAS_STATUS_ARCH_MISMATCH   HIPBLAS_STATUS_ARCH_MISMATCH
+#define CUBLAS_STATUS_EXECUTION_FAILED HIPBLAS_STATUS_EXECUTION_FAILED
+#define CUBLAS_STATUS_NOT_SUPPORTED   HIPBLAS_STATUS_NOT_SUPPORTED
diff --git a/hip_compat/cuda.h b/hip_compat/cuda.h
new file mode 100644
index 0000000..f637b4d
--- /dev/null
+++ b/hip_compat/cuda.h
@@ -0,0 +1,2 @@
+#pragma once
+#include "zhilight_cuda_to_hip.h"
diff --git a/hip_compat/cuda_bf16.h b/hip_compat/cuda_bf16.h
new file mode 100644
index 0000000..ea38caa
--- /dev/null
+++ b/hip_compat/cuda_bf16.h
@@ -0,0 +1,7 @@
+#pragma once
+#include "zhilight_cuda_to_hip.h"
+// hip_bf16.h is clang-only (clang elementwise builtins); host g++ TUs that
+// include <cuda_bf16.h> never use the bf16 type, so skip it there.
+#if defined(__clang__)
+#include <hip/hip_bf16.h>
+#endif
diff --git a/hip_compat/cuda_bf16.hpp b/hip_compat/cuda_bf16.hpp
new file mode 100644
index 0000000..cac8084
--- /dev/null
+++ b/hip_compat/cuda_bf16.hpp
@@ -0,0 +1,2 @@
+#pragma once
+#include "cuda_bf16.h"
diff --git a/hip_compat/cuda_fp16.h b/hip_compat/cuda_fp16.h
new file mode 100644
index 0000000..4ed6931
--- /dev/null
+++ b/hip_compat/cuda_fp16.h
@@ -0,0 +1,3 @@
+#pragma once
+#include "zhilight_cuda_to_hip.h"
+#include <hip/hip_fp16.h>
diff --git a/hip_compat/cuda_fp16.hpp b/hip_compat/cuda_fp16.hpp
new file mode 100644
index 0000000..81591bc
--- /dev/null
+++ b/hip_compat/cuda_fp16.hpp
@@ -0,0 +1,2 @@
+#pragma once
+#include "cuda_fp16.h"
diff --git a/hip_compat/cuda_runtime.h b/hip_compat/cuda_runtime.h
new file mode 100644
index 0000000..f637b4d
--- /dev/null
+++ b/hip_compat/cuda_runtime.h
@@ -0,0 +1,2 @@
+#pragma once
+#include "zhilight_cuda_to_hip.h"
diff --git a/hip_compat/cuda_runtime_api.h b/hip_compat/cuda_runtime_api.h
new file mode 100644
index 0000000..f637b4d
--- /dev/null
+++ b/hip_compat/cuda_runtime_api.h
@@ -0,0 +1,2 @@
+#pragma once
+#include "zhilight_cuda_to_hip.h"
diff --git a/hip_compat/curand.h b/hip_compat/curand.h
new file mode 100644
index 0000000..5782e72
--- /dev/null
+++ b/hip_compat/curand.h
@@ -0,0 +1,22 @@
+#pragma once
+#include "zhilight_cuda_to_hip.h"
+#include <hiprand/hiprand.h>
+
+#define curandGenerator_t        hiprandGenerator_t
+#define curandStatus_t           hiprandStatus_t
+#define curandRngType_t          hiprandRngType_t
+#define curandOrdering_t         hiprandOrdering_t
+
+#define CURAND_STATUS_SUCCESS    HIPRAND_STATUS_SUCCESS
+#define CURAND_RNG_PSEUDO_MRG       HIPRAND_RNG_PSEUDO_MRG32K3A
+#define CURAND_RNG_PSEUDO_MRG32K3A  HIPRAND_RNG_PSEUDO_MRG32K3A
+#define CURAND_ORDERING_PSEUDO_BEST HIPRAND_ORDERING_PSEUDO_BEST
+
+#define curandCreateGenerator    hiprandCreateGenerator
+#define curandDestroyGenerator   hiprandDestroyGenerator
+#define curandGenerateNormal     hiprandGenerateNormal
+#define curandGenerateUniform    hiprandGenerateUniform
+#define curandSetGeneratorOffset hiprandSetGeneratorOffset
+#define curandSetGeneratorOrdering hiprandSetGeneratorOrdering
+#define curandSetPseudoRandomGeneratorSeed hiprandSetPseudoRandomGeneratorSeed
+#define curandSetStream          hiprandSetStream
diff --git a/hip_compat/curand_kernel.h b/hip_compat/curand_kernel.h
new file mode 100644
index 0000000..082467c
--- /dev/null
+++ b/hip_compat/curand_kernel.h
@@ -0,0 +1,3 @@
+#pragma once
+#include "zhilight_cuda_to_hip.h"
+#include <hiprand/hiprand_kernel.h>
diff --git a/hip_compat/library_types.h b/hip_compat/library_types.h
new file mode 100644
index 0000000..f637b4d
--- /dev/null
+++ b/hip_compat/library_types.h
@@ -0,0 +1,2 @@
+#pragma once
+#include "zhilight_cuda_to_hip.h"
diff --git a/hip_compat/mma.h b/hip_compat/mma.h
new file mode 100644
index 0000000..617561c
--- /dev/null
+++ b/hip_compat/mma.h
@@ -0,0 +1,2 @@
+#pragma once
+// tensor-core wmma path disabled on HIP (deferred); see attention_kernel.cu
diff --git a/hip_compat/nccl.h b/hip_compat/nccl.h
new file mode 100644
index 0000000..5169c43
--- /dev/null
+++ b/hip_compat/nccl.h
@@ -0,0 +1,3 @@
+#pragma once
+#include "zhilight_cuda_to_hip.h"
+#include <rccl/rccl.h>
diff --git a/hip_compat/vector_types.h b/hip_compat/vector_types.h
new file mode 100644
index 0000000..f637b4d
--- /dev/null
+++ b/hip_compat/vector_types.h
@@ -0,0 +1,2 @@
+#pragma once
+#include "zhilight_cuda_to_hip.h"
diff --git a/hip_compat/zhilight_cuda_to_hip.h b/hip_compat/zhilight_cuda_to_hip.h
new file mode 100644
index 0000000..a36ccd7
--- /dev/null
+++ b/hip_compat/zhilight_cuda_to_hip.h
@@ -0,0 +1,153 @@
+// ZhiLight CUDA -> HIP compatibility shim (ROCm/HIP port).
+//
+// Force-included on every HIP translation unit (CMAKE_HIP_FLAGS -include) and
+// pulled in by the toolkit-named forwarding shims (cuda_runtime.h, cublas_v2.h,
+// nccl.h, ...) placed on the HIP-only include path. On a CUDA build this file
+// and the hip_compat/ shim dir are never on the include path, so the NVIDIA
+// build is unchanged.
+//
+// Authored with assistance from Claude (Anthropic).
+#pragma once
+
+#if defined(USE_HIP) || defined(__HIP_PLATFORM_AMD__)
+
+// libc host decls must win over HIP's __device__ memcpy/memset overloads once
+// <hip/hip_runtime.h> is in scope, so include them first.
+#include <cstring>
+#include <cstdlib>
+
+// Host C++ TUs (compiled by g++/clang++, not hipcc) need the AMD platform macro
+// for <hip/hip_runtime.h> to expose the runtime API (hipMalloc, hipStream_t...).
+// hipcc defines it already; defining it twice is harmless.
+#ifndef __HIP_PLATFORM_AMD__
+#define __HIP_PLATFORM_AMD__ 1
+#endif
+
+#include <hip/hip_runtime.h>
+#include <hip/hip_fp16.h>
+// hip_bf16.h uses clang elementwise builtins (__builtin_elementwise_fma etc.)
+// that g++ lacks, so it only parses under the clang/HIP compiler. Host C++ TUs
+// (g++) never use the bf16 type directly; skip it there.
+#if defined(__clang__)
+#include <hip/hip_bf16.h>
+#endif
+
+// ---- per-arch warp size (C10_WARP_SIZE pattern) -------------------------
+// __GFX9__ is defined only in the device compilation pass; host code that
+// needs the value must query hipGetDeviceProperties(...).warpSize at runtime.
+#if defined(__GFX9__)
+#define ZL_WARP_SIZE 64   // CDNA: gfx90a, gfx94x
+#else
+#define ZL_WARP_SIZE 32   // RDNA: gfx10xx, gfx11xx (and CUDA)
+#endif
+
+// ---- bf16 type spellings ------------------------------------------------
+// HIP's hip_bf16.h provides __hip_bfloat16 / __hip_bfloat162 with the full set
+// of arithmetic operators and the __float2bfloat16 / __bfloat162float helpers.
+// Only available under clang (the bf16 header is clang-only, see above).
+#if defined(__clang__)
+using nv_bfloat16 = __hip_bfloat16;
+using nv_bfloat162 = __hip_bfloat162;
+#define __nv_bfloat16 __hip_bfloat16
+#define __nv_bfloat162 __hip_bfloat162
+#endif
+
+// ---- cuda data-type enum (cublas/library_types) -------------------------
+// hipblaslt/hipblas pull in hip_common's hipDataType; mirror the CUDA names.
+#define cudaDataType_t hipDataType
+#define cudaDataType   hipDataType
+#define CUDA_R_32F  HIP_R_32F
+#define CUDA_R_16F  HIP_R_16F
+#define CUDA_R_16BF HIP_R_16BF
+#define CUDA_R_32I  HIP_R_32I
+#define CUDA_R_8I   HIP_R_8I
+#define CUDA_R_8F_E4M3 HIP_R_8F_E4M3
+#define CUDA_R_8F_E5M2 HIP_R_8F_E5M2
+
+// ---- runtime: types -----------------------------------------------------
+#define cudaError_t            hipError_t
+#define cudaSuccess            hipSuccess
+#define cudaErrorPeerAccessAlreadyEnabled hipErrorPeerAccessAlreadyEnabled
+#define cudaStream_t           hipStream_t
+#define cudaEvent_t            hipEvent_t
+#define cudaDeviceProp         hipDeviceProp_t
+#define cudaFuncAttributes     hipFuncAttributes
+
+// ---- runtime: memcpy kinds ----------------------------------------------
+#define cudaMemcpyHostToDevice   hipMemcpyHostToDevice
+#define cudaMemcpyDeviceToHost   hipMemcpyDeviceToHost
+#define cudaMemcpyDeviceToDevice hipMemcpyDeviceToDevice
+#define cudaMemcpyHostToHost     hipMemcpyHostToHost
+
+// ---- runtime: functions -------------------------------------------------
+#define cudaMalloc             hipMalloc
+#define cudaFree               hipFree
+#define cudaMallocHost         hipHostMalloc
+#define cudaHostAlloc          hipHostMalloc
+#define cudaFreeHost           hipHostFree
+#define cudaMemcpy             hipMemcpy
+#define cudaMemcpyAsync        hipMemcpyAsync
+#define cudaMemcpyPeer         hipMemcpyPeer
+#define cudaMemsetAsync        hipMemsetAsync
+#define cudaMemGetInfo         hipMemGetInfo
+#define cudaGetDevice          hipGetDevice
+#define cudaSetDevice          hipSetDevice
+#define cudaGetDeviceCount     hipGetDeviceCount
+#define cudaGetDeviceProperties hipGetDeviceProperties
+#define cudaDeviceSynchronize  hipDeviceSynchronize
+#define cudaDeviceCanAccessPeer hipDeviceCanAccessPeer
+#define cudaDeviceEnablePeerAccess hipDeviceEnablePeerAccess
+#define cudaGetLastError       hipGetLastError
+#define cudaGetErrorString     hipGetErrorString
+#define cudaDeviceGetAttribute hipDeviceGetAttribute
+
+// ---- runtime: device attributes -----------------------------------------
+#define cudaDevAttrMaxSharedMemoryPerBlockOptin hipDeviceAttributeSharedMemPerBlockOptin
+
+// ---- runtime: streams ---------------------------------------------------
+#define cudaStreamCreate            hipStreamCreate
+#define cudaStreamCreateWithFlags   hipStreamCreateWithFlags
+#define cudaStreamCreateWithPriority hipStreamCreateWithPriority
+#define cudaStreamDestroy           hipStreamDestroy
+#define cudaStreamSynchronize       hipStreamSynchronize
+#define cudaStreamWaitEvent         hipStreamWaitEvent
+#define cudaStreamNonBlocking       hipStreamNonBlocking
+
+// ---- runtime: events ----------------------------------------------------
+#define cudaEventCreate          hipEventCreate
+#define cudaEventCreateWithFlags hipEventCreateWithFlags
+#define cudaEventDestroy         hipEventDestroy
+#define cudaEventRecord          hipEventRecord
+#define cudaEventQuery           hipEventQuery
+#define cudaEventSynchronize     hipEventSynchronize
+#define cudaEventElapsedTime     hipEventElapsedTime
+#define cudaEventDefault         hipEventDefault
+#define cudaEventDisableTiming   hipEventDisableTiming
+
+// ---- runtime: kernel func attributes ------------------------------------
+// HIP takes the kernel entry as `const void*`; CUDA takes the typed function
+// pointer. Wrap so call sites pass a typed kernel pointer unchanged.
+#define cudaFuncAttributeMaxDynamicSharedMemorySize hipFuncAttributeMaxDynamicSharedMemorySize
+#define cudaFuncSetAttribute(func, attr, value) \
+    hipFuncSetAttribute(reinterpret_cast<const void*>(func), (attr), (value))
+#define cudaFuncGetAttributes(attrptr, func) \
+    hipFuncGetAttributes((attrptr), reinterpret_cast<const void*>(func))
+
+// ---- device cache-hint load ---------------------------------------------
+// __ldcs (streaming/last-use cache hint) has no HIP equivalent; map to the
+// read-only __ldg load (same result, just a different cache hint). Used in
+// device code only, so this is inert in host TUs.
+#define __ldcs __ldg
+
+// ---- kernel-parameter qualifier -----------------------------------------
+// __grid_constant__ (CUDA sm70+ by-value kernel-param hint) has no HIP keyword;
+// it is purely an optimization, so drop it on AMD.
+#define __grid_constant__
+
+// ---- L2 access-policy window: unsupported on CDNA, stub out -------------
+// gfx90a has no programmable L2 persisting-cache window. The only caller
+// (setL2AccessPolicyWindow) is already commented out at its call sites, and
+// the device-limit set is best-effort, so map these to no-ops / drop them.
+#define ZL_HIP_NO_L2_ACCESS_POLICY 1
+
+#endif // USE_HIP
diff --git a/src/generator/random_util.cu b/src/generator/random_util.cu
index f9c1a5e..646d8e0 100644
--- a/src/generator/random_util.cu
+++ b/src/generator/random_util.cu
@@ -11,6 +11,34 @@
 #include <cub/cub.cuh>
 
 namespace beam_utility {
+
+// rocPRIM (hipCUB) radix sort lacks a built-in codec for __hip_bfloat16 keys
+// (half and float are fine). bf16 top-p sampling is not on the validated path,
+// so on HIP its cub::DeviceRadixSort instantiation is skipped (would fail to
+// compile via rocPRIM's identity_decomposer assertion) and the call throws at
+// runtime. CUDA keeps the original cub call for every dtype.
+template<typename KeyT>
+struct HipBf16Key { static constexpr bool value = false; };
+#if defined(USE_HIP)
+template<> struct HipBf16Key<nv_bfloat16> { static constexpr bool value = true; };
+#endif
+
+template<typename KeyT, typename ValueT>
+static inline hipError_t radixSortPairsDescending(
+    void* d_temp, size_t& temp_bytes,
+    const KeyT* keys_in, KeyT* keys_out,
+    const ValueT* vals_in, ValueT* vals_out,
+    int num_items, int begin_bit, int end_bit, hipStream_t stream) {
+    if constexpr (HipBf16Key<KeyT>::value) {
+        BM_EXCEPTION("top-p sampling with bfloat16 probs is not supported on ROCm/HIP "
+                     "(rocPRIM radix sort has no bf16 key codec). Use float/half probs.");
+        return hipSuccess;
+    } else {
+        return cub::DeviceRadixSort::SortPairsDescending(
+            d_temp, temp_bytes, keys_in, keys_out, vals_in, vals_out,
+            num_items, begin_bit, end_bit, stream);
+    }
+}
 // gridDim (n / 1024, 1, 1),    blockDim(1024, 1, 1)
 template<typename T>
 static __global__ void BM_KERNEL(random_repetition_penalty)(
@@ -139,14 +167,14 @@ void random_sampler_gpu(
 
     BM_DTYPE_DISPATCH_FLOAT(probs.dtype(), {
         size_t temp_buffer_size1, temp_buffer_size2;
-        cub::DeviceRadixSort::SortPairsDescending(
-            nullptr,
+        radixSortPairsDescending(
+            (void*) nullptr,
             temp_buffer_size1,
             values_out.data<scalar_t>(),
             values_out.data<scalar_t>(),
             indicies_in.data<int32_t>(),
             indicies_out.data<int32_t>(),
-            n_classes);
+            n_classes, 0, sizeof(scalar_t) * 8, (hipStream_t) 0);
         cub::DeviceScan::InclusiveSum(
             nullptr,
             temp_buffer_size2,
@@ -161,7 +189,7 @@ void random_sampler_gpu(
         CURAND_CHECK(curandGenerateUniform(gen, p_random.data<float>(), p_random.size(0)));
         for (int i = 0; i < batch; i++) {
             scalar_t* offset_prob = (probs.data<scalar_t>() + i * n_classes);
-            BM_CUDART_ASSERT(cub::DeviceRadixSort::SortPairsDescending(
+            BM_CUDART_ASSERT(radixSortPairsDescending(
                 temp.data(),
                 temp_buffer_size,
                 offset_prob,
diff --git a/src/nn/attention/attention_kernel.cu b/src/nn/attention/attention_kernel.cu
index 47ab8d3..2f9cabc 100644
--- a/src/nn/attention/attention_kernel.cu
+++ b/src/nn/attention/attention_kernel.cu
@@ -23,6 +23,28 @@ using bmengine::core::DataType;
 using bmengine::core::Tensor;
 
 #define WARP_SIZE 32
+
+// Width-32 raw shuffles for the scalar attention paths, which tile a 32-lane
+// NVIDIA warp (128-dim head as lane*4). On HIP the width argument confines the
+// permute to 32-lane sub-groups so it is correct on wave64; on CUDA it is the
+// usual masked warp shuffle.
+template<typename T>
+static __inline__ __device__ T attnShflDown32(T x, int offset) {
+#if defined(__HIP_PLATFORM_AMD__) || defined(USE_HIP)
+    return __shfl_down(x, offset, 32);
+#else
+    return __shfl_down_sync(0xFFFFFFFF, x, offset, 32);
+#endif
+}
+template<typename T>
+static __inline__ __device__ T attnShfl32(T x, int src) {
+#if defined(__HIP_PLATFORM_AMD__) || defined(USE_HIP)
+    return __shfl(x, src, 32);
+#else
+    return __shfl_sync(0xFFFFFFFF, x, src, 32);
+#endif
+}
+
 template<typename T, typename T2 = T, int DIM_HEAD = 128>
 static __inline__ __device__ void multiply_q_k_block(
     const T* __restrict__ g_q, // (dim_head)
@@ -43,7 +65,7 @@ static __inline__ __device__ void multiply_q_k_block(
         res += float(*reinterpret_cast<T*>(&a4.y)) * float(*reinterpret_cast<T*>(&b4.y));
         res += float(*reinterpret_cast<T*>(&a4.z)) * float(*reinterpret_cast<T*>(&b4.z));
         res += float(*reinterpret_cast<T*>(&a4.w)) * float(*reinterpret_cast<T*>(&b4.w));
-        res = functions::warpReduceSum<float>(res);
+        res = functions::warpReduceSumWidth<float, WARP_SIZE>(res);
         if (laneId == 0)
             logit[col] = T2(res);
     }
@@ -93,7 +115,7 @@ static __inline__ __device__ void multiply_q_k_block_d64(
         short2 b2 = *reinterpret_cast<const short2*>(g_k + col * stride_k + laneId * 2);
         res += float(*reinterpret_cast<T*>(&a2.x)) * float(*reinterpret_cast<T*>(&b2.x));
         res += float(*reinterpret_cast<T*>(&a2.y)) * float(*reinterpret_cast<T*>(&b2.y));
-        res = functions::warpReduceSum<float>(res);
+        res = functions::warpReduceSumWidth<float, WARP_SIZE>(res);
         if (laneId == 0)
             logit[col] = T2(res);
     }
@@ -137,7 +159,7 @@ static __inline__ __device__ void multiply_q_k_block_mq1(
             res += float(*reinterpret_cast<T*>(&a4.y)) * float(*reinterpret_cast<T*>(&b4.y));
             res += float(*reinterpret_cast<T*>(&a4.z)) * float(*reinterpret_cast<T*>(&b4.z));
             res += float(*reinterpret_cast<T*>(&a4.w)) * float(*reinterpret_cast<T*>(&b4.w));
-            res = functions::warpReduceSum<float>(res);
+            res = functions::warpReduceSumWidth<float, WARP_SIZE>(res);
             if (laneId == 0)
                 logit[q * len_buf + col] = T2(res);
         }
@@ -190,12 +212,12 @@ static __device__ void multiply_score_v_block2(
     __syncthreads();
     float x = tmp[threadIdx.x];
     if (NUM_SPLIT == 16)
-        x += __shfl_down_sync(0xFFFFFFFF, x, 8);
+        x += attnShflDown32(x, 8);
     if (NUM_SPLIT >= 8)
-        x += __shfl_down_sync(0xFFFFFFFF, x, 4);
+        x += attnShflDown32(x, 4);
     if (NUM_SPLIT >= 4)
-        x += __shfl_down_sync(0xFFFFFFFF, x, 2);
-    x += __shfl_down_sync(0xFFFFFFFF, x, 1);
+        x += attnShflDown32(x, 2);
+    x += attnShflDown32(x, 1);
     if ((threadIdx.x % NUM_SPLIT) == 0) {
         output[threadIdx.x / NUM_SPLIT] = T2(x);
     }
@@ -230,6 +252,11 @@ static __device__ void multiply_score_v_block_mq1(
     }
 }
 
+// Tensor-core (wmma / nvcuda) attention path: NVIDIA-only, reimplement-not-port.
+// Disabled on the first AMD pass; the dispatch below routes to the portable
+// scalar multiply_q_k_block / multiply_score_v_block path under USE_HIP. An
+// AMD-native rocWMMA/MFMA decode path is a deferred follow-up.
+#if !defined(USE_HIP)
 using namespace nvcuda;
 
 template<typename T, int M_Q = 8, int N_BUF = 32, int K_DIM = 16, int DIM_HEAD = 128>
@@ -390,45 +417,58 @@ static __forceinline__ __device__ void multiply_score_v_block_wmma2(
     if (M_Q == 16)
         output[q_head_offset + threadIdx.x] = T(v1);
 }
+#endif // !USE_HIP  (tensor-core wmma path)
 
-// use outer shared memory
+// use outer shared memory. Wave-size driven (BM_WARP_SIZE = 64 on CDNA, 32 on
+// RDNA/CUDA) so the per-warp partials match functions::warpReduce*; shared[]
+// here must be sized >= num_warps+1 by the caller (max 1024/32 + 1 = 33).
 template<typename T>
 __inline__ __device__ T blockReduceMax(T x, T* shared) {
-    int lane = threadIdx.x % 32;
-    int wid = threadIdx.x / 32;
+    int lane = threadIdx.x % BM_WARP_SIZE;
+    int wid = threadIdx.x / BM_WARP_SIZE;
     x = functions::warpReduceMax<T>(x);
     if (lane == 0)
         shared[wid] = x;
     __syncthreads();
 
     if (wid == 0) {
-        x = (threadIdx.x < blockDim.x / 32) ? shared[lane] : -functions::Inf<T>();
+#if defined(__HIP_PLATFORM_AMD__) || defined(USE_HIP)
+        int num_warps = (blockDim.x + BM_WARP_SIZE - 1) / BM_WARP_SIZE;
+        x = (lane < num_warps) ? shared[lane] : -functions::Inf<T>();
+#else
+        x = (threadIdx.x < blockDim.x / BM_WARP_SIZE) ? shared[lane] : -functions::Inf<T>();
+#endif
         x = functions::warpReduceMax<T>(x);
         if (lane == 0)
-            shared[32] = x;
+            shared[BM_BLOCK_REDUCE_MAX_WARPS - 1] = x;
     }
     __syncthreads();
-    return shared[32]; // avoid RAW hazard
+    return shared[BM_BLOCK_REDUCE_MAX_WARPS - 1]; // avoid RAW hazard
 }
 
 // use outer shared memory
 template<typename T>
 __inline__ __device__ T blockReduceSum(T x, T* shared) {
-    int lane = threadIdx.x % 32;
-    int wid = threadIdx.x / 32;
+    int lane = threadIdx.x % BM_WARP_SIZE;
+    int wid = threadIdx.x / BM_WARP_SIZE;
     x = functions::warpReduceSum<T>(x);
     if (lane == 0)
         shared[wid] = x;
     __syncthreads();
 
     if (wid == 0) {
-        x = (threadIdx.x < blockDim.x / 32) ? shared[lane] : T(0.);
+#if defined(__HIP_PLATFORM_AMD__) || defined(USE_HIP)
+        int num_warps = (blockDim.x + BM_WARP_SIZE - 1) / BM_WARP_SIZE;
+        x = (lane < num_warps) ? shared[lane] : T(0.);
+#else
+        x = (threadIdx.x < blockDim.x / BM_WARP_SIZE) ? shared[lane] : T(0.);
+#endif
         x = functions::warpReduceSum<T>(x);
         if (lane == 0)
-            shared[32] = x;
+            shared[BM_BLOCK_REDUCE_MAX_WARPS - 1] = x;
     }
     __syncthreads();
-    return shared[32]; // avoid RAW hazard
+    return shared[BM_BLOCK_REDUCE_MAX_WARPS - 1]; // avoid RAW hazard
 }
 
 template<typename T>
@@ -800,6 +840,8 @@ static __global__ void KERNEL_mqa_rag_buffer_split_kv(
 }
 
 // gridDim (len_q * split, num_kv_heads * m_query, batch),  blockDim (dim_head)
+// int8-quantized-KV attention; NVIDIA-PTX dequant, deferred on the first AMD pass.
+#if !defined(USE_HIP)
 template<typename T, int DIM_HEAD = 128, int O_DIM_HEAD = DIM_HEAD>
 static __global__ void KERNEL_mqa_rag_buffer_split_kv_quant(
     const half* __restrict__ g_q,      // (batch, len_q, num_kv_heads * m_query, dim_head）
@@ -876,6 +918,7 @@ static __global__ void KERNEL_mqa_rag_buffer_split_kv_quant(
         DEV_mul_score_v_v1<T, O_DIM_HEAD, float>(smem, val_buf, cache, len_buf2, stride_kv);
     }
 }
+#endif // !USE_HIP  (int8-quantized-KV attention kernel)
 
 // gridDim (num_virtual_heads / batch, batch),  blockDim (dim_head)
 template<typename T>
@@ -897,12 +940,12 @@ static __global__ void KERNEL_mqa_combine(
         local_max = g_local_max[h * num_split + laneId];
         local_sum_exp = g_local_sum_exp[h * num_split + laneId];
     }
-    float global_max = functions::warpReduceMax<float>(local_max);
-    global_max = __shfl_sync(0xFFFFFFFF, global_max, 0);
+    float global_max = functions::warpReduceMaxWidth<float, WARP_SIZE>(local_max);
+    global_max = attnShfl32(global_max, 0);
     float scale1 = laneId < num_split ? expf(local_max - global_max) : 0.;
     float local_sum_exp2 = local_sum_exp * scale1;
-    float global_sum_exp = functions::warpReduceSum<float>(local_sum_exp2);
-    global_sum_exp = __shfl_sync(0xFFFFFFFF, global_sum_exp, 0);
+    float global_sum_exp = functions::warpReduceSumWidth<float, WARP_SIZE>(local_sum_exp2);
+    global_sum_exp = attnShfl32(global_sum_exp, 0);
     float scale2 = local_sum_exp / global_sum_exp * scale1;
     //    if (threadIdx.x == 1 && h == 1)
     //        printf("h=%d local_max=%f, global_max=%f, local_sum_exp2=%f, global_sum_exp=%f\n", h,
@@ -914,7 +957,7 @@ static __global__ void KERNEL_mqa_combine(
     for (int i = 0; i < num_split; ++i) {
         // float v = cache[(h * num_split + i) * DIM_HEAD + threadIdx.x];
         float v = cache[i * DIM_HEAD];
-        float scale = __shfl_sync(0xFFFFFFFF, scale2, i);
+        float scale = attnShfl32(scale2, i);
         //        if (threadIdx.x == 0 && h == 1)
         //            printf("h=%d v=%f, scale=%f\n", h, v, scale);
         res += v * scale;
@@ -1054,10 +1097,15 @@ static __global__ void KERNEL_mqa_self(
     const unsigned int len_buf_cut = (q + 32) / 32 * 32; // little speed up after cut
     // const unsigned int len_buf_cut = len_buf;
     // Q * K  (m_query, dim_head) * (len_buf, dim_head) => (m_query, len_buf)
+#if defined(USE_HIP)
+    // wmma tensor-core path disabled on AMD; always use the scalar path.
+    multiply_q_k_block_mq1(s_q, g_k, score, len_buf_cut);
+#else
     if (high_precision)
         multiply_q_k_block_mq1(s_q, g_k, score, len_buf_cut);
     else
         multiply_q_k_block_wmma<T>(s_q, g_k, score, len_buf_cut);
+#endif
     __syncthreads();
 
     // Softmax
@@ -1070,6 +1118,9 @@ static __global__ void KERNEL_mqa_self(
     __syncthreads();
 
     // Score * V (m_query, len_buf) * (len_buf, dim_head) => (m_query, dim_head)
+#if defined(USE_HIP)
+    multiply_score_v_block_mq1(score, g_v, output, len_buf_cut);
+#else
     if (high_precision > 0) {
         multiply_score_v_block_mq1(score, g_v, output, len_buf_cut);
     } else {
@@ -1084,6 +1135,7 @@ static __global__ void KERNEL_mqa_self(
                 score_t, g_v, output, len_buf_cut, reinterpret_cast<float*>(s_q));
         }
     }
+#endif
 }
 
 void mul_qk_rag_buffer(
@@ -1301,7 +1353,13 @@ void multi_query_attention_rag_buffer(
     }
     // std::cout << "max_len_buf=" << max_len_buf << ", num_split=" << num_split << ", split_kv_thres=" << split_kv_thres << endl;
 
+#if defined(USE_HIP)
+    // The wmma (tensor-core) MQA path is disabled on AMD; force the portable
+    // scalar KERNEL_mqa_rag_buffer1 path.
+    static bool use_mma = false;
+#else
     static bool use_mma = utils::get_int_env("CPM_MQ_ATTN_MMA", 0) && m_query == 8;
+#endif
     if (algo_id == -1 || batch_q.dtype() == DataType::kBFloat16)
         algo_id = 1; // default to 1
     if (num_split > 1 && algo_id == 1) {
@@ -1314,6 +1372,10 @@ void multi_query_attention_rag_buffer(
 
         auto out_dtype = batch_q.dtype();
         if (scale_key_addrs.numel()) {
+#if defined(USE_HIP)
+            BM_EXCEPTION("int8-quantized KV-cache attention is not supported on ROCm/HIP "
+                         "(NVIDIA-PTX dequant path deferred). Use an unquantized KV cache.");
+#else
             BM_ASSERT(dim_head == 128, "Unsupported dim_head");
             BM_ASSERT_EQ(batch_q.dtype(), core::DataType::kHalf, "input must be half");
             out_dtype = dequant_dtype;
@@ -1339,6 +1401,7 @@ void multi_query_attention_rag_buffer(
                 );
             });
             BM_CUDART_ASSERT(cudaGetLastError());
+#endif
         } else {
             BM_DTYPE_DISPATCH_HALF(out_dtype, {
                 auto kernel = KERNEL_mqa_rag_buffer_split_kv<scalar_t, 128>;
@@ -1485,8 +1548,9 @@ void multi_query_self_attention(
         dim3 gridDim(len_q, num_kv_heads);
         size_t dynamic_size = 8 * (dim_head + std::max(2048, len_buf)) * sizeof(float) + 32;
 
+        auto mqa_self_kernel = KERNEL_mqa_self<__half, 128, 8>;
         BM_CUDART_ASSERT(cudaFuncSetAttribute(
-            KERNEL_mqa_self<__half, 128, 8>,
+            mqa_self_kernel,
             cudaFuncAttributeMaxDynamicSharedMemorySize,
             dynamic_size));
         KERNEL_mqa_self<__half, 128, 8><<<gridDim, 1024, dynamic_size, stream>>>(
diff --git a/src/nn/attention/flash_decoding.cpp b/src/nn/attention/flash_decoding.cpp
index 45e0569..36b8f55 100644
--- a/src/nn/attention/flash_decoding.cpp
+++ b/src/nn/attention/flash_decoding.cpp
@@ -16,6 +16,17 @@ using bmengine::functions::transpose_2_1;
 using namespace flash;
 #endif
 
+#if defined(USE_HIP)
+// The flash-attn .so is a CUDA binary not linked on ROCm (ABI/symbols differ);
+// its run_mha_fwd entry is unavailable. The in-tree attention path covers
+// prefill (ENABLE_FLASH_MHA is undefined). Any call into the flash-attn prefill
+// here is a configuration error on AMD, so fail loudly instead of referencing
+// the missing symbol. ROCm flash-attn is a deferred follow-up.
+#define run_mha_fwd(...) \
+    throw std::runtime_error("flash-attn prefill is not available on ROCm/HIP (deferred); " \
+                             "use the in-tree attention path")
+#endif
+
 FlashDecoding::FlashDecoding(const Context& ctx) {
     BM_CUDART_ASSERT(cudaGetDeviceProperties(&dprops, ctx.active_device()));
     params = new Flash_fwd_params();
diff --git a/src/nn/attention/quant_attention.cuh b/src/nn/attention/quant_attention.cuh
index 7ab534b..7ecff8e 100644
--- a/src/nn/attention/quant_attention.cuh
+++ b/src/nn/attention/quant_attention.cuh
@@ -7,6 +7,12 @@
 #include <vector_types.h>
 #include <bmengine/functions/reduce.cuh>
 
+// int8-quantized-KV attention: the dequant uses NVIDIA prmt.b32 / sub.f16x2
+// PTX (reimplement-not-port). Disabled on the first AMD pass; the int8 KV-cache
+// attention kernel and its launcher are gated off under USE_HIP and route to an
+// explicit unsupported error. An AMD-native int8-KV path is a deferred follow-up.
+#if !defined(USE_HIP)
+
 static __inline__ __device__ void DEV_dequant_int8(
     uint32_t const i8s, half2 (&result)[2]
 ) {
@@ -121,3 +127,5 @@ static __device__ void DEV_mul_score_v_v1(
         output[threadIdx.x / NUM_SPLIT] = T2(x);
     }
 }
+
+#endif // !USE_HIP  (int8-quantized-KV attention PTX path)
diff --git a/src/nn/feedforward/feedforward.cpp b/src/nn/feedforward/feedforward.cpp
index 0369596..8d08924 100644
--- a/src/nn/feedforward/feedforward.cpp
+++ b/src/nn/feedforward/feedforward.cpp
@@ -97,12 +97,18 @@ class FeedForward::impl::NormalImpl : public FeedForward::impl {
     }
 
     Tensor gptq_fused_up(const core::Context& ctx, const Tensor& input) {
+#if defined(USE_HIP)
+        // GPTQ fused gate-up uses the deferred gptq kernels; never reached on
+        // AMD (support_fuse_gptq_gate_in returns false here).
+        BM_EXCEPTION("GPTQ fused gate-up is not supported on ROCm/HIP (deferred).");
+#else
         // Fuse 'w_in' and 'w_gated' and activation into a kernel.
         // i.e. activate(input X 'w_in') * (input * 'w_gated')
         auto[qw1, qz1, scales1, sym1] = w_in.get_gptq_weights();
         auto[qw2, qz2, scales2, sym2] = w_gated.get_gptq_weights();
         return nn::gptq::gemm_fuse_gate_in(
             ctx, input, qw1, qz1, scales1, Tensor(), qw2, qz2, scales2, Tensor(), sym1);
+#endif
     }
 
     void try_fuse_up_weights(const core::Context& ctx) {
@@ -868,6 +874,9 @@ class FeedForward::impl::FusedMOE : public FeedForward::impl::MOEImpl {
 
 };
 
+// GPTQ-fused and fp8-block MOE experts use the deferred gptq / fp8 quant
+// kernels (excluded from the HIP build); their factory branches throw on AMD.
+#if !defined(USE_HIP)
 class FeedForward::impl::GPTQMOE : public FeedForward::impl::FusedMOE {
 public:
     GPTQMOE(
@@ -1239,6 +1248,7 @@ class FeedForward::impl::FP8BlockMOE : public FeedForward::impl::FusedMOE {
         return with_share(ctx, input, ret);
     }
 };
+#endif // !USE_HIP  (GPTQ / fp8-block MOE experts)
 
 FeedForward::FeedForward(
     const core::Context& ctx,
@@ -1261,6 +1271,15 @@ FeedForward::FeedForward(
 //                << ", dyn_shared=" << dyn_shared << endl;
 //        }
         impl::MOEImpl* ptr;
+#if defined(USE_HIP)
+        // GPTQ-fused and fp8-block MOE use deferred quant kernels on AMD; the
+        // generic MOEImpl path (fp16/bf16) is used instead.
+        if (grouped_fp8_gemm || (fuse_moe && gptq_kernel_algo == 1)) {
+            BM_EXCEPTION("FUSE_GPTQ_MOE / GROUPED_FP8_GEMM MOE is not supported on ROCm/HIP "
+                         "(deferred); unset these env vars to use the generic MOE path.");
+        }
+        ptr = new impl::MOEImpl(ctx, cfg, quant_config, parallel);
+#else
         if (grouped_fp8_gemm) {
             ptr = new impl::FP8BlockMOE(ctx, cfg, quant_config, parallel, dyn_shared);
         } else if (fuse_moe && gptq_kernel_algo == 1) {
@@ -1268,6 +1287,7 @@ FeedForward::FeedForward(
         } else {
             ptr = new impl::MOEImpl(ctx, cfg, quant_config, parallel);
         }
+#endif
         add_submodule("router", ptr->router.get());
         for (int i = 0; i < ptr->num_local_experts; ++i) {
             auto p = ptr->experts[i];
diff --git a/src/nn/feedforward/ff_kernel.cu b/src/nn/feedforward/ff_kernel.cu
index 09803b0..62cd554 100644
--- a/src/nn/feedforward/ff_kernel.cu
+++ b/src/nn/feedforward/ff_kernel.cu
@@ -136,7 +136,11 @@ static __device__ __forceinline__ void DEV_softmax_inplace(const T* logits, floa
     if (blockDim.x > 32)
         local_max = functions::blockReduceMax(local_max);
     else
+#if defined(__HIP_PLATFORM_AMD__) || defined(USE_HIP)
+        local_max = functions::warpReduceMaxWidthB<float, 32>(local_max);
+#else
         local_max = functions::warpReduceMaxB(local_max);
+#endif
 
     float local_sum = 1e-20;
     for (int i = threadIdx.x; i < n; i += blockDim.x) {
@@ -146,7 +150,11 @@ static __device__ __forceinline__ void DEV_softmax_inplace(const T* logits, floa
     if (blockDim.x > 32)
         local_sum = functions::blockReduceSum(local_sum);
     else
+#if defined(__HIP_PLATFORM_AMD__) || defined(USE_HIP)
+        local_sum = functions::warpReduceSumWidthB<float, 32>(local_sum);
+#else
         local_sum = functions::warpReduceSumB(local_sum);
+#endif
 
     for (int i = threadIdx.x; i < n; i += blockDim.x) {
         data[i] /= local_sum;
@@ -284,8 +292,13 @@ static __device__ inline void warpBitonicSort(T& v1, int& pos, bool asc = false)
         bool desc = ((lane_id & k) == 0) ^ asc;
 #pragma unroll
         for (int j = k / 2; j > 0; j /= 2) {
+#if defined(__HIP_PLATFORM_AMD__) || defined(USE_HIP)
+            T v2 = __shfl_xor(v1, j);
+            int pos2 = __shfl_xor(pos, j);
+#else
             T v2 = __shfl_xor_sync(0xFFFFFFFF, v1, j);
             int pos2 = __shfl_xor_sync(0xFFFFFFFF, pos, j);
+#endif
             bool upper = (lane_id & j) != 0;
 
             if (desc ^ (v1 > v2 || (v1 == v2 && pos < pos2)) ^ upper) {
diff --git a/src/nn/linear/cuda_helper.cpp b/src/nn/linear/cuda_helper.cpp
index 10d8178..2fbf411 100644
--- a/src/nn/linear/cuda_helper.cpp
+++ b/src/nn/linear/cuda_helper.cpp
@@ -1,6 +1,11 @@
 #include <bmengine/core/exception.h>
 #include <cuda.h>
 
+#if defined(USE_HIP)
+// gfx90a/CDNA has no programmable L2 access-policy window; this is a no-op on
+// AMD (the only caller is commented out at its call sites in linear.cpp).
+void setL2AccessPolicyWindow(cudaStream_t stream, void* data, int window_size, float hitRatio) {}
+#else
 // https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#l2-cache-access-window
 void setL2AccessPolicyWindow(cudaStream_t stream, void* data, int window_size, float hitRatio) {
     cudaStreamAttrValue stream_attribute;                                         // Stream level attributes data structure
@@ -17,3 +22,4 @@ void setL2AccessPolicyWindow(cudaStream_t stream, void* data, int window_size, f
 //    if (window_size == 0)
 //        BM_CUDART_ASSERT(cudaCtxResetPersistingL2Cache());
 }
+#endif
diff --git a/src/nn/linear/linear.cpp b/src/nn/linear/linear.cpp
index 0e6c150..54c0ce7 100644
--- a/src/nn/linear/linear.cpp
+++ b/src/nn/linear/linear.cpp
@@ -635,6 +635,11 @@ class Linear::impl::Int8Linear : public Linear::impl {
     }
 };
 
+// GPTQ / Marlin / AWQ / fp8 quantized Linear implementations are deferred on the
+// first AMD pass: they call NVIDIA tensor-core/PTX kernels (marlin mma.sync,
+// awq ldmatrix, gptq prmt/__vsub4, fp8 cvt.e4m3x2) excluded from the HIP build.
+// The Linear factory above throws for these quant types under USE_HIP.
+#if !defined(USE_HIP)
 class Linear::impl::Int4GPTQ : public Linear::impl {
     const int group_size;
 public:
@@ -1948,6 +1953,7 @@ class Linear::impl::Fp8Block : public Linear::impl {
 #endif
     }
 };
+#endif // !USE_HIP  (GPTQ/Marlin/AWQ/fp8 quantized Linear impls)
 
 Linear::Linear(
     const core::Context& ctx,
@@ -1962,6 +1968,19 @@ Linear::Linear(
     core::DataType dtype)
     : Layer() {
     int quant = static_cast<int>(quant_config.quant_type);
+#if defined(USE_HIP)
+    // GPTQ/AWQ/Marlin/fp8 quant linears use NVIDIA tensor-core / PTX kernels
+    // that are deferred on the first AMD pass. fp16/bf16 (NormalLinear) and the
+    // cublasLt-int8 path (Int8Linear) are supported; quantized weights are not.
+    if (quant_config.quant_type == model::QuantType::FP8
+        || quant_config.quant_type == model::QuantType::FP8_Block
+        || quant_config.quant_type == model::QuantType::AWQ
+        || quant_config.quant_type == model::QuantType::GPTQ
+        || quant_config.quant_type == model::QuantType::GPTQ_Marlin) {
+        BM_EXCEPTION("GPTQ/AWQ/Marlin/fp8 quantized Linear is not supported on ROCm/HIP "
+                     "on the first AMD pass (tensor-core/PTX kernels deferred).");
+    } else if (quant) {
+#else
     if (quant_config.quant_type == model::QuantType::FP8) {
         pimpl.reset(new impl::Fp8Linear(
             ctx, dim_in, dim_out, act_fn_type, quant, dtype, parallel, dist_layout));
@@ -1989,6 +2008,7 @@ Linear::Linear(
             ctx, dim_in, dim_out, act_fn_type, quant, dtype, parallel, dist_layout, act_order, group_size, sym);
         pimpl = std::unique_ptr<impl>((impl *) tmp);
     } else if (quant) {
+#endif
         auto tmp = new impl::Int8Linear(
             ctx, dim_in, dim_out, act_fn_type, quant, scale_weights, weight_transposed, dtype, parallel, dist_layout);
         add_parameter("weight_quant", tmp->weight);
@@ -2094,6 +2114,7 @@ void Linear::load_state_dict(
     this->prefix = prefix;
     pimpl->load_state_dict(ctx, state_dict, prefix, allow_missing);
 
+#if !defined(USE_HIP)
     bool dequant_desc_act = utils::get_int_env("DEQUANT_DESC_ACT", 0) > 0;
     impl::Int4GPTQ* q = dynamic_cast<impl::Int4GPTQ*>(pimpl.get());
     if (dequant_desc_act && q && q->parallel && q->act_order && !q->dim_out_parallel) {
@@ -2104,6 +2125,7 @@ void Linear::load_state_dict(
         *new_p->weight = w;
         pimpl.reset(new_p);
     }
+#endif
 }
 
 Linear* Linear::fuse(const core::Context& ctx, Linear& q, Linear& k) {
@@ -2165,6 +2187,9 @@ std::vector<Linear*> Linear::split(const core::Context& ctx, size_t n_split, boo
 }
 
 bool Linear::support_fuse_gptq_gate_in(const Tensor& input) {
+#if defined(USE_HIP)
+    return false;  // GPTQ quant Linear is deferred on AMD.
+#else
     static int fuse_w_in = utils::get_int_env("CPM_FUSE_FF_IN", 0);
     auto ptr = dynamic_cast<impl::Int4GPTQ*>(pimpl.get());
     return ptr
@@ -2174,14 +2199,19 @@ bool Linear::support_fuse_gptq_gate_in(const Tensor& input) {
             && !ptr->act_order
             && !ptr->trt_kernel
             && ptr->qweight.numel() > 0;
+#endif
 }
 
 std::tuple<core::Tensor, core::Tensor, core::Tensor, bool> Linear::get_gptq_weights() {
+#if defined(USE_HIP)
+    return std::make_tuple(Tensor(), Tensor(), Tensor(), false);  // GPTQ deferred on AMD.
+#else
     auto ptr = dynamic_cast<impl::Int4GPTQ*>(pimpl.get());
     if (!ptr) {
         return {Tensor(), Tensor(), Tensor(), false};
     }
     return {ptr->qweight, ptr->qzeros, ptr->scales, ptr->sym};
+#endif
 }
 
 void Linear::set_has_bias(bool b) {
@@ -2189,11 +2219,13 @@ void Linear::set_has_bias(bool b) {
 }
 
 void Linear::dequant_cache_weight(core::Context& ctx, const core::Tensor& fake_input) {
+#if !defined(USE_HIP)
     auto gptq_ptr = dynamic_cast<impl::Int4GPTQ*>(pimpl.get());
     model::ModelContext* m_ctx = model::ModelContext::cast(ctx);
     if (gptq_ptr && m_ctx) {
         gptq_ptr->dequant_cache_weight(*m_ctx, fake_input);
     }
+#endif
 }
 
 core::Tensor Linear::grouped_gemm_fp8_block(
@@ -2201,9 +2233,13 @@ core::Tensor Linear::grouped_gemm_fp8_block(
     const core::Tensor& input,
     const core::Tensor& m_indices,
     int num_groups) {
+#if defined(USE_HIP)
+    BM_EXCEPTION("fp8-block grouped GEMM is not supported on ROCm/HIP (deferred).");
+#else
     auto ptr = dynamic_cast<impl::Fp8Block*>(pimpl.get());
     BM_ASSERT(ptr, "Not Fp8Block impl");
     return ptr->grouped_gemm(ctx, input, m_indices, num_groups);
+#endif
 }
 
 // When num_kv_heads < TP; we load parameter of specified partition.
diff --git a/src/nn/quant/int8/quant_reduce_kernel.cu b/src/nn/quant/int8/quant_reduce_kernel.cu
index 461636a..e18ab26 100644
--- a/src/nn/quant/int8/quant_reduce_kernel.cu
+++ b/src/nn/quant/int8/quant_reduce_kernel.cu
@@ -23,7 +23,7 @@ static __global__ void KERNEL_quant_group_32(
         size_t offset = m * 32;
         int i = threadIdx.x;
         float v = inp[offset + i];
-        float abs_max = functions::warpReduceMaxB<T>(fabsf(v));
+        float abs_max = functions::warpReduceMaxWidthB<float, 32>(fabsf(v));
 
         out[offset + i] = int8_t(nearbyintf(v * 127.0f / abs_max));
         if (threadIdx.x == 0) {
@@ -53,7 +53,7 @@ static __global__ void KERNEL_quant_group_32_v2(
             size_t offset = m * 32;
             int i = threadIdx.x;
             float v = inp[offset + i];
-            float abs_max = functions::warpReduceMaxB<T>(fabsf(v));
+            float abs_max = functions::warpReduceMaxWidthB<float, 32>(fabsf(v));
 
             out[offset + i] = int8_t(nearbyintf(v * 127.0f / abs_max));
             if (threadIdx.x == 0) {
@@ -89,7 +89,7 @@ std::tuple<core::Tensor, core::Tensor> quant_group_32(
     cudaStream_t stream = ctx.current_stream()->ptr;
 
     BM_DTYPE_DISPATCH_HALF(input.dtype(), {
-        KERNEL_quant_group_32<scalar_t><<<round_up(M, 32) / 32, {32, 32}, 0, stream>>>(
+        KERNEL_quant_group_32<scalar_t><<<round_up(M, 32) / 32, dim3(32, 32), 0, stream>>>(
         input.data<scalar_t>(),
         output.mutable_data<int8_t>(),
         output_scale.mutable_data<scalar_t>(),
@@ -130,7 +130,7 @@ void dequant_group_32(
 
     cudaStream_t stream = ctx.current_stream()->ptr;
     BM_DTYPE_DISPATCH_FLOAT(scale.dtype(), {
-        KERNEL_dequant_group_32<scalar_t><<<round_up(M, 32) / 32, {32, 32}, 0, stream>>>(
+        KERNEL_dequant_group_32<scalar_t><<<round_up(M, 32) / 32, dim3(32, 32), 0, stream>>>(
             q.data<int8_t>(),
             scale.data<scalar_t>(),
             output->mutable_data<scalar_t>(),
@@ -227,7 +227,7 @@ core::Tensor dequant_group_fuse_add(
 
     cudaStream_t stream = ctx.current_stream()->ptr;
     BM_DTYPE_DISPATCH_FLOAT(scale.dtype(), {
-        KERNEL_dequant_group_fuse_add<scalar_t><<<round_up(M, 32) / 32, {32, 32}, 0, stream>>>(
+        KERNEL_dequant_group_fuse_add<scalar_t><<<round_up(M, 32) / 32, dim3(32, 32), 0, stream>>>(
             q.data<int8_t>(),
             scale.data<scalar_t>(),
             c.data<scalar_t>(),
@@ -264,7 +264,7 @@ static __global__ void KERNEL_dequant_sum_quant_g32(
 //        sum += q * s;
     }
 
-    float abs_max = functions::warpReduceMaxB<T>(fabsf(sum));
+    float abs_max = functions::warpReduceMaxWidthB<float, 32>(fabsf(sum));
 
     out_q[offset_i] = int8_t(nearbyintf(sum * 127.0f / abs_max));
 
diff --git a/tests/py_export_internal/CMakeLists.txt b/tests/py_export_internal/CMakeLists.txt
index c92ebf5..bc1dfa0 100644
--- a/tests/py_export_internal/CMakeLists.txt
+++ b/tests/py_export_internal/CMakeLists.txt
@@ -1,10 +1,18 @@
-enable_language(CUDA)
+if(USE_HIP)
+    enable_language(HIP)
+else()
+    enable_language(CUDA)
+    set(TORCH_CUDA_ARCH_LIST "8.0+PTX;8.9")
+endif()
 
 project(internals_)
-set(CMAKE_CXX_STANDARD 17)
+if(USE_HIP)
+    set(CMAKE_CXX_STANDARD 20)  # ROCm torch 2.13 headers need C++20 concepts
+else()
+    set(CMAKE_CXX_STANDARD 17)
+endif()
 set(CMAKE_CXX_VISIBILITY_PRESET default)
 set(CMAKE_CUDA_VISIBILITY_PRESET default)
-set(TORCH_CUDA_ARCH_LIST "8.0+PTX;8.9")
 
 execute_process(COMMAND ${PYTHON_EXECUTABLE} -c "import torch; print(torch.utils.cmake_prefix_path)" OUTPUT_VARIABLE PYTORCH_CMAKE_PREFIX_PATH OUTPUT_STRIP_TRAILING_WHITESPACE)
 list(APPEND CMAKE_PREFIX_PATH "${PYTORCH_CMAKE_PREFIX_PATH}/Torch")
@@ -18,7 +26,14 @@ file(GLOB_RECURSE TEST_SRC_NN "*.cpp")
 list(APPEND TEST_SRC_NN "../../src/py_export/py_utils.cpp")
 
 message("ADD pybind11 module for unittest ${PROJECT_NAME}")
-pybind11_add_module(${PROJECT_NAME} ${TEST_SRC_NN})
+if(USE_HIP)
+    # NO_EXTRAS drops pybind11's -flto (HIP link does not finalize LTO -> the
+    # .so loses PyInit_*).
+    pybind11_add_module(${PROJECT_NAME} NO_EXTRAS ${TEST_SRC_NN})
+    set_target_properties(${PROJECT_NAME} PROPERTIES INTERPROCEDURAL_OPTIMIZATION OFF)
+else()
+    pybind11_add_module(${PROJECT_NAME} ${TEST_SRC_NN})
+endif()
 
 #target_include_directories(${PROJECT_NAME} PRIVATE ${Python_SITELIB}/torch/include)
 
diff --git a/tests/py_export_internal/layer_feedforward.cpp b/tests/py_export_internal/layer_feedforward.cpp
index c5013c2..f12bd9f 100644
--- a/tests/py_export_internal/layer_feedforward.cpp
+++ b/tests/py_export_internal/layer_feedforward.cpp
@@ -200,20 +200,36 @@ class PyFeedForward {
     void load_state_dict(const std::map<std::string, py::array>& state_dict)
         __attribute__((visibility("hidden"))) {
 
-        std::vector<std::thread> threads;
-
-        for (int i = 0; i < engine->num_gpus(); ++i) {
+        // bind::load_state_dict touches py::array (numpy buffer protocol = Python
+        // C-API), which requires the GIL. The main thread holds the GIL and then
+        // join()s the worker, so a worker that needs the GIL deadlocks. For the
+        // single-GPU case (the validation path) run inline on the GIL-holding
+        // main thread; only fan out to worker threads when there are multiple
+        // GPUs, and have each worker acquire the GIL before the numpy access.
+        int n = engine->num_gpus();
+        if (n <= 1) {
+            auto ctx = engine->create_context({ 0 });
+            bmengine::core::WithDevice device(ctx, 0);
+            auto named_params = mds[0]->named_parameters("ff", true);
+            bind::load_state_dict(ctx, state_dict, named_params);
+            return;
+        }
 
+        std::vector<std::thread> threads;
+        for (int i = 0; i < n; ++i) {
             threads.emplace_back([this, i, &state_dict] {
+                py::gil_scoped_acquire gil;
                 auto ctx = engine->create_context({ i });
                 bmengine::core::WithDevice device(ctx, 0);
                 auto named_params = mds[i]->named_parameters("ff", true);
                 bind::load_state_dict(ctx, state_dict, named_params);
             });
         }
-
-        for (auto it = threads.begin(); it != threads.end(); ++it) {
-            it->join();
+        {
+            py::gil_scoped_release release;  // let the workers acquire the GIL
+            for (auto it = threads.begin(); it != threads.end(); ++it) {
+                it->join();
+            }
         }
     }