From 7e397dfc18b21d47d54e0d7c5adc1a6963105804 Mon Sep 17 00:00:00 2001 From: Skyler Ruiter Date: Sat, 25 Jul 2026 08:21:16 -0500 Subject: [PATCH] adds small preset for h200 --- CHANGELOG.md | 3 +++ CMakePresets.json | 25 +++++++++++++++++++++++++ 2 files changed, 28 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 921041de..a9b55e1b 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -9,6 +9,9 @@ Version numbers follow [Semantic Versioning](https://semver.org/). ## [Unreleased] — 2.0.0 +### Added +- **`cuda-h200` CMake preset** (configure/build/test): CUDA backend, `CMAKE_CUDA_ARCHITECTURES=90` — the project default (86, A100) doesn't cover Hopper. Verified via `ctest --preset cuda-h200` on a gpuH200x8 node (NCSA Delta): 40/40 tests pass. + ### Fixed - **`test_ginterp`'s 4 `_Double3D` failures on HIP root-caused: not a code bug, a real AMD CDNA hardware ceiling.** `GInterpStage`'s 3-D spline kernel carves a `17^3 * 2 tiles * 8 bytes ≈ 76.8 KB` dynamic shared-memory allocation per block — CUDA Volta+ opts into that via `cudaFuncSetAttribute(cudaFuncAttributeMaxDynamicSharedMemorySize)`, which scales up to ~163-227 KB on modern NVIDIA GPUs, but MI100 (confirmed via `hipDeviceGetAttribute(hipDeviceAttributeSharedMemPerBlockOptin)`, returns 65536) has a *fixed* 64 KB LDS with no opt-in region at all — no software fix can make 76.8 KB fit in 64 KB. Previously this surfaced as an opaque `cudaGetLastError() → invalid argument` from the kernel launch. `ginterpRaiseSmemIfNeeded()` (`modules/fused/ginterp/ginterp_kernels.cu`) now queries the device's actual opt-in ceiling before attempting to raise it and throws a clear, actionable `std::runtime_error` when the requested size exceeds hardware capability, instead of silently mis-configuring the launch. The 4 affected tests (`GID1/3/4/5`, `tests/stages/test_ginterp.cpp`) now query the same device attribute up front and `GTEST_SKIP()` with the same explanation when the hardware can't support the configuration, rather than crashing. Float 3-D (~38.3 KB) and double 2-D (~4.5 KB) are unaffected and continue to pass — confirmed via full `ctest` run on MI100: **40/40 test binaries pass** (`test_ginterp`: 50/54 pass, 4 clean skips). Reducing the anchor tile size for double-precision 3-D specifically (so it fits within 64 KB on this class of hardware) is a real algorithmic option but changes the spline's tile geometry (currently tied to `kLevel=4` via `16 = 2^4`) and was left as a follow-up rather than attempted under this session's time constraints. - **`ANSStage.PartialBlock` (and any other small/single-block ANS input) segfaulted deterministically on real (non-vGPU) hardware.** Root cause was an `nvc++` host-code-generation bug, not a memory-safety issue (`compute-sanitizer --tool memcheck` reported 0 errors on the crashing binary): `ANSStage::execute()` handled both the forward and inverse paths — many locals, several templated `<<<>>>` kernel launches, and vGPU-only conditional branches — in one large function, and the compiler emitted a 16-byte-aligned `movaps` store into a CUDA kernel-launch argument slot that was only 8-byte aligned on the stack, faulting with SIGSEGV on that specific alignment mismatch. Confirmed via `ncu`-free `gdb` backtrace + register inspection (`$r15+0x10 ≡ 8 mod 16`), and confirmed the mechanism (not just the symptom) by reproducing a clean pass when recompiling `ans_stage.cu`'s host partition with `g++` instead of `nvc++` as `CMAKE_CUDA_HOST_COMPILER` — not shipped as the fix (toolchain swap), used only to verify the diagnosis. Fixed by splitting `ANSStage::execute()` into `executeForward()`/`executeInverse()` (`modules/coders/ans/ans_stage.{h,cu}`), so neither function's stack frame holds both branches' locals and launches at once; this changes the compiler's stack layout enough to avoid the misalignment, with no logic change. Verified deterministic (5/5 clean runs pre-fix reproduced the crash, 5/5 post-fix pass), `compute-sanitizer --tool memcheck` clean, and the full `ctest` suite (40/40, including `test_concurrency`'s multi-slot ANS case) passes with the original `nvc++` toolchain — no build/toolchain change shipped. Pre-existing since the ANS stage's introduction; unrelated to this session's `AdaptiveBitpackStage`/`TiledLorenzoStage` optimization work (confirmed via `git blame` — different files, different history) and distinct from the earlier, already-fixed vGPU-only spurious-SIGSEGV issue (that one is gated on `MemoryPool::isFallbackMode()`, which is `false` on real hardware and was confirmed `false` on this machine). diff --git a/CMakePresets.json b/CMakePresets.json index 01c453cf..6516511b 100644 --- a/CMakePresets.json +++ b/CMakePresets.json @@ -54,6 +54,18 @@ } }, + { + "name": "cuda-h200", + "inherits": "base", + "displayName": "CUDA (H200 / Hopper, sm_90)", + "description": "Release build against the CUDA backend for compute capability 9.0 (H100/H200). The project default (86, A100) does not cover Hopper, so this preset overrides CMAKE_CUDA_ARCHITECTURES explicitly.", + "cacheVariables": { + "CMAKE_BUILD_TYPE": "Release", + "FZGMOD_BACKEND": "CUDA", + "CMAKE_CUDA_ARCHITECTURES": "90" + } + }, + { "name": "compute-san", "inherits": "base", @@ -88,6 +100,11 @@ "configurePreset": "hip", "displayName": "Build HIP (AMD ROCm)" }, + { + "name": "cuda-h200", + "configurePreset": "cuda-h200", + "displayName": "Build CUDA (H200 / Hopper)" + }, { "name": "compute-san", "configurePreset": "compute-san", @@ -130,6 +147,14 @@ "output": { "outputOnFailure": true } }, + { + "name": "cuda-h200", + "configurePreset": "cuda-h200", + "displayName": "All tests (CUDA H200)", + "description": "Run the full test suite against the CUDA sm_90 (H200/Hopper) build.", + "output": { "outputOnFailure": true } + }, + { "name": "asan", "configurePreset": "asan",