Fix HIP build errors in unified kernel files

- Add hip/hip_bfloat16.h include to compat.cuh (bnb_bfloat16 type alias
  requires hip_bfloat16 to be defined)
- Add __syncwarp() no-op macro for HIP (AMD warps are always in lockstep)
- Add hipblas version check (#if hipblasVersionMajor >= 3) for GemmEx calls
  (ROCm 6.1 ships hipblas v2 which uses HIPBLAS_R_* not HIPBLAS_COMPUTE_*)
- Fix include in ops.cuh: common.h -> common.cuh (BNB_WARP_SIZE visibility)

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>

Author: TimDettmers

csrc/compat.cuh

@@ -26,6 +26,7 @@
 #if BNB_HIP
 
 #include <hip/hip_fp16.h>
+#include <hip/hip_bfloat16.h>
 #include <hip/hip_math_constants.h>
 #include <hip/hip_runtime.h>
 #include <hipblas/hipblas.h>
@@ -85,6 +86,17 @@ using bnb_error_t = cudaError_t;
 // Keep backward compat for existing code during migration
 #define CUDA_CHECK_RETURN(value) BNB_CHECK_RETURN(value)
 
+// ============================================================================
+// Warp synchronization
+//
+// HIP warps are always in lockstep (no independent thread scheduling),
+// so __syncwarp() is a no-op. CUDA needs it for warp convergence.
+// ============================================================================
+
+#if BNB_HIP
+#define __syncwarp() do {} while(0)
+#endif
+
 // ============================================================================
 // BFloat16 type alias
 //

csrc/ops.cu

@@ -270,10 +270,18 @@ void gemmex(
 #if BNB_HIP
     hipblasStatus_t status;
 
+#if hipblasVersionMajor >= 3
     status = hipblasGemmEx(
         context->m_handle, transposeA ? HIPBLAS_OP_T : HIPBLAS_OP_N, transposeB ? HIPBLAS_OP_T : HIPBLAS_OP_N, m, n, k,
         alpha, A, HIP_R_8I, lda, B, HIP_R_8I, ldb, beta, C, HIP_R_32I, ldc, HIPBLAS_COMPUTE_32I, HIPBLAS_GEMM_DEFAULT
     );
+#else
+    status = hipblasGemmEx(
+        context->m_handle, transposeA ? HIPBLAS_OP_T : HIPBLAS_OP_N, transposeB ? HIPBLAS_OP_T : HIPBLAS_OP_N, m, n, k,
+        alpha, A, HIPBLAS_R_8I, lda, B, HIPBLAS_R_8I, ldb, beta, C, HIPBLAS_R_32I, ldc, HIPBLAS_R_32I,
+        HIPBLAS_GEMM_DEFAULT
+    );
+#endif
 
     if (status != HIPBLAS_STATUS_SUCCESS) {
         std::cout << "HIPBLAS ERROR: Status " << status << std::endl;
@@ -304,11 +312,19 @@ void strided_gemmex(
 #if BNB_HIP
     hipblasStatus_t status;
 
+#if hipblasVersionMajor >= 3
     status = hipblasGemmStridedBatchedEx(
         context->m_handle, transposeA ? HIPBLAS_OP_T : HIPBLAS_OP_N, transposeB ? HIPBLAS_OP_T : HIPBLAS_OP_N, m, n, k,
         alpha, A, HIP_R_8I, lda, (long long int)strideA, B, HIP_R_8I, ldb, (long long int)strideB, beta, C, HIP_R_32I,
         ldc, (long long int)strideC, batchCount, HIPBLAS_COMPUTE_32I, HIPBLAS_GEMM_DEFAULT
     );
+#else
+    status = hipblasGemmStridedBatchedEx(
+        context->m_handle, transposeA ? HIPBLAS_OP_T : HIPBLAS_OP_N, transposeB ? HIPBLAS_OP_T : HIPBLAS_OP_N, m, n, k,
+        alpha, A, HIPBLAS_R_8I, lda, (long long int)strideA, B, HIPBLAS_R_8I, ldb, (long long int)strideB, beta, C,
+        HIPBLAS_R_32I, ldc, (long long int)strideC, batchCount, HIPBLAS_R_32I, HIPBLAS_GEMM_DEFAULT
+    );
+#endif
 
     if (status != HIPBLAS_STATUS_SUCCESS) {
         std::cout << "HIPBLAS ERROR: Status " << status << std::endl;

csrc/ops.cuh

@@ -14,7 +14,7 @@
 #include <vector>
 
 #include "compat.cuh"
-#include <common.h>
+#include "common.cuh"
 
 // ============================================================================
 // Error checking helpers