fix: verify_submission.sh — strip semicolons from path extraction grep

- patch_ops.sh: call cccl_preload/build_cccl_preload.sh (new CCCL deps)
  instead of old build_cccl_preload_allocator.sh (mock)
- computility-run.yaml: add LD_PRELOAD + CCCL_ALLOC_DISABLE env vars
- Remove old mock files: cccl_preload_allocator.cu, build script, test
- .dockerignore: exclude cccl_upstream/ upstream_ref/ vllm/ *.zip
- verify_submission.sh: 31-point pre-submission check
  (file structure, CCCL chain, path matching, prebuilt integrity,
   corex imports, docker context, GPU smoke test)

computility-run.yaml

@@ -48,14 +48,8 @@ env:
     - name: BI100_MOE_COREX_TOPK_SOFTMAX
       value: '1'
     - name: PYTORCH_CUDA_ALLOC_CONF
-      value: expandable_segments:True
+      value: max_split_size_mb:512
     - name: LD_PRELOAD
-      value: /workspace/qwen3_6_scripts/cccl_preload_allocator.so
-    - name: CCCL_ALLOC_BIN_GROWTH
-      value: '8'
-    - name: CCCL_ALLOC_MIN_BIN
-      value: '3'
-    - name: CCCL_ALLOC_MAX_BIN
-      value: '13'
-    - name: CCCL_ALLOC_MAX_CACHED_MB
-      value: '4096'
+      value: /workspace/qwen3_6_scripts/libcccl_allocator.so
+    - name: CCCL_ALLOC_DISABLE
+      value: '0'

qwen3_6_scripts/build_cccl_preload_allocator.sh

@@ -1,99 +0,0 @@
-#!/usr/bin/env bash
-set -euo pipefail
-
-# Build the CCCL CachingDeviceAllocator LD_PRELOAD .so
-#
-# Usage: bash build_cccl_preload_allocator.sh [output_dir]
-# Default output: ./cccl_preload_allocator.so
-#
-# Test:  LD_PRELOAD=./cccl_preload_allocator.so python3 -c "import torch; x=torch.zeros(1024,device='cuda'); del x; print('OK')"
-
-SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
-OUTPUT_DIR=${1:-${SCRIPT_DIR}}
-OUTPUT=${OUTPUT_DIR}/cccl_preload_allocator.so
-SOURCE=${SCRIPT_DIR}/cccl_preload_allocator.cu
-
-COREX_ROOT=${COREX_ROOT:-/usr/local/corex-3.2.3}
-
-# Check if we have the corex compiler
-if [[ -x "${COREX_ROOT}/bin/clang++" ]]; then
-    COMPILER="${COREX_ROOT}/bin/clang++"
-    echo "[build] Using corex clang++: ${COMPILER}"
-
-    "${COMPILER}" \
-        -std=c++17 -O3 -shared -fPIC \
-        --cuda-path="${COREX_ROOT}" --cuda-gpu-arch=ivcore10 \
-        --no-cuda-version-check -D_GLIBCXX_USE_CXX11_ABI=0 \
-        -I"${COREX_ROOT}/include" \
-        "${SOURCE}" \
-        -L"${COREX_ROOT}/lib64" -lcudart -ldl \
-        -Wl,-rpath,"${COREX_ROOT}/lib64" \
-        -o "${OUTPUT}"
-else
-    # Fallback: try to compile as pure C++ (no CUDA kernels needed)
-    # The allocator is entirely host-side code
-    echo "[build] corex clang++ not found, trying system g++"
-    echo "[build] Note: this is HOST-only code, no GPU kernels involved"
-
-    # Find cuda include path
-    CUDA_INC=""
-    for p in /usr/local/corex-3.2.3/include /usr/local/cuda/include /usr/local/corex/include; do
-        if [[ -d "$p" ]]; then CUDA_INC="$p"; break; fi
-    done
-
-    CUDA_LIB=""
-    for p in /usr/local/corex-3.2.3/lib64 /usr/local/cuda/lib64 /usr/local/corex/lib64; do
-        if [[ -d "$p" ]]; then CUDA_LIB="$p"; break; fi
-    done
-
-    if [[ -z "$CUDA_INC" || -z "$CUDA_LIB" ]]; then
-        echo "[build] ERROR: Cannot find CUDA headers/libs" >&2
-        exit 2
-    fi
-
-    # Rename .cu → .cpp for g++ (it's all host code anyway)
-    TMP_CPP=$(mktemp /tmp/cccl_alloc_XXXXXX.cpp)
-    cp "${SOURCE}" "${TMP_CPP}"
-
-    g++ -std=c++17 -O3 -shared -fPIC \
-        -D_GLIBCXX_USE_CXX11_ABI=0 \
-        -I"${CUDA_INC}" \
-        "${TMP_CPP}" \
-        -L"${CUDA_LIB}" -lcudart -ldl \
-        -Wl,-rpath,"${CUDA_LIB}" \
-        -o "${OUTPUT}"
-
-    rm -f "${TMP_CPP}"
-fi
-
-# Verify
-if [[ ! -s "${OUTPUT}" ]]; then
-    echo "[build] ERROR: output is empty" >&2
-    exit 2
-fi
-
-# Check it's a proper shared library with our symbols
-if command -v nm &>/dev/null; then
-    HAS_MALLOC=$(nm -D "${OUTPUT}" 2>/dev/null | grep -c "T cudaMalloc" || true)
-    HAS_FREE=$(nm -D "${OUTPUT}" 2>/dev/null | grep -c "T cudaFree" || true)
-    if [[ "$HAS_MALLOC" -gt 0 && "$HAS_FREE" -gt 0 ]]; then
-        echo "[build] OK: cudaMalloc and cudaFree symbols exported"
-    else
-        echo "[build] WARNING: symbol check inconclusive (nm output may differ)"
-    fi
-fi
-
-echo "[build] Built: ${OUTPUT} ($(stat -c%s "${OUTPUT}" 2>/dev/null || stat -f%z "${OUTPUT}") bytes)"
-echo ""
-echo "Test command:"
-echo "  LD_PRELOAD=${OUTPUT} python3 -c \"import torch; x=torch.zeros(1024,device='cuda'); del x; print('OK')\""
-echo ""
-echo "Production usage (add to computility-run.yaml or Dockerfile CMD):"
-echo "  LD_PRELOAD=${OUTPUT} python3 -m vllm.entrypoints.openai.api_server ..."
-echo ""
-echo "Tuning env vars:"
-echo "  CCCL_ALLOC_BIN_GROWTH=8       # Geometric growth factor"
-echo "  CCCL_ALLOC_MIN_BIN=3          # Min bin (growth^3 = 512B)"
-echo "  CCCL_ALLOC_MAX_BIN=13         # Max bin (8^13 = ~550MB)"
-echo "  CCCL_ALLOC_MAX_CACHED_MB=4096 # Max 4GB cached per device"
-echo "  CCCL_ALLOC_DEBUG=1            # Print every alloc/free"

qwen3_6_scripts/cccl_preload/build_cccl_preload.sh

@@ -1,41 +1,43 @@
 #!/usr/bin/env bash
-# Build libcccl_allocator.so — LD_PRELOAD .so for CUB CachingDeviceAllocator
+# Build libcccl_allocator.so
+#
+# Full CCCL dependency chain (288 headers) in ./include/
+# Source: cccl_upstream/cub/cub/util_allocator.cuh + transitive deps
 #
 # Usage:
 #   bash build_cccl_preload.sh [output_dir]
-#
-# On BI-V100 with CoreX SDK:
-#   bash build_cccl_preload.sh /workspace/qwen3_6_scripts/cccl_preload
-#
-# The .so intercepts cudaMalloc/cudaFree and routes through CUB's
-# caching allocator, bypassing CoreX's "expandable segment not supported"
-# ASSERT in CUDACachingAllocator.cpp:545.
 
 set -euo pipefail
 
 SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
 OUTPUT_DIR="${1:-${SCRIPT_DIR}}"
 SRC="${SCRIPT_DIR}/cccl_allocator_preload.cu"
+INC="${SCRIPT_DIR}/include"
 OUT="${OUTPUT_DIR}/libcccl_allocator.so"
 
-# Find CoreX clang++ (preferred) or system g++
-if [[ -x /usr/local/corex-3.2.3/bin/clang++ ]]; then
-    CXX=/usr/local/corex-3.2.3/bin/clang++
-    echo "[build] Using CoreX clang++: ${CXX}"
-elif [[ -x /usr/local/corex/bin/clang++ ]]; then
-    CXX=/usr/local/corex/bin/clang++
-    echo "[build] Using CoreX clang++ (alt): ${CXX}"
-else
-    CXX=g++
-    echo "[build] CoreX clang++ not found, falling back to g++"
-fi
+[[ -d "${INC}/cub" ]] || { echo "CCCL include tree missing: ${INC}/cub"; exit 2; }
+[[ -d "${INC}/cuda" ]] || { echo "CCCL include tree missing: ${INC}/cuda"; exit 2; }
 
-# Find CUDA include path
+# Find compiler
+CXX=""
+for candidate in \
+    /usr/local/corex-3.2.3/bin/clang++ \
+    /usr/local/corex/bin/clang++ \
+    /usr/local/corex/lib64/clang/16/bin/clang++ \
+    ; do
+    if [[ -x "${candidate}" ]]; then
+        CXX="${candidate}"
+        break
+    fi
+done
+[[ -n "${CXX}" ]] || { CXX=g++; echo "[build] no CoreX clang++, falling back to g++"; }
+echo "[build] CXX=${CXX}"
+
+# Find CUDA headers (for cuda_runtime_api.h)
 CUDA_INC=""
 for candidate in \
     /usr/local/corex/include \
     /usr/local/cuda/include \
-    /usr/local/corex/lib64/clang/16/include \
     ; do
     if [[ -f "${candidate}/cuda_runtime_api.h" ]]; then
         CUDA_INC="${candidate}"
@@ -43,7 +45,7 @@ for candidate in \
     fi
 done
 
-# Find CUDA lib path for linking
+# Find CUDA libs
 CUDA_LIB=""
 for candidate in \
     /usr/local/corex/lib64 \
@@ -55,57 +57,43 @@ for candidate in \
     fi
 done
 
-if [[ -z "${CUDA_INC}" ]]; then
-    echo "[WARN] cuda_runtime_api.h not found — trying compile anyway"
-fi
-
 echo "[build] CUDA include: ${CUDA_INC:-system}"
-echo "[build] CUDA lib: ${CUDA_LIB:-system}"
-echo "[build] Source: ${SRC}"
-echo "[build] Output: ${OUT}"
+echo "[build] CUDA lib:     ${CUDA_LIB:-system}"
+echo "[build] CCCL include: ${INC} ($(find "${INC}" -type f | wc -l) files)"
+echo "[build] Source:        ${SRC}"
+echo "[build] Output:        ${OUT}"
+
+COMMON_FLAGS=(
+    -shared -fPIC -O2 -std=c++17
+    -I"${INC}"
+    ${CUDA_INC:+-I"${CUDA_INC}"}
+    ${CUDA_LIB:+-L"${CUDA_LIB}"}
+    -lcudart -ldl
+    # Suppress CCCL warnings that don't affect correctness
+    -Wno-unused-function
+    -Wno-unknown-pragmas
+    # 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
+)
 
-# Build as shared library
-# -x cuda or -x c++ depending on compiler
 if [[ "${CXX}" == *clang++* ]]; then
-    # CoreX clang++ can compile .cu natively
-    ${CXX} \
-        -shared -fPIC \
-        -O2 \
-        ${CUDA_INC:+-I"${CUDA_INC}"} \
-        ${CUDA_LIB:+-L"${CUDA_LIB}"} \
-        -lcudart \
-        -ldl \
-        -std=c++17 \
-        -o "${OUT}" \
-        "${SRC}"
+    "${CXX}" "${COMMON_FLAGS[@]}" -x c++ -o "${OUT}" "${SRC}" 2>&1
 else
-    # g++ needs .cu renamed or treated as C++
-    # cuda_runtime_api.h should still work with host compiler
-    ${CXX} \
-        -shared -fPIC \
-        -O2 \
-        ${CUDA_INC:+-I"${CUDA_INC}"} \
-        ${CUDA_LIB:+-L"${CUDA_LIB}"} \
-        -lcudart \
-        -ldl \
-        -std=c++17 \
-        -x c++ \
-        -o "${OUT}" \
-        "${SRC}"
+    "${CXX}" "${COMMON_FLAGS[@]}" -x c++ -o "${OUT}" "${SRC}" 2>&1
 fi
 
 if [[ -f "${OUT}" ]]; then
-    SIZE=$(stat -c%s "${OUT}" 2>/dev/null || stat -f%z "${OUT}" 2>/dev/null || echo "?")
+    SIZE=$(stat -c%s "${OUT}" 2>/dev/null || echo "?")
+    echo ""
     echo "[build] SUCCESS: ${OUT} (${SIZE} bytes)"
     echo ""
-    echo "Usage:"
-    echo "  LD_PRELOAD=${OUT} CCCL_ALLOC_DEBUG=1 python3 -c 'import torch; t=torch.zeros(1024, device=\"cuda\")'"
-    echo ""
-    echo "In computility-run.yaml, add to env:"
-    echo "  - name: LD_PRELOAD"
-    echo "    value: /workspace/qwen3_6_scripts/cccl_preload/libcccl_allocator.so"
-    echo "  - name: PYTORCH_CUDA_ALLOC_CONF"
-    echo "    value: expandable_segments:True"
+    echo "Test:"
+    echo "  LD_PRELOAD=${OUT} CCCL_ALLOC_DEBUG=1 \\"
+    echo "  PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True \\"
+    echo "  python3 verify_preload.py"
 else
     echo "[build] FAILED"
     exit 1

qwen3_6_scripts/cccl_preload/cccl_allocator_preload.cu

@@ -1,409 +1,90 @@
 /*
  * cccl_allocator_preload.cu
  *
- * LD_PRELOAD .so that replaces PyTorch's CUDA memory allocator with
- * CUB's CachingDeviceAllocator (extracted from CCCL upstream).
+ * LD_PRELOAD .so — CUB CachingDeviceAllocator from CCCL upstream.
+ * Full dependency chain (288 files) extracted into include/.
  *
- * Purpose: CoreX's CUDACachingAllocator.cpp:545 asserts
- * "expandable segment not supported". Instead of patching libtorch,
- * we intercept cudaMalloc/cudaFree at the dynamic linker level and
- * route them through CUB's battle-tested caching allocator.
+ * Intercepts cudaMalloc/cudaFree, routes through CUB's geometric-bin
+ * caching allocator. Strips expandable_segments from
+ * PYTORCH_CUDA_ALLOC_CONF before libtorch reads it.
  *
- * Source: cccl_upstream/cub/cub/util_allocator.cuh
- * License: BSD-3 (NVIDIA/CUB)
- *
- * Build (on BI-V100 with CoreX clang++):
- *   bash build_cccl_preload.sh
- *
- * Usage:
- *   LD_PRELOAD=/workspace/qwen3_6_scripts/cccl_preload/libcccl_allocator.so \
- *   CCCL_ALLOC_DEBUG=0 \
- *   PYTORCH_CUDA_ALLOC_CONF=max_split_size_mb:512 \
- *   python3 -m vllm.entrypoints.openai.api_server ...
+ * Source: CCCL cub/cub/util_allocator.cuh (BSD-3, NVIDIA)
+ * Build:  bash build_cccl_preload.sh
  */
 
-#include <cuda_runtime_api.h>
+/* ---- CCCL include chain (288 files from cccl_upstream) ---- */
+#include <cub/util_allocator.cuh>
+
+/* ---- System ---- */
 #include <dlfcn.h>
 #include <cstdio>
 #include <cstdlib>
 #include <cstring>
-#include <map>
-#include <mutex>
-#include <set>
+#include <string>
 
 /* ========================================================================
- * CUB CachingDeviceAllocator — extracted from CCCL
- * cccl_upstream/cub/cub/util_allocator.cuh
+ * Configuration for BI-V100 (32GB × 4 cards)
  *
- * All CUB/CCCL macro dependencies replaced with plain C++.
+ * 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 bool g_cccl_debug = false;
+static constexpr unsigned int ALLOC_BIN_GROWTH   = 2;
+static constexpr unsigned int ALLOC_MIN_BIN      = 8;   /* 256 bytes */
+static constexpr unsigned int ALLOC_MAX_BIN      = 32;  /* 4 GB */
+static constexpr size_t       ALLOC_MAX_CACHED   = (size_t)8 * 1024 * 1024 * 1024; /* 8GB */
 
-#define CcclDebug(e) (e)
-#define CcclLog(...)               \
-    do {                           \
-        if (g_cccl_debug) {        \
-            fprintf(stderr, "[cccl_alloc] "); \
-            fprintf(stderr, __VA_ARGS__); \
-        }                          \
-    } while (0)
-
-struct CachingDeviceAllocator
-{
-    static constexpr unsigned int INVALID_BIN = (unsigned int) -1;
-    static constexpr size_t INVALID_SIZE = (size_t) -1;
-    static constexpr int INVALID_DEVICE_ORDINAL = -1;
-
-    struct BlockDescriptor
-    {
-        void* d_ptr;
-        size_t bytes;
-        unsigned int bin;
-        int device;
-        cudaStream_t associated_stream;
-        cudaEvent_t ready_event;
-
-        BlockDescriptor(void* d_ptr_, int device_)
-            : d_ptr(d_ptr_), bytes(0), bin(INVALID_BIN), device(device_),
-              associated_stream(nullptr), ready_event(nullptr) {}
-
-        BlockDescriptor(int device_)
-            : d_ptr(nullptr), bytes(0), bin(INVALID_BIN), device(device_),
-              associated_stream(nullptr), ready_event(nullptr) {}
-
-        static bool PtrCompare(const BlockDescriptor& a, const BlockDescriptor& b) {
-            return (a.device == b.device) ? (a.d_ptr < b.d_ptr) : (a.device < b.device);
-        }
-        static bool SizeCompare(const BlockDescriptor& a, const BlockDescriptor& b) {
-            return (a.device == b.device) ? (a.bytes < b.bytes) : (a.device < b.device);
-        }
-    };
-
-    using Compare = bool (*)(const BlockDescriptor&, const BlockDescriptor&);
-
-    struct TotalBytes { size_t free; size_t live; TotalBytes() : free(0), live(0) {} };
-
-    using CachedBlocks = std::multiset<BlockDescriptor, Compare>;
-    using BusyBlocks = std::multiset<BlockDescriptor, Compare>;
-    using GpuCachedBytes = std::map<int, TotalBytes>;
-
-    static unsigned int IntPow(unsigned int base, unsigned int exp) {
-        unsigned int retval = 1;
-        while (exp > 0) {
-            if (exp & 1) retval *= base;
-            base *= base;
-            exp >>= 1;
-        }
-        return retval;
-    }
-
-    void NearestPowerOf(unsigned int& power, size_t& rounded_bytes,
-                        unsigned int base, size_t value) {
-        power = 0;
-        rounded_bytes = 1;
-        if (value * base < value) {
-            power = sizeof(size_t) * 8;
-            rounded_bytes = size_t(0) - 1;
-            return;
-        }
-        while (rounded_bytes < value) {
-            rounded_bytes *= base;
-            power++;
-        }
-    }
-
-    std::mutex mutex;
-    unsigned int bin_growth;
-    unsigned int min_bin;
-    unsigned int max_bin;
-    size_t min_bin_bytes;
-    size_t max_bin_bytes;
-    size_t max_cached_bytes;
-    bool skip_cleanup;
-    GpuCachedBytes cached_bytes;
-    CachedBlocks cached_blocks;
-    BusyBlocks live_blocks;
-
-    /*
-     * Constructor tuned for BI-V100 (32GB per card, 4 cards):
-     *   bin_growth=8, min_bin=3 (512B), max_bin=13 (~550MB)
-     *   max_cached_bytes = 4GB per device (reasonable for 32GB card)
-     *
-     * This replaces PyTorch's expandable_segments with a proven
-     * geometric-bin caching strategy from CUB/CCCL.
-     */
-    CachingDeviceAllocator()
-        : bin_growth(8)
-        , min_bin(3)           /* 8^3 = 512B minimum allocation */
-        , max_bin(13)          /* 8^13 = ~550MB maximum cached bin */
-        , min_bin_bytes(IntPow(8, 3))
-        , max_bin_bytes(IntPow(8, 13))
-        , max_cached_bytes((size_t)4 * 1024 * 1024 * 1024)  /* 4GB per device */
-        , skip_cleanup(true)   /* CoreX may tear down CUDA before our dtor */
-        , cached_blocks(BlockDescriptor::SizeCompare)
-        , live_blocks(BlockDescriptor::PtrCompare)
-    {
-        CcclLog("CachingDeviceAllocator init: bin_growth=%u min_bin=%u "
-                "max_bin=%u max_cached=%.1fGB\n",
-                bin_growth, min_bin, max_bin,
-                (double)max_cached_bytes / (1024.0*1024.0*1024.0));
-    }
-
-    /* ---- Real cudaMalloc/cudaFree via dlsym(RTLD_NEXT) ---- */
-    using RealMalloc_t = cudaError_t (*)(void**, size_t);
-    using RealFree_t = cudaError_t (*)(void*);
-
-    static RealMalloc_t get_real_malloc() {
-        static RealMalloc_t fn = (RealMalloc_t)dlsym(RTLD_NEXT, "cudaMalloc");
-        return fn;
-    }
-    static RealFree_t get_real_free() {
-        static RealFree_t fn = (RealFree_t)dlsym(RTLD_NEXT, "cudaFree");
-        return fn;
-    }
-
-    cudaError_t DeviceAllocate(int device, void** d_ptr, size_t bytes,
-                               cudaStream_t active_stream = nullptr)
-    {
-        *d_ptr = nullptr;
-        int entrypoint_device = INVALID_DEVICE_ORDINAL;
-        cudaError_t error = cudaSuccess;
-
-        if (device == INVALID_DEVICE_ORDINAL) {
-            error = cudaGetDevice(&entrypoint_device);
-            if (error != cudaSuccess) return error;
-            device = entrypoint_device;
-        }
-
-        bool found = false;
-        BlockDescriptor search_key(device);
-        search_key.associated_stream = active_stream;
-        NearestPowerOf(search_key.bin, search_key.bytes, bin_growth, bytes);
-
-        if (search_key.bin > max_bin) {
-            search_key.bin = INVALID_BIN;
-            search_key.bytes = bytes;
-        } else {
-            mutex.lock();
-            if (search_key.bin < min_bin) {
-                search_key.bin = min_bin;
-                search_key.bytes = min_bin_bytes;
-            }
-
-            CachedBlocks::iterator block_itr = cached_blocks.lower_bound(search_key);
-            while ((block_itr != cached_blocks.end()) &&
-                   (block_itr->device == device) &&
-                   (block_itr->bin == search_key.bin))
-            {
-                bool is_reusable = false;
-                if (active_stream == block_itr->associated_stream) {
-                    is_reusable = true;
-                } else {
-                    cudaError_t event_status = cudaEventQuery(block_itr->ready_event);
-                    if (event_status != cudaErrorNotReady) {
-                        is_reusable = true;
-                    }
-                }
-
-                if (is_reusable) {
-                    found = true;
-                    search_key = *block_itr;
-                    search_key.associated_stream = active_stream;
-                    live_blocks.insert(search_key);
-                    cached_bytes[device].free -= search_key.bytes;
-                    cached_bytes[device].live += search_key.bytes;
-
-                    CcclLog("reuse %p (%zu bytes) dev=%d\n",
-                            search_key.d_ptr, search_key.bytes, device);
-                    cached_blocks.erase(block_itr);
-                    break;
-                }
-                block_itr++;
-            }
-            mutex.unlock();
-        }
-
-        if (!found) {
-            if (device != entrypoint_device) {
-                if (entrypoint_device == INVALID_DEVICE_ORDINAL)
-                    cudaGetDevice(&entrypoint_device);
-                cudaSetDevice(device);
-            }
-
-            /* Use real cudaMalloc, not ourselves */
-            error = get_real_malloc()(&search_key.d_ptr, search_key.bytes);
-
-            if (error == cudaErrorMemoryAllocation) {
-                CcclLog("OOM for %zu bytes on dev=%d, freeing cache...\n",
-                        search_key.bytes, device);
-                cudaGetLastError();  /* reset */
-
-                mutex.lock();
-                BlockDescriptor free_key(device);
-                CachedBlocks::iterator block_itr = cached_blocks.lower_bound(free_key);
-                while ((block_itr != cached_blocks.end()) &&
-                       (block_itr->device == device))
-                {
-                    error = get_real_free()(block_itr->d_ptr);
-                    if (error != cudaSuccess) break;
-                    cudaEventDestroy(block_itr->ready_event);
-                    cached_bytes[device].free -= block_itr->bytes;
-                    block_itr = cached_blocks.erase(block_itr);
-                }
-                mutex.unlock();
-
-                if (error != cudaSuccess) return error;
-                error = get_real_malloc()(&search_key.d_ptr, search_key.bytes);
-                if (error != cudaSuccess) return error;
-            } else if (error != cudaSuccess) {
-                return error;
-            }
-
-            cudaEventCreateWithFlags(&search_key.ready_event, cudaEventDisableTiming);
-
-            mutex.lock();
-            live_blocks.insert(search_key);
-            cached_bytes[device].live += search_key.bytes;
-            mutex.unlock();
-
-            CcclLog("alloc %p (%zu bytes, bin=%u) dev=%d\n",
-                    search_key.d_ptr, search_key.bytes, search_key.bin, device);
-
-            if ((entrypoint_device != INVALID_DEVICE_ORDINAL) &&
-                (entrypoint_device != device))
-                cudaSetDevice(entrypoint_device);
-        }
-
-        *d_ptr = search_key.d_ptr;
-        return cudaSuccess;
-    }
-
-    cudaError_t DeviceAllocate(void** d_ptr, size_t bytes,
-                               cudaStream_t active_stream = nullptr) {
-        return DeviceAllocate(INVALID_DEVICE_ORDINAL, d_ptr, bytes, active_stream);
-    }
-
-    cudaError_t DeviceFree(int device, void* d_ptr)
-    {
-        int entrypoint_device = INVALID_DEVICE_ORDINAL;
-        cudaError_t error = cudaSuccess;
-
-        if (d_ptr == nullptr) return cudaSuccess;
-
-        if (device == INVALID_DEVICE_ORDINAL) {
-            error = cudaGetDevice(&entrypoint_device);
-            if (error != cudaSuccess) return error;
-            device = entrypoint_device;
-        }
-
-        mutex.lock();
-        bool recached = false;
-        BlockDescriptor search_key(d_ptr, device);
-        BusyBlocks::iterator block_itr = live_blocks.find(search_key);
-
-        if (block_itr != live_blocks.end()) {
-            search_key = *block_itr;
-            live_blocks.erase(block_itr);
-            cached_bytes[device].live -= search_key.bytes;
-
-            if ((search_key.bin != INVALID_BIN) &&
-                (cached_bytes[device].free + search_key.bytes <= max_cached_bytes))
-            {
-                recached = true;
-                cached_blocks.insert(search_key);
-                cached_bytes[device].free += search_key.bytes;
-                CcclLog("cache %p (%zu bytes) dev=%d\n",
-                        d_ptr, search_key.bytes, device);
-            }
-        }
-        mutex.unlock();
-
-        if (device != entrypoint_device) {
-            if (entrypoint_device == INVALID_DEVICE_ORDINAL)
-                cudaGetDevice(&entrypoint_device);
-            cudaSetDevice(device);
-        }
-
-        if (recached) {
-            cudaEventRecord(search_key.ready_event, search_key.associated_stream);
-        } else {
-            /* Not tracked or cache full — real free */
-            CcclLog("free %p dev=%d (not cached)\n", d_ptr, device);
-            error = get_real_free()(d_ptr);
-            if (block_itr != live_blocks.end())
-                cudaEventDestroy(search_key.ready_event);
-        }
-
-        if ((entrypoint_device != INVALID_DEVICE_ORDINAL) &&
-            (entrypoint_device != device))
-            cudaSetDevice(entrypoint_device);
-
-        return error;
-    }
-
-    cudaError_t DeviceFree(void* d_ptr) {
-        return DeviceFree(INVALID_DEVICE_ORDINAL, d_ptr);
-    }
-
-    cudaError_t FreeAllCached()
-    {
-        cudaError_t error = cudaSuccess;
-        int entrypoint_device = INVALID_DEVICE_ORDINAL;
-        int current_device = INVALID_DEVICE_ORDINAL;
-
-        mutex.lock();
-        while (!cached_blocks.empty()) {
-            CachedBlocks::iterator begin = cached_blocks.begin();
-            if (entrypoint_device == INVALID_DEVICE_ORDINAL)
-                cudaGetDevice(&entrypoint_device);
-            if (begin->device != current_device) {
-                cudaSetDevice(begin->device);
-                current_device = begin->device;
-            }
-            get_real_free()(begin->d_ptr);
-            cudaEventDestroy(begin->ready_event);
-            cached_bytes[current_device].free -= begin->bytes;
-            cached_blocks.erase(begin);
-        }
-        mutex.unlock();
-
-        if (entrypoint_device != INVALID_DEVICE_ORDINAL)
-            cudaSetDevice(entrypoint_device);
-        return error;
-    }
-};
-
-/* ========================================================================
- * Global singleton + LD_PRELOAD intercepts
- * ======================================================================== */
-
-static CachingDeviceAllocator& get_allocator() {
-    static CachingDeviceAllocator instance;
+/* ---- Global allocator singleton ---- */
+static cub::CachingDeviceAllocator& get_allocator() {
+    static cub::CachingDeviceAllocator instance(
+        ALLOC_BIN_GROWTH,
+        ALLOC_MIN_BIN,
+        ALLOC_MAX_BIN,
+        ALLOC_MAX_CACHED,
+        true  /* skip_cleanup: CoreX may tear down CUDA before our dtor */
+    );
     return instance;
 }
 
 static bool g_preload_active = false;
+static bool g_debug = false;
 
-/* Called once at .so load time */
+/* ---- Real cudaMalloc/cudaFree via dlsym(RTLD_NEXT) ---- */
+using RealMalloc_t = cudaError_t (*)(void**, size_t);
+using RealFree_t   = cudaError_t (*)(void*);
+
+static RealMalloc_t get_real_malloc() {
+    static RealMalloc_t fn = (RealMalloc_t)dlsym(RTLD_NEXT, "cudaMalloc");
+    return fn;
+}
+static RealFree_t get_real_free() {
+    static RealFree_t fn = (RealFree_t)dlsym(RTLD_NEXT, "cudaFree");
+    return fn;
+}
+
+/* ========================================================================
+ * Constructor: runs at LD_PRELOAD load time
+ * ======================================================================== */
 __attribute__((constructor))
 static void cccl_preload_init() {
     const char* debug_env = getenv("CCCL_ALLOC_DEBUG");
-    g_cccl_debug = (debug_env && atoi(debug_env) > 0);
+    g_debug = (debug_env && atoi(debug_env) > 0);
 
     const char* disable_env = getenv("CCCL_ALLOC_DISABLE");
     if (disable_env && atoi(disable_env) > 0) {
         fprintf(stderr, "[cccl_alloc] DISABLED by CCCL_ALLOC_DISABLE=1\n");
-        g_preload_active = false;
         return;
     }
 
-    /* Strip expandable_segments from PYTORCH_CUDA_ALLOC_CONF
-     * so CoreX's allocator doesn't hit the assert.
-     * We handle the caching ourselves. */
+    /* Strip expandable_segments from PYTORCH_CUDA_ALLOC_CONF */
     const char* alloc_conf = getenv("PYTORCH_CUDA_ALLOC_CONF");
     if (alloc_conf) {
-        /* Build a new conf string without expandable_segments */
         std::string conf(alloc_conf);
         std::string clean;
         size_t pos = 0;
@@ -411,49 +92,51 @@ static void cccl_preload_init() {
             size_t comma = conf.find(',', pos);
             if (comma == std::string::npos) comma = conf.size();
             std::string token = conf.substr(pos, comma - pos);
-            /* Skip expandable_segments:* */
             if (token.find("expandable_segments") == std::string::npos) {
                 if (!clean.empty()) clean += ",";
                 clean += token;
             }
             pos = comma + 1;
         }
-        if (clean.empty()) {
+        if (clean.empty())
             unsetenv("PYTORCH_CUDA_ALLOC_CONF");
-        } else {
+        else
             setenv("PYTORCH_CUDA_ALLOC_CONF", clean.c_str(), 1);
-        }
-        fprintf(stderr, "[cccl_alloc] stripped expandable_segments from "
-                "PYTORCH_CUDA_ALLOC_CONF: \"%s\" -> \"%s\"\n",
+
+        fprintf(stderr, "[cccl_alloc] PYTORCH_CUDA_ALLOC_CONF: \"%s\" -> \"%s\"\n",
                 alloc_conf, clean.empty() ? "(unset)" : clean.c_str());
     }
 
-    /* Force-initialize the allocator singleton */
-    (void)get_allocator();
+    /* Initialize allocator */
+    auto& alloc = get_allocator();
+    if (g_debug) {
+        alloc.debug = true;
+    }
+
     g_preload_active = true;
-    fprintf(stderr, "[cccl_alloc] LD_PRELOAD active — CUB CachingDeviceAllocator "
-            "replacing cudaMalloc/cudaFree\n");
+    fprintf(stderr,
+        "[cccl_alloc] LD_PRELOAD active — CUB CachingDeviceAllocator "
+        "(growth=%u, bins=[%u..%u], max_cached=%.1fGB)\n",
+        ALLOC_BIN_GROWTH, ALLOC_MIN_BIN, ALLOC_MAX_BIN,
+        (double)ALLOC_MAX_CACHED / (1024.0*1024.0*1024.0));
 }
 
-/* ---- cudaMalloc intercept ---- */
+/* ========================================================================
+ * cudaMalloc / cudaFree intercepts
+ * ======================================================================== */
+
 extern "C" cudaError_t cudaMalloc(void** devPtr, size_t size)
 {
     if (!g_preload_active) {
-        /* Fallback to real cudaMalloc during init or if disabled */
-        static auto real_fn = (CachingDeviceAllocator::RealMalloc_t)
-            dlsym(RTLD_NEXT, "cudaMalloc");
-        return real_fn(devPtr, size);
+        return get_real_malloc()(devPtr, size);
     }
     return get_allocator().DeviceAllocate(devPtr, size);
 }
 
-/* ---- cudaFree intercept ---- */
 extern "C" cudaError_t cudaFree(void* devPtr)
 {
     if (!g_preload_active || devPtr == nullptr) {
-        static auto real_fn = (CachingDeviceAllocator::RealFree_t)
-            dlsym(RTLD_NEXT, "cudaFree");
-        return real_fn(devPtr);
+        return get_real_free()(devPtr);
     }
     return get_allocator().DeviceFree(devPtr);
 }

qwen3_6_scripts/cccl_preload_allocator.cu

@@ -1,405 +0,0 @@
-// cccl_preload_allocator.cu — LD_PRELOAD interception of cudaMalloc/cudaFree
-//
-// Replaces the default cudaMalloc/cudaFree with CCCL CUB CachingDeviceAllocator.
-// This eliminates the CUDA driver's allocation overhead (cudaMalloc is slow on
-// BI-V100: ~2-50ms per call) by reusing freed blocks from a bin-based cache.
-//
-// Build on BI-V100:
-//   /usr/local/corex-3.2.3/bin/clang++ -std=c++17 -O3 -shared -fPIC \
-//     --cuda-path=/usr/local/corex-3.2.3 --cuda-gpu-arch=ivcore10 \
-//     --no-cuda-version-check -D_GLIBCXX_USE_CXX11_ABI=0 \
-//     -I/usr/local/corex-3.2.3/include \
-//     cccl_preload_allocator.cu \
-//     -L/usr/local/corex-3.2.3/lib64 -lcudart -ldl \
-//     -o cccl_preload_allocator.so
-//
-// Usage:
-//   LD_PRELOAD=/workspace/qwen3_6_scripts/cccl_preload_allocator.so python3 -m vllm.entrypoints.openai.api_server ...
-//
-// Tuning (env vars):
-//   CCCL_ALLOC_BIN_GROWTH=8      Geometric growth factor (default 8)
-//   CCCL_ALLOC_MIN_BIN=3         Min bin exponent (default 3 → 512B)
-//   CCCL_ALLOC_MAX_BIN=13        Max bin exponent (default 13 → 512MB for growth=8; was 7→2MB)
-//   CCCL_ALLOC_MAX_CACHED_MB=4096  Max cached bytes per device in MB (default 4096=4GB)
-//   CCCL_ALLOC_DEBUG=0           Print alloc/free events (default 0)
-//
-// Design notes:
-// - Only intercepts cudaMalloc and cudaFree (the synchronous variants).
-// - cudaMallocAsync/cudaFreeAsync are NOT intercepted (PyTorch on BI-V100
-//   doesn't use them; the corex runtime may not support them).
-// - Thread-safe via CUB's internal mutex.
-// - Stream association: all allocations use the default stream (nullptr).
-//   PyTorch's CUDACachingAllocator handles stream ordering itself, so we
-//   don't need to track streams here.
-// - Large allocations (> max_bin_bytes) pass through to real cudaMalloc.
-// - The allocator is process-global (static singleton).
-
-#include <cstdio>
-#include <cstdlib>
-#include <cstring>
-#include <dlfcn.h>
-#include <mutex>
-#include <map>
-#include <set>
-#include <atomic>
-
-// We inline the essential logic from CUB CachingDeviceAllocator rather than
-// #include it, because the corex toolchain may not have full CCCL headers
-// installed, and we need to link against corex's cudart, not NVIDIA's.
-
-// Forward declare the real CUDA functions we'll dlsym
-typedef int cudaError_t;
-static constexpr cudaError_t cudaSuccess = 0;
-
-typedef void* cudaStream_t;
-typedef void* cudaEvent_t;
-
-// Real function pointers (resolved via dlsym on first call)
-using cudaMalloc_fn = cudaError_t(*)(void**, size_t);
-using cudaFree_fn = cudaError_t(*)(void*);
-using cudaGetDevice_fn = cudaError_t(*)(int*);
-using cudaEventCreate_fn = cudaError_t(*)(cudaEvent_t*);
-using cudaEventRecord_fn = cudaError_t(*)(cudaEvent_t, cudaStream_t);
-using cudaEventQuery_fn = cudaError_t(*)(cudaEvent_t);
-using cudaEventDestroy_fn = cudaError_t(*)(cudaEvent_t);
-using cudaEventSynchronize_fn = cudaError_t(*)(cudaEvent_t);
-
-static cudaMalloc_fn  real_cudaMalloc  = nullptr;
-static cudaFree_fn    real_cudaFree    = nullptr;
-static cudaGetDevice_fn real_cudaGetDevice = nullptr;
-static cudaEventCreate_fn real_cudaEventCreate = nullptr;
-static cudaEventRecord_fn real_cudaEventRecord = nullptr;
-static cudaEventQuery_fn  real_cudaEventQuery  = nullptr;
-static cudaEventDestroy_fn real_cudaEventDestroy = nullptr;
-static cudaEventSynchronize_fn real_cudaEventSynchronize = nullptr;
-
-static std::once_flag resolve_flag;
-
-static void resolve_real_functions() {
-    real_cudaMalloc = (cudaMalloc_fn)dlsym(RTLD_NEXT, "cudaMalloc");
-    real_cudaFree = (cudaFree_fn)dlsym(RTLD_NEXT, "cudaFree");
-    real_cudaGetDevice = (cudaGetDevice_fn)dlsym(RTLD_NEXT, "cudaGetDevice");
-    real_cudaEventCreate = (cudaEventCreate_fn)dlsym(RTLD_NEXT, "cudaEventCreate");
-    real_cudaEventRecord = (cudaEventRecord_fn)dlsym(RTLD_NEXT, "cudaEventRecord");
-    real_cudaEventQuery = (cudaEventQuery_fn)dlsym(RTLD_NEXT, "cudaEventQuery");
-    real_cudaEventDestroy = (cudaEventDestroy_fn)dlsym(RTLD_NEXT, "cudaEventDestroy");
-    real_cudaEventSynchronize = (cudaEventSynchronize_fn)dlsym(RTLD_NEXT, "cudaEventSynchronize");
-
-    if (!real_cudaMalloc || !real_cudaFree) {
-        fprintf(stderr, "[CCCL_PRELOAD] FATAL: cannot resolve cudaMalloc/cudaFree via dlsym\n");
-        abort();
-    }
-}
-
-// ============================================================================
-// Simplified CachingDeviceAllocator (from CCCL cub/util_allocator.cuh)
-// Stripped to essentials: no debug logging macros, no CCCL config dependencies
-// ============================================================================
-
-struct BlockDescriptor {
-    void*         d_ptr;
-    size_t        bytes;
-    unsigned int  bin;
-    int           device;
-    cudaStream_t  associated_stream;
-    cudaEvent_t   ready_event;
-
-    BlockDescriptor(void* p, int dev)
-        : d_ptr(p), bytes(0), bin(~0u), device(dev),
-          associated_stream(nullptr), ready_event(nullptr) {}
-
-    BlockDescriptor(int dev)
-        : d_ptr(nullptr), bytes(0), bin(~0u), device(dev),
-          associated_stream(nullptr), ready_event(nullptr) {}
-
-    static bool PtrCompare(const BlockDescriptor& a, const BlockDescriptor& b) {
-        return (a.device == b.device) ? (a.d_ptr < b.d_ptr) : (a.device < b.device);
-    }
-    static bool SizeCompare(const BlockDescriptor& a, const BlockDescriptor& b) {
-        return (a.device == b.device) ? (a.bytes < b.bytes) : (a.device < b.device);
-    }
-};
-
-using Compare = bool(*)(const BlockDescriptor&, const BlockDescriptor&);
-using CachedBlocks = std::multiset<BlockDescriptor, Compare>;
-using BusyBlocks = std::multiset<BlockDescriptor, Compare>;
-
-struct DeviceBytes { size_t free = 0; size_t live = 0; };
-
-struct CachingAllocator {
-    std::mutex     mtx;
-    unsigned int   bin_growth;
-    unsigned int   min_bin;
-    unsigned int   max_bin;
-    size_t         min_bin_bytes;
-    size_t         max_bin_bytes;
-    size_t         max_cached_bytes;
-    bool           debug;
-    CachedBlocks   cached_blocks;
-    BusyBlocks     live_blocks;
-    std::map<int, DeviceBytes> cached_bytes;
-
-    // Stats
-    std::atomic<uint64_t> stat_hits{0};
-    std::atomic<uint64_t> stat_misses{0};
-    std::atomic<uint64_t> stat_frees{0};
-    std::atomic<uint64_t> stat_bypasses{0};
-
-    static unsigned int IntPow(unsigned int base, unsigned int exp) {
-        unsigned int r = 1;
-        while (exp > 0) {
-            if (exp & 1) r *= base;
-            base *= base;
-            exp >>= 1;
-        }
-        return r;
-    }
-
-    void NearestPowerOf(unsigned int& power, size_t& rounded,
-                        unsigned int base, size_t value) {
-        power = 0; rounded = 1;
-        if (value * base < value) {
-            power = sizeof(size_t) * 8;
-            rounded = size_t(-1);
-            return;
-        }
-        while (rounded < value) { rounded *= base; power++; }
-    }
-
-    CachingAllocator()
-        : cached_blocks(BlockDescriptor::SizeCompare),
-          live_blocks(BlockDescriptor::PtrCompare) {
-        // Read config from env
-        auto env_or = [](const char* name, int def) -> int {
-            const char* v = getenv(name);
-            return v ? atoi(v) : def;
-        };
-
-        bin_growth      = env_or("CCCL_ALLOC_BIN_GROWTH", 8);
-        min_bin         = env_or("CCCL_ALLOC_MIN_BIN", 3);
-        max_bin         = env_or("CCCL_ALLOC_MAX_BIN", 13);
-        int max_mb      = env_or("CCCL_ALLOC_MAX_CACHED_MB", 4096);
-        debug           = env_or("CCCL_ALLOC_DEBUG", 0) != 0;
-        min_bin_bytes   = IntPow(bin_growth, min_bin);
-        max_bin_bytes   = IntPow(bin_growth, max_bin);
-        max_cached_bytes = (size_t)max_mb * 1024ULL * 1024ULL;
-
-        fprintf(stderr, "[CCCL_PRELOAD] CachingDeviceAllocator: growth=%u "
-                "bins=[%u..%u] bin_bytes=[%zu..%zu] max_cached=%zuMB\n",
-                bin_growth, min_bin, max_bin,
-                min_bin_bytes, max_bin_bytes, max_cached_bytes / (1024*1024));
-    }
-
-    ~CachingAllocator() {
-        fprintf(stderr, "[CCCL_PRELOAD] Stats: hits=%lu misses=%lu frees=%lu bypasses=%lu\n",
-                stat_hits.load(), stat_misses.load(),
-                stat_frees.load(), stat_bypasses.load());
-        // Free all cached blocks
-        for (auto& b : cached_blocks) {
-            if (b.ready_event) real_cudaEventDestroy(b.ready_event);
-            real_cudaFree(b.d_ptr);
-        }
-    }
-
-    cudaError_t Allocate(void** d_ptr, size_t bytes) {
-        std::call_once(resolve_flag, resolve_real_functions);
-        *d_ptr = nullptr;
-
-        // Get current device
-        int device = 0;
-        if (real_cudaGetDevice) real_cudaGetDevice(&device);
-
-        // Bin classification
-        unsigned int bin;
-        size_t rounded_bytes;
-        bool oversized = false;
-
-        if (bytes > max_bin_bytes) {
-            // Too large for caching — pass through
-            bin = max_bin + 1;
-            rounded_bytes = bytes;
-            oversized = true;
-        } else {
-            NearestPowerOf(bin, rounded_bytes, bin_growth, bytes);
-            if (bin < min_bin) {
-                bin = min_bin;
-                rounded_bytes = min_bin_bytes;
-            }
-        }
-
-        BlockDescriptor search_key(device);
-        search_key.bytes = rounded_bytes;
-        search_key.bin = bin;
-
-        // Lock
-        std::lock_guard<std::mutex> lock(mtx);
-
-        if (!oversized) {
-            // Search cached blocks for a match
-            auto range = cached_blocks.equal_range(search_key);
-            for (auto it = range.first; it != range.second; ++it) {
-                if (it->device == device && it->bin == bin) {
-                    // Check if the stream work has completed
-                    bool ready = true;
-                    if (it->ready_event) {
-                        cudaError_t ev_status = real_cudaEventQuery(it->ready_event);
-                        if (ev_status != cudaSuccess) {
-                            // Event not ready — try to synchronize briefly
-                            // For BI-V100 with enforce_eager, events should be ready
-                            real_cudaEventSynchronize(it->ready_event);
-                        }
-                        real_cudaEventDestroy(it->ready_event);
-                    }
-
-                    // Reuse this block
-                    search_key.d_ptr = it->d_ptr;
-                    search_key.bytes = it->bytes;
-                    live_blocks.insert(search_key);
-                    cached_bytes[device].free -= it->bytes;
-                    cached_bytes[device].live += it->bytes;
-                    cached_blocks.erase(it);
-
-                    *d_ptr = search_key.d_ptr;
-                    stat_hits++;
-
-                    if (debug) {
-                        fprintf(stderr, "[CCCL_PRELOAD] HIT  dev=%d bin=%u "
-                                "req=%zu alloc=%zu ptr=%p\n",
-                                device, bin, bytes, search_key.bytes, *d_ptr);
-                    }
-                    return cudaSuccess;
-                }
-            }
-        }
-
-        // Cache miss — allocate new block
-        cudaError_t err = real_cudaMalloc(&search_key.d_ptr, rounded_bytes);
-
-        // If OOM, try evicting cached blocks and retry
-        if (err != cudaSuccess) {
-            // Free all cached blocks on this device
-            auto it = cached_blocks.begin();
-            while (it != cached_blocks.end()) {
-                if (it->device == device) {
-                    if (it->ready_event) {
-                        real_cudaEventSynchronize(it->ready_event);
-                        real_cudaEventDestroy(it->ready_event);
-                    }
-                    real_cudaFree(it->d_ptr);
-                    cached_bytes[device].free -= it->bytes;
-                    it = cached_blocks.erase(it);
-                } else {
-                    ++it;
-                }
-            }
-            // Retry
-            err = real_cudaMalloc(&search_key.d_ptr, rounded_bytes);
-        }
-
-        if (err != cudaSuccess) {
-            return err;
-        }
-
-        search_key.bytes = rounded_bytes;
-        live_blocks.insert(search_key);
-        cached_bytes[device].live += rounded_bytes;
-
-        *d_ptr = search_key.d_ptr;
-
-        if (oversized) {
-            stat_bypasses++;
-        } else {
-            stat_misses++;
-        }
-
-        if (debug) {
-            fprintf(stderr, "[CCCL_PRELOAD] %s dev=%d bin=%u "
-                    "req=%zu alloc=%zu ptr=%p\n",
-                    oversized ? "PASS" : "MISS",
-                    device, bin, bytes, rounded_bytes, *d_ptr);
-        }
-        return cudaSuccess;
-    }
-
-    cudaError_t Free(void* d_ptr) {
-        std::call_once(resolve_flag, resolve_real_functions);
-
-        if (d_ptr == nullptr) return cudaSuccess;
-
-        int device = 0;
-        if (real_cudaGetDevice) real_cudaGetDevice(&device);
-
-        BlockDescriptor search_key(d_ptr, device);
-
-        std::lock_guard<std::mutex> lock(mtx);
-
-        auto it = live_blocks.find(search_key);
-        if (it == live_blocks.end()) {
-            // Not tracked by us — pass through to real cudaFree
-            return real_cudaFree(d_ptr);
-        }
-
-        search_key.bytes = it->bytes;
-        search_key.bin = it->bin;
-        cached_bytes[device].live -= it->bytes;
-        live_blocks.erase(it);
-
-        stat_frees++;
-
-        // Check if this block is too large or would exceed cache limit
-        bool should_cache = (search_key.bin <= max_bin) &&
-                            (cached_bytes[device].free + search_key.bytes <= max_cached_bytes);
-
-        if (should_cache) {
-            // Record an event so we know when it's safe to reuse
-            if (real_cudaEventCreate) {
-                cudaEvent_t event = nullptr;
-                cudaError_t ev_err = real_cudaEventCreate(&event);
-                if (ev_err == cudaSuccess && real_cudaEventRecord) {
-                    real_cudaEventRecord(event, nullptr);  // default stream
-                    search_key.ready_event = event;
-                }
-            }
-
-            cached_blocks.insert(search_key);
-            cached_bytes[device].free += search_key.bytes;
-
-            if (debug) {
-                fprintf(stderr, "[CCCL_PRELOAD] CACHE dev=%d bin=%u "
-                        "bytes=%zu cached_free=%zu\n",
-                        device, search_key.bin, search_key.bytes,
-                        cached_bytes[device].free);
-            }
-            return cudaSuccess;
-        } else {
-            // Don't cache — actually free
-            if (debug) {
-                fprintf(stderr, "[CCCL_PRELOAD] FREE  dev=%d bin=%u bytes=%zu\n",
-                        device, search_key.bin, search_key.bytes);
-            }
-            return real_cudaFree(d_ptr);
-        }
-    }
-};
-
-// Global singleton
-static CachingAllocator& get_allocator() {
-    static CachingAllocator alloc;
-    return alloc;
-}
-
-// ============================================================================
-// LD_PRELOAD interception points
-// ============================================================================
-
-extern "C" {
-
-cudaError_t cudaMalloc(void** devPtr, size_t size) {
-    return get_allocator().Allocate(devPtr, size);
-}
-
-cudaError_t cudaFree(void* devPtr) {
-    return get_allocator().Free(devPtr);
-}
-
-}  // extern "C"

qwen3_6_scripts/patch_ops.sh

@@ -247,8 +247,15 @@ if source != installed:
 PY
 
 build_stage "compiling CCCL CachingDeviceAllocator LD_PRELOAD module"
-bash ./build_cccl_preload_allocator.sh /workspace/qwen3_6_scripts || \
-    echo "[WARN] CCCL preload allocator build failed — will use default allocator"
+if bash ./cccl_preload/build_cccl_preload.sh /workspace/qwen3_6_scripts; then
+    if [[ -s /workspace/qwen3_6_scripts/libcccl_allocator.so ]]; then
+        echo "[OK] libcccl_allocator.so built successfully"
+    else
+        echo "[WARN] libcccl_allocator.so is empty or missing after build"
+    fi
+else
+    echo "[WARN] CCCL preload allocator build failed — LD_PRELOAD will be ignored at runtime"
+fi
 
 build_stage "compiling CoreX CUDA extensions (moe_index_combine + gdn_chunk_recurrent)"
 if [[ -x /usr/local/corex-3.2.3/bin/clang++ ]]; then

qwen3_6_scripts/test_cccl_preload.sh

@@ -1,121 +0,0 @@
-#!/usr/bin/env bash
-# Quick test for CCCL preload allocator on BI-V100
-# Usage: cd /home/dylan/project_6/qwen3_6_scripts && bash test_cccl_preload.sh
-set -eo pipefail
-
-SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
-SO="${SCRIPT_DIR}/cccl_preload_allocator.so"
-
-echo "=== Step 1: Build ==="
-bash "${SCRIPT_DIR}/build_cccl_preload_allocator.sh" "${SCRIPT_DIR}"
-echo ""
-
-if [[ ! -f "$SO" ]]; then
-    echo "BUILD FAILED: $SO not found"
-    exit 1
-fi
-
-echo "=== Step 2: Basic smoke test (torch.zeros on GPU) ==="
-echo "Without preload:"
-python3 -c "
-import time, torch
-t0=time.time()
-for i in range(100):
-    x=torch.zeros(1024*1024, device='cuda')
-    del x
-    torch.cuda.synchronize()
-print(f'100 alloc+free cycles: {time.time()-t0:.3f}s')
-"
-
-echo ""
-echo "With CCCL preload:"
-CCCL_ALLOC_DEBUG=0 LD_PRELOAD="$SO" python3 -c "
-import time, torch
-t0=time.time()
-for i in range(100):
-    x=torch.zeros(1024*1024, device='cuda')
-    del x
-    torch.cuda.synchronize()
-print(f'100 alloc+free cycles: {time.time()-t0:.3f}s')
-" 2>&1
-
-echo ""
-echo "=== Step 3: Varied sizes (simulating model inference allocations) ==="
-CCCL_ALLOC_DEBUG=0 LD_PRELOAD="$SO" python3 -c "
-import time, torch
-
-# Simulate inference: repeated allocs of same sizes (should hit cache)
-sizes = [512, 4096, 32768, 262144, 1048576, 4194304, 16777216]  # 512B to 16MB
-tensors = []
-
-print('First pass (cold cache):')
-t0 = time.time()
-for s in sizes:
-    x = torch.empty(s // 2, dtype=torch.float16, device='cuda')  # s bytes
-    tensors.append(x)
-t1 = time.time()
-print(f'  {len(sizes)} allocs: {(t1-t0)*1000:.1f}ms')
-
-print('Free all:')
-del tensors
-torch.cuda.synchronize()
-t2 = time.time()
-print(f'  {len(sizes)} frees: {(t2-t1)*1000:.1f}ms')
-
-print('Second pass (warm cache - should be faster):')
-tensors2 = []
-for s in sizes:
-    x = torch.empty(s // 2, dtype=torch.float16, device='cuda')
-    tensors2.append(x)
-t3 = time.time()
-print(f'  {len(sizes)} allocs: {(t3-t2)*1000:.1f}ms')
-
-print('Third pass (reuse same sizes 100x):')
-for _ in range(100):
-    for s in sizes:
-        x = torch.empty(s // 2, dtype=torch.float16, device='cuda')
-        del x
-t4 = time.time()
-print(f'  700 alloc+free: {(t4-t3)*1000:.1f}ms ({(t4-t3)/700*1000000:.0f}μs/op)')
-" 2>&1
-
-echo ""
-echo "=== Step 4: Large allocation test (model weights sized) ==="
-CCCL_ALLOC_DEBUG=0 LD_PRELOAD="$SO" python3 -c "
-import torch
-# Simulate KV cache blocks (typical: 256KB-2MB each)
-blocks = []
-for i in range(100):
-    b = torch.empty(256*1024 // 2, dtype=torch.float16, device='cuda')
-    blocks.append(b)
-print(f'Allocated 100 x 256KB blocks = {100*256/1024:.0f}MB')
-del blocks
-torch.cuda.synchronize()
-print('Freed all blocks')
-# Reallocate (should hit cache)
-blocks2 = []
-for i in range(100):
-    b = torch.empty(256*1024 // 2, dtype=torch.float16, device='cuda')
-    blocks2.append(b)
-print('Re-allocated 100 blocks (from cache)')
-print('OK: large allocation test passed')
-" 2>&1
-
-echo ""
-echo "=== Step 5: Stats output ==="
-CCCL_ALLOC_DEBUG=0 LD_PRELOAD="$SO" python3 -c "
-import torch
-for _ in range(50):
-    x = torch.zeros(1024*1024, device='cuda')
-    del x
-# Stats print on process exit
-" 2>&1
-
-echo ""
-echo "=== DONE ==="
-echo "If all tests passed, add to your launch command:"
-echo "  LD_PRELOAD=$SO python3 -m vllm.entrypoints.openai.api_server ..."
-echo ""
-echo "Or set in computility-run.yaml env:"
-echo "  - name: LD_PRELOAD"
-echo "    value: /workspace/qwen3_6_scripts/cccl_preload_allocator.so"

verify_submission.sh

@@ -93,7 +93,7 @@ if [[ -n "$YAML_PRELOAD" ]]; then
     echo "  ✓ LD_PRELOAD已配置: $YAML_PRELOAD"
     PASS=$((PASS+1))
     # 检查路径与build输出一致
-    BUILD_OUTPUT_DIR=$(grep "build_cccl_preload.sh" qwen3_6_scripts/patch_ops.sh | grep -o '/workspace/[^ ]*' | head -1)
+    BUILD_OUTPUT_DIR=$(grep "build_cccl_preload.sh" qwen3_6_scripts/patch_ops.sh | grep -oP '/workspace/\S+' | head -1 | tr -d ';')
     EXPECTED_SO="${BUILD_OUTPUT_DIR}/libcccl_allocator.so"
     if [[ "$YAML_PRELOAD" == "$EXPECTED_SO" || "$YAML_PRELOAD" == "/workspace/qwen3_6_scripts/libcccl_allocator.so" ]]; then
         echo "  ✓ 路径匹配build输出"