dylanyunlong/project_6
1
0
Fork 0
Code Issues Pull Requests Actions Projects Releases Wiki Activity

Compare commits

base: dylanyunlong:19946d4c720cac6cca4a350e1e83fc5feaf0ec92
Branches Tags
dylanyunlong:main
...
compare: dylanyunlong:ffe764f8b2241154236d9a3640f1c8aa23fa8c0c
Branches Tags
dylanyunlong:main

8 Commits
19946d4c72 ... ffe764f8b2

Author SHA1 Message Date
root
ffe764f8b2 Merge branch 'main' of https://dev.modelhub.org.cn/dylanyunlong/project_6 2026-08-13 09:46:52 +00:00
dylanyunlon
2cbae09d5e feat: CCCL CachingDeviceAllocator preload — 完整依赖链 288 files 
从 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
 2026-08-13 09:45:45 +00:00
dylanyunlon
5449456534 feat: CCCL CachingDeviceAllocator LD_PRELOAD — bypass CoreX expandable_segments ASSERT 
从 CCCL upstream cub/cub/util_allocator.cuh 提取 CachingDeviceAllocator
核心算法，去掉所有 CUB/CCCL 宏依赖，编译为独立 .so。

用 LD_PRELOAD 拦截 cudaMalloc/cudaFree，路由到 CUB 的 geometric-bin
缓存分配器。同时在 constructor 中 strip PYTORCH_CUDA_ALLOC_CONF 里的
expandable_segments 配置，避免 CoreX CUDACachingAllocator.cpp:545 ASSERT。

BI-V100 调优参数:
  bin_growth=8, min_bin=3 (512B), max_bin=13 (~550MB)
  max_cached_bytes=4GB per device (32GB卡的合理上限)

真机测试步骤:
  1. bash build_cccl_preload.sh
  2. LD_PRELOAD=./libcccl_allocator.so CCCL_ALLOC_DEBUG=1 \
     PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True \
     python3 verify_preload.py
 2026-08-13 09:45:45 +00:00
dylanyunlon
1e8d1ddfa3 fix: remove expandable_segments — CoreX CUDACachingAllocator不支持 
BI-V100 CoreX PyTorch的CUDACachingAllocator.cpp:545没有实现
expandable segment特性，导致模型加载阶段(VocabParallelEmbedding)
直接INTERNAL ASSERT FAILED崩溃。

替换为max_split_size_mb:512减少内存碎片化。
 2026-08-13 09:43:35 +00:00
Claude
bf0ad625c3 feat(CCCL): LD_PRELOAD CachingDeviceAllocator — intercept cudaMalloc/cudaFree 
Route C: replace PyTorch's cudaMalloc/cudaFree with CCCL CUB's
CachingDeviceAllocator via LD_PRELOAD. Eliminates driver-level allocation
overhead by reusing freed GPU memory from a bin-based cache.

Based on cccl_upstream/cub/cub/util_allocator.cuh (901 lines).
Self-contained .so with no CCCL header dependencies at compile time.

Files:
- cccl_preload_allocator.cu: the allocator (405 lines)
- build_cccl_preload_allocator.sh: build script (corex clang++ or g++ fallback)
- test_cccl_preload.sh: smoke test suite for BI-V100
- patch_ops.sh: build during docker build
- computility-run.yaml: LD_PRELOAD env var for runtime

Config via env:
  CCCL_ALLOC_BIN_GROWTH=8, MIN_BIN=3, MAX_BIN=13, MAX_CACHED_MB=4096

Test on real machine:
  cd qwen3_6_scripts && bash test_cccl_preload.sh
 2026-08-13 09:43:35 +00:00
root
1aa6e16d4e build: trigger rebuild after merge reconciliation 2026-08-13 09:43:18 +00:00
root
93fb9bd8b7 build: trigger rebuild after merge reconciliation 2026-08-13 09:43:18 +00:00
dylanyunlon
93e498197a fix: thread_local reentrant guard — prevent cudaMalloc infinite recursion 
CUB CachingDeviceAllocator 内部在 cache miss 时调 cudaMalloc，
被我们的 LD_PRELOAD 再次拦截 → DeviceAllocate → cudaMalloc → 无限递归 → segfault。

加 thread_local bool inside_cub 标志:
  外部调用 → CUB allocator (带缓存)
  CUB 内部调用 → 直接走 dlsym(RTLD_NEXT) 的真实 cudaMalloc
 2026-08-13 09:42:26 +00:00
2 changed files with 4 additions and 5 deletions
Expand all files Collapse all files

3
qwen3_6_scripts/cccl_preload/build_cccl_preload.sh
View File

@@ -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

6
qwen3_6_scripts/cccl_preload/cccl_allocator_preload.cu
View File

@@ -41,10 +41,8 @@ static constexpr unsigned int ALLOC_MAX_BIN = 32; /* 4 GB */
static constexpr size_t ALLOC_MAX_CACHED = (size_t)8 * 1024 * 1024 * 1024; /* 8GB */
/* ---- Global allocator singleton ---- */
using CubAllocator = CUB_NS_QUALIFIER::CachingDeviceAllocator;
static CubAllocator& get_allocator() {
static CubAllocator instance(
static cub::CachingDeviceAllocator& get_allocator() {
static cub::CachingDeviceAllocator instance(
ALLOC_BIN_GROWTH,
ALLOC_MIN_BIN,
ALLOC_MAX_BIN,