diff --git a/qwen3_6_scripts/cccl_preload/build_cccl_preload.sh b/qwen3_6_scripts/cccl_preload/build_cccl_preload.sh index b97fde7c..70838bff 100755 --- a/qwen3_6_scripts/cccl_preload/build_cccl_preload.sh +++ b/qwen3_6_scripts/cccl_preload/build_cccl_preload.sh @@ -72,8 +72,7 @@ COMMON_FLAGS=( # 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 + # Non-NVCC: disable arch magic, wrap namespace to avoid ODR conflicts -D__CUDA_ARCH_LIST__=700 -DCUB_DISABLE_NAMESPACE_MAGIC -DCUB_WRAPPED_NAMESPACE=cccl_preload diff --git a/qwen3_6_scripts/cccl_preload/cccl_allocator_preload.cu b/qwen3_6_scripts/cccl_preload/cccl_allocator_preload.cu index 50061632..8d3de4ec 100644 --- a/qwen3_6_scripts/cccl_preload/cccl_allocator_preload.cu +++ b/qwen3_6_scripts/cccl_preload/cccl_allocator_preload.cu @@ -41,8 +41,10 @@ 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 ---- */ -static cub::CachingDeviceAllocator& get_allocator() { - static cub::CachingDeviceAllocator instance( +using CubAllocator = CUB_NS_QUALIFIER::CachingDeviceAllocator; + +static CubAllocator& get_allocator() { + static CubAllocator instance( ALLOC_BIN_GROWTH, ALLOC_MIN_BIN, ALLOC_MAX_BIN,