From bdac8be75ba9ddfc9ac9fc1bf11b6365b2c65a4b Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Sat, 8 Aug 2026 23:48:58 +0000 Subject: [PATCH 1/3] 3P-ADMM-PC2: keep this branch's project state across the trunk merge --- projects/3P-ADMM-PC2/notes.md | 366 +++++++++++++++++++++++++++++++ projects/3P-ADMM-PC2/plan.md | 186 ++++++++++++++++ projects/3P-ADMM-PC2/stats.jsonl | 14 ++ projects/3P-ADMM-PC2/status.json | 174 +++++++++++++++ 4 files changed, 740 insertions(+) create mode 100644 projects/3P-ADMM-PC2/notes.md create mode 100644 projects/3P-ADMM-PC2/plan.md create mode 100644 projects/3P-ADMM-PC2/stats.jsonl create mode 100644 projects/3P-ADMM-PC2/status.json diff --git a/projects/3P-ADMM-PC2/notes.md b/projects/3P-ADMM-PC2/notes.md new file mode 100644 index 00000000..b096c68c --- /dev/null +++ b/projects/3P-ADMM-PC2/notes.md @@ -0,0 +1,366 @@ +# 3P-ADMM-PC2 notes + +Privacy-preserving distributed ADMM for LASSO. CPU does the optimization +(Paillier homomorphic encryption + quantized ADMM, gmpy2). The GPU's only job +is to accelerate the Paillier modular exponentiation `g^m mod n` over a batch +of big integers, via cuFFT-based large-integer polynomial multiplication + +Barrett reduction. + +## Build classification +ext_type `nvcc-shared`: ad-hoc nvcc commands in README.md sec.1 (no CMake, no +Makefile, no setup.py). The `.cu` is compiled + a generated C wrapper, linked +with `g++ -shared` into `/tmp/lib_cufft.so`. The Python side +(`crypto/paillier_gpu.py`, `protocol/edge_crt_helper.py`) loads that .so by +absolute path via `ctypes.CDLL` and calls `init_gpu(int)` / `run_modexp(...)`. + +## Port scope (LIVE GPU path only) +Ported ONLY `gpu/cufft_modexp.cu` (cuFFT Z2Z + scalar bigint kernels). This is +the only GPU code wired into the protocol. + +Deliberately EXCLUDED: +- Dead register/warp-NTT files: `gpu/reg_ntt_modexp*.cu`, `gpu/ntt_modexp*.cu`, + `gpu/ntt_funcs.cu`, `gpu/ntt_kernel2.cu`, `gpu/ntt_wrapper*.cu`, + `gpu/modexp_final.cu`. Referenced by nothing in the repo (no Python import, + no launch site, not built by the README). The `reg_ntt_*` and `ntt_funcs` + files DO use `__shfl_xor_sync(0xffffffff, ..., delta)` warp-NTT butterflies + with blockDim=32 (one warp/task) -- classic wave64 hazards -- but they are + not on any live path, so porting them would expose the wave64 fault class for + no validatable benefit. A future contributor reviving the register-NTT path + on AMD would need the `__GFX9__`->64 / RDNA->32 + width-aware ballot + treatment. +- Standalone PyCUDA self-test: `gpu/gpu_modexp.py` (JITs `gpu/modexp.cu` / + `gpu/modexp_v2.cu`). PyCUDA is not in requirements.txt and this is a + `__main__` self-test, not called by the protocol. Out of scope. + +## What changed +- ADD `gpu/cuda_to_hip.h`: single compat shim. Under `__HIP__`/ + `__HIP_PLATFORM_AMD__`/`USE_HIP` it includes `` + + `` and aliases the exact CUDA symbols the kernel uses + (`cudaMalloc/Free/Memcpy` + dirs; `cufftHandle`/`cufftDoubleComplex`, + `cufftPlanMany`/`cufftExecZ2Z`/`cufftDestroy`, `CUFFT_Z2Z`/`CUFFT_FORWARD`, + `CUFFT_INVERSE`->`HIPFFT_BACKWARD`); else it includes the original CUDA + headers. NVIDIA build stays byte-identical. + - Gotcha: hipcc compiling `-x hip` defines `__HIP__`, NOT + `__HIP_PLATFORM_AMD__`, at preprocess time. Key the `#if` off `__HIP__` + too, or the shim falls through to `` and fails to find it. + - libc headers (``) are included BEFORE the HIP + runtime (gpuRIR lesson): host `calloc`/`free`/`round` in the .cu would + otherwise risk binding to HIP `__device__` overloads. +- EDIT `gpu/cufft_modexp.cu`: replaced ``/``/`` + with `#include "cuda_to_hip.h"`. No body changes; all cuda*/cufft* spellings + preserved. +- ADD `gpu/build_hip.sh`: hipcc recipe mirroring the README nvcc steps, + producing `/tmp/lib_cufft.so` with the SAME `init_gpu`/`run_modexp` C ABI + (Python untouched), linking `-lhipfft`. Arch from `HIP_ARCH` (default + `gfx90a`, never a literal). hipcc does host+device link in one step (no + separate `-dlink`/`-fgpu-rdc` needed for this non-RDC build). + +## Build (gfx90a) +``` +HIP_VISIBLE_DEVICES=1 \ +BUILD= OUT=/tmp/lib_cufft.so \ + bash gpu/build_hip.sh +``` +Followers: `HIP_ARCH=gfx1100` / `gfx1151`. Multi-arch fat binary: +`HIP_ARCH="gfx90a,gfx1100"` builds clean (verified). The only warnings are +benign `-Wunused-value` on the cuda*Memcpy nodiscard return (present in the +NVIDIA build too). hipcc/ROCm 7.2.1, hipfft present at /opt/rocm. + +## hipFFT rounding note (the one real risk -- PASSED) +The bigint multiply is a double-precision Z2Z FFT then `round()` to recover +exact integer limbs (`norm_round` + `carry_prop`), relying on a < 0.5 rounding +margin. hipFFT twiddles differ bit-for-bit from cuFFT, so a thin margin could +in principle round a limb the wrong way. Validated statistically: 6144 random +modexp cases across moduli up to 2046-bit n^2 -- all exact-match gmpy2. The +double-FFT margin holds on gfx90a (full-rate fp64). No code change was needed. + +## Validation (real gfx90a, GCD 1, HIP_VISIBLE_DEVICES=1) +- modexp-vs-gmpy2 gold: 6144/6144 exact match. `gpu_batch_modexp(g,m,n)` vs + `gmpy2.powmod(g,m,n)`, n at 255/510/511/1022/768/1535/1023/2046 bits, varied + m_bits, 256/batch. Harness: agent_space/3p-admm/validate.py (not committed). +- ALGORITHM PRECONDITION: the kernel assumes base g < modulus n (no initial + reduction). A test with g >= n mismatches gmpy2 -- this is expected (contract + violation), NOT a HIP bug. Real Paillier always has g < modulus. Reducing + g %= n before the call makes it match exactly. +- encrypt round-trip: `crypto/paillier_gpu.encrypt_batch_gpu` (GPU g^m mod n^2) + -> CPU `crypto/paillier.decrypt` = 64/64 plaintexts recovered. +- CPU regression (no GPU touched): `crypto/test_paillier.py`, + `test_quantization.py`, `test_full_chain.py` all pass. +- gfx90a dispatch confirmed via AMD_LOG_LEVEL=3 (hipFFT `twiddle_gen_radices_dp` + + our kernels on device; rocminfo Name gfx90a). + +## Wave size +Not exposed in the live path: no `__shfl`/`__ballot`/`warpSize`/`__syncwarp`, +no cub/thrust/curand, no textures/atomics. The `<<>>` launches +(`cond_copy`, `load_complex` grid.x) are independent-thread strided copies, not +warp collectives. Multi-arch is trivially fine; gfx1100/gfx1151 deltas are +expected to be no-op revalidations (same .so rebuilt for the arch). hipFFT is +arch-agnostic. + +## Upstream warts (NOT touched -- out of port scope) +- Hardcoded `/mnt/3p-admm-pc2/...` absolute paths in Python (lib + kernel source + paths) and `sys.path.append('/mnt/3p-admm-pc2')` in the test files. For our + build+validate we point the loader at the real /tmp/lib_cufft.so and set + PYTHONPATH. We do not refactor upstream's path scheme. +- `/tmp` is non-persistent (README notes "recompile each boot"). +- Distributed end-to-end tests (`experiments/*.py`, `protocol/master_node.py`) + need multi-node SSH / matpool cloud; not runnable here. The modexp gold gate + + encrypt round-trip are the representative validatable proxy for the GPU work + those drive. + +## Validation 2026-06-02 (validator, linux-gfx90a, GCD 1) + +validated_sha: 6ef301f3204579779bbaa1f32a466934f720903a + +GPU arch: gfx90a (AMD Instinct MI250X, ROCm 7.2.1, HIP_VISIBLE_DEVICES=1) + +Commands run: +``` +# gfx90a build +HIP_VISIBLE_DEVICES=1 BUILD=/var/lib/jenkins/moat/agent_space/3p-admm OUT=/tmp/lib_cufft.so \ + bash utils/timeit.sh 3P-ADMM-PC2 compile -- bash projects/3P-ADMM-PC2/src/gpu/build_hip.sh + +# fat binary (gfx90a + gfx1100) +HIP_VISIBLE_DEVICES=1 HIP_ARCH="gfx90a,gfx1100" \ + BUILD=/var/lib/jenkins/moat/agent_space/3p-admm OUT=/tmp/lib_cufft_multi.so \ + bash utils/timeit.sh 3P-ADMM-PC2 compile -- bash projects/3P-ADMM-PC2/src/gpu/build_hip.sh + +# modexp gold match +HIP_VISIBLE_DEVICES=1 AMD_LOG_LEVEL=0 \ + bash utils/timeit.sh 3P-ADMM-PC2 test -- python3 agent_space/3p-admm/validate.py + +# encrypt round-trip, CPU regression tests +HIP_VISIBLE_DEVICES=1 PYTHONPATH= python3 crypto/test_paillier.py +HIP_VISIBLE_DEVICES=1 PYTHONPATH= python3 crypto/test_quantization.py +HIP_VISIBLE_DEVICES=1 PYTHONPATH= python3 crypto/test_full_chain.py +``` + +Results: +- Build gfx90a: PASS (only benign -Wunused-value on nodiscard hipMemcpy/hipFree, matches CUDA build) +- Fat binary gfx90a+gfx1100: PASS; roc-obj-ls confirms both code objects embedded + (hipv4-amdgcn-amd-amdhsa--gfx1100 @ offset=12288, hipv4-amdgcn-amd-amdhsa--gfx90a @ offset=40960) +- modexp vs gmpy2 gold: 6144/6144 exact match + n at 255/510/511/768/1022/1023/1535/2046 bits, 256/batch, 3 reps each, varied m_bits +- Encrypt round-trip: 64/64 plaintexts recovered (GPU g^m mod n^2 -> CPU decrypt) +- CPU regression: test_paillier PASS, test_quantization PASS, test_full_chain PASS (max error 1.33e-09) +- Native gfx90a dispatch confirmed (AMD_LOG_LEVEL=3): + "Using native code object for device: amdgcn-amd-amdhsa--gfx90a:sramecc+:xnack-" + hipFFT twiddle_gen_radices_dp dispatched; all our kernels dispatched via hipLaunchKernel +- hipFFT double-FFT rounding margin holds: 0 mismatches across all 6144 crypto-significant cases + +VERDICT: PASS -- review-passed -> completed + +## Review 2026-06-02 (reviewer, gfx90a) +review-passed. No problems found; no changes requested. + +Verified independently on real gfx90a (HIP_VISIBLE_DEVICES=1, GCD 1, ROCm 7.2.1): +- Scope correct: diff touches only gpu/cuda_to_hip.h (new), gpu/build_hip.sh (new), + gpu/cufft_modexp.cu (one include-line swap). No Python, no dead reg_ntt_*/ntt_funcs + warp-NTT files, no PyCUDA self-test touched -- so the live path genuinely has no + warp/wave64 surface (grep confirms no __shfl/__ballot/warpSize/__syncwarp in the .cu; + the <<>> launches are independent-thread strided copies). Multi-arch trivially fine. +- Shim symbol coverage complete: every cuda*/cufft*/CUFFT_* token used by cufft_modexp.cu + (cudaMalloc/Free/Memcpy + H2D/D2H, cufftHandle/DoubleComplex/PlanMany/ExecZ2Z/Destroy, + CUFFT_Z2Z/FORWARD/INVERSE->HIPFFT_BACKWARD) is aliased. libc (cstdlib/cstring/cmath) + included before . #if keys off __HIP__ -- independently confirmed + `hipcc -x hip -dM -E` defines __HIP__ but NOT __HIP_PLATFORM_AMD__ at preprocess, so the + __HIP__ key is load-bearing and correct. NVIDIA #else path falls through to original CUDA + headers (math.h -> cmath is a superset; round() still resolves), byte-identical build. +- build_hip.sh wrapper (init_gpu/run_modexp -> cufft_init/cufft_modexp) is identical to the + README's wr_cufft.cu, preserving the ctypes C ABI; arch from HIP_ARCH (no literal); no + -fgpu-rdc needed (no cross-TU __device__ funcs). +- Built the gfx90a,gfx1100 fat binary clean (only benign -Wunused-value on nodiscard + hipMemcpy, present on CUDA too); roc-obj-ls confirms both code objects embedded. +- Reproduced the correctness gate: 6144/6144 exact vs gmpy2.powmod across n at + 255/510/511/768/1022/1023/1535/2046 bits, varied m_bits, 256/batch. hipFFT double-FFT + rounding margin holds. +- g>=n: confirmed the kernel has NO initial g%=n reduction (identical to upstream CUDA), so + the precondition is pre-existing, not a port defect; after g%=n the batch matches gmpy2. +- Commit hygiene: title "[ROCm] Port live cuFFT Paillier modexp to hipFFT" (47 chars), + mentions Claude, no noreply trailer, no ghstack, no em-dash. Fork features == upstream + features @ dd96d5d (clean mirror); single port commit on moat-port. Actions disabled. + +Non-blocking: the __HIP__-vs-__HIP_PLATFORM_AMD__ compat-shim-keying trap (an #include'd shim +keyed only on __HIP_PLATFORM_AMD__ silently falls through to because hipcc +defines __HIP__, not __HIP_PLATFORM_AMD__, at preprocess) is a generalizable lesson and was +appended to PORTING_GUIDE.md. + +## Validation 2026-06-02 (gfx1100, linux-gfx1100, HIP_VISIBLE_DEVICES=0) + +validated_sha: 6ef301f3204579779bbaa1f32a466934f720903a + +GPU arch: gfx1100 (AMD Radeon Pro W7800 48GB, ROCm 7.2.1, HIP_VISIBLE_DEVICES=0, wave32) + +No code change from gfx90a lead -- validate-first follower, fork untouched at 6ef301f. + +Commands run: +``` +# gfx1100 build (HIP_ARCH=gfx1100, no literal) +HIP_VISIBLE_DEVICES=0 HIP_ARCH=gfx1100 \ + BUILD=/var/lib/jenkins/moat/agent_space/3p-admm-gfx1100 OUT=/tmp/lib_cufft_gfx1100.so \ + bash utils/timeit.sh 3P-ADMM-PC2 compile -- bash projects/3P-ADMM-PC2/src/gpu/build_hip.sh + +# modexp gold match (run twice for determinism) +HIP_VISIBLE_DEVICES=0 AMD_LOG_LEVEL=0 LIB_CUFFT=/tmp/lib_cufft_gfx1100.so \ + bash utils/timeit.sh 3P-ADMM-PC2 test -- python3 agent_space/3p-admm-gfx1100/validate.py + +# gfx1100 dispatch confirmation +HIP_VISIBLE_DEVICES=0 AMD_LOG_LEVEL=3 LIB_CUFFT=/tmp/lib_cufft_gfx1100.so \ + python3 agent_space/3p-admm-gfx1100/validate.py 2>&1 | grep "native code\|gfx1100\|0x1016" + +# CPU regression +PYTHONPATH=projects/3P-ADMM-PC2/src python3 crypto/test_paillier.py +PYTHONPATH=projects/3P-ADMM-PC2/src python3 crypto/test_quantization.py +PYTHONPATH=projects/3P-ADMM-PC2/src python3 crypto/test_full_chain.py +``` + +Results: +- Build gfx1100: PASS -- only benign -Wunused-value on nodiscard hipMemcpy/hipFree (identical to gfx90a) +- roc-obj-ls confirms gfx1100 code object: hipv4-amdgcn-amd-amdhsa--gfx1100 @ offset=12288 + (single-arch gfx1100 only .so; no gfx90a object present in this build) +- Native gfx1100 dispatch confirmed (AMD_LOG_LEVEL=3): + "Using native code object for device: amdgcn-amd-amdhsa--gfx1100" + hipFFT twiddle_gen_radices_dp dispatched; all kernels hipLaunchKernel hipSuccess; no HSA 0x1016 +- modexp vs gmpy2 gold (run 1): 6144/6144 exact match + n at 255/510/511/768/1022/1023/1535/2046 bits, 256/batch, 3 reps each, varied m_bits +- modexp vs gmpy2 gold (run 2, determinism): 6144/6144 exact match +- hipFFT double-FFT rounding margin holds on wave32 (gfx1100 fp64 full-rate): 0 mismatches +- CPU regression: test_paillier PASS, test_quantization PASS, test_full_chain PASS (max error 1.33e-09) +- Fork working tree: clean (only __pycache__ .pyc touched by running tests; no source change) +- No fork push (zero-churn follower validation; no code change needed) + +Wave32 note: live path has no warp intrinsics (__shfl/__ballot/warpSize/__syncwarp absent), +no cub/thrust/curand. The <<>> launches are independent-thread strided copies -- wave32 +behavior is identical to wave64. Dead warp-NTT files remain excluded (no live path, wave64 hazard). + +VERDICT: PASS -- port-ready -> completed +gfx1100 hipFFT Z2Z + Barrett reduction correct at wave32; matches gfx90a@6ef301f exactly. + +## Validation 2026-06-04 (windows-gfx1151, gfx1151, AMD Radeon 8060S, TheRock 7.13) + +validated_sha: 6ef301f3204579779bbaa1f32a466934f720903a + +GPU arch: gfx1151 (AMD Radeon 8060S, RDNA3.5, wave32, TheRock ROCm 7.13 pip wheels) + +No code change from gfx90a lead -- zero-churn follower validation, fork untouched at 6ef301f. + +### Windows-specific build notes + +On Windows the build_hip.sh Linux recipe needs two adaptations: +1. Remove `-fPIC` (unsupported on MSVC target; position-independent code is implicit for Windows DLLs). +2. Add a `.def` file to the linker so `init_gpu` and `run_modexp` are exported. On Windows with + clang/MSVC-ABI, `extern "C"` functions in a DLL are NOT auto-exported (unlike Linux .so); + `__declspec(dllexport)` is silently ignored in device-code TUs compiled with hipcc, so a + `/DEF:lib_cufft.def` passed via `-Wl,/DEF:...` is the clean solution. +3. DLL runtime env: copy TheRock DLLs (amdhip64_7.dll, amd_comgr0713.dll, rocm_kpack.dll, + hipfft.dll, rocfft.dll, hiprtc07013.dll, hiprtc-builtins07013.dll) from + `_rocm_sdk_devel/bin/` + `_rocm_sdk_libraries_gfx1151/bin/` to the same directory as + lib_cufft.dll. Windows loader searches the DLL's own directory first, so this ensures + TheRock's runtime loads instead of the broken System32 Adrenalin amdhip64_7.dll. In the + Python ctypes loader, call `os.add_dll_directory(dll_dir)` before loading. + +The hipfft.dll and rocfft.dll from `_rocm_sdk_libraries_gfx1151/bin/` (the gfx1151-specific +package) contain the precompiled gfx1151 device kernels; no separate .kpack file needed. + +Commands run: +``` +# Build (agent_space/3p-admm-win/build_win.sh) +HIP_ARCH=gfx1151 bash utils/timeit.sh 3P-ADMM-PC2 compile -- \ + bash agent_space/3p-admm-win/build_win.sh +# -> lib_cufft.dll (150 KB, exports: init_gpu, run_modexp) + +# modexp gold match +bash utils/timeit.sh 3P-ADMM-PC2 test -- \ + C:/Users//AppData/Local/Programs/Python/Python313/python.exe \ + agent_space/3p-admm-win/validate_win.py + +# CPU regression (PYTHONUTF8=1 works around upstream Chinese print() on cp1252 terminal) +PYTHONPATH=projects/3P-ADMM-PC2/src PYTHONUTF8=1 python crypto/test_paillier.py +PYTHONPATH=projects/3P-ADMM-PC2/src PYTHONUTF8=1 python crypto/test_quantization.py +PYTHONPATH=projects/3P-ADMM-PC2/src PYTHONUTF8=1 python crypto/test_full_chain.py +``` + +Results: +- Build gfx1151: PASS -- 1 benign warning (`--ld-path` unused during linking); DLL exports + confirmed (dumpbin: init_gpu @ 0x28A0, run_modexp @ 0x28B0) +- modexp vs gmpy2 gold: 6144/6144 exact match + n at 255/510/511/768/1022/1023/1535/2046 bits, 256/batch, 3 reps, varied m_bits +- hipFFT double-FFT rounding margin holds on gfx1151 (RDNA3.5, wave32, fp64 via f64 ALU): + 0 mismatches across all 6144 crypto-significant cases +- CPU regression: test_paillier PASS, test_quantization PASS, test_full_chain PASS + (max error 1.33e-09, identical to gfx90a/gfx1100) +- No fork code change needed (zero-delta follower; same commit as gfx90a) + +LOW-NUMERIC-RISK target confirmed: the gfx1151 RDNA3.5 FP-divergence class does NOT apply here +because the correctness gate is BIT-EXACT integer modexp against gmpy2 (not FP tolerance), and +the Z2Z double-FFT rounding margin holds at < 0.5 limb error even on gfx1151. + +VERDICT: PASS -- port-ready -> completed + +## Validation 2026-06-04 (windows-gfx1101, AMD Radeon PRO V710, RDNA3, ROCm 7.14 / TheRock) + +validated_sha: 6ef301f3204579779bbaa1f32a466934f720903a (zero-churn follower; fork untouched) + +GPU arch: gfx1101 (AMD Radeon PRO V710, RDNA3, wave32). Host = the gfx1101+gfx1201 Windows +workstation (memory windows-gfx1101-gfx1201-host); HIP_VISIBLE_DEVICES=0 pins gfx1101 (the +only visible device this process). TheRock pip ROCm SDK 7.14 (rocm-sdk-devel @ 60850-d34cbb64). + +### Windows gfx1101 build recipe (REUSABLE; scripts in agent_space/3p-admm-win/, not committed) +Single ROCm root = the pip rocm-sdk-devel tree +`.../site-packages/_rocm_sdk_devel` (clang++, hipcc, hip cmake, hipfft.lib all present). +build_win.sh adapts gpu/build_hip.sh for Windows, same two adaptations the gfx1151 note found: +1. NO -fPIC (clang rejects it for the x86_64-pc-windows-msvc target; PIC is implicit on Windows). +2. Export init_gpu/run_modexp via a `.def` (`-Wl,/DEF:lib_cufft.def`) -- extern "C" funcs are + not auto-exported from a Windows DLL. +Link `-L$ROCM/lib -lhipfft` (resolves hipfft.lib). hipcc all-clang, `--offload-arch=gfx1101`. +Output lib_cufft.dll (143 KB; only the benign `--ld-path unused` warning). PATH must include +`$ROCM/bin:$ROCM/lib/llvm/bin` so hipcc finds clang + lld-link. + +### Windows DLL runtime (no per-arch library package needed) +There is NO _rocm_sdk_libraries_gfx1101 package on this host; the multi-arch `_rocm_sdk_libraries` +plus rocFFT runtime kernel generation cover gfx1101 (AMD_LOG_LEVEL=3 shows rocFFT RTC-compiling +its FFT kernels at run time). The validate harness calls `os.add_dll_directory()` for +`_rocm_sdk_core/bin`, `_rocm_sdk_devel/bin`, and `_rocm_sdk_libraries/bin` before loading +lib_cufft.dll, so TheRock's amdhip64_7.dll / amd_comgr.dll / rocm_kpack.dll / hiprtc0714.dll / +hipfft.dll / rocfft.dll load instead of the broken System32 Adrenalin amdhip64. gmpy2 was +pip-installed into the TheRock venv. The repo's crypto/paillier_gpu.py hardcodes +`/tmp/lib_cufft.so`; the harness monkeypatches `paillier_gpu._lib` to the built DLL. + +### Result +``` +# build (HIP_ARCH=gfx1101) +bash utils/timeit.sh 3P-ADMM-PC2 compile -- bash agent_space/3p-admm-win/build_win.sh +# modexp gold (HIP_VISIBLE_DEVICES=0 = gfx1101) +HIP_VISIBLE_DEVICES=0 .../venv/Scripts/python.exe agent_space/3p-admm-win/validate_win.py +``` +- Build gfx1101: PASS (lib_cufft.dll, exports init_gpu/run_modexp). +- modexp vs gmpy2 gold: 6144/6144 EXACT match -- n at 255/510/511/768/1022/1023/1535/2046 bits, + 256/batch, 3 reps/size, varied m_bits (64 / nb/2 / nb-1), g < n precondition honored. +- hipFFT double-FFT rounding margin holds on gfx1101 (RDNA3, wave32, fp64): 0 mismatches. +- Native gfx1101 dispatch: only gfx1101 visible (HIP_VISIBLE_DEVICES=0), hipGetDevice success, + rocFFT RTC-compiled device kernels; bit-exact GPU results (the .cu kernels are GPU-only, so a + silent CPU fallback is impossible). No fork code change (zero-churn follower; same 6ef301f). + +VERDICT: PASS -- port-ready -> completed. gfx1101 matches gfx90a/gfx1100/gfx1151 exactly. +This is the first real GPU validation on the new Windows host -- it establishes the reusable +Windows ROCm build + DLL-runtime recipe for the rest of the gfx1101/gfx1201 sweep. + +## Validation 2026-06-04 (windows-gfx1201, AMD Radeon RX 9070 XT, RDNA4, ROCm 7.14 / TheRock) + +validated_sha: 6ef301f3204579779bbaa1f32a466934f720903a (zero-churn follower; fork untouched) + +GPU arch: gfx1201 (AMD Radeon RX 9070 XT, RDNA4, wave32). HIP_VISIBLE_DEVICES=1 pins gfx1201 +(detect_arch confirms device 1 = gfx1201). Same host/recipe as the gfx1101 run above; only +`--offload-arch=gfx1201` and the device pin change. +``` +HIP_ARCH=gfx1201 OUT=.../lib_cufft_gfx1201.dll bash agent_space/3p-admm-win/build_win.sh +HIP_VISIBLE_DEVICES=1 LIB_CUFFT=.../lib_cufft_gfx1201.dll .../python.exe agent_space/3p-admm-win/validate_win.py +``` +- Build gfx1201: PASS (lib_cufft_gfx1201.dll, exports init_gpu/run_modexp). +- modexp vs gmpy2 gold: 6144/6144 EXACT match (same 8 moduli x 3 reps x 256/batch as gfx1101). +- hipFFT double-FFT rounding margin holds on gfx1201 (RDNA4, wave32, fp64): 0 mismatches. +- No fork code change (zero-churn follower; same 6ef301f). RDNA4 FP behavior irrelevant -- the + gate is bit-exact integer modexp, no wave-size surface in the live path. + +VERDICT: PASS -- port-ready -> completed. Both Windows archs (gfx1101 + gfx1201) now match +gfx90a/gfx1100/gfx1151 exactly; all five platforms terminal -> PR-ready. diff --git a/projects/3P-ADMM-PC2/plan.md b/projects/3P-ADMM-PC2/plan.md new file mode 100644 index 00000000..9a5a418f --- /dev/null +++ b/projects/3P-ADMM-PC2/plan.md @@ -0,0 +1,186 @@ +# Port plan: 3P-ADMM-PC2 (linux-gfx90a lead) + +## Project +- Name: 3P-ADMM-PC2 +- Upstream: https://github.com/Samarvivian/3P-ADMM-PC2 +- Default branch: `features` (per upstream.json; README also lists `releases`, `master`) +- What it is: a privacy-preserving distributed ADMM framework for LASSO. CPU does the + optimization (Paillier homomorphic encryption + quantized ADMM, gmpy2). The GPU's ONLY + job is to accelerate the Paillier modular exponentiation `g^m mod n` over a batch of + big integers, via large-integer polynomial multiplication. The ADMM math itself + (`admm/*.py`, `protocol/*.py`) is pure Python/numpy/gmpy2 -- there are NO GPU + linear-algebra kernels, no sparse/dense matvec, no cuBLAS/cuSPARSE/cuSOLVER. The GPU + surface is crypto big-integer arithmetic. + +## Existing AMD support +- None. NVIDIA-only: README targets CUDA 12.1, `-arch=sm_86` (RTX A4000/A2000), links + `-lcudart -lcufft`. No HIP path, no OpenCL/Vulkan/SYCL path. A ROCm/HIP port of the + CUDA ModExp kernel adds clear value. +- Decision: PROCEED with a mechanical CUDA->HIP port of the live GPU path. This is NOT a + CUTLASS/Hopper/perf-rewrite case (confirmed: no CUTLASS, no CuTe, no wgmma, no warp + specialization anywhere). Correctness-first mechanical port is the right and sufficient + first step. + +## Build classification: ad-hoc nvcc (NOT CMake, NOT a torch extension, NOT a Makefile) +- Evidence: `find` shows no CMakeLists.txt, no Makefile, no setup.py, no pyproject.toml. + `requirements.txt` is `fastapi`, `uvicorn`, `numpy` only -- no torch, no pycuda listed. +- The build is hand-rolled `nvcc` commands embedded in README.md sec. "1. 每次开机重新编译GPU库" + (lines 161-198): compile `gpu/cufft_modexp.cu` with a generated C wrapper `wr_cufft.cu` + (exposing `init_gpu(int)` / `run_modexp(...)`), device-link, then `g++ -shared` into + `/tmp/lib_cufft.so` linking `-lcuda -lcudart -lcufft`. +- The Python side (`crypto/paillier_gpu.py`, `protocol/edge_crt_helper.py`) loads + `/tmp/lib_cufft.so` by absolute path via `ctypes.CDLL` and calls `run_modexp`/`init_gpu`. +- ext_type recorded as `nvcc-shared` (custom; closest to Strategy A, "only the `.cu` sees + the GPU toolchain"). Set in upstream.json + status.json. + +## Port strategy: A-flavored compat header + a HIP build recipe (no CMake to gate) +Rationale: there is no CMake `enable_language(HIP)` to flip and no torch hipify to lean on. +The minimal-footprint analogue of Strategy A here is: +1. Add ONE compat header `gpu/cuda_to_hip.h` that, under `__HIP_PLATFORM_AMD__`/`USE_HIP`, + `#include ` + `` and aliases the exact CUDA symbols + the live kernel uses to their HIP spellings; else `#include `+``. + Include ``/`` BEFORE the HIP runtime (PORTING_GUIDE gpuRIR lesson: + inside a .cu compiled as HIP, host memcpy/memset/free/calloc can otherwise resolve to + HIP `__device__` overloads). `cufft_modexp.cu` uses `calloc`/`free`/`round` on the host. +2. `#include "cuda_to_hip.h"` at the top of `gpu/cufft_modexp.cu` in place of the bare + `cuda_runtime.h`/`cufft.h` includes. Keep all `cudaXxx`/`cufftXxx` spellings in the body; + the header aliases them. This keeps the NVIDIA build byte-identical. +3. Provide a ROCm build recipe alongside the README's nvcc one: `hipcc` compiling the same + `cufft_modexp.cu` + wrapper, linking `-lhipfft`, producing `/tmp/lib_cufft.so` with the + SAME `init_gpu`/`run_modexp` C ABI so the Python ctypes side is UNCHANGED. Add it as a + small `gpu/build_hip.sh` (and document in notes.md). Default arch from an env/arg, never + a literal (so gfx1100/gfx1151 followers reuse it with only an arch change). +4. The hardcoded `/mnt/3p-admm-pc2/...` absolute paths in the Python (lib path, kernel + source paths) are an upstream portability wart, not a HIP issue; for our build+validate + we point the loader at our actual `/tmp/lib_cufft.so` path. Do not refactor upstream's + path scheme as part of the port (out of scope; flag in notes). + +## CUDA surface inventory +LIVE GPU code (what actually runs in the protocol), `gpu/cufft_modexp.cu` (204 lines): +- Library: cuFFT. `cufftHandle`, `cufftPlanMany(...CUFFT_Z2Z...)`, `cufftExecZ2Z` (FORWARD/ + INVERSE), `cufftDestroy`, `cufftDoubleComplex`. -> hipFFT: `hipfftHandle`, + `hipfftPlanMany`, `hipfftExecZ2Z` (`HIPFFT_FORWARD`/`HIPFFT_BACKWARD`), `hipfftDestroy`, + `hipfftDoubleComplex`. APIs are ~1:1; watch the FORWARD/INVERSE enum spelling + (CUFFT_INVERSE -> HIPFFT_BACKWARD) and that hipFFT complex type aliases `double2`. +- Runtime: `cudaMalloc`, `cudaFree`, `cudaMemcpy` (H2D/D2H), `cudaMemcpyHostToDevice`, + `cudaMemcpyDeviceToHost`. -> `hipMalloc`/`hipFree`/`hipMemcpy` 1:1. +- Kernels (all plain, no intrinsics): `cmul`, `norm_round`, `load_complex`, `carry_prop`, + `shr_kernel`, `sub_correct`, `extract_bit`, `cond_copy`. Launches use `dim3`/blockDim + 32 or 256; `cond_copy<<>>` is a strided memcpy (`for i=j;i PyCUDA JITs `gpu/modexp.cu` / `gpu/modexp_v2.cu`: +- `import pycuda.autoinit / pycuda.driver / pycuda.compiler.SourceModule`. PyCUDA is NOT in + requirements.txt and this module is a standalone `__main__` self-test, NOT called by the + protocol/experiments. `modexp.cu`: one-thread-per-task scalar bigint, no intrinsics. + `modexp_v2.cu`: 128-thread block, all `__syncthreads()` (block barrier, wave-size + agnostic), no warp intrinsics. Hardcodes `/mnt/3p-admm-pc2/gpu/*.cu` source paths. + DECISION: out of scope for the lead port. PyCUDA-on-ROCm is non-standard and this path is + not exercised by any validatable workload. Note it; do not port unless validation needs it. + +DEAD/experimental files (in `gpu/`, referenced by NOTHING in the repo -- no Python import, +no `<<<...>>>` launch site, not built by the README): `reg_ntt_modexp.cu`, +`reg_ntt_modexp_v2.cu`, `reg_ntt_modexp_v3.cu`, `ntt_modexp.cu`, `ntt_modexp_v2.cu`, +`ntt_modexp_fast.cu`, `ntt_funcs.cu`, `ntt_kernel2.cu`, `ntt_wrapper.cu`, +`ntt_wrapper2.cu`, `modexp_final.cu`, `cufft_modexp` siblings. The `reg_ntt_*` and +`ntt_funcs` files DO contain `__shfl_xor_sync(0xffffffff, ..., delta)` warp-NTT butterflies +launched (conceptually) with blockDim=32 (one warp/task) and `for(i=lane;i<...;i+=32)` +strides -- i.e. classic wave64 hazards. But they are NOT in any live path. DECISION: DO NOT +port the dead files. Porting them would be effort against unused code AND would expose the +wave64 fault class for no validatable benefit. If a future maintainer wants the register-NTT +path on AMD, that is a separate task (see Open questions). + +## Risk list +- LOW overall: the live kernel is the easy class (cuFFT + scalar bigint, no warp ops). +- hipFFT enum/type drift: `CUFFT_INVERSE` -> `HIPFFT_BACKWARD`, `CUFFT_FORWARD` -> + `HIPFFT_FORWARD`, `cufftDoubleComplex` -> `hipfftDoubleComplex`. Verify `hipfftPlanMany` + batched (the `N` batch arg) semantics match cuFFT's (they do in ROCm 7.2.x). Alias these + in the compat header so the body is untouched. +- hipFFT precision: this is a NUMERIC-CORRECTNESS-VIA-FFT bigint multiply. The kernel does + Z2Z (double-precision) FFT then `round()` to recover exact integer coefficients + (`norm_round` + `carry_prop`). gfx90a has full-rate fp64, so double FFT round-off should + stay within the < 0.5 rounding margin the algorithm relies on, same as CUDA. BUT hipFFT's + twiddle/rounding differs from cuFFT's bit-for-bit, so the recovered coefficients could + occasionally round the wrong way if the design margin is thin. This is the one real risk + to watch in validation -- check `g^m mod n` against gmpy2 over many random inputs, not + just one. (No bit-exactness is required vs CUDA; only that the final integer matches the + gmpy2 gold.) +- Wave size: NOT exposed in the live path (no warp intrinsics, no hardcoded-32 warp logic; + `<<>>` blocks are independent-thread strided copies). The MULTI-ARCH per-arch + warpSize standard therefore needs NO code change for the lead port: there is no device + warp constant to set and no host warpSize query to add. The follower deltas (gfx1100/ + gfx1151, wave32) are expected to be no-op revalidations (same .so rebuilt for the arch). + Flag the dead `reg_ntt_*` shfl files only so a later contributor knows they would need the + `__GFX9__`->64 / RDNA->32 + width-aware ballot treatment if ever revived. +- Build wart (not a HIP bug): hardcoded `/mnt/3p-admm-pc2` absolute paths in Python; `/tmp` + non-persistent .so. We control these for our build; do not refactor upstream's scheme. +- Include order: host `calloc`/`free`/`round` in `cufft_modexp.cu` -- put ``/ + ``/`` before `` in the compat header (gpuRIR lesson). +- No rule-of-five/texture/OOB/pitch/atomicMin classes apply (no such constructs in the live + kernel). `sub_correct`/`carry_prop`/`shr_kernel` index within per-task strided buffers; + bounds are guarded by `if(t>=N) return` and `src`, ``) + with `#include "cuda_to_hip.h"`. No body changes expected (aliases cover it). If hipFFT + needs `HIPFFT_BACKWARD` where the body has `CUFFT_INVERSE`, alias `CUFFT_INVERSE` -> + `HIPFFT_BACKWARD` in the header rather than editing the body. +- ADD `gpu/build_hip.sh` -- hipcc recipe mirroring the README nvcc steps, arch from + `${HIP_ARCH:-gfx90a}`, links `-lhipfft`, emits `/tmp/lib_cufft.so` with the same C ABI. +- (notes.md) record the build + the `/mnt` path caveat + the dead-file decision. +- DO NOT touch: any Python, the dead `gpu/*.cu` files, `admm/`, `protocol/`, `crypto/` + (except possibly a thin validate harness in agent_space, not committed to the fork). + +## Build commands (gfx90a) +ROCm build (analogue of README sec.1, via build_hip.sh): +``` +hipcc -O2 --offload-arch=gfx90a -fPIC -c gpu/cufft_modexp.cu -o /tmp/cufft_modexp.o +# generated wrapper wr_cufft.cu exposing init_gpu/run_modexp (same as README) +hipcc -O2 --offload-arch=gfx90a -fPIC -c /tmp/wr_cufft.cu -o /tmp/wr_cufft.o +hipcc -shared -fPIC /tmp/cufft_modexp.o /tmp/wr_cufft.o -lhipfft -o /tmp/lib_cufft.so +``` +(hipcc does host+device link in one step; no separate `-dlink` needed for a non-RDC build. +If RDC is wanted, add `-fgpu-rdc` to compiles and a `--hip-link` device-link step.) +Followers: same script with `HIP_ARCH=gfx1100` / `gfx1151`. + +## Test plan +GPU-validatable slice (real GPU, the load-bearing correctness gate): +- The GPU's contract is `out[i] == (g^m mod n)` over a batch. Validate the ported + `/tmp/lib_cufft.so` by calling `run_modexp` (through `crypto/paillier_gpu.py:gpu_batch_modexp` + or `protocol/edge_crt_helper.py:gpu_modexp_diff_g`) on many random (g, m, n=p*q) triples + and comparing against `gmpy2.powmod(g, m, n)` -- the exact CPU gold the repo already uses + in `gpu/gpu_modexp.py:__main__` and throughout paillier_gpu.py. Pass = all batches match + gmpy2 (with the final `% n`). Use realistic Paillier sizes (n^2 up to ~2048-bit, LEN=128 + base-65536 limbs) AND small sizes; vary m_bits. A single-input check is insufficient + given the hipFFT rounding-margin risk -- run a few thousand. +- Higher-level GPU functional check: `crypto/paillier_gpu.py:encrypt_batch_gpu` / + `encrypt_batch_gpu_fast` produce ciphertexts; decrypt with the CPU `crypto/paillier.py` + and confirm round-trip equals the plaintext batch. This exercises the GPU g^m path inside + the real encryption routine. Buildable as a small local harness (agent_space), not + committed to the fork. +Non-GPU regression set (must not regress; pure CPU, runnable as-is up to the /mnt path): +- `crypto/test_paillier.py` (keygen/enc/dec/homomorphic add+mul), `crypto/test_quantization.py` + (gamma1/gamma2 quantization), `crypto/test_full_chain.py`. These are CPU/gmpy2 only and + the port does not touch them; run them to confirm no collateral breakage. +- NOT runnable in CI here (need multi-node SSH / matpool cloud, out of scope): the + `experiments/test_distributed_pc2.py`, `test_large_scale.py`, `test_cen_vs_dis.py`, + `monitor_gpu.py`, and `protocol/master_node.py` end-to-end protocol. The GPU-slice + + encrypt round-trip above is the representative validatable proxy for the GPU work these + drive. + +## Open questions +- Does upstream want the dead `reg_ntt_*` register-NTT variants ported too? They are the + faster-claimed path in the README perf table but are not wired into the live protocol and + carry the wave64 warp-NTT hazard. Lead port deliberately excludes them; revisit only if + the user/upstream asks (would be a separate AMD-native or width-aware effort). +- hipFFT batched-plan numeric margin: confirm in validation that double-FFT bigint multiply + recovers exact limbs across the full input range on gfx90a; if a thin-margin miss appears, + options are a larger FFT length or a wider integer accumulation -- but expect it to pass + given fp64 parity. +- PyInstaller `dist/3P-ADMM-PC2` binary and the FastAPI web app are NVIDIA/x86 packaging + artifacts; not part of the GPU port. diff --git a/projects/3P-ADMM-PC2/stats.jsonl b/projects/3P-ADMM-PC2/stats.jsonl new file mode 100644 index 00000000..fb576081 --- /dev/null +++ b/projects/3P-ADMM-PC2/stats.jsonl @@ -0,0 +1,14 @@ +{"kind":"phase","ts":"2026-06-02T06:08:44Z","phase":"compile","seconds":0.005,"exit":127,"cmd":"bash gpu/build_hip.sh"} +{"kind":"phase","ts":"2026-06-02T06:08:50Z","phase":"compile","seconds":0.529,"exit":1,"cmd":"bash /var/lib/jenkins/moat/projects/3P-ADMM-PC2/src/gpu/build_hip.sh"} +{"kind":"phase","ts":"2026-06-02T06:09:13Z","phase":"compile","seconds":3.183,"exit":0,"cmd":"bash /var/lib/jenkins/moat/projects/3P-ADMM-PC2/src/gpu/build_hip.sh"} +{"kind":"phase","ts":"2026-06-02T06:19:43Z","phase":"compile","seconds":3.244,"exit":0,"cmd":"bash projects/3P-ADMM-PC2/src/gpu/build_hip.sh"} +{"kind":"phase","ts":"2026-06-02T06:19:58Z","phase":"compile","seconds":4.799,"exit":0,"cmd":"bash projects/3P-ADMM-PC2/src/gpu/build_hip.sh"} +{"kind":"phase","ts":"2026-06-02T06:21:24Z","phase":"test","seconds":70.009,"exit":0,"cmd":"python3 agent_space/3p-admm/validate.py"} +{"kind":"phase","ts":"2026-06-04T17:58:18Z","phase":"compile","seconds":0.813,"exit":1,"cmd":"bash agent_space/3p-admm-win/build_win.sh"} +{"kind":"phase","ts":"2026-06-04T17:58:42Z","phase":"compile","seconds":4.596,"exit":0,"cmd":"bash agent_space/3p-admm-win/build_win.sh"} +{"kind":"phase","ts":"2026-06-04T18:00:46Z","phase":"test","seconds":0.326,"exit":1,"cmd":"C:/Users//AppData/Local/Programs/Python/Python313/python.exe agent_space/3p-admm-win/validate_win.py"} +{"kind":"phase","ts":"2026-06-04T18:01:06Z","phase":"test","seconds":1.029,"exit":1,"cmd":"C:/Users//AppData/Local/Programs/Python/Python313/python.exe agent_space/3p-admm-win/validate_win.py"} +{"kind":"phase","ts":"2026-06-04T18:02:27Z","phase":"test","seconds":12.169,"exit":0,"cmd":"C:/Users//AppData/Local/Programs/Python/Python313/python.exe agent_space/3p-admm-win/validate_win.py"} +{"kind": "tokens", "ts": "2026-06-04T18:05:34Z", "tokens": 74493, "source": "validator"} +{"kind":"phase","ts":"2026-06-02T06:26:57Z","phase":"compile","seconds":2.378,"exit":0,"cmd":"bash projects/3P-ADMM-PC2/src/gpu/build_hip.sh"} +{"kind":"phase","ts":"2026-06-02T06:28:26Z","phase":"test","seconds":13.720,"exit":0,"cmd":"python3 /var/lib/jenkins/moat/agent_space/3p-admm-gfx1100/validate.py"} diff --git a/projects/3P-ADMM-PC2/status.json b/projects/3P-ADMM-PC2/status.json new file mode 100644 index 00000000..42393324 --- /dev/null +++ b/projects/3P-ADMM-PC2/status.json @@ -0,0 +1,174 @@ +{ + "schema_version": 3, + "name": "3P-ADMM-PC2", + "upstream_url": "https://github.com/Samarvivian/3P-ADMM-PC2", + "fork_url": "https://github.com/AMD-Ecosystem/3P-ADMM-PC2", + "fork_default_branch": "features", + "priority": 4.906, + "ext_type": "nvcc-shared", + "adopted_at": "2026-05-30T00:46:51Z", + "updated_at": "2026-08-07T07:04:24Z", + "head_sha": "2c1cafae0204fc8a728e046caae153e9ee1df834", + "pr_url": "https://github.com/Samarvivian/3P-ADMM-PC2/pull/10", + "pr_state": "merged", + "pr_number": 10, + "pr_opened_at": "2026-06-08T16:22:00Z", + "pr_merged_at": "2026-06-10T14:23:02Z", + "porting": null, + "waivers": {}, + "license_clearance": { + "approved_by": "jeffdaily", + "at": "2026-08-06T02:41:26Z", + "tier": 4, + "note": "carries the org review of 2026-08-06: the risk is in USING these projects, not in contributing to them; covers this project only" + }, + "license_spdx": "no licence file", + "upstream_repo_id": 1092826294, + "stage": "review-passed", + "platforms": { + "linux-gfx90a": { + "state": "completed", + "blocked": false, + "blocked_reason": null, + "validated_sha": "096cc8f37bbd3cb974f159416579aae93cb3ea9a", + "started_at": "2026-06-02T06:08:10Z", + "completed_at": "2026-06-07T15:15:22Z", + "updated_at": "2026-06-10T14:23:02Z", + "stats": { + "tokens_total": 0, + "tokens_approx": true, + "wall_seconds": { + "thinking": 0, + "compile": 0, + "test": 0, + "misc": 0 + }, + "session_count": 0, + "first_session_at": null, + "last_session_at": null + }, + "last_agent": "moat-checkup", + "carry_forward": { + "to": "096cc8f37bbd3cb974f159416579aae93cb3ea9a", + "method": "source-class", + "detail": "bash comment-only edit in gpu/build_hip.sh (jargon scrub); comments are never executed so the build recipe runs byte-identical commands and produces an identical .so on every arch -- inert", + "at": "2026-06-07T15:15:22Z" + } + }, + "linux-gfx1100": { + "state": "completed", + "blocked": false, + "blocked_reason": null, + "validated_sha": "2c1cafae0204fc8a728e046caae153e9ee1df834", + "started_at": "2026-06-02T06:40:00Z", + "completed_at": "2026-06-08T18:11:05Z", + "updated_at": "2026-06-08T18:11:05Z", + "stats": { + "tokens_total": 0, + "tokens_approx": true, + "wall_seconds": { + "thinking": 0, + "compile": 0, + "test": 0, + "misc": 0 + }, + "session_count": 1, + "first_session_at": "2026-06-02T06:40:00Z", + "last_session_at": "2026-06-02T06:50:00Z" + }, + "last_agent": "validator", + "carry_forward": { + "to": "2c1cafae0204fc8a728e046caae153e9ee1df834", + "method": "source-class", + "detail": "README: document ROCm/HIP build (doc-only)", + "at": "2026-06-08T18:11:05Z" + } + }, + "windows-gfx1101": { + "state": "completed", + "blocked": false, + "blocked_reason": null, + "validated_sha": "2c1cafae0204fc8a728e046caae153e9ee1df834", + "started_at": null, + "completed_at": "2026-06-08T18:11:05Z", + "updated_at": "2026-06-08T18:11:05Z", + "stats": { + "tokens_total": 0, + "tokens_approx": true, + "wall_seconds": { + "thinking": 0, + "compile": 0, + "test": 0, + "misc": 0 + }, + "session_count": 0, + "first_session_at": null, + "last_session_at": null + }, + "last_agent": "validator", + "carry_forward": { + "to": "2c1cafae0204fc8a728e046caae153e9ee1df834", + "method": "source-class", + "detail": "README: document ROCm/HIP build (doc-only)", + "at": "2026-06-08T18:11:05Z" + } + }, + "windows-gfx1201": { + "state": "completed", + "blocked": false, + "blocked_reason": null, + "validated_sha": "2c1cafae0204fc8a728e046caae153e9ee1df834", + "started_at": null, + "completed_at": "2026-06-08T18:11:05Z", + "updated_at": "2026-06-08T18:11:05Z", + "stats": { + "tokens_total": 0, + "tokens_approx": true, + "wall_seconds": { + "thinking": 0, + "compile": 0, + "test": 0, + "misc": 0 + }, + "session_count": 0, + "first_session_at": null, + "last_session_at": null + }, + "last_agent": "validator", + "carry_forward": { + "to": "2c1cafae0204fc8a728e046caae153e9ee1df834", + "method": "source-class", + "detail": "README: document ROCm/HIP build (doc-only)", + "at": "2026-06-08T18:11:05Z" + } + }, + "windows-gfx1151": { + "state": "completed", + "blocked": false, + "blocked_reason": null, + "validated_sha": "2c1cafae0204fc8a728e046caae153e9ee1df834", + "started_at": null, + "completed_at": "2026-06-08T18:11:05Z", + "updated_at": "2026-06-08T18:11:05Z", + "stats": { + "tokens_total": 0, + "tokens_approx": true, + "wall_seconds": { + "thinking": 0, + "compile": 0, + "test": 0, + "misc": 0 + }, + "session_count": 0, + "first_session_at": null, + "last_session_at": null + }, + "carry_forward": { + "to": "2c1cafae0204fc8a728e046caae153e9ee1df834", + "method": "source-class", + "detail": "README: document ROCm/HIP build (doc-only)", + "at": "2026-06-08T18:11:05Z" + } + } + } +} From ff55bb8fa69263ea0da9ab97876beba8550eee3f Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Sat, 8 Aug 2026 23:52:39 +0000 Subject: [PATCH 2/3] 3P-ADMM-PC2: revalidate linux-gfx90a (carry-forward, doc-only delta) + record CUDA no-regression gate --- projects/3P-ADMM-PC2/notes.md | 58 ++++++++++++++++++++++++++++++++ projects/3P-ADMM-PC2/stats.jsonl | 2 ++ projects/3P-ADMM-PC2/status.json | 14 ++++---- 3 files changed, 67 insertions(+), 7 deletions(-) diff --git a/projects/3P-ADMM-PC2/notes.md b/projects/3P-ADMM-PC2/notes.md index b096c68c..d1d3add9 100644 --- a/projects/3P-ADMM-PC2/notes.md +++ b/projects/3P-ADMM-PC2/notes.md @@ -364,3 +364,61 @@ HIP_VISIBLE_DEVICES=1 LIB_CUFFT=.../lib_cufft_gfx1201.dll .../python.exe agent_s VERDICT: PASS -- port-ready -> completed. Both Windows archs (gfx1101 + gfx1201) now match gfx90a/gfx1100/gfx1151 exactly; all five platforms terminal -> PR-ready. + +## Validation 2026-08-08 (revalidate, linux-gfx90a, GCD 1) + +linux-gfx90a was still `completed` at validated_sha `096cc8f37bbd3cb974f159416579aae93cb3ea9a` +while fork head had moved to `2c1cafae0204fc8a728e046caae153e9ee1df834` (the other four platforms +had already been carried forward at 2026-06-08T18:11:05Z for the same delta; gfx90a's +`advance_head` pass apparently ran without the fork cloned locally, so `_classify_safe` returned +None and it fell back to conservative revalidate -- lagged rather than actually regressed). + +Delta (validated_sha..head_sha) is exactly one commit: +``` +2c1cafa [ROCm] Document the AMD GPU (ROCm/HIP) build in the README +``` +``` +git diff --stat 096cc8f37bbd3cb974f159416579aae93cb3ea9a..2c1cafae0204fc8a728e046caae153e9ee1df834 + README.md | 6 ++++++ + 1 file changed, 6 insertions(+) +``` + +Classification: +``` +python3 utils/moatlib.py classify 3P-ADMM-PC2 096cc8f37bbd3cb974f159416579aae93cb3ea9a 2c1cafae0204fc8a728e046caae153e9ee1df834 +class=doc-only arch_independent=True inert=True +``` +Doc-only, arch-independent, inert -- the carry-forward shortcut applies (validator.md step 1-2). +No GPU re-run needed; carried forward without touching the .so or re-running the modexp gold +gate that already passed at 096cc8f. + +``` +python3 utils/moatlib.py carry-forward 3P-ADMM-PC2 linux-gfx90a 2c1cafae0204fc8a728e046caae153e9ee1df834 source-class "README: document ROCm/HIP build (doc-only); classify verdict class=doc-only arch_independent=True inert=True" +``` + +### CUDA no-regression gate (not previously recorded at this head_sha) + +nvcc 12.8 via `/opt/conda/envs/cuda-12.8/bin/nvcc`, arch pinned `-arch=sm_80` (no NVIDIA GPU on +this host; native autodetection would silently degrade). `gpu/cuda_to_hip.h` falls through to +the original CUDA headers when `__HIP__`/`__HIP_PLATFORM_AMD__`/`USE_HIP` are undefined, so the +CUDA branch of the shim is exercised unmodified. +``` +bash utils/timeit.sh 3P-ADMM-PC2 cuda-compile -- \ + /opt/conda/envs/cuda-12.8/bin/nvcc -arch=sm_80 -O2 --compiler-options '-fPIC' \ + -dc projects/3P-ADMM-PC2/src/gpu/cufft_modexp.cu -o /tmp/cufft_modexp_cudacheck.o +# wrapper + device link (mirrors README steps 2-4) +/opt/conda/envs/cuda-12.8/bin/nvcc -arch=sm_80 -O2 --compiler-options '-fPIC' -dc /tmp/wr_cufft_cudacheck.cu -o /tmp/wr_cufft_cudacheck.o +/opt/conda/envs/cuda-12.8/bin/nvcc -arch=sm_80 --compiler-options '-fPIC' -dlink /tmp/cufft_modexp_cudacheck.o /tmp/wr_cufft_cudacheck.o -o /tmp/dl_cufft_cudacheck.o +``` +Result: PASS, no errors or warnings on any of the three nvcc invocations -- pure passthrough +confirmed, no CUDA regression. (This gate is per head_sha, not per arch; recorded here so no +other arch re-runs it for `2c1cafae0204fc8a728e046caae153e9ee1df834`.) + +### Other gates +- `python3 utils/jargon.py --port 3P-ADMM-PC2` -> clean. +- ROCm build documented in README.md sec.1 alongside the CUDA nvcc recipe (house style; already + present from the 2c1cafa commit that motivated this revalidation). +- Fork working tree: clean after classification/CUDA-compile checks (compile artifacts written + only to /tmp, nothing under `projects/3P-ADMM-PC2/src` modified). + +VERDICT: PASS (carried forward, no GPU re-run) -- linux-gfx90a completed @ 2c1cafae0204fc8a728e046caae153e9ee1df834. diff --git a/projects/3P-ADMM-PC2/stats.jsonl b/projects/3P-ADMM-PC2/stats.jsonl index fb576081..a2328ee7 100644 --- a/projects/3P-ADMM-PC2/stats.jsonl +++ b/projects/3P-ADMM-PC2/stats.jsonl @@ -12,3 +12,5 @@ {"kind": "tokens", "ts": "2026-06-04T18:05:34Z", "tokens": 74493, "source": "validator"} {"kind":"phase","ts":"2026-06-02T06:26:57Z","phase":"compile","seconds":2.378,"exit":0,"cmd":"bash projects/3P-ADMM-PC2/src/gpu/build_hip.sh"} {"kind":"phase","ts":"2026-06-02T06:28:26Z","phase":"test","seconds":13.720,"exit":0,"cmd":"python3 /var/lib/jenkins/moat/agent_space/3p-admm-gfx1100/validate.py"} +{"kind":"phase","ts":"2026-08-08T23:51:37Z","phase":"cuda-compile","seconds":1.166,"exit":0,"cmd":"/opt/conda/envs/cuda-12.8/bin/nvcc -arch=sm_80 -O2 --compiler-options -fPIC -dc projects/3P-ADMM-PC2/src/gpu/cufft_modexp.cu -o /tmp/cufft_modexp_cudacheck.o"} +{"kind":"phase","ts":"2026-08-08T23:51:50Z","phase":"cuda-compile","seconds":0.998,"exit":0,"cmd":"bash -c \n/opt/conda/envs/cuda-12.8/bin/nvcc -arch=sm_80 -O2 --compiler-options \"-fPIC\" -dc /tmp/wr_cufft_cudacheck.cu -o /tmp/wr_cufft_cudacheck.o &&\n/opt/conda/envs/cuda-12.8/bin/nvcc -arch=sm_80 --compiler-options \"-fPIC\" -dlink /tmp/cufft_modexp_cudacheck.o /tmp/wr_cufft_cudacheck.o -o /tmp/dl_cufft_cudacheck.o\n"} diff --git a/projects/3P-ADMM-PC2/status.json b/projects/3P-ADMM-PC2/status.json index 42393324..e178df75 100644 --- a/projects/3P-ADMM-PC2/status.json +++ b/projects/3P-ADMM-PC2/status.json @@ -7,7 +7,7 @@ "priority": 4.906, "ext_type": "nvcc-shared", "adopted_at": "2026-05-30T00:46:51Z", - "updated_at": "2026-08-07T07:04:24Z", + "updated_at": "2026-08-08T23:52:05Z", "head_sha": "2c1cafae0204fc8a728e046caae153e9ee1df834", "pr_url": "https://github.com/Samarvivian/3P-ADMM-PC2/pull/10", "pr_state": "merged", @@ -30,10 +30,10 @@ "state": "completed", "blocked": false, "blocked_reason": null, - "validated_sha": "096cc8f37bbd3cb974f159416579aae93cb3ea9a", + "validated_sha": "2c1cafae0204fc8a728e046caae153e9ee1df834", "started_at": "2026-06-02T06:08:10Z", - "completed_at": "2026-06-07T15:15:22Z", - "updated_at": "2026-06-10T14:23:02Z", + "completed_at": "2026-08-08T23:52:05Z", + "updated_at": "2026-08-08T23:52:05Z", "stats": { "tokens_total": 0, "tokens_approx": true, @@ -49,10 +49,10 @@ }, "last_agent": "moat-checkup", "carry_forward": { - "to": "096cc8f37bbd3cb974f159416579aae93cb3ea9a", + "to": "2c1cafae0204fc8a728e046caae153e9ee1df834", "method": "source-class", - "detail": "bash comment-only edit in gpu/build_hip.sh (jargon scrub); comments are never executed so the build recipe runs byte-identical commands and produces an identical .so on every arch -- inert", - "at": "2026-06-07T15:15:22Z" + "detail": "README: document ROCm/HIP build (doc-only); classify verdict class=doc-only arch_independent=True inert=True", + "at": "2026-08-08T23:52:05Z" } }, "linux-gfx1100": { From 991a8dc5272b1a98c2a49863fd8e89c919dfdfb2 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Sat, 8 Aug 2026 23:55:33 +0000 Subject: [PATCH 3/3] 3P-ADMM-PC2: record the validator's token cost --- projects/3P-ADMM-PC2/stats.jsonl | 1 + 1 file changed, 1 insertion(+) diff --git a/projects/3P-ADMM-PC2/stats.jsonl b/projects/3P-ADMM-PC2/stats.jsonl index a2328ee7..781a6275 100644 --- a/projects/3P-ADMM-PC2/stats.jsonl +++ b/projects/3P-ADMM-PC2/stats.jsonl @@ -14,3 +14,4 @@ {"kind":"phase","ts":"2026-06-02T06:28:26Z","phase":"test","seconds":13.720,"exit":0,"cmd":"python3 /var/lib/jenkins/moat/agent_space/3p-admm-gfx1100/validate.py"} {"kind":"phase","ts":"2026-08-08T23:51:37Z","phase":"cuda-compile","seconds":1.166,"exit":0,"cmd":"/opt/conda/envs/cuda-12.8/bin/nvcc -arch=sm_80 -O2 --compiler-options -fPIC -dc projects/3P-ADMM-PC2/src/gpu/cufft_modexp.cu -o /tmp/cufft_modexp_cudacheck.o"} {"kind":"phase","ts":"2026-08-08T23:51:50Z","phase":"cuda-compile","seconds":0.998,"exit":0,"cmd":"bash -c \n/opt/conda/envs/cuda-12.8/bin/nvcc -arch=sm_80 -O2 --compiler-options \"-fPIC\" -dc /tmp/wr_cufft_cudacheck.cu -o /tmp/wr_cufft_cudacheck.o &&\n/opt/conda/envs/cuda-12.8/bin/nvcc -arch=sm_80 --compiler-options \"-fPIC\" -dlink /tmp/cufft_modexp_cudacheck.o /tmp/wr_cufft_cudacheck.o -o /tmp/dl_cufft_cudacheck.o\n"} +{"kind": "tokens", "ts": "2026-08-08T23:55:32Z", "tokens": 81908, "source": "validator"}