fix: wire CCCL preload into build+launch chain + pre-submission verification
- 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)
This commit is contained in:
Claude
@@ -1,3 +1,18 @@
|
|||||||
**/__pycache__
|
**/__pycache__
|
||||||
**/*.pyc
|
**/*.pyc
|
||||||
**/.git
|
**/.git
|
||||||
|
|
||||||
|
# Large directories not needed in Docker image
|
||||||
|
cccl_upstream/
|
||||||
|
upstream_ref/
|
||||||
|
vllm/
|
||||||
|
ixformer_sdk/
|
||||||
|
muh/
|
||||||
|
ex_engine/
|
||||||
|
docs/
|
||||||
|
*.zip
|
||||||
|
*.txt
|
||||||
|
*.md
|
||||||
|
*.json
|
||||||
|
|
||||||
|
# Keep only qwen3_6_scripts/ computility-run.yaml Dockerfile
|
||||||
|
|||||||
@@ -49,3 +49,7 @@ env:
|
|||||||
value: '1'
|
value: '1'
|
||||||
- name: PYTORCH_CUDA_ALLOC_CONF
|
- name: PYTORCH_CUDA_ALLOC_CONF
|
||||||
value: max_split_size_mb:512
|
value: max_split_size_mb:512
|
||||||
|
- name: LD_PRELOAD
|
||||||
|
value: /workspace/qwen3_6_scripts/libcccl_allocator.so
|
||||||
|
- name: CCCL_ALLOC_DISABLE
|
||||||
|
value: '0'
|
||||||
|
|||||||
@@ -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"
|
|
||||||
@@ -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"
|
|
||||||
@@ -247,8 +247,15 @@ if source != installed:
|
|||||||
PY
|
PY
|
||||||
|
|
||||||
build_stage "compiling CCCL CachingDeviceAllocator LD_PRELOAD module"
|
build_stage "compiling CCCL CachingDeviceAllocator LD_PRELOAD module"
|
||||||
bash ./build_cccl_preload_allocator.sh /workspace/qwen3_6_scripts || \
|
if bash ./cccl_preload/build_cccl_preload.sh /workspace/qwen3_6_scripts; then
|
||||||
echo "[WARN] CCCL preload allocator build failed — will use default allocator"
|
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)"
|
build_stage "compiling CoreX CUDA extensions (moe_index_combine + gdn_chunk_recurrent)"
|
||||||
if [[ -x /usr/local/corex-3.2.3/bin/clang++ ]]; then
|
if [[ -x /usr/local/corex-3.2.3/bin/clang++ ]]; then
|
||||||
|
|||||||
@@ -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"
|
|
||||||
235
verify_submission.sh
Executable file
235
verify_submission.sh
Executable file
@@ -0,0 +1,235 @@
|
|||||||
|
#!/usr/bin/env bash
|
||||||
|
# verify_submission.sh — 竞赛提交前的完整验证
|
||||||
|
#
|
||||||
|
# 在真机上运行:
|
||||||
|
# cd /home/dylan/project_6
|
||||||
|
# bash verify_submission.sh
|
||||||
|
#
|
||||||
|
# 检查项:
|
||||||
|
# 1. Dockerfile语法
|
||||||
|
# 2. patch_ops.sh可执行 + 引用的文件全部存在
|
||||||
|
# 3. CCCL preload编译链完整(头文件+源文件+build脚本)
|
||||||
|
# 4. computility-run.yaml路径与build输出匹配
|
||||||
|
# 5. prebuilt .so文件完整性(SHA256)
|
||||||
|
# 6. qwen3_5.py imports的corex模块全部有对应.so或build脚本
|
||||||
|
# 7. Docker context大小(不要超过平台限制)
|
||||||
|
# 8. 单卡冒烟测试(如果有GPU)
|
||||||
|
|
||||||
|
PASS=0
|
||||||
|
FAIL=0
|
||||||
|
WARN=0
|
||||||
|
|
||||||
|
check() {
|
||||||
|
local name=$1; shift
|
||||||
|
if "$@" >/dev/null 2>&1; then
|
||||||
|
echo " ✓ $name"
|
||||||
|
PASS=$((PASS+1))
|
||||||
|
else
|
||||||
|
echo " ✗ $name"
|
||||||
|
FAIL=$((FAIL+1))
|
||||||
|
fi
|
||||||
|
}
|
||||||
|
|
||||||
|
warn() {
|
||||||
|
echo " △ $1"
|
||||||
|
WARN=$((WARN+1))
|
||||||
|
}
|
||||||
|
|
||||||
|
ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
|
||||||
|
cd "$ROOT"
|
||||||
|
|
||||||
|
echo "=== 1. 文件结构 ==="
|
||||||
|
check "Dockerfile存在" test -f Dockerfile
|
||||||
|
check "computility-run.yaml存在" test -f computility-run.yaml
|
||||||
|
check "patch_ops.sh存在" test -f qwen3_6_scripts/patch_ops.sh
|
||||||
|
check "patch_ops.sh可执行" test -x qwen3_6_scripts/patch_ops.sh
|
||||||
|
check ".dockerignore存在" test -f .dockerignore
|
||||||
|
|
||||||
|
echo ""
|
||||||
|
echo "=== 2. CCCL Preload编译链 ==="
|
||||||
|
check "build脚本存在" test -f qwen3_6_scripts/cccl_preload/build_cccl_preload.sh
|
||||||
|
check "源文件存在" test -f qwen3_6_scripts/cccl_preload/cccl_allocator_preload.cu
|
||||||
|
check "CCCL头文件目录存在" test -d qwen3_6_scripts/cccl_preload/include/cub
|
||||||
|
check "libcudacxx头文件存在" test -d qwen3_6_scripts/cccl_preload/include/cuda
|
||||||
|
|
||||||
|
CCCL_HEADERS=$(find qwen3_6_scripts/cccl_preload/include -type f 2>/dev/null | wc -l)
|
||||||
|
if [[ "$CCCL_HEADERS" -ge 280 ]]; then
|
||||||
|
echo " ✓ CCCL头文件数量: $CCCL_HEADERS (≥280)"
|
||||||
|
PASS=$((PASS+1))
|
||||||
|
else
|
||||||
|
echo " ✗ CCCL头文件数量: $CCCL_HEADERS (期望≥280)"
|
||||||
|
FAIL=$((FAIL+1))
|
||||||
|
fi
|
||||||
|
|
||||||
|
# 检查关键头文件
|
||||||
|
check "cub/util_allocator.cuh" test -f qwen3_6_scripts/cccl_preload/include/cub/util_allocator.cuh
|
||||||
|
check "cub/config.cuh" test -f qwen3_6_scripts/cccl_preload/include/cub/config.cuh
|
||||||
|
check "cuda/__cccl_config" test -f qwen3_6_scripts/cccl_preload/include/cuda/__cccl_config
|
||||||
|
check "nv/target" test -f qwen3_6_scripts/cccl_preload/include/nv/target
|
||||||
|
|
||||||
|
echo ""
|
||||||
|
echo "=== 3. patch_ops.sh → CCCL preload 调用链 ==="
|
||||||
|
if grep -q "cccl_preload/build_cccl_preload.sh" qwen3_6_scripts/patch_ops.sh; then
|
||||||
|
echo " ✓ patch_ops.sh调用新版build脚本"
|
||||||
|
PASS=$((PASS+1))
|
||||||
|
else
|
||||||
|
echo " ✗ patch_ops.sh未调用cccl_preload/build_cccl_preload.sh"
|
||||||
|
FAIL=$((FAIL+1))
|
||||||
|
fi
|
||||||
|
|
||||||
|
# 检查旧文件是否残留
|
||||||
|
if [[ -f qwen3_6_scripts/cccl_preload_allocator.cu ]]; then
|
||||||
|
echo " ✗ 旧mock文件残留: cccl_preload_allocator.cu"
|
||||||
|
FAIL=$((FAIL+1))
|
||||||
|
else
|
||||||
|
echo " ✓ 旧mock文件已清理"
|
||||||
|
PASS=$((PASS+1))
|
||||||
|
fi
|
||||||
|
|
||||||
|
echo ""
|
||||||
|
echo "=== 4. computility-run.yaml 路径匹配 ==="
|
||||||
|
YAML_PRELOAD=$(grep -A1 "LD_PRELOAD" computility-run.yaml | grep "value:" | awk '{print $2}')
|
||||||
|
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)
|
||||||
|
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输出"
|
||||||
|
PASS=$((PASS+1))
|
||||||
|
else
|
||||||
|
echo " ✗ 路径不匹配: yaml=$YAML_PRELOAD expected=$EXPECTED_SO"
|
||||||
|
FAIL=$((FAIL+1))
|
||||||
|
fi
|
||||||
|
else
|
||||||
|
echo " ✗ computility-run.yaml缺少LD_PRELOAD"
|
||||||
|
FAIL=$((FAIL+1))
|
||||||
|
fi
|
||||||
|
|
||||||
|
# 检查CCCL_ALLOC_DISABLE
|
||||||
|
if grep -q "CCCL_ALLOC_DISABLE" computility-run.yaml; then
|
||||||
|
echo " ✓ CCCL_ALLOC_DISABLE可控"
|
||||||
|
PASS=$((PASS+1))
|
||||||
|
fi
|
||||||
|
|
||||||
|
echo ""
|
||||||
|
echo "=== 5. Prebuilt .so 完整性 ==="
|
||||||
|
PREBUILT_DIR="qwen3_6_scripts/prebuilt/corex-3.2.3-ivcore10"
|
||||||
|
if [[ -f "${PREBUILT_DIR}/SHA256SUMS" ]]; then
|
||||||
|
EXPECTED=$(wc -l < "${PREBUILT_DIR}/SHA256SUMS")
|
||||||
|
ACTUAL=$(ls "${PREBUILT_DIR}"/*.so 2>/dev/null | wc -l)
|
||||||
|
if [[ "$ACTUAL" -ge "$EXPECTED" ]]; then
|
||||||
|
echo " ✓ prebuilt .so数量: $ACTUAL (manifest expects $EXPECTED)"
|
||||||
|
PASS=$((PASS+1))
|
||||||
|
else
|
||||||
|
echo " ✗ prebuilt .so数量不足: $ACTUAL < $EXPECTED"
|
||||||
|
FAIL=$((FAIL+1))
|
||||||
|
fi
|
||||||
|
# 如果有sha256sum工具, 验证checksum
|
||||||
|
if command -v sha256sum &>/dev/null; then
|
||||||
|
if (cd "$PREBUILT_DIR" && sha256sum --status --check SHA256SUMS 2>/dev/null); then
|
||||||
|
echo " ✓ SHA256校验通过"
|
||||||
|
PASS=$((PASS+1))
|
||||||
|
else
|
||||||
|
echo " ✗ SHA256校验失败"
|
||||||
|
FAIL=$((FAIL+1))
|
||||||
|
fi
|
||||||
|
fi
|
||||||
|
else
|
||||||
|
echo " △ SHA256SUMS不存在, 跳过校验"
|
||||||
|
WARN=$((WARN+1))
|
||||||
|
fi
|
||||||
|
|
||||||
|
echo ""
|
||||||
|
echo "=== 6. qwen3_5.py corex imports vs prebuilt ==="
|
||||||
|
IMPORTS=$(grep "from vllm import corex_" qwen3_6_scripts/qwen3_5.py 2>/dev/null | sed 's/.*import //' | sed 's/ as.*//' | sort -u)
|
||||||
|
for mod in $IMPORTS; do
|
||||||
|
SO_FILE="${PREBUILT_DIR}/${mod}.so"
|
||||||
|
BUILD_SCRIPT="qwen3_6_scripts/build_${mod}.sh"
|
||||||
|
if [[ -f "$SO_FILE" ]]; then
|
||||||
|
echo " ✓ $mod → prebuilt .so"
|
||||||
|
PASS=$((PASS+1))
|
||||||
|
elif [[ -f "$BUILD_SCRIPT" ]]; then
|
||||||
|
echo " △ $mod → build脚本存在 (docker内编译)"
|
||||||
|
WARN=$((WARN+1))
|
||||||
|
else
|
||||||
|
echo " ✗ $mod → 无.so也无build脚本"
|
||||||
|
FAIL=$((FAIL+1))
|
||||||
|
fi
|
||||||
|
done
|
||||||
|
|
||||||
|
echo ""
|
||||||
|
echo "=== 7. Docker context大小 ==="
|
||||||
|
# 排除.git和大文件
|
||||||
|
if [[ -f .dockerignore ]]; then
|
||||||
|
# 粗略估算
|
||||||
|
CONTEXT_SIZE=$(du -sh --exclude='.git' --exclude='cccl_upstream' --exclude='upstream_ref' --exclude='*.zip' --exclude='vllm' --exclude='ixformer_sdk' qwen3_6_scripts/ computility-run.yaml Dockerfile 2>/dev/null | tail -1 | awk '{print $1}')
|
||||||
|
echo " 核心文件大小: ~$CONTEXT_SIZE"
|
||||||
|
echo " (.dockerignore应排除cccl_upstream/, upstream_ref/, vllm/, *.zip等)"
|
||||||
|
fi
|
||||||
|
|
||||||
|
# 检查.dockerignore是否排除大目录
|
||||||
|
if [[ -f .dockerignore ]]; then
|
||||||
|
for dir in cccl_upstream upstream_ref vllm ixformer_sdk "*.zip"; do
|
||||||
|
if grep -q "$dir" .dockerignore 2>/dev/null; then
|
||||||
|
echo " ✓ .dockerignore排除: $dir"
|
||||||
|
else
|
||||||
|
warn ".dockerignore未排除: $dir (可能导致docker context过大)"
|
||||||
|
fi
|
||||||
|
done
|
||||||
|
fi
|
||||||
|
|
||||||
|
echo ""
|
||||||
|
echo "=== 8. 单卡冒烟测试 ==="
|
||||||
|
if command -v ixsmi &>/dev/null || command -v nvidia-smi &>/dev/null; then
|
||||||
|
echo " 检测到GPU, 运行基础验证..."
|
||||||
|
if python3 -c "import torch; assert torch.cuda.is_available(); print(f' ✓ PyTorch CUDA: {torch.cuda.get_device_name(0)}')" 2>/dev/null; then
|
||||||
|
PASS=$((PASS+1))
|
||||||
|
else
|
||||||
|
echo " ✗ PyTorch CUDA不可用"
|
||||||
|
FAIL=$((FAIL+1))
|
||||||
|
fi
|
||||||
|
|
||||||
|
# 测试CCCL preload编译
|
||||||
|
if [[ -f qwen3_6_scripts/cccl_preload/build_cccl_preload.sh ]]; then
|
||||||
|
echo " 尝试编译CCCL preload .so..."
|
||||||
|
if bash qwen3_6_scripts/cccl_preload/build_cccl_preload.sh /tmp 2>/dev/null; then
|
||||||
|
if [[ -s /tmp/libcccl_allocator.so ]]; then
|
||||||
|
echo " ✓ CCCL preload编译成功"
|
||||||
|
PASS=$((PASS+1))
|
||||||
|
# 测试加载
|
||||||
|
if LD_PRELOAD=/tmp/libcccl_allocator.so python3 -c "import torch; x=torch.zeros(1024,device='cuda'); del x; print(' ✓ LD_PRELOAD加载正常')" 2>/dev/null; then
|
||||||
|
PASS=$((PASS+1))
|
||||||
|
else
|
||||||
|
echo " ✗ LD_PRELOAD加载失败"
|
||||||
|
FAIL=$((FAIL+1))
|
||||||
|
fi
|
||||||
|
rm -f /tmp/libcccl_allocator.so
|
||||||
|
else
|
||||||
|
echo " ✗ 编译产物为空"
|
||||||
|
FAIL=$((FAIL+1))
|
||||||
|
fi
|
||||||
|
else
|
||||||
|
echo " ✗ CCCL preload编译失败"
|
||||||
|
FAIL=$((FAIL+1))
|
||||||
|
fi
|
||||||
|
fi
|
||||||
|
else
|
||||||
|
echo " (无GPU, 跳过)"
|
||||||
|
fi
|
||||||
|
|
||||||
|
echo ""
|
||||||
|
echo "==========================================="
|
||||||
|
echo " PASS: $PASS FAIL: $FAIL WARN: $WARN"
|
||||||
|
echo "==========================================="
|
||||||
|
|
||||||
|
if [[ "$FAIL" -gt 0 ]]; then
|
||||||
|
echo ""
|
||||||
|
echo "⚠ 有 $FAIL 个检查失败, 提交前请修复!"
|
||||||
|
exit 1
|
||||||
|
else
|
||||||
|
echo ""
|
||||||
|
echo "✓ 全部检查通过, 可以提交"
|
||||||
|
exit 0
|
||||||
|
fi
|
||||||
Reference in New Issue
Block a user