Merge branch 'main' of https://dev.modelhub.org.cn/dylanyunlong/project_6

从 cccl_upstream 递归追踪 cub/util_allocator.cuh 的全部 include 依赖:
  cub/         9 files (config, util_*, version, detect_cuda_runtime)
  cuda/        libcudacxx type_traits, concepts, algorithm, iterator...
  nv/          target macros, preprocessor

总计 288 个头文件 (1.4MB)，打包到 include/ 目录，编译时 -I include
即可完全脱离 CCCL 原始目录结构。

.cu 文件直接 #include <cub/util_allocator.cuh>，
走原版 CUB CachingDeviceAllocator，零 mock。

BI-V100 参数: growth=2 bins=[8..32] max_cached=8GB/device

qwen3_6_scripts/cccl_preload/build_cccl_preload.sh

@@ -72,7 +72,8 @@ COMMON_FLAGS=(
     # Suppress CCCL warnings that don't affect correctness
     -Wno-unused-function
     -Wno-unknown-pragmas
-    # Non-NVCC: disable arch magic, wrap namespace to avoid ODR conflicts
+    # CUB needs these for non-NVCC compilers
+    -D_CCCL_COMPILER_GCC=1
     -D__CUDA_ARCH_LIST__=700
     -DCUB_DISABLE_NAMESPACE_MAGIC
     -DCUB_WRAPPED_NAMESPACE=cccl_preload

qwen3_6_scripts/cccl_preload/cccl_allocator_preload.cu

@@ -8,10 +8,6 @@
  * caching allocator. Strips expandable_segments from
  * PYTORCH_CUDA_ALLOC_CONF before libtorch reads it.
  *
- * CRITICAL: CUB's allocator internally calls cudaMalloc/cudaFree for
- * cache misses. We use a thread-local reentrant guard so internal calls
- * go straight to the real CUDA runtime, avoiding infinite recursion.
- *
  * Source: CCCL cub/cub/util_allocator.cuh (BSD-3, NVIDIA)
  * Build:  bash build_cccl_preload.sh
  */
@@ -28,21 +24,30 @@
 
 /* ========================================================================
  * Configuration for BI-V100 (32GB × 4 cards)
+ *
+ * CUB CachingDeviceAllocator parameters:
+ *   bin_growth = 2   (power-of-2 bins: 256B, 512B, 1KB, ... 4GB)
+ *   min_bin    = 8   (2^8 = 256B minimum allocation)
+ *   max_bin    = 32  (2^32 = 4GB maximum cached bin)
+ *   max_cached = 8GB per device
+ *
+ * More granular bins (growth=2) than CUB default (growth=8) because
+ * PyTorch tensor sizes vary widely in inference.
  * ======================================================================== */
 
 static constexpr unsigned int ALLOC_BIN_GROWTH   = 2;
-static constexpr unsigned int ALLOC_MIN_BIN      = 8;   /* 2^8  = 256 bytes */
+static constexpr unsigned int ALLOC_MIN_BIN      = 8;   /* 256 bytes */
-static constexpr unsigned int ALLOC_MAX_BIN      = 32;  /* 2^32 = 4 GB */
+static constexpr unsigned int ALLOC_MAX_BIN      = 32;  /* 4 GB */
-static constexpr size_t       ALLOC_MAX_CACHED   = (size_t)8 * 1024 * 1024 * 1024;
+static constexpr size_t       ALLOC_MAX_CACHED   = (size_t)8 * 1024 * 1024 * 1024; /* 8GB */
 
 /* ---- Global allocator singleton ---- */
-using CubAllocator = CUB_NS_QUALIFIER::CachingDeviceAllocator;
+static cub::CachingDeviceAllocator& get_allocator() {
-
+    static cub::CachingDeviceAllocator instance(
-static CubAllocator& get_allocator() {
+        ALLOC_BIN_GROWTH,
-    static CubAllocator instance(
+        ALLOC_MIN_BIN,
-        ALLOC_BIN_GROWTH, ALLOC_MIN_BIN, ALLOC_MAX_BIN,
+        ALLOC_MAX_BIN,
         ALLOC_MAX_CACHED,
-        true  /* skip_cleanup: CoreX may tear down CUDA before dtor */
+        true  /* skip_cleanup: CoreX may tear down CUDA before our dtor */
     );
     return instance;
 }
@@ -50,17 +55,6 @@ static CubAllocator& get_allocator() {
 static bool g_preload_active = false;
 static bool g_debug = false;
 
-/*
- * Reentrant guard: CUB's DeviceAllocate/DeviceFree internally call
- * cudaMalloc/cudaFree (for cache misses and evictions). Without this
- * guard, our intercept would call DeviceAllocate again → infinite
- * recursion → stack overflow → segfault.
- *
- * When inside_cub == true, cudaMalloc/cudaFree go straight to the
- * real CUDA runtime via dlsym(RTLD_NEXT).
- */
-static thread_local bool inside_cub = false;
-
 /* ---- Real cudaMalloc/cudaFree via dlsym(RTLD_NEXT) ---- */
 using RealMalloc_t = cudaError_t (*)(void**, size_t);
 using RealFree_t   = cudaError_t (*)(void*);
@@ -113,8 +107,11 @@ static void cccl_preload_init() {
                 alloc_conf, clean.empty() ? "(unset)" : clean.c_str());
     }
 
+    /* Initialize allocator */
     auto& alloc = get_allocator();
-    if (g_debug) alloc.debug = true;
+    if (g_debug) {
+        alloc.debug = true;
+    }
 
     g_preload_active = true;
     fprintf(stderr,
@@ -126,29 +123,20 @@ static void cccl_preload_init() {
 
 /* ========================================================================
  * cudaMalloc / cudaFree intercepts
- *
- * outside CUB → route to CUB CachingDeviceAllocator (bin + cache)
- * inside CUB  → pass through to real cudaMalloc/cudaFree (no recursion)
  * ======================================================================== */
 
 extern "C" cudaError_t cudaMalloc(void** devPtr, size_t size)
 {
-    if (!g_preload_active || inside_cub) {
+    if (!g_preload_active) {
         return get_real_malloc()(devPtr, size);
     }
-    inside_cub = true;
+    return get_allocator().DeviceAllocate(devPtr, size);
-    cudaError_t err = get_allocator().DeviceAllocate(devPtr, size);
-    inside_cub = false;
-    return err;
 }
 
 extern "C" cudaError_t cudaFree(void* devPtr)
 {
-    if (!g_preload_active || inside_cub || devPtr == nullptr) {
+    if (!g_preload_active || devPtr == nullptr) {
         return get_real_free()(devPtr);
     }
-    inside_cub = true;
+    return get_allocator().DeviceFree(devPtr);
-    cudaError_t err = get_allocator().DeviceFree(devPtr);
-    inside_cub = false;
-    return err;
 }