1fae917c22
- Migrate from NVRTC JIT to cached nvcc -O3 for all CUDA kernels - Fuse guard kernels, increase prefetch chunk, eliminate per-step GPU alloc - H100-specific: fused Adam, warp reductions, shmem tiling, PPO occupancy - Vectorize gather_states with __ldg() and 4x unroll - sincosf() Box-Muller + paired Gaussian generation in noisy nets - Shared-memory tiled branching DQN forward pass for sm_<90 - GPU-resident training guard kernel replacing Candle tensor ops - Eliminate all to_vec1/to_vec2 CPU roundtrips, DtoD weight copy - GPU PER mandatory everywhere — kill CPU replay path on CUDA - Full GPU action masking — eliminate CPU fallback path - Fix cuBLAS handle sharing via OnceLock (root cause of 49 cascade failures) - Fix ILLEGAL_ADDRESS: scratch1_dist buffer overflow, stack sizing, curand determinism - Fix CudaStream lifetime: bind before .context() to extend lifetime - Keep raw cudarc buffers alive across epochs - Add gpu-hotpath-guard.sh (37 patterns) and ptx-cache-invalidate.sh Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
53 lines
1.9 KiB
Bash
Executable File
53 lines
1.9 KiB
Bash
Executable File
#!/usr/bin/env bash
|
|
# ptx-cache-invalidate.sh — Invalidate PTX cache when CUDA kernels change.
|
|
#
|
|
# The Rust runtime caches compiled PTX in $CARGO_TARGET_DIR/.ptx_cache/ (or
|
|
# /tmp/.ptx_cache/). The cache key includes a SHA-256 of the kernel source,
|
|
# so changed .cu files naturally miss. However, if the CI PVC preserves the
|
|
# cache across runs and the build binary is stale, old PTX can survive.
|
|
#
|
|
# This script computes a content hash of all .cu/.cuh kernel source files and
|
|
# compares it against a stamp file in the cache directory. If the hash differs
|
|
# (any kernel source changed), the entire PTX cache is purged, forcing nvcc
|
|
# recompilation on the next run.
|
|
#
|
|
# Usage: scripts/ptx-cache-invalidate.sh [cache_dir]
|
|
# cache_dir defaults to $CARGO_TARGET_DIR/.ptx_cache, then /tmp/.ptx_cache
|
|
|
|
set -euo pipefail
|
|
|
|
KERNEL_DIR="crates/ml/src/cuda_pipeline"
|
|
CACHE_DIR="${1:-${CARGO_TARGET_DIR:+${CARGO_TARGET_DIR}/.ptx_cache}}"
|
|
CACHE_DIR="${CACHE_DIR:-/tmp/.ptx_cache}"
|
|
|
|
# Compute content hash of all CUDA source files
|
|
HASH=$(find "$KERNEL_DIR" -type f \( -name '*.cu' -o -name '*.cuh' \) \
|
|
-exec sha256sum {} + 2>/dev/null | sort | sha256sum | cut -d' ' -f1)
|
|
|
|
STAMP_FILE="${CACHE_DIR}/.kernel_hash"
|
|
|
|
if [ -f "$STAMP_FILE" ]; then
|
|
OLD_HASH=$(cat "$STAMP_FILE")
|
|
if [ "$HASH" = "$OLD_HASH" ]; then
|
|
echo "PTX cache valid — kernel hash unchanged ($HASH)"
|
|
exit 0
|
|
fi
|
|
echo "PTX cache STALE — kernel hash changed"
|
|
echo " old: $OLD_HASH"
|
|
echo " new: $HASH"
|
|
else
|
|
echo "PTX cache stamp missing — first run or cache cleared"
|
|
fi
|
|
|
|
# Purge stale PTX cache
|
|
if [ -d "$CACHE_DIR" ]; then
|
|
PTX_COUNT=$(find "$CACHE_DIR" -name '*.ptx' -type f 2>/dev/null | wc -l)
|
|
echo "Purging ${PTX_COUNT} cached PTX files from ${CACHE_DIR}"
|
|
rm -f "${CACHE_DIR}"/*.ptx
|
|
fi
|
|
|
|
# Write new stamp
|
|
mkdir -p "$CACHE_DIR"
|
|
echo "$HASH" > "$STAMP_FILE"
|
|
echo "PTX cache stamp updated: $HASH"
|