Activate unified CUDA/HIP kernel files from csrc/examples/

Move the unified portability-header-based files from csrc/examples/
into csrc/, replacing the duplicated CUDA and HIP kernel files.

- Add compat.cuh and compat_device.cuh (portability headers)
- Replace common.cuh, kernels.cu, ops.cu, ops.cuh, pythonInterface.cpp,
  CMakeLists.txt with unified versions
- Update kernels.cuh: rename kQuantizeBlockwise32 -> kQuantizeBlockwiseSmall
- Delete HIP-only files: common_hip.cuh, kernels.hip, kernels_hip.cuh,
  ops.hip, ops_hip.cuh
- Delete csrc/examples/ (files are now in their final locations)

Net: 10 source files -> 7, ~3300 fewer lines of duplicated code.
Same .cu files compiled by both nvcc (CUDA) and hipcc (HIP).

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

Author: TimDettmers

CMakeLists.txt

@@ -24,8 +24,7 @@ endif()
 
 # Define included source files
 set(CPP_FILES csrc/cpu_ops.cpp csrc/pythonInterface.cpp)
-set(CUDA_FILES csrc/ops.cu csrc/kernels.cu)
-set(HIP_FILES csrc/ops.hip csrc/kernels.hip)
+set(GPU_FILES csrc/ops.cu csrc/kernels.cu)
 set(MPS_FILES csrc/mps_ops.mm)
 set(METAL_FILES csrc/mps_kernels.metal)
 set(XPU_FILES csrc/xpu_ops.cpp csrc/xpu_kernels.cpp)
@@ -195,7 +194,7 @@ if(BUILD_CUDA)
     message(STATUS "CUDA Targets: ${CMAKE_CUDA_ARCHITECTURES}")
     message(STATUS "CUDA NVCC Flags: ${CMAKE_CUDA_FLAGS}")
 
-    list(APPEND SRC_FILES ${CUDA_FILES})
+    list(APPEND SRC_FILES ${GPU_FILES})
 
     string(APPEND BNB_OUTPUT_NAME "_cuda${CUDA_VERSION_SHORT}")
     add_compile_definitions(BUILD_CUDA)
@@ -213,7 +212,7 @@ elseif(BUILD_HIP)
     endif()
     message(STATUS "HIP Targets: ${CMAKE_HIP_ARCHITECTURES}")
 
-    list(APPEND SRC_FILES ${HIP_FILES})
+    list(APPEND SRC_FILES ${GPU_FILES})
 
     string(APPEND BNB_OUTPUT_NAME "_rocm")
 
@@ -339,7 +338,7 @@ if(BUILD_HIP)
     target_link_libraries(bitsandbytes PUBLIC roc::hipblas hip::hiprand roc::hipsparse)
 
     target_compile_definitions(bitsandbytes PUBLIC BNB_USE_HIP)
-    set_source_files_properties(${HIP_FILES} PROPERTIES LANGUAGE HIP)
+    set_source_files_properties(${GPU_FILES} PROPERTIES LANGUAGE HIP)
     set_target_properties(bitsandbytes PROPERTIES LINKER_LANGUAGE CXX)
 
     if(HIP_VERSION VERSION_LESS "6.1")

csrc/common.cuh

@@ -1,6 +1,41 @@
+// common.cuh — Merged architecture constants for CUDA and HIP
+//
+// This replaces both the old csrc/common.cuh and csrc/common_hip.cuh.
+// Platform detection uses compat.cuh's BNB_HIP macro.
+
 #pragma once
 
-// TODO: Let's make some of these constexpr and put in a namespace.
+#include "compat.cuh"
+
+// ============================================================================
+// Warp size
+// ============================================================================
+
+#if BNB_HIP
+// AMD GFX9 (CDNA) uses 64-wide warps; RDNA uses 32-wide
+#ifdef __GFX9__
+#define BNB_WARP_SIZE 64
+#else
+#define BNB_WARP_SIZE 32
+#endif
+#else
+#define BNB_WARP_SIZE 32
+#endif
+
+// ============================================================================
+// BF16 availability
+// ============================================================================
+
+#if BNB_HIP
+// BF16 is available on all currently-supported ROCm architectures (CDNA2+, RDNA3+)
+#define BNB_BF16_AVAILABLE true
+#else
+#define BNB_BF16_AVAILABLE (__CUDA_ARCH__ >= BNB_CC_AMPERE)
+#endif
+
+// ============================================================================
+// CUDA compute capability constants (CUDA-only, but harmless to define on HIP)
+// ============================================================================
 
 #define BNB_CC_PASCAL 600
 #define BNB_CC_PASCAL_X2 620
@@ -14,31 +49,45 @@
 #define BNB_CC_HOPPER 900
 #define BNB_CC_BLACKWELL 1000
 
+// ============================================================================
+// Feature availability based on arch (CUDA uses __CUDA_ARCH__, HIP is simpler)
+// ============================================================================
+
+#if BNB_HIP
+// HIP: MMA not supported via mma.h; FP8 support varies by arch
+#define BNB_FP16_MMA_AVAILABLE 0
+#define BNB_INT8_MMA_AVAILABLE 0
+#define BNB_FP8_AVAILABLE 0
+#else
 #define BNB_FP16_MMA_AVAILABLE (__CUDA_ARCH__ >= BNB_CC_VOLTA)
 #define BNB_INT8_MMA_AVAILABLE (__CUDA_ARCH__ >= BNB_CC_VOLTA_XAVIER)
-#define BNB_BF16_AVAILABLE (__CUDA_ARCH__ >= BNB_CC_AMPERE)
 #define BNB_FP8_AVAILABLE (__CUDA_ARCH__ >= BNB_CC_ADA)
+#endif
 
-#define BNB_WARP_SIZE 32
+// ============================================================================
+// Maximum threads per SM/CU
+// ============================================================================
 
-// The maximum number of resident threads per SM varies by arch.
-// For A100/H100 and all prior to Turing, it is 2048, which allows
-// for 2 full blocks of 1024 threads per SM.
-// Reference:
-// https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#features-and-technical-specifications-technical-specifications-per-compute-capability
+#if BNB_HIP
+// For currently supported ROCm architectures (CDNA2, RDNA3)
+#define BNB_MAX_THREADS_PER_SM 2048
+#else
+// The maximum number of resident threads per SM varies by NVIDIA arch.
+// Reference: CUDA Programming Guide, Technical Specifications per Compute Capability
 #if __CUDA_ARCH__ == 750
 #define BNB_MAX_THREADS_PER_SM 1024
 #elif __CUDA_ARCH__ >= 860 && __CUDA_ARCH__ <= 890
 #define BNB_MAX_THREADS_PER_SM 1536
 #else
 #define BNB_MAX_THREADS_PER_SM 2048
 #endif
+#endif
 
-// Maximum resident warps per SM is always directly related to the number of threads.
+// Maximum resident warps per SM/CU
 #define BNB_MAX_WARPS_PER_SM ((BNB_MAX_THREADS_PER_SM) / (BNB_WARP_SIZE))
 
-// Maximum resident blocks per SM may vary.
-#if __CUDA_ARCH__ == 860 || __CUDA_ARCH__ == 870
+// Maximum resident blocks per SM/CU
+#if !BNB_HIP && (defined(__CUDA_ARCH__)) && (__CUDA_ARCH__ == 860 || __CUDA_ARCH__ == 870)
 #define BNB_MAX_BLOCKS_PER_SM 16
 #else
 #define BNB_MAX_BLOCKS_PER_SM ((BNB_MAX_WARPS_PER_SM) / 2)