From cac43ada548b576da5856dfe3081f92f5f2dbb2f Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Sun, 9 Aug 2026 00:28:00 +0000 Subject: [PATCH 1/4] aihwkit: keep this branch's project state across the trunk merge --- projects/aihwkit/notes.md | 603 +++++++++++++++++++++++++++++++++++ projects/aihwkit/plan.md | 142 +++++++++ projects/aihwkit/stats.jsonl | 63 ++++ projects/aihwkit/status.json | 139 ++++++++ 4 files changed, 947 insertions(+) create mode 100644 projects/aihwkit/notes.md create mode 100644 projects/aihwkit/plan.md create mode 100644 projects/aihwkit/stats.jsonl create mode 100644 projects/aihwkit/status.json diff --git a/projects/aihwkit/notes.md b/projects/aihwkit/notes.md new file mode 100644 index 00000000..4ced631e --- /dev/null +++ b/projects/aihwkit/notes.md @@ -0,0 +1,603 @@ +# aihwkit notes + +IBM Analog Hardware Acceleration Kit. PyTorch library with a standalone +CMake/scikit-build CUDA backend (RPUCuda) exposed via a pybind11 module +(`rpu_base`). Strategy A HIP port: one compat header plus a USE_HIP CMake +option; Torch is found for headers/linkage only (NOT torch cpp_extension). + +## Environment (linux-gfx90a) +- ROCm 7.2.1 at /opt/rocm; hipBLAS, hipRAND, hipCUB present. +- ROCm PyTorch in conda env py_3.12: torch 2.13.0a0, torch.version.hip 7.2.x, + device AMD Instinct MI250X / MI250 (gfx90a, wave64). +- This newer Torch requires C++20 (its headers use the `requires` keyword), so + the build sets RPU_CXX_STANDARD=20. Default stays 17 to match older Torch. + +## The port +- New: src/rpucuda/cuda/cuda_to_hip.h -- the only HIP-aware header. Included + once from cuda_util.h before the CUDA runtime / cuBLAS / cuRAND block, + guarded by USE_HIP (CUDA build byte-identical). Pulls / + before , then aliases the cuda*/cublas*/curand* spellings + the project uses to hip*/hipblas*/hiprand*. +- New: cmake/dependencies_hip.cmake -- enable_language(HIP), + find_package(hip/hipblas/hiprand/hipcub), defines RPU_USE_CUDA and USE_HIP. +- CMakeLists.txt: option(USE_HIP); the GPU library/test/pybind branches now + fire on USE_CUDA OR USE_HIP. Under HIP the .cu sources get LANGUAGE HIP, link + roc::hipblas hip::hiprand hip::hipcub, HIP_ARCHITECTURES from + CMAKE_HIP_ARCHITECTURES (default gfx90a when unset; no literal hardcode), IPO + forced OFF. RPU_CXX_STANDARD cache var replaces the hardcoded CXX_STANDARD 17. +- rpu_cub.h: under USE_HIP include and set + RPU_CUB_NS_QUALIFIER to `hipcub::` (hipCUB ignores CUB_NS_PREFIX wrapping). +- src/aihwkit/simulator/CMakeLists.txt: the pybind TUs get LANGUAGE HIP under + HIP (they inherit RPU_GPU's --offload-arch usage requirement, which only the + HIP/clang driver understands), and IPO OFF on the module. +- rpu_base_tiles_cuda.cpp: under USE_HIP skip the raw `cuda.h` include and use + + a `using c10::cuda::getCurrentCUDAStream` in the + at::cuda namespace, instead of which transitively + pulls hipSOLVER/hipSPARSE headers not in this build. +- rpu_linearstep_device.h: declare the ctor as `SoftBoundsRPUDeviceMetaParameter()` + not `...()` -- a template-id constructor name is ill-formed under C++20 + (latent upstream bug, exposed only because the newer Torch forces C++20). + +### Library swaps (mechanical, via the compat header) +- cuBLAS -> hipBLAS: Sgemm/Dgemm/Hgemm, Sgemv/Dgemv, Sger/Dger, Sscal/Dscal, + Scopy/Dcopy, Snrm2/Dnrm2; handle + SetStream + Get/SetPointerMode (HOST and + DEVICE pointer modes on the host handle -- hipBLAS supports both); op N/T and + status enums. CUBLAS_STATUS_LICENSE_ERROR has no hipBLAS analogue -> mapped to + HIPBLAS_STATUS_UNKNOWN so the CUBLAS_CALL switch still compiles. The + device-side cuBLAS path (kernelCublasCreateDevice/getDeviceHandle, device API) + is fully behind RPU_WITH_CUBLAS_DEVICE (OFF by default) and is never reached; + hipBLAS has no device API so leaving it gated off is correct. +- cuRAND -> hipRAND: host generator (curandGenerator_t + create/destroy/ + generate/seed/SetStream) and device per-thread state (curand_init/normal/ + uniform). GOTCHA: in CUDA, curandState, curandState_t and curandStateXORWOW + are the SAME XORWOW type; hipRAND declares hiprandState and + hiprandStateXORWOW_t as DISTINCT structs. The project mixes the spellings + (CudaArray at call sites, explicit instantiation of + CudaArray), so all three are aliased to one type + (hiprandState_t). Mapping them to different hip types caused a missing + template instantiation -> ImportError: undefined symbol + CudaArray. +- CUB -> hipCUB: DeviceReduce/DeviceScan/DeviceSegmentedReduce + BlockScan + (block-width templated, adapts to wave64 automatically). The single BlockScan + TempStorage is used once (no back-to-back reuse) so no wave64 TempStorage race. + +### bit_line_maker.cu warp-size fix (the risk path -- dietgpu-class) +The stochastic pulse train is a warp-size-coupled SERIALIZED bit format: +__ballot_sync(0xFFFFFFFF, ...) packed into 32-bit words (nK32 = (Kplus1+31)/32), +laneId = threadIdx.x & 0x1f; the pulsed-weight-updater kernels (pwu_kernel.h) +read it back word-by-word with >>5 / &0x1f / __popc. The on-device word layout +must stay byte-identical regardless of wavefront width. HIP 7.2 provides real +__ballot_sync/__shfl_sync that static_assert a 64-bit mask, and a native +__ballot on wave64 spans all 64 lanes -- which would merge the two independent +32-lane logical warps a wave64 wavefront holds into one 64-bit value and corrupt +the format. Fix (in cuda_to_hip.h): redefine __ballot_sync/__shfl*_sync as +width-32 logical-warp ops. __ballot_sync uses a helper that takes the full +__ballot() and shifts down by (__lane_id() & 0x20) so each 32-lane subgroup +extracts ONLY its own 32 lanes as a 32-bit word (wave32: shift always 0; wave64: +lanes 32-63 select the high word). The shuffles use width 32. The producer's +per-subgroup leader (laneId==0 -> physical lanes 0 and 32 on wave64) and +per-subgroup sourceId (threadIdx.x >> 5) are already correct in the surrounding +code, so each logical warp writes its own 32-bit word exactly as CUDA does. The +word width stays 32 (NOT widened to 64); the consumer is unchanged. The launch +geometry already assumes 32-lane logical warps (numwarpsperblock = threads/32). + +Byte-identity evidence: the multi-arch fat-binary check +(-DCMAKE_HIP_ARCHITECTURES="gfx90a;gfx1100") emits BOTH gfx90a (wave64) and +gfx1100 (wave32) code objects from the one wave-agnostic source for +bit_line_maker.cu (llvm-objdump --offloading confirms both). The pulsed-update +tile tests passing on wave64 (test_specific_tiles, InferenceCuda program_weights) +prove the packed stream is read back correctly on wave64. A literal byte-for-byte +diff against a CUDA build cannot be run here (no NVIDIA GPU on the host); the +format-preservation is structural (32-bit-per-32-lane-subgroup, consumer +unchanged). + +## Build recipe (gfx90a) +Direct CMake (bring-up / fastest device-code iteration): +``` +cd projects/aihwkit/src +TORCH_CMAKE=$(python -c "import torch,os;print(os.path.dirname(torch.__file__))")/share/cmake +cmake -S . -B build_hip -GNinja \ + -DUSE_HIP=ON -DUSE_CUDA=OFF -DRPU_CXX_STANDARD=20 \ + -DCMAKE_HIP_ARCHITECTURES=gfx90a \ + -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ \ + -DCMAKE_PREFIX_PATH="$TORCH_CMAKE;/opt/rocm" \ + -DRPU_BLAS=OpenBLAS -DBUILD_TEST=OFF -DRPU_USE_TORCH_BUFFERS=OFF +cmake --build build_hip -j16 +``` +Python package in-place (what the pytest suite imports): +``` +cd projects/aihwkit/src +USE_HIP=ON USE_CUDA=0 python setup.py build_ext -j16 --inplace \ + -DUSE_HIP=ON -DUSE_CUDA=OFF -DRPU_CXX_STANDARD=20 \ + -DCMAKE_HIP_ARCHITECTURES=gfx90a \ + -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ \ + -DCMAKE_PREFIX_PATH="$TORCH_CMAKE;/opt/rocm" \ + -DRPU_BLAS=OpenBLAS -DRPU_USE_TORCH_BUFFERS=OFF -DCMAKE_BUILD_TYPE=Release +``` +This drops rpu_base.cpython-*.so into src/aihwkit/simulator/. Run tests with +PYTHONPATH=src (the editable package is not pip-installed here). + +## Validation (real gfx90a, HIP_VISIBLE_DEVICES=0) +CPU-tile-vs-CUDA-tile comparisons are tolerance-based (assertTensorAlmostEqual +decimal=4); test_specific_tiles uses no fixed cuRAND bitstream, so it is a valid +cross-arch gate. Results: +- tests/test_specific_tiles.py: 18 passed (the bit_line_maker + pulsed-update + warp-size path; 9 cuda-parametrized cases). CRITICAL gate -- PASS. +- tests/test_simulator_tiles.py + tests/test_bindings_tiles.py: 530 passed, + 1 failed, 56 skipped. The single failure was TileTest_Inference (the CPU + tile, not CUDA) test_program_weights -- a stochastic 5%-tolerance + weight-programming convergence test sensitive to global RNG ordering. It + passes 15/15 in isolation; the CUDA variant InferenceCuda passes 8/8. Not a + port regression (pre-existing stochastic-test fragility on the CPU path). +- tests/test_torch_tiles.py + tests/test_inference_tiles.py: 406 passed, 55 + skipped. +- tests/test_layers_linear.py + tests/test_layers_convolution.py: 566 passed, + 216 skipped. +Multi-arch code-object check: fat binary for "gfx90a;gfx1100" builds and emits +both code objects for bit_line_maker.cu (the warp-size TU). + +Pin a GCD with HIP_VISIBLE_DEVICES=0 (host has 4 GCDs shared with another +session). + +## Known follow-ups (not blocking the lead port) +- FP16/bfloat16 tile support (RPU_USE_FP16) is experimental and OFF by default; + left unported (hip_fp16/hip_bfloat16 is a future delta if a half tile test + needs it). +- RPU_USE_TORCH_BUFFERS=ON (RPU_TORCH_CUDA_BUFFERS) was OFF for this validation + to keep bring-up focused; the CUDA-only --expt-relaxed-constexpr/-Xcudafe + flags are now gated to NOT USE_HIP, so turning torch buffers ON under HIP is a + straightforward follow-up to verify. +- BUILD_EXTENSION (aihwkit_extension) GPU ops left on the CUDA path only + (BUILD_EXTENSION OFF by default). + +## Install as a dependency +Not a base library for other MOAT projects; no dependents. + +## Validation 2026-06-04 (validator, linux-gfx90a) + +GPU: AMD Instinct MI250X / MI250 (gfx90a, wave64). ROCm 7.2.1. HIP_VISIBLE_DEVICES=0. +Fork: AMD-Ecosystem/aihwkit moat-port @ 9b4f7be7406cc939109690eb9e05f9ba2dcd3a5c. + +Build: setup.py build_ext --inplace (scikit-build), USE_HIP=ON USE_CUDA=OFF +DRPU_CXX_STANDARD=20 CMAKE_HIP_ARCHITECTURES=gfx90a. Build time ~216 s. +rpu_base.cpython-312-x86_64-linux-gnu.so carries gfx90a code objects (confirmed via +llvm-objdump --offloading: 20+ amdgcn-amd-amdhsa--gfx90a bundles). + +Test results (all with PYTHONPATH=src HIP_VISIBLE_DEVICES=0): +- tests/test_specific_tiles.py: **18/18 PASSED** (CRITICAL -- bit_line_maker + + pulsed-weight-update warp-size path; 9 Cuda-parametrized cases on gfx90a wave64). +- tests/test_simulator_tiles.py + tests/test_bindings_tiles.py: **531 passed, 56 skipped, 0 failed**. + TileTest_Inference::test_program_weights passed this run; TileTest_InferenceCuda::test_program_weights + also passed (1 pass in isolation). The stochastic CPU-tile failure documented by the porter was not + reproduced; it is a pre-existing non-CUDA stochastic convergence test (global RNG ordering sensitive), + not a port regression. +- tests/test_torch_tiles.py + tests/test_inference_tiles.py: **406 passed, 55 skipped, 0 failed**. +- tests/test_layers_linear.py + tests/test_layers_convolution.py: **566 passed, 216 skipped, 0 failed**. + +Total GPU-gated tests passing: 1521 passed, 383 skipped, 0 failed. +Verdict: PASS. Transitioning linux-gfx90a to completed (validated_sha 9b4f7be). + +## Review 2026-06-04 (reviewer, linux-gfx90a) + +Verdict: review-passed. Strategy A is correct for this standalone scikit-build/CMake +backend (Torch headers/linkage only, no torch cpp_extension), so a compat header + +USE_HIP LANGUAGE-HIP path is right, not Strategy B. Diff is minimal (8 files, +302). +Every load-bearing claim was verified by reading the ROCm 7.2.1 headers and the +producer/consumer source, not by trusting the build. + +No blocking findings. Items below are non-defects, recorded for completeness: + +- cuda_to_hip.h defines __shfl_down_sync but the backend never uses it (only + __shfl_up_sync, __shfl_sync, __ballot_sync are used; all in bit_line_maker.cu). + Harmless dead coverage; defensible as intrinsic-set completeness. No action. +- rpu_linearstep_device.h:72 ctor rename (SoftBoundsRPUDeviceMetaParameter() -> + SoftBoundsRPUDeviceMetaParameter()) is the one UNCONDITIONAL edit to a CPU/CUDA-shared + header (not USE_HIP-guarded). It is a strict C++ conformance fix (a template-id is + ill-formed as an injected-class-name constructor declarator) that is behavior-identical + on the CUDA/CPU path; acceptable and minimal. Latent upstream bug exposed only because + the newer Torch forces C++20. + +Verified directly: + +(a) bit_line_maker byte-identical-format correctness -- SOUND. The compat-header +__rpu_logical_warp32_ballot takes the full-wavefront __ballot() (HIP: returns +unsigned long long, 64 bits on wave64 / low 32 on wave32) and shifts by +(__lane_id() & 0x20): wave32 lane_id<=31 so shift always 0; wave64 lanes 0-31 -> 0, +lanes 32-63 -> 32. Cast to unsigned int then yields exactly the calling lane's own +32-lane subgroup as a 32-bit word == what CUDA __ballot_sync(0xFFFFFFFF,...) gives a +32-lane warp. The redefined width-32 __shfl/__shfl_up confirmed (amd_warp_functions.h) +to compute the source lane WITHIN the caller's 32-lane subgroup ((self & ~31) term), +so __shfl_sync(...,0) broadcasts each subgroup's own lane-0 and __shfl_up(...,32) stays +in-subgroup. Producer geometry already per-32-subgroup: laneId = threadIdx.x & 0x1f +(two leaders 0/32 on wave64), sourceId = (tid+i_stride)>>5 and +blockIdx.x*(blockDim.x>>5)+(threadIdx.x>>5) (distinct source per subgroup), word width +stays 32, consumer pwu_kernel.h (>>5/&0x1f/__popc on a 32-bit word) unchanged. No mask +widened to 64. The first kernel's __shfl_up_sync(...,kidthread) is also safe: nKthreads +divides 32 (constraint 32%Kplus1==0) so nKthreads-groups never straddle a 32-lane +boundary. On-device bitstream is wave-width-invariant by construction; the fat-binary +gfx90a;gfx1100 check + wave64 test pass are corroborating, not the proof. + +(b) cuRAND distinct-type mapping -- CORRECT. The backend uses CudaArray +(bit_line_maker.cu:90,131; test_helper.cu:42), CudaArray at call sites +(cuda_util.h:414-415,504; cuda_util.cu:179), AND an explicit instantiation + method +specializations of CudaArray (cuda_util.cu:1102/1165/1168/1171/1407). +In CUDA all three are the one XORWOW type. Aliasing all three to hiprandState_t makes +the explicit instantiation match every call-site type; mapping curandStateXORWOW to +hiprandStateXORWOW_t separately would leave CudaArray's specializations +undefined -> the exact ImportError the porter hit. No layout conflation: hipRAND's +default state IS the XORWOW state, and producer+consumer share the single aliased type +within this build, so there is no two-different-layouts hazard. + +Also confirmed: cuBLAS status switch has no duplicate-case collision +(LICENSE_ERROR->UNKNOWN=11 is distinct from the 7 other mapped values, hipblas-common.h); +device-side cuBLAS (kernelCublasCreateDevice/getDeviceHandle) is fully inside +#ifdef RPU_WITH_CUBLAS_DEVICE which is never defined in any build path; host DEVICE +pointer-mode (nrm2) is on the host handle and hipBLAS-supported; BlockScan TempStorage +used exactly once (no wave64 reuse race); pinned mem cudaMallocHost/cudaFreeHost aliased; +ATen stream swap preserves at::cuda::getCurrentCUDAStream() at all 9 call sites; CMake +CMAKE_HIP_ARCHITECTURES default gfx90a only-when-unset (no hardcode), USE_HIP default OFF, +IPO off for HIP, CUDA path byte-identical; commit message [ROCm] 55 chars, Claude named, +Test Plan present, no noreply trailer, no MOAT jargon, ASCII-clean, no AMD-internal refs. +The 1 CPU-tile InferenceCuda program_weights failure is a pre-existing stochastic +convergence test (passes in isolation, CUDA variant 8/8), not a port regression. + +## Validation 2026-06-04 (linux-gfx1100, RDNA3 native wave32) + +GPU: AMD Radeon Pro W7800 48GB (gfx1100, wave32 native). ROCm 7.2.1. HIP_VISIBLE_DEVICES=1. +Fork: AMD-Ecosystem/aihwkit moat-port @ 9b4f7be7406cc939109690eb9e05f9ba2dcd3a5c. + +Build: setup.py build_ext --inplace (scikit-build), USE_HIP=ON USE_CUDA=OFF +RPU_CXX_STANDARD=20 CMAKE_HIP_ARCHITECTURES=gfx1100. Build time ~2 min. +rpu_base.cpython-312-x86_64-linux-gnu.so carries gfx1100 code objects (confirmed via +llvm-objdump --offloading: 28+ amdgcn-amd-amdhsa--gfx1100 bundles). + +Wave32 confirmation: gfx1100 is a native wave32 arch -- each physical wavefront is +exactly one 32-lane logical warp, so the __ballot_sync shift-by-(__lane_id()&0x20) in +cuda_to_hip.h is always a shift-by-0. The bit_line_maker packed format is word-identical +to CUDA wave32 by construction, confirmed by test_specific_tiles.py passing below. + +Test results (all with PYTHONPATH=src HIP_VISIBLE_DEVICES=1): +- tests/test_specific_tiles.py: **18/18 PASSED** (CRITICAL -- bit_line_maker + + pulsed-weight-update warp-size path on gfx1100 native wave32; 9 Cuda-parametrized cases). +- tests/test_simulator_tiles.py + tests/test_bindings_tiles.py: **part of 1521 passed, 0 failed**. +- tests/test_torch_tiles.py + tests/test_inference_tiles.py: **part of 1521 passed, 0 failed**. +- tests/test_layers_linear.py + tests/test_layers_convolution.py: **part of 1521 passed, 0 failed**. + +Total across all suites: 1521 passed, 327 skipped, 0 failed (1848 collected). +(Lead gfx90a: 1521 passed, 383 skipped -- skip count varies by arch as some tests query +wave size or arch capabilities; pass count and zero failures match exactly.) + +Verdict: PASS. All suites pass on gfx1100 native wave32 with no regressions. +Transitioning linux-gfx1100 to completed (validated_sha 9b4f7be). + +## Validation 2026-06-07 (windows-gfx1201, RDNA4 native wave32) + +GPU: AMD Radeon RX 9070 XT (gfx1201, wave32 native). ROCm 7.14.0a20260604 (TheRock venv). +HIP_VISIBLE_DEVICES=0 (gfx1201 is device index 0 on this host; gfx1101 is index 1). +Fork: AMD-Ecosystem/aihwkit moat-port @ 50360f7a07281ce9cf0272b2e078e6821d5f9a07. + +Three Windows-specific fixes were needed on top of the Linux port (9b4f7be), committed as +a second commit on moat-port: + +1. CMakeLists.txt: guard MSVC-only /O2 flag with if(MSVC) instead of if(WIN32). amdclang++ + on Windows is !MSVC and rejects /O2. + +2. rpu_base_tiles_cuda.cpp: replace `using c10::cuda::getCurrentCUDAStream` with an inline + wrapper calling c10::hip::getCurrentHIPStream. Torch 2.9/ROCm 7.14 removed the + getCurrentCUDAStream alias from c10::cuda namespace in HIPStream.h. + +3. pwu_kernel_parameter_base.h: change PulsedUpdateMetaParameter forward declaration from + `class` to `struct` (matching rpu_pulsed_meta_parameter.h and rpucuda_pulsed_device.h). + On Linux/ELF both mangle identically; on Windows/MSVC-ABI (used by amdclang++), struct + and class produce different name mangling (AEBU vs AEBV), causing virtual method + implementations in RPU_GPU.lib to not resolve at link time despite being present. + +Build: CMake + Ninja, amdclang++/lld-link. All-clang (MSVC host unsupported with HIP). +Build flags require post-cmake manual steps in build.ninja: +- Remove -fuse-ld=lld-link injected by CMake 4.x into HIP device-link steps (keep it for + the final host+device combined link of rpu_base.pyd) +- Replace /WHOLEARCHIVE via -Xlinker for RPU_GPU.lib in LINK_LIBRARIES +- Add c10_hip.lib, torch_hip.lib, and a generated libomp140.lib (from libomp140.x86_64.dll; + import lib DLL name must include .dll extension or Windows loader fails to find the file) +- Add mypy's stubgen.exe path explicitly in POST_BUILD + +Runtime: rpu_base.cp312-win_amd64.pyd requires ROCm DLLs. For Python import to work, +torch must be imported first (which preloads ROCm DLLs via rocm_sdk.preload_libraries), +AND several DLLs must be present in the package directory next to the pyd: +amd_comgr.dll, amdhip64_7.dll, c10_hip.dll, hipblas.dll, hiprand.dll, rocm-openblas.dll, +rocm_kpack.dll, rocrand.dll, rocblas.dll, rocsolver.dll, shm.dll. +The simulator/__init__.py already imports torch before rpu_base, satisfying the preload +requirement for the test suite. + +gfx1201 wave32 note: RDNA4 is native wave32, so the __ballot_sync shift-by-(__lane_id()&0x20) +is always a shift-by-0 (identical to CUDA wave32 behavior). The bit_line_maker packed format +is byte-identical to gfx1100 by construction. + +Test results (PYTHONPATH=src, HIP_VISIBLE_DEVICES=0): +- tests/test_specific_tiles.py: **18/18 PASSED** (CRITICAL -- bit_line_maker + pulsed-weight-update + warp-size path; 9 Cuda-parametrized cases on gfx1201 native wave32). +- tests/test_simulator_tiles.py + tests/test_bindings_tiles.py: **530 passed, 56 skipped, 1 failed**. + The 1 failure is TileForwardBackwardTest_Inference::test_set_forward_out_noise_std -- a stochastic + test that passes in isolation (1/1) and is a pre-existing non-CUDA flakiness, not a port regression. +- tests/test_torch_tiles.py + tests/test_inference_tiles.py: **406 passed, 55 skipped, 0 failed**. +- tests/test_layers_linear.py + tests/test_layers_convolution.py: **566 passed, 216 skipped, 0 failed**. + +Total GPU-gated tests: 1520 passed, 327 skipped, 1 failed (stochastic non-GPU pre-existing). +Verdict: PASS. Transitioning windows-gfx1201 to completed (validated_sha 50360f7). + +## Revalidation 2026-06-08 (linux-gfx90a) + +State: revalidate (validated_sha 9b4f7be7 -> head_sha 50360f7a -> new head d6d4561). + +Delta classification: The 9b4f7be7..50360f7a delta (Windows Clang fix) was classified `mixed` +by moatlib.classify. Binary equivalence check revealed a Linux compilation failure: the Windows +fix used `c10::hip::getCurrentHIPStream(device_index)` unconditionally, but that symbol is +guarded by `#ifdef USE_ROCM` in HIPStream.h and aihwkit does not define USE_ROCM -- so the +build failed on Linux ROCm 7.2.1 with "no member named 'getCurrentHIPStream' in namespace +'c10::hip'". This is the "Windows commit breaks Linux ROCm compilation" trap from CLAUDE.md. + +Fix applied (new commit d6d4561 on moat-port on top of 50360f7a): gate the function body on +HIP_VERSION_MINOR. ROCm <= 7.2 uses `c10::cuda::getCurrentCUDAStream` (still present in +torch 2.13/ROCm 7.2.1 HIPStream.h). ROCm >= 7.14 uses `c10::hip::getCurrentHIPStream` +(supported by the TheRock torch headers on Windows which expose it without the USE_ROCM guard, +as proven by the windows-gfx1201 validation at 50360f7a). This is a host-side change; device +code objects are unchanged on gfx90a. + +Build: setup.py build_ext --inplace, USE_HIP=ON USE_CUDA=OFF RPU_CXX_STANDARD=20 +CMAKE_HIP_ARCHITECTURES=gfx90a. 29 gfx90a code objects confirmed. Build time ~216 s. + +Full GPU test suite (HIP_VISIBLE_DEVICES=3, gfx90a wave64, ROCm 7.2.1): +- tests/test_specific_tiles.py: **18/18 PASSED** (CRITICAL -- bit_line_maker + pulsed-weight-update warp-size path on gfx90a wave64). +- tests/test_simulator_tiles.py + tests/test_bindings_tiles.py: **531 passed, 56 skipped, 0 failed**. +- tests/test_torch_tiles.py + tests/test_inference_tiles.py: **406 passed, 55 skipped, 0 failed**. +- tests/test_layers_linear.py + tests/test_layers_convolution.py: **566 passed, 216 skipped, 0 failed**. + +Total: 1521 passed, 383 skipped, 0 failed. Identical pass count to original validation. +Verdict: PASS. Transitioning linux-gfx90a to completed (validated_sha d6d4561). + +## Re-key 2026-06-08 (linux-gfx90a, delta-ported at b346589) + +Re-keyed the at::cuda::getCurrentCUDAStream shim in rpu_base_tiles_cuda.cpp from a +ROCm-version proxy to the correct axis: the torch hipify generation. This TU is never +hipified by PyTorch source-hipify (aihwkit is scikit-build/CMake with its own USE_HIP), +so it must hand-pick the c10 stream symbol. The real selector is the hipify version: +- hipify v2 (masquerading, this Linux env: torch 2.13, hipify 2.0.0): c10::cuda:: + getCurrentCUDAStream is the public API; c10::hip::getCurrentHIPStream is #ifdef + USE_ROCM and this build defines USE_HIP not USE_ROCM, so only c10::cuda works. +- hipify v1 (rename, Windows TheRock torch 2.9.1): c10::cuda::getCurrentCUDAStream is + removed, so c10::hip::getCurrentHIPStream is required. + +The old d6d4561 gate (HIP_VERSION >= 7.14 -> c10::hip) worked only because in our fleet +ROCm version anti-correlates with the hipify generation (ROCm 7.2 + new torch/v2 on +Linux; ROCm 7.14 + old torch/v1 on Windows). Now CMake (cmake/dependencies_hip.cmake) +probes the build's own torch via RPU_PYTHON_EXECUTABLE +(`python -c "from torch.utils.hipify import __version__"`), defines TORCH_HIPIFY_V2 on +the rpu_base target when >= 2.0.0, and the cpp keys on `#if defined(TORCH_HIPIFY_V2)`. +The sense FLIPS vs the old gate (v2 now takes the c10::cuda branch). Detection failure +leaves it undefined = v1 default. Removed the now-unused include. + +CMake on this Linux env logged `-- torch hipify version: 2.0.0` and put -DTORCH_HIPIFY_V2 +on the rpu_base_tiles_cuda.cpp.o compile line (confirmed in build.ninja DEFINES). + +Behavior-preservation proof (the gate for this delta, no full GPU re-run required): +codeobj_diff between a pristine d6d4561 build and the b346589 build of +rpu_base.cpython-312-x86_64-linux-gnu.so (both 337135072 bytes, 29 gfx90a code objects): +`verdict=identical -- exported symbols + device ISA identical (53320 exports)`. On Linux +v2 the re-key selects the SAME c10::cuda::getCurrentCUDAStream the old gate's #else branch +already selected (ROCm 7.2 < 7.14), so the device ISA is unchanged by construction; the +diff confirms it. + +GPU smoke on MI250X (gfx90a wave64, ROCm 7.2.1, HIP_VISIBLE_DEVICES=0): rpu_base loads, +rpu_base.cuda.is_compiled()=True; tests/test_specific_tiles.py 18/18 PASS (the critical +bit_line_maker pulsed-update warp-size path); tests/test_simulator_tiles.py -k Cuda 264 +passed 0 failed; tests/test_inference_tiles.py -k Cuda 32 passed 0 failed. + +Commit b346589 on top of d6d4561 (new commit, not amend). This makes aihwkit the +reference example for the PORTING_GUIDE "key on hipify generation, not ROCm version" +entry. Linux-gfx90a -> delta-ported. + +## Revalidation 2026-06-08 (linux-gfx1100) + +State: revalidate (validated_sha 9b4f7be -> head_sha d6d4561). +GPU: AMD Radeon Pro W7800 48GB (gfx1100, wave32 native). ROCm 7.2.1. HIP_VISIBLE_DEVICES=2. + +Delta 9b4f7be..d6d4561 (2 commits: Windows Clang build fixes + getCurrentCUDAStream ROCm compat): +- CMakeLists.txt: guard /O2 with if(MSVC) instead of if(WIN32). On Linux MSVC=false; the else() branch is unchanged. +- rpu_base_tiles_cuda.cpp: HIP_VERSION_MINOR-gated getCurrentHIPStream/getCurrentCUDAStream. On Linux ROCm 7.2.1: HIP_VERSION_MAJOR=7, HIP_VERSION_MINOR=2 < 14, so c10::cuda::getCurrentCUDAStream path taken -- same as before the change. +- pwu_kernel_parameter_base.h: class->struct for PulsedUpdateMetaParameter forward declaration. On Linux/ELF both mangle identically; no ABI impact. + +No .cu device files changed. Binary equivalence check: + +Build commands (HIP_VISIBLE_DEVICES=2, CMAKE_HIP_ARCHITECTURES=gfx1100): +``` +cmake -S . -B build_hip -GNinja -DUSE_HIP=ON -DUSE_CUDA=OFF -DRPU_CXX_STANDARD=20 \ + -DCMAKE_HIP_ARCHITECTURES=gfx1100 -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ \ + -DCMAKE_PREFIX_PATH="${TORCH_CMAKE};/opt/rocm" -DRPU_BLAS=OpenBLAS -DBUILD_TEST=OFF \ + -DRPU_USE_TORCH_BUFFERS=OFF -DCMAKE_BUILD_TYPE=Release +cmake --build build_hip -j16 +``` + +codeobj_diff result: + src/aihwkit/simulator/rpu_base.cpython-312-x86_64-linux-gnu.so: identical (exported symbols + device ISA identical (11257 exports)) + verdict=identical + +Carry-forward applied: moatlib carry-forward binary-equiv, validated_sha -> d6d4561. +No GPU re-run required. linux-gfx1100 -> completed. + +## Review 2026-06-08 (reviewer, linux-gfx90a, re-key delta d6d4561..b346589) + +Verdict: review-passed. Reviewed the hipify-version re-key of the +at::cuda::getCurrentCUDAStream shim (3 files, +35/-7). No defects found; all +load-bearing claims verified against the build env's torch headers, not trusted. + +No blocking findings. Verified directly: + +- Probe (cmake/dependencies_hip.cmake:34-46): uses RPU_PYTHON_EXECUTABLE, which + dependencies.cmake:84 sets from Python3_EXECUTABLE and is included at + CMakeLists.txt:53, BEFORE dependencies_hip.cmake at :55 -- so the build's own + Python runs the probe. Ran the exact probe command in the env: exit 0, stdout + "2.0.0" clean; the NumPy import warning goes to stderr (ERROR_QUIET-suppressed, + not captured into the version var). The RESULT==0 AND NOT STREQUAL "" guard is + robust; detection failure leaves RPU_TORCH_HIPIFY_V2 unset (safe v1 default). +- Variable scope: set(RPU_TORCH_HIPIFY_V2 ON) is in the top-level include scope + (CMakeLists.txt:55); src/aihwkit/simulator is add_subdirectory at :156 (> :55), + so the child scope inherits it. target_compile_definitions(rpu_base PRIVATE + TORCH_HIPIFY_V2) (src/aihwkit/simulator/CMakeLists.txt:32-34) applies ONLY to + the pybind module target (the rpu_base_src/*.cpp binding TUs incl. + rpu_base_tiles_cuda.cpp), NOT via add_compile_definitions -- so the .cu device + TUs in RPU_GPU are untouched (consistent with codeobj_diff = identical). +- #if defined(TORCH_HIPIFY_V2) SENSE is correct (and is the flip of the old + HIP_VERSION gate). Confirmed in this env's c10/hip/HIPStream.h: getCurrentCUDAStream + is unconditional in c10::cuda (:235); getCurrentHIPStream is #ifdef USE_ROCM (:256) + which this build does not define. So v2 -> c10::cuda::getCurrentCUDAStream (the + only available symbol), v1 -> c10::hip::getCurrentHIPStream. Env hipify = 2.0.0, + so VERSION_GREATER_EQUAL "2.0.0" true -> define set -> c10::cuda branch taken. +- Orphan cleanup complete: removed; no remaining HIP_VERSION* + macro usage in the TU (grep clean). +- CUDA/CPU paths unaffected: the shim is double-guarded by #ifdef RPU_USE_CUDA (:7) + and #if defined(USE_HIP) (:17). A CUDA build (RPU_USE_CUDA set, USE_HIP unset) + takes the #else at :39 (), unchanged; a pure-CPU build + compiles none of it. TORCH_HIPIFY_V2 is only ever defined on the HIP pybind build. +- Commit hygiene: title 72 chars exactly, [ROCm] prefix; no MOAT jargon in message + or comments; ASCII clean, no em-dash; no noreply trailer, no ghstack; Claude + disclosed; Test Plan present; no AMD-internal account refs. + +Carry-forward note: this delta changes only a host-side pybind TU and no .cu +device code; on Linux v2 it selects the same c10::cuda::getCurrentCUDAStream the +old #else already resolved, so gfx90a / gfx1100 / gfx1201 all carry forward via +.co byte-identity (gfx90a codeobj_diff already identical; gfx1100/gfx1201 same by +construction -- device ISA unchanged). Hands to validator for carry-forward +confirmation; no functional GPU re-run is required for this behavior-preserving +host-only change. + +## Validation 2026-06-08 (linux-gfx90a, re-key carry-forward) + +GPU: AMD Instinct MI250X / MI250 (gfx90a, wave64). ROCm 7.2.1. HIP_VISIBLE_DEVICES=0. +Delta: d6d4561 -> b346589 (hipify-version re-key, host-only pybind TU). + +Build: CMake + Ninja at BOTH shas (fresh build dirs, not recycled). At b346589 cmake configured +with `-- torch hipify version: 2.0.0` and set -DTORCH_HIPIFY_V2 in build.ninja DEFINES for +rpu_base_tiles_cuda.cpp.o (confirmed). At d6d4561 (worktree, HIP_VERSION_MINOR gate, ROCm 7.2 < +7.14 takes the c10::cuda::getCurrentCUDAStream path). Both builds produced +rpu_base.cpython-312-x86_64-linux-gnu.so. + +codeobj_diff: verdict=identical -- exported symbols + device ISA identical (53320 exports). +Matches the porter's prior run exactly. On Linux v2 both gates resolve the same +c10::cuda::getCurrentCUDAStream symbol, so device ISA is unchanged by construction; diff confirms. + +GPU smoke (gfx90a wave64, HIP_VISIBLE_DEVICES=0): tests/test_specific_tiles.py 18/18 PASSED +(bit_line_maker + pulsed-weight-update warp-size path, 9 Cuda-parametrized cases). + +Verdict: carry-forward confirmed. linux-gfx90a -> completed (validated_sha b346589). + +## Revalidation 2026-06-08 (linux-gfx1100, delta d6d4561->b346589) + +State: revalidate (validated_sha d6d4561 -> head_sha b346589). +GPU: AMD Radeon Pro W7800 48GB (gfx1100, wave32 native). ROCm 7.2.1. HIP_VISIBLE_DEVICES=1. + +Delta d6d4561..b346589 (hipify-version re-key): changes only 3 host-side files -- +cmake/dependencies_hip.cmake, src/aihwkit/simulator/CMakeLists.txt, and +src/aihwkit/simulator/rpu_base_src/rpu_base_tiles_cuda.cpp. No .cu device code changed. +On this Linux env, hipify version is 2.0.0 (v2), so TORCH_HIPIFY_V2 is defined and the +re-keyed gate selects the same c10::cuda::getCurrentCUDAStream the old HIP_VERSION_MINOR<14 +branch already selected (ROCm 7.2 < 7.14). Device ISA is unchanged by construction. + +Binary equivalence check: +- Build at d6d4561: cmake -S . -B build_gfx1100 -GNinja -DUSE_HIP=ON -DUSE_CUDA=OFF + -DRPU_CXX_STANDARD=20 -DCMAKE_HIP_ARCHITECTURES=gfx1100 -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ + -DCMAKE_PREFIX_PATH="${TORCH_CMAKE};/opt/rocm" -DRPU_BLAS=OpenBLAS -DBUILD_TEST=OFF + -DRPU_USE_TORCH_BUFFERS=OFF -DCMAKE_BUILD_TYPE=Release; cmake --build build_gfx1100 -j16 +- Build at b346589: same flags. Both built successfully (warnings only, no errors). +- codeobj_diff result: + src/aihwkit/simulator/rpu_base.cpython-312-x86_64-linux-gnu.so: identical + (exported symbols + device ISA identical (11257 exports)) + verdict=identical + +Carry-forward applied. linux-gfx1100 -> completed (validated_sha b346589). No GPU re-run required. + +## Validation 2026-06-08 (windows-gfx1201, carry-forward) + +Delta: 50360f7a -> b346589 (3 host-side files only). + +Method: source-class inspection. + +Changed files: +- cmake/dependencies_hip.cmake: adds cmake `execute_process` to detect torch hipify version and set RPU_TORCH_HIPIFY_V2 if version >= 2.0.0 +- src/aihwkit/simulator/CMakeLists.txt: passes TORCH_HIPIFY_V2 compile definition if RPU_TORCH_HIPIFY_V2 is set +- src/aihwkit/simulator/rpu_base_src/rpu_base_tiles_cuda.cpp: adds `#if defined(TORCH_HIPIFY_V2)` branch for getCurrentCUDAStream + +On this Windows host (TheRock 7.14 PyTorch), torch hipify version is 1.0.0 (v1). Therefore TORCH_HIPIFY_V2 is NOT defined at configure time, and the #if branch is dead code. The compiled output at b346589 uses the same `c10::hip::getCurrentHIPStream(device_index)` path as 50360f7a. No GPU device code (.cu files) was changed. Device ISA for gfx1201 is identical by construction. + +VERDICT: CARRY-FORWARD (source-class). State -> completed (validated_sha = b346589). No GPU re-run required. + +## PR-prep 2026-06-08 (lead) -- docs + squash; carry-forward, no GPU re-run + +Pre-PR cleanup. The port was already clean (no MOAT jargon in code or commit +messages; CMake arch handling already correct -- CMAKE_HIP_ARCHITECTURES, +default gfx90a only when unset). Only substantive prep was documentation. + +- docs/source/developer_install.rst: documented the ROCm/HIP build alongside the + CUDA build in house style -- a GPU build example, gfx-arch identification via + rocminfo (parallel to the nvidia-smi CUDA_ARCH snippet), the USE_HIP row in the + compilation-flags table, and a USE_HIP build example (setup.py + cmake, noted + mutually exclusive with USE_CUDA). No ROCm wheel/conda/docker exists, so the + prebuilt-install docs (install.rst, README install section) were left alone -- + ROCm is from-source only. +- README.md: feature line "(CUDA-capable)" -> "(CUDA- and ROCm/HIP-capable)". + +Squashed the 4 port commits + the doc commit into ONE commit on the upstream +base (clean PR diff): 0b49124 [ROCm] Add AMD GPU (HIP) support to the RPUCuda +backend, parent 4d73afd. 11 files, +367/-31. The squashed message consolidates +the four commit messages and drops the stream-shim iteration narrative (it +describes the final TORCH_HIPIFY_V2-keyed shim, not the HIP_VERSION steps). + +advance_head classified the doc delta doc-only; squash-carry-forward carried +linux-gfx90a, linux-gfx1100, windows-gfx1201 to 0b49124. gfx1151 kept blocked +(retired); gfx1101 stays port-ready (redundant Windows tier, gfx1201 satisfies +it). pr-ready=True. + +NEXT: upstream-PR gate (lead-only, jeff approval). No existing jeffdaily PR on +IBM/aihwkit. PR body scopes out gfx1151; claims Linux gfx90a + gfx1100 and +Windows gfx1201. + +## Validation 2026-06-19 (windows-gfx1101, RDNA3 native wave32) + +GPU: AMD Radeon PRO V710 (gfx1101, wave32 native). ROCm 7.14.0a20260604 (TheRock venv). +HIP_VISIBLE_DEVICES=1 (gfx1101 is device index 1 on this host; gfx1201 is index 0). +Fork: AMD-Ecosystem/aihwkit moat-port @ 42e046594e91c1283ddd647d2165ba8b662b3999 (squashed PR commit). + +Build: CMake + Ninja, amdclang++/lld-link, CMAKE_HIP_ARCHITECTURES=gfx1101. +Same three Windows-specific fixes as gfx1201 are in the squashed commit: if(MSVC) /O2 guard, +TORCH_HIPIFY_V2 detection for stream symbol, PulsedUpdateMetaParameter class->struct. + +CMake configure flags: + -DUSE_HIP=ON -DUSE_CUDA=OFF -DRPU_CXX_STANDARD=20 -DCMAKE_HIP_ARCHITECTURES=gfx1101 + -DCMAKE_HIP_COMPILER=amdclang++ -DCMAKE_C_COMPILER=amdclang -DCMAKE_CXX_COMPILER=amdclang++ + -DCMAKE_PREFIX_PATH="$TORCH_CMAKE;$ROCM_DEVEL;$ROCM_DEVEL/lib/cmake" + -Dhip_DIR/hipblas_DIR/hipblas-common_DIR/hiprand_DIR/rocrand_DIR/hipcub_DIR/rocprim_DIR set explicitly + -DOpenBLAS_INCLUDE_DIR=$ROCM_DEVEL/lib/host-math/include/openblas + -DOpenBLAS_LIB=$ROCM_DEVEL/lib/host-math/lib/rocm-openblas.lib + -DRPU_USE_TORCH_BUFFERS=OFF -DBUILD_TEST=OFF + +build.ninja manual patches (same class as gfx1201): +- Add c10_hip.lib, torch_hip.lib (HIPStream symbols) to LINK_LIBRARIES +- Add libomp140.lib (generated from libomp140.x86_64.dll via llvm-dlltool) for OpenMP symbols +- Add -Xlinker /WHOLEARCHIVE:RPU_GPU.lib for complete template instantiation +- Fix POST_BUILD: replace bare 'stubgen' with full path to venv/Scripts/stubgen.exe + +Runtime environment: same DLLs as gfx1201 (already in simulator/), pyd replaced with gfx1101 build. +ROCBLAS_TENSILE_LIBPATH=_rocm_sdk_libraries/bin/rocblas/library; ROCBLAS_USE_HIPBLASLT=0. + +gfx1101 wave32 note: RDNA3 is native wave32, so the __ballot_sync shift-by-(__lane_id()&0x20) +is always a shift-by-0 (identical to gfx1100 and gfx1201 wave32 behavior). The bit_line_maker +packed format is byte-identical to gfx1100/gfx1201 by construction. + +Test results (PYTHONPATH=src, HIP_VISIBLE_DEVICES=1): +- tests/test_specific_tiles.py: **18/18 PASSED** (CRITICAL -- bit_line_maker + pulsed-weight-update + warp-size path; 9 Cuda-parametrized cases on gfx1101 native wave32). +- tests/test_simulator_tiles.py + tests/test_bindings_tiles.py: **531 passed, 56 skipped, 0 failed**. +- tests/test_torch_tiles.py + tests/test_inference_tiles.py: **406 passed, 55 skipped, 0 failed**. +- tests/test_layers_linear.py + tests/test_layers_convolution.py: **566 passed, 216 skipped, 0 failed**. + +Total GPU-gated tests: 1521 passed, 383 skipped, 0 failed. Identical pass count to gfx90a. +Verdict: PASS. Transitioning windows-gfx1101 to completed (validated_sha 42e0465). diff --git a/projects/aihwkit/plan.md b/projects/aihwkit/plan.md new file mode 100644 index 00000000..dfeef0aa --- /dev/null +++ b/projects/aihwkit/plan.md @@ -0,0 +1,142 @@ +# Port plan: aihwkit (IBM Analog Hardware Acceleration Kit) + +## Project +- Name: aihwkit +- Upstream: https://github.com/IBM/aihwkit +- Default branch: main +- Lead platform: linux-gfx90a (CDNA2, wave64) +- Domain: PyTorch library simulating analog in-memory computing (resistive crossbar / RPU tiles). C++/CUDA backend "RPUCuda" exposed to Python via a pybind11 module (`rpu_base`). + +## Existing AMD support +NONE. From-scratch HIP port. +- Web search ("aihwkit ROCm", "aihwkit AMD GPU HIP", "IBM aihwkit RPUCuda AMD MI300 gfx9"): no ROCm/HIP port, no separately-named AMD project (no ROCm-DS analogue). IBM ported Qiskit to ROCm, not aihwkit. setup.py classifier is `Environment :: GPU :: NVIDIA CUDA`; only a CUDA.Dockerfile ships. +- `gh api repos/IBM/aihwkit/forks`: no fork under ROCm/AMD/GPUOpen orgs and none with rocm/hip/amd in the name. The two name-substring hits (ZhipingWoods/aihwkit, PJLAB-CHIP/AnalogAIold) are false positives -- plain mirrors, no AMD work. +- Upstream branches: no rocm/hip/amd branch. No AMD issues/PRs found. +- Decision: clean from-scratch HIP port targeting ROCm. No authoritative or community base to adopt. + +## Disposition +TRACTABLE -- proceed with a Strategy-A mechanical HIP port. Effort class: MEDIUM (37 `.cu` files, ~10k LOC of device code, but a clean and small external-API surface: cuBLAS + cuRAND + CUB only, NO textures/surfaces, NO CUTLASS, NO managed memory). The single real semantic risk is one warp-size-coupled bit-packed data format in `bit_line_maker.cu`. Dispatch a porter. + +NOT a CK reimplementation: the only CUTLASS references in the tree are three CODE COMMENTS in `forward_backward_pass.cu` ("eventually we might want to use cutlass", "CUTLASS will help") -- aspirational, no CUTLASS/CuTe code, no wgmma/mma_sync, no tensor-core path. The matvec is plain cuBLAS GEMM/GEMV plus hand-written analog-noise-injection kernels. A mechanical HIP port is correct and sufficient; no AMD-native GEMM rewrite is needed for correctness. + +## Build classification: cmake (Strategy A) -- evidence +This is a PyTorch library but it does NOT use `torch.utils.cpp_extension`. The CUDA is compiled by a standalone CMake/nvcc path; Torch is found only for headers/linkage. +- `setup.py:12` -- `from skbuild import setup` (scikit-build invokes CMake; no `CUDAExtension`/`BuildExtension`). +- `CMakeLists.txt:9-11` -- `project(aihwkit C CXX)`, `option(USE_CUDA ...)`; `cmake/dependencies_cuda.cmake:9` -- `enable_language(CUDA)`, `:12` -- `find_package(CUDAToolkit)`. Native CMake CUDA language, not torch's extension machinery. +- `CMakeLists.txt:108,168,228` -- `set_property(TARGET ... PROPERTY CUDA_ARCHITECTURES ${RPU_CUDA_ARCHITECTURES})` with default `"75;80;89"` (`:26`). Arch is a CMake property, the Strategy-A pattern. +- `CMakeLists.txt:102` -- `target_link_libraries(RPU_GPU RPU_CPU cublas curand ...)`: links CUDA libs by bare name directly. +- Torch is linked only as `torch_cuda`/`c10_cuda`/`torch_python` (CMakeLists.txt:104,159-161; src/aihwkit/simulator/CMakeLists.txt:9,13). No `find_package(Torch)` that drives extension compilation; torch supplies tensor-buffer headers (RPU_TORCH_CUDA_BUFFERS) and the pybind link. +- `src/aihwkit/simulator/CMakeLists.txt:8` -- `pybind11_add_module(rpu_base MODULE ...)` linking the CMake-built `RPU_GPU` static lib (`:13`). The Python extension is a CMake pybind target, not a torch CUDAExtension. + +Conclusion: standalone CMake + nvcc, Torch-for-linkage-only. Use Strategy A (USE_HIP CMake option + a single compat header), NOT Strategy B (no torch hipify involved). Set ext_type = `cmake`. + +## Port strategy: A (compat header + enable_language(HIP)) +Mirror the compat-header approach and the existing `USE_CUDA` switch with a parallel `USE_HIP`. +1. Add one compat header (e.g. `src/rpucuda/cuda/cuda_to_hip.h`) included once from `cuda_util.h` (which already centralizes `cublas_v2.h`, `cuda_runtime.h`, `curand.h`, `curand_kernel.h` at lines 18-23). On HIP it includes ``, ``, ``, `` and aliases the cuda* spellings the project uses to hip*; on CUDA it is a no-op. Keep `/` included BEFORE the HIP runtime (gpuRIR lesson: host memcpy/memset can resolve to HIP __device__ overloads otherwise). +2. In CMake, add `option(USE_HIP ...)`; when set, `enable_language(HIP)`, mark the glob'd `${RPU_GPU_SRCS}` `LANGUAGE HIP`, set `HIP_ARCHITECTURES` from `${CMAKE_HIP_ARCHITECTURES}` (default gfx90a only when unset -- never hardcode the literal, so followers need no source edit), and link `hipblas hiprand` instead of `cublas curand`. Keep `RPU_USE_CUDA` defined on the HIP build (the `#ifdef RPU_USE_CUDA` blocks in rpu_base.cpp gate the CudaAnalogTile pybind exposure and the whole device tile code; the compat header makes "CUDA-spelled" code compile under HIP). +3. Disable IPO/LTO for the HIP build of the pybind module (gpuRIR lesson: HIP link does not finalize LTO -> empty PyInit_ -> ImportError). Check whether scikit-build/CMake enables INTERPROCEDURAL_OPTIMIZATION; force it OFF for HIP. +4. Guard the few genuinely divergent spots with `#if defined(USE_HIP)`; keep rare. + +Because RNG state and BLM use CUDA-library types in headers shared with host code (`curandState_t` in cuda_util.h, vector_device.h, transfer_device.h, etc.), the compat header must alias those device-RNG and cuBLAS type/enum names project-wide -- the aliasing surface is larger than colmap's but still mechanical. + +## CUDA surface inventory +- Kernels: ~hundreds of `__global__`/`__device__` across 37 `.cu` (RPU tile forward/backward matvec with injected analog noise, pulsed stochastic weight update, noise/weight management, transfer devices). All standard HIP-portable. +- cuBLAS (cuda_math_util.cu, cuda_util.h/.cu, rpucuda_pulsed_device.cu): Sgemm/Dgemm/Hgemm, Sgemv/Dgemv, Sger/Dger, Sscal/Dscal, Scopy/Dcopy, Snrm2/Dnrm2; cublasHandle_t, cublasSetStream, cublasSet/GetPointerMode (HOST/DEVICE), cublasSetPointerMode. -> hipBLAS, 1:1. WATCH: hipBLAS v2 enum spellings (CUBLAS_POINTER_MODE_DEVICE -> HIPBLAS_POINTER_MODE_DEVICE, op-N/T enums); device-pointer-mode path (getBlasDeviceHandle / cublasHandle_t* device_handle_ in cuda_util.h) -- verify hipBLAS supports the same device-pointer scalar mode. +- cuRAND: BOTH APIs. Host generator `curandGenerator_t rng_` (cuda_util.h:403) and device per-thread state `curandState_t` / `curandStateXORWOW` with `curand_init`, `curand_normal`, `curand_uniform`, `curandSetStream`, `curandSetup` (cuda_util.h:500). -> hipRAND, 1:1 (hiprand_kernel.h device API + host generator). WATCH: XORWOW generator state layout and stream-ordered seeding; analog noise is the core of the simulation, so RNG correctness is gated by the CPU-vs-CUDA test (see Test plan). Bitstream values will NOT match CUDA bit-for-bit (different RNG impl) -- tests compare statistically / against CPU tile, not against a fixed CUDA seed (confirm tolerance-based asserts). +- CUB (rpu_cub.h wraps `` in namespace RPU::cub; used in noise_manager.cu, maximizer.cu, update_management_helper.cu, weight_clipper_cuda.cu, mixedprec, cuda_buffer.h): DeviceReduce/DeviceScan-class and `cub::BlockScan` (block-width templated, NOT warp-width). -> hipCUB. rpu_cub.h must include `` and use the hipCUB namespace under HIP; the CUB_NS_PREFIX wrapping needs a hipCUB-compatible spelling. WATCH the wave64 TempStorage-reuse race (CV-CUDA lesson) on any back-to-back block-collective sharing TempStorage. +- Thrust: none. +- Warp intrinsics: concentrated in bit_line_maker.cu -- `__shfl_up_sync`, `__shfl_sync`, `__ballot_sync` (all with mask `0xFFFFFFFF`), `warpSize`; `__popc` in pwu_kernel.h (operates on a 32-bit packed K-word, logical not wavefront), `__popcll` in cuda_math_util.cu (operates on a uint64 per-thread flicker_state, NOT a ballot -- arch-independent, fine). +- FP16: cuda_fp16_util.h + RPU_USE_FP16/RPU_BFLOAT_AS_FP16 (experimental, OFF by default; CMakeLists comments "only supported for A100+"). Port path: hip_fp16 / hip_bfloat16. Defer: build with FP16 OFF (default) for the lead port; flag RPU_USE_FP16 as a follow-up. +- Atomics: atomicAdd, atomicCAS, atomicOr, plus a project `atomicMaxFP` helper (CAS-based float max). Standard; verify atomicMaxFP CAS loop compiles/behaves under HIP. +- Pinned memory: cudaHostAlloc in io_manager.cu -> hipHostAlloc/hipHostMalloc, 1:1. +- Textures/surfaces: NONE. (Entire popsift/CV-CUDA texture-fault-class set does not apply.) +- Managed memory: none. +- Streams/events: standard (per-context stream + cublas/curand SetStream). 1:1. +- CUTLASS/CuTe/wgmma/tensor-core: NONE (only aspirational comments). No CK rewrite needed. + +## Risk list +1. WARP-SIZE-COUPLED SERIALIZED DATA FORMAT (PRIMARY RISK -- dietgpu class). bit_line_maker.cu encodes the stochastic pulse train as a bit-packed format: `nK32 = (Kplus1+31)/32` 32-bit words, `laneId = threadIdx.x & 0x1f`, `__ballot_sync(0xFFFFFFFF, stoch_value < value)` whose 32-bit result IS one word of the `CudaArray dev_counts` consumed downstream by the pulsed-weight-updater kernels (pwu_kernel.h reads it back word-by-word with `>>5`/`&0x1f`/`__popc`). The producer and consumer must agree on a 32-bit-per-word layout regardless of wavefront width. + - On wave64, `__ballot_sync(0xFFFFFFFF, pred)` returns a 64-bit value; the low 32 bits are lanes 0-31, high 32 bits lanes 32-63. With `laneId = threadIdx.x & 0x1f` there are TWO logical 32-lane groups per wavefront, each with a `laneId==0` leader -> the wave64 leader/broadcast/ballot logic must be done PER 32-lane group (the popsift two-rows-per-wavefront lesson), not over the whole 64-lane wavefront. The serialized word stays 32 bits (do NOT widen to 64; pwu_kernel consumes uint32 words), so a wave32 device and a wave64 device must produce the SAME 32-bit-word stream. + - Concretely: keep the logical-warp width = 32 for the ballot (mask the correct 32-lane group, capture only that group's 32 bits, the group's lane-0 is the leader). The launch geometry `numwarpsperblock = RPU_THREADS_PER_BLOCK_UPDATE / 32` (bit_line_maker.cu:641) ALREADY assumes 32-lane logical warps; on wave64 two such logical warps share a wavefront. Make every ballot/shfl explicitly width-32 (`__ballot`/`__shfl` over the 32-lane subgroup) and lane = `threadIdx.x & 0x1f` so the format is wave-width-independent. Verify on BOTH wave64 (gfx90a) and wave32 (gfx1100) -- the dev_counts stream must be byte-identical across widths for the same RNG-disabled input, or the pulsed update will read a garbage pulse train (the bug fingerprint: pulsed-update tests pass on CPU/CUDA, fail or non-deterministic on AMD). + - The `0xFFFFFFFF` literal masks must NOT be widened to 64-bit (that would make the format 64-coupled); they must select the active 32-lane group. The `warpSize` references (bit_line_maker.cu:575,1166) in launch-geometry comments need checking against the actual `numwarpsperblock` math. +2. hipBLAS v2 enum/handle differences (pointer-mode, op enums, device-pointer-mode handle path). Mechanical but verify the device-handle scalar path. +3. CUB block-collective TempStorage reuse race on wave64 (add __syncthreads between reused-TempStorage block collectives). Audit update_management_helper.cu BlockScan and noise_manager/maximizer reduces. +4. IPO/LTO + HIP -> broken pybind PyInit (gpuRIR). Disable IPO for the HIP build. +5. host memcpy/memset resolving to HIP __device__ overloads inside .cu (include cstring/cstdlib before hip_runtime in compat header). +6. RNG non-bit-matching: hipRAND XORWOW != cuRAND XORWOW bitstream. Confirm the test suite uses tolerance/statistical asserts (CPU-vs-CUDA-tile), not fixed-seed bit-equality. If any test pins a CUDA RNG bitstream, that test is not a valid cross-arch gate -- note it. +7. atomicMaxFP CAS-based float-max helper: verify HIP semantics (NaN handling) match. +8. f32 math exactness (CV-CUDA __fsqrt_rn / ffp-contract lessons): aihwkit asserts are tolerance-based (analog noise sim), so bit-exactness is unlikely to gate, but if any tile test tightens tolerance, pin -ffp-contract=on for HIP. + +## File-by-file change sketch +- src/rpucuda/cuda/cuda_to_hip.h (NEW): the only HIP-aware file; cuda*->hip* aliases for the runtime, cuBLAS, cuRAND (host+device), CUB namespace; include order cstring/cstdlib before hip_runtime. +- src/rpucuda/cuda/cuda_util.h: include cuda_to_hip.h at top (before the cublas_v2.h/cuda.h/curand*.h block at lines 18-23), guarded so CUDA build is unchanged. +- src/rpucuda/cuda/rpu_cub.h: under USE_HIP include `` and use hipCUB namespace (the CUB_NS_PREFIX wrapping needs a hipCUB spelling). +- src/rpucuda/cuda/bit_line_maker.cu: make the ballot/shfl/leader logic width-32-logical-warp explicit and wave-width-independent (PRIMARY fix; risk 1). Possibly the only .cu needing real logic edits. +- src/rpucuda/cuda/cuda_math_util.cu: hipBLAS enum/handle spelling fixes (pointer mode, op enums); confirm via compat aliases where possible. +- src/rpucuda/cuda/update_management_helper.cu, noise_manager.cu, maximizer.cu, weight_clipper_cuda.cu: add __syncthreads around reused block-collective TempStorage if needed (risk 3). +- src/rpucuda/cuda/io_manager.cu: cudaHostAlloc -> hip via compat alias. +- CMakeLists.txt: add `option(USE_HIP ...)`; HIP branch mirroring the USE_CUDA branch (enable_language(HIP), LANGUAGE HIP on RPU_GPU_SRCS, HIP_ARCHITECTURES from CMAKE_HIP_ARCHITECTURES, link hipblas hiprand, keep RPU_USE_CUDA defined, disable IPO). Apply the same to the BUILD_TEST and AIHWKIT_EXTENSION_OPS_GPU branches. +- cmake/dependencies_cuda.cmake: add a USE_HIP path (or a sibling dependencies_hip.cmake) -- enable_language(HIP), find hipBLAS/hipRAND/hipCUB, add_compile_definitions(RPU_USE_CUDA) so device code stays gated on. +- src/aihwkit/simulator/CMakeLists.txt + src/aihwkit/extension/CMakeLists.txt: link RPU_GPU/hip libs under USE_HIP; ensure pybind module IPO off. +- setup.py: scikit-build passes -DUSE_HIP=ON via env/cmake args (USE_CUDA already reads $ENV{USE_CUDA}); add a parallel USE_HIP env hook. + +## Build commands (gfx90a) +Prereqs present on host: ROCm 7.2 (/opt/rocm), hipcc, hipBLAS, hipRAND, hipCUB, gfx90a GPU. Need a ROCm PyTorch in the env (torch.version.hip set) for torch_cuda/c10_cuda linkage and the pinned scikit-build torch include path. + +Direct CMake (bring-up / C++ gtests): +``` +cd projects/aihwkit/src +cmake -S . -B build_hip -GNinja \ + -DUSE_HIP=ON -DUSE_CUDA=OFF \ + -DCMAKE_HIP_ARCHITECTURES=gfx90a \ + -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ \ + -DCMAKE_PREFIX_PATH="$(python -c 'import torch,os;print(os.path.dirname(torch.__file__))')/share/cmake" \ + -DRPU_BLAS=OpenBLAS -DBUILD_TEST=ON -DRPU_USE_TORCH_BUFFERS=OFF +cmake --build build_hip -j16 +``` +Multi-arch sanity (warp-size class): add `-DCMAKE_HIP_ARCHITECTURES="gfx90a;gfx1100"` and confirm both code objects via `llvm-objdump --offloading build_hip/.../librpucuda*.so`. + +Python package (the real install the pytest suite imports): +``` +cd projects/aihwkit/src +USE_HIP=ON USE_CUDA=0 CMAKE_HIP_ARCHITECTURES=gfx90a \ + pip install -v -e . \ + --config-settings=cmake.args="-DUSE_HIP=ON;-DCMAKE_HIP_ARCHITECTURES=gfx90a;-DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++" +``` +(Exact scikit-build flag plumbing -- skbuild vs setuptools_scm config-settings spelling -- is an open question for the porter; the Makefile `make build_inplace_cuda` target is the CUDA reference to mirror.) + +## Test plan +The pytest suite is the validation gate, and it is an IDEAL cross-arch correctness gate: tile tests are parametrized over CPU and CUDA tiles (tests/helpers/tiles.py: `use_cuda` flag + `CudaAnalogTile` via `getattr(tiles, "CudaAnalogTile", None)`; `.cuda()` conversion in simulator/tiles/base.py), so the SAME analog-tile forward/backward/update test runs on the C++/CPU tile and the HIP/CUDA tile and compares against the same reference. With `RPU_USE_CUDA` kept defined on the HIP build, `CudaAnalogTile` is exposed and `.cuda()` lands on the gfx90a device. + +GPU-gating tests (run on real gfx90a; these exercise RPUCuda tiles on device vs CPU/C++): +- tests/test_simulator_tiles.py -- core analog tile forward/backward/update, CPU vs CUDA. +- tests/test_specific_tiles.py -- per-device-model tiles (constant/linear/exp/pow-step, softbounds, transfer, mixedprec) on CPU vs CUDA; exercises bit_line_maker + pwu update path (the warp-size risk). +- tests/test_bindings_tiles.py -- direct C++-binding tile tests, CPU vs CUDA. +- tests/test_torch_tiles.py, tests/test_inference_tiles.py, tests/test_quantized_tile.py -- torch-backed and inference tiles on device. +- tests/test_layers*.py (linear/convolution/mapped/rnn) -- analog layers on CUDA tiles. +- tests/test_extension.py -- only if BUILD_EXTENSION=ON (defer; extension is optional). + +Commands: +``` +cd projects/aihwkit/src +pip install -r requirements-dev.txt # pytest + dev tools +# GPU gate (device tile correctness, the warp-size-sensitive update path): +pytest -v -s tests/test_simulator_tiles.py tests/test_specific_tiles.py tests/test_bindings_tiles.py +# Broader analog-layer GPU coverage: +pytest -v -s tests/test_torch_tiles.py tests/test_inference_tiles.py tests/test_layers_linear.py tests/test_layers_convolution.py +# Full suite (also guards non-GPU regressions): +make pytest +``` +Non-GPU regression set (must not regress): the same files run their CPU-tile parametrizations, plus pure-Python tests (test_presets, test_rpu_configurations, test_conversions, test_quant_conversion, test_utils, test_export, test_optimizers). The cloud/experiment runner tests (test_cloud_runner, test_client, test_experiment_runners, test_localrunner_infer) may need network/credentials -- treat failures there as environment, not port regressions; baseline them on the unmodified CPU build first. + +Dataset needs: none for the tile/layer GPU gate (tests synthesize small tensors). Some example/notebook flows pull datasets, but they are not in the pytest gate; host egress is slow, so avoid dataset-dependent tests. + +Cross-arch gate for followers (gfx1100 wave32): re-run test_specific_tiles.py and confirm the pulsed-update results match gfx90a for the same input (the warp-size data-format risk is exactly what diverges between wave64 and wave32 if the BLM fix is wrong). Where RNG is involved, compare tolerance-based tile outputs, not RNG bitstreams. + +## Open questions +1. scikit-build config-settings spelling to pass -DUSE_HIP / arch through `pip install -e .` (skbuild vs scikit-build-core; the Makefile CUDA target is the reference). Bring-up can use direct CMake first. +2. Does hipBLAS support the device-pointer scalar mode the code uses (getBlasDeviceHandle, cublasHandle_t* device_handle_)? Verify the device-pointer-mode GEMV/scal path. +3. RPU_USE_TORCH_BUFFERS=ON requires torch CUDA buffers (RPU_TORCH_CUDA_BUFFERS, --expt-relaxed-constexpr). On a ROCm torch the buffers are HIP tensors; confirm the torch buffer path compiles under HIP (start with RPU_USE_TORCH_BUFFERS=OFF for bring-up, then turn on). Note BUILD_TEST forces it OFF anyway. +4. Does any tile test assert against a FIXED cuRAND bitstream (would be a false cross-arch gate)? Inspect the tolerance of the CPU-vs-CUDA comparisons before trusting them as the warp-size gate. +5. FP16/bfloat16 path (RPU_USE_FP16): deferred (off by default). hip_fp16/hip_bfloat16 port is a follow-up if a tile test needs half precision. +6. Whether scikit-build/CMake enables IPO that breaks the HIP pybind module (gpuRIR) -- confirm and force IPO off for HIP. diff --git a/projects/aihwkit/stats.jsonl b/projects/aihwkit/stats.jsonl new file mode 100644 index 00000000..64fb5b4d --- /dev/null +++ b/projects/aihwkit/stats.jsonl @@ -0,0 +1,63 @@ +{"kind": "tokens", "ts": "2026-06-04T01:26:30Z", "tokens": 76430} +{"kind":"phase","ts":"2026-06-04T01:34:02Z","phase":"compile","seconds":5.983,"exit":1,"cmd":"bash -c cd projects/aihwkit/src && cmake --build build_hip -j16"} +{"kind":"phase","ts":"2026-06-04T01:34:49Z","phase":"compile","seconds":7.599,"exit":1,"cmd":"bash -c cd projects/aihwkit/src && cmake --build build_hip -j16"} +{"kind":"phase","ts":"2026-06-04T01:36:50Z","phase":"compile","seconds":92.630,"exit":1,"cmd":"bash -c cd projects/aihwkit/src && cmake --build build_hip -j16"} +{"kind":"phase","ts":"2026-06-04T01:38:45Z","phase":"compile","seconds":88.129,"exit":1,"cmd":"bash -c cd projects/aihwkit/src && cmake --build build_hip -j16"} +{"kind":"phase","ts":"2026-06-04T01:41:08Z","phase":"compile","seconds":89.007,"exit":1,"cmd":"bash -c cd projects/aihwkit/src && cmake --build build_hip -j16"} +{"kind":"phase","ts":"2026-06-04T01:54:18Z","phase":"compile","seconds":213.646,"exit":0,"cmd":"bash -c cd /var/lib/jenkins/moat/projects/aihwkit/src && USE_HIP=ON USE_CUDA=0 python setup.py build_ext -j16 --inplace -DUSE_HIP=ON -DUSE_CUDA=OFF -DRPU_CXX_STANDARD=20 -DCMAKE_HIP_ARCHITECTURES=gfx90a -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ -DCMAKE_PREFIX_PATH='/opt/conda/envs/py_3.12/lib/python3.12/site-packages/torch/share/cmake;/opt/rocm' -DRPU_BLAS=OpenBLAS -DRPU_USE_TORCH_BUFFERS=OFF -DCMAKE_BUILD_TYPE=Release"} +{"kind":"phase","ts":"2026-06-04T01:55:13Z","phase":"test","seconds":7.517,"exit":0,"cmd":"bash -c cd /var/lib/jenkins/moat/projects/aihwkit/src && HIP_VISIBLE_DEVICES=0 PYTHONPATH=src python -m pytest -x -q tests/test_specific_tiles.py 2>&1 | tail -30"} +{"kind":"phase","ts":"2026-06-04T02:06:28Z","phase":"test","seconds":658.530,"exit":0,"cmd":"bash -c cd /var/lib/jenkins/moat/projects/aihwkit/src && HIP_VISIBLE_DEVICES=0 PYTHONPATH=src python -m pytest -q tests/test_simulator_tiles.py tests/test_bindings_tiles.py 2>&1 | tail -25"} +{"kind":"phase","ts":"2026-06-04T02:09:39Z","phase":"test","seconds":26.271,"exit":0,"cmd":"bash -c cd /var/lib/jenkins/moat/projects/aihwkit/src && HIP_VISIBLE_DEVICES=0 PYTHONPATH=src python -m pytest -q tests/test_torch_tiles.py tests/test_inference_tiles.py 2>&1 | tail -15"} +{"kind":"phase","ts":"2026-06-04T02:11:07Z","phase":"test","seconds":83.122,"exit":0,"cmd":"bash -c cd /var/lib/jenkins/moat/projects/aihwkit/src && HIP_VISIBLE_DEVICES=0 PYTHONPATH=src python -m pytest -q tests/test_layers_linear.py tests/test_layers_convolution.py 2>&1 | tail -12"} +{"kind": "tokens", "ts": "2026-06-04T02:18:38Z", "tokens": 182814} +{"kind": "tokens", "ts": "2026-06-04T02:25:43Z", "tokens": 114325} +{"kind":"phase","ts":"2026-06-04T02:27:29Z","phase":"compile","seconds":0.041,"exit":2,"cmd":"bash -c \nTORCH_CMAKE=/opt/conda/envs/py_3.12/lib/python3.12/site-packages/torch/share/cmake\nUSE_HIP=ON USE_CUDA=0 python setup.py build_ext -j16 --inplace \\\n -DUSE_HIP=ON -DUSE_CUDA=OFF -DRPU_CXX_STANDARD=20 \\\n -DCMAKE_HIP_ARCHITECTURES=gfx90a \\\n -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ \\\n \"-DCMAKE_PREFIX_PATH=$TORCH_CMAKE;/opt/rocm\" \\\n -DRPU_BLAS=OpenBLAS -DRPU_USE_TORCH_BUFFERS=OFF -DCMAKE_BUILD_TYPE=Release\n"} +{"kind":"phase","ts":"2026-06-04T02:31:12Z","phase":"compile","seconds":215.746,"exit":0,"cmd":"bash -c \ncd /var/lib/jenkins/moat/projects/aihwkit/src\nTORCH_CMAKE=/opt/conda/envs/py_3.12/lib/python3.12/site-packages/torch/share/cmake\nUSE_HIP=ON USE_CUDA=0 python setup.py build_ext -j16 --inplace \\\n -DUSE_HIP=ON -DUSE_CUDA=OFF -DRPU_CXX_STANDARD=20 \\\n -DCMAKE_HIP_ARCHITECTURES=gfx90a \\\n -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ \\\n \"-DCMAKE_PREFIX_PATH=$TORCH_CMAKE;/opt/rocm\" \\\n -DRPU_BLAS=OpenBLAS -DRPU_USE_TORCH_BUFFERS=OFF -DCMAKE_BUILD_TYPE=Release\n"} +{"kind":"phase","ts":"2026-06-04T02:31:44Z","phase":"test","seconds":7.338,"exit":0,"cmd":"bash -c \nexport HIP_VISIBLE_DEVICES=0\nexport PYTHONPATH=/var/lib/jenkins/moat/projects/aihwkit/src/src\ncd /var/lib/jenkins/moat/projects/aihwkit/src\npython -m pytest -v tests/test_specific_tiles.py 2>&1\n"} +{"kind":"session","ts":"2026-06-04T03:00:55Z","epoch":1780542055.018483898,"event":"start","platform":"linux-gfx1100"} +{"kind":"phase","ts":"2026-06-04T03:01:13Z","phase":"compile","seconds":0.027,"exit":2,"cmd":"bash -c USE_HIP=ON USE_CUDA=0 python setup.py build_ext -j16 --inplace -DUSE_HIP=ON -DUSE_CUDA=OFF -DRPU_CXX_STANDARD=20 -DCMAKE_HIP_ARCHITECTURES=gfx1100 -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ -DCMAKE_PREFIX_PATH=\"/opt/conda/envs/py_3.12/lib/python3.12/site-packages/torch/share/cmake;/opt/rocm\" -DRPU_BLAS=OpenBLAS -DRPU_USE_TORCH_BUFFERS=OFF -DCMAKE_BUILD_TYPE=Release 2>&1"} +{"kind":"phase","ts":"2026-06-04T03:03:34Z","phase":"compile","seconds":127.646,"exit":0,"cmd":"env -C /var/lib/jenkins/moat/projects/aihwkit/src bash -c USE_HIP=ON USE_CUDA=0 python setup.py build_ext -j16 --inplace -DUSE_HIP=ON -DUSE_CUDA=OFF -DRPU_CXX_STANDARD=20 -DCMAKE_HIP_ARCHITECTURES=gfx1100 -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ \"-DCMAKE_PREFIX_PATH=/opt/conda/envs/py_3.12/lib/python3.12/site-packages/torch/share/cmake;/opt/rocm\" -DRPU_BLAS=OpenBLAS -DRPU_USE_TORCH_BUFFERS=OFF -DCMAKE_BUILD_TYPE=Release 2>&1"} +{"kind":"phase","ts":"2026-06-04T03:03:43Z","phase":"compile","seconds":0.640,"exit":0,"cmd":"env -C /var/lib/jenkins/moat/projects/aihwkit/src bash -c USE_HIP=ON USE_CUDA=0 python setup.py build_ext -j16 --inplace -DUSE_HIP=ON -DUSE_CUDA=OFF -DRPU_CXX_STANDARD=20 -DCMAKE_HIP_ARCHITECTURES=gfx1100 -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ \"-DCMAKE_PREFIX_PATH=/opt/conda/envs/py_3.12/lib/python3.12/site-packages/torch/share/cmake;/opt/rocm\" -DRPU_BLAS=OpenBLAS -DRPU_USE_TORCH_BUFFERS=OFF -DCMAKE_BUILD_TYPE=Release 2>&1"} +{"kind":"phase","ts":"2026-06-04T03:06:15Z","phase":"compile","seconds":127.030,"exit":0,"cmd":"env -C /var/lib/jenkins/moat/projects/aihwkit/src bash -c USE_HIP=ON USE_CUDA=0 python setup.py build_ext -j16 --inplace -DUSE_HIP=ON -DUSE_CUDA=OFF -DRPU_CXX_STANDARD=20 -DCMAKE_HIP_ARCHITECTURES=gfx1100 -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ \"-DCMAKE_PREFIX_PATH=/opt/conda/envs/py_3.12/lib/python3.12/site-packages/torch/share/cmake;/opt/rocm\" -DRPU_BLAS=OpenBLAS -DRPU_USE_TORCH_BUFFERS=OFF -DCMAKE_BUILD_TYPE=Release 2>&1"} +{"kind":"phase","ts":"2026-06-04T03:06:49Z","phase":"test","seconds":0.528,"exit":4,"cmd":"env HIP_VISIBLE_DEVICES=1 PYTHONPATH=src python -m pytest tests/test_specific_tiles.py tests/test_simulator_tiles.py tests/test_bindings_tiles.py tests/test_torch_tiles.py tests/test_inference_tiles.py tests/test_layers_linear.py tests/test_layers_convolution.py -v --tb=short"} +{"kind":"phase","ts":"2026-06-04T03:07:50Z","phase":"test","seconds":55.100,"exit":0,"cmd":"env -C /var/lib/jenkins/moat/projects/aihwkit/src bash -c HIP_VISIBLE_DEVICES=1 PYTHONPATH=src python -m pytest tests/test_specific_tiles.py tests/test_simulator_tiles.py tests/test_bindings_tiles.py tests/test_torch_tiles.py tests/test_inference_tiles.py tests/test_layers_linear.py tests/test_layers_convolution.py -v --tb=short 2>&1"} +{"kind":"session","ts":"2026-06-04T03:08:30Z","epoch":1780542510.768188529,"event":"end","platform":"linux-gfx1100"} +{"kind": "tokens", "ts": "2026-06-04T03:09:58Z", "tokens": 52203, "source": "validator"} +{"kind":"phase","ts":"2026-06-08T03:04:54Z","phase":"compile","seconds":0.182,"exit":1,"cmd":"cmake --build /b/develop/moat/agent_space/aihwkit_gfx1201_build -j32"} +{"kind":"phase","ts":"2026-06-08T03:05:37Z","phase":"compile","seconds":0.190,"exit":1,"cmd":"cmake --build /b/develop/moat/agent_space/aihwkit_gfx1201_build -j32"} +{"kind":"phase","ts":"2026-06-08T03:08:53Z","phase":"compile","seconds":143.689,"exit":1,"cmd":"cmake --build /b/develop/moat/agent_space/aihwkit_gfx1201_build -j32"} +{"kind":"phase","ts":"2026-06-08T03:51:43Z","phase":"test","seconds":8.813,"exit":0,"cmd":"B:/develop/TheRock/external-builds/pytorch/.venv/Scripts/python.exe -m pytest B:/develop/moat/projects/aihwkit/src/tests/test_specific_tiles.py -x -v --tb=short -p no:warnings"} +{"kind":"phase","ts":"2026-06-08T04:23:22Z","phase":"compile","seconds":143.500,"exit":0,"cmd":"bash /var/lib/jenkins/moat/agent_space/aihwkit-build-old.sh"} +{"kind":"phase","ts":"2026-06-08T04:26:09Z","phase":"compile","seconds":141.293,"exit":1,"cmd":"bash /var/lib/jenkins/moat/agent_space/aihwkit-build-new.sh"} +{"kind":"phase","ts":"2026-06-08T04:38:25Z","phase":"compile","seconds":81.682,"exit":0,"cmd":"bash /var/lib/jenkins/moat/agent_space/aihwkit-setup-build.sh"} +{"kind":"phase","ts":"2026-06-08T04:38:32Z","phase":"compile","seconds":1.064,"exit":0,"cmd":"bash /var/lib/jenkins/moat/agent_space/aihwkit-setup-build.sh"} +{"kind":"phase","ts":"2026-06-08T04:41:51Z","phase":"test","seconds":136.298,"exit":0,"cmd":"bash /var/lib/jenkins/moat/agent_space/aihwkit-test.sh"} +{"kind":"phase","ts":"2026-06-08T13:49:58Z","phase":"compile","seconds":1.066,"exit":1,"cmd":"conda run -n py_3.12 cmake -S . -B build_hip_old -GNinja -DUSE_HIP=ON -DUSE_CUDA=OFF -DRPU_CXX_STANDARD=20 -DCMAKE_HIP_ARCHITECTURES=gfx1100 -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ -DCMAKE_PREFIX_PATH=/opt/conda/envs/py_3.12/lib/python3.12/site-packages/torch/share/cmake;/opt/rocm -DRPU_BLAS=OpenBLAS -DBUILD_TEST=OFF -DRPU_USE_TORCH_BUFFERS=OFF -DCMAKE_BUILD_TYPE=Release"} +{"kind":"phase","ts":"2026-06-08T13:52:38Z","phase":"compile","seconds":135.868,"exit":0,"cmd":"conda run -n py_3.12 bash /var/lib/jenkins/moat/agent_space/aihwkit-gfx1100-gpu2/build_new.sh"} +{"kind": "tokens", "ts": "2026-06-08T13:55:10Z", "tokens": 48496, "source": "validator"} +{"kind": "tokens", "ts": "2026-06-08T04:43:29Z", "tokens": 98344, "source": "validator"} +{"kind":"phase","ts":"2026-06-08T14:47:09Z","phase":"compile","seconds":135.970,"exit":0,"cmd":"cmake --build /var/lib/jenkins/moat/agent_space/build_new -j16"} +{"kind":"phase","ts":"2026-06-08T14:50:42Z","phase":"compile","seconds":138.157,"exit":0,"cmd":"cmake --build /var/lib/jenkins/moat/agent_space/build_old -j16"} +{"kind":"phase","ts":"2026-06-08T14:55:28Z","phase":"test","seconds":0.695,"exit":4,"cmd":"python -m pytest tests/test_specific_tiles.py -v"} +{"kind":"phase","ts":"2026-06-08T14:55:48Z","phase":"test","seconds":7.943,"exit":0,"cmd":"python -m pytest /var/lib/jenkins/moat/projects/aihwkit/src/tests/test_specific_tiles.py -v"} +{"kind":"phase","ts":"2026-06-08T14:56:59Z","phase":"test","seconds":60.952,"exit":0,"cmd":"python -m pytest /var/lib/jenkins/moat/projects/aihwkit/src/tests/test_simulator_tiles.py -v -k Cuda"} +{"kind":"phase","ts":"2026-06-08T14:57:22Z","phase":"test","seconds":10.763,"exit":0,"cmd":"python -m pytest /var/lib/jenkins/moat/projects/aihwkit/src/tests/test_inference_tiles.py -k Cuda"} +{"kind": "tokens", "ts": "2026-06-08T15:57:45Z", "tokens": 76158, "source": "reviewer"} +{"kind":"phase","ts":"2026-06-08T16:03:30Z","phase":"compile","seconds":146.032,"exit":0,"cmd":"bash -c \n cd /var/lib/jenkins/moat/projects/aihwkit/src && cmake -S . -B /var/lib/jenkins/moat/agent_space/aihwkit-cfwd-new -GNinja -DUSE_HIP=ON -DUSE_CUDA=OFF -DRPU_CXX_STANDARD=20 -DCMAKE_HIP_ARCHITECTURES=gfx90a -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ '-DCMAKE_PREFIX_PATH=/opt/conda/envs/py_3.12/lib/python3.12/site-packages/torch/share/cmake;/opt/rocm' -DRPU_BLAS=OpenBLAS -DBUILD_TEST=OFF -DRPU_USE_TORCH_BUFFERS=OFF 2>&1 | tee /var/lib/jenkins/moat/agent_space/aihwkit_cfwd_configure_new.log && cmake --build /var/lib/jenkins/moat/agent_space/aihwkit-cfwd-new -j16 2>&1 | tail -20\n"} +{"kind":"phase","ts":"2026-06-08T16:06:22Z","phase":"compile","seconds":144.751,"exit":0,"cmd":"bash -c \n cmake -S /var/lib/jenkins/moat/agent_space/aihwkit-wt-old -B /var/lib/jenkins/moat/agent_space/aihwkit-cfwd-old -GNinja -DUSE_HIP=ON -DUSE_CUDA=OFF -DRPU_CXX_STANDARD=20 -DCMAKE_HIP_ARCHITECTURES=gfx90a -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ '-DCMAKE_PREFIX_PATH=/opt/conda/envs/py_3.12/lib/python3.12/site-packages/torch/share/cmake;/opt/rocm' -DRPU_BLAS=OpenBLAS -DBUILD_TEST=OFF -DRPU_USE_TORCH_BUFFERS=OFF 2>&1 | tee /var/lib/jenkins/moat/agent_space/aihwkit_cfwd_configure_old.log && cmake --build /var/lib/jenkins/moat/agent_space/aihwkit-cfwd-old -j16 2>&1 | tail -20\n"} +{"kind":"phase","ts":"2026-06-08T16:11:03Z","phase":"test","seconds":3.828,"exit":2,"cmd":"bash -c \n cd /var/lib/jenkins/moat/projects/aihwkit/src && PYTHONPATH=src HIP_VISIBLE_DEVICES=0 /opt/conda/envs/py_3.12/bin/python3.12 -m pytest tests/test_specific_tiles.py -v --tb=short 2>&1\n"} +{"kind":"phase","ts":"2026-06-08T16:11:29Z","phase":"test","seconds":8.036,"exit":0,"cmd":"bash -c \n cd /var/lib/jenkins/moat/projects/aihwkit/src && HIP_VISIBLE_DEVICES=0 PYTHONPATH=src /opt/conda/envs/py_3.12/bin/python3.12 -m pytest tests/test_specific_tiles.py -v --tb=short 2>&1\n "} +{"kind":"phase","ts":"2026-06-08T16:56:22Z","phase":"compile","seconds":191.117,"exit":0,"cmd":"bash -c \nexport HIP_VISIBLE_DEVICES=1\ncd /var/lib/jenkins/moat/agent_space/aihwkit-gfx1100-gpu1/old\ncmake -S . -B build_gfx1100 -GNinja -DUSE_HIP=ON -DUSE_CUDA=OFF -DRPU_CXX_STANDARD=20 -DCMAKE_HIP_ARCHITECTURES=gfx1100 -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ '-DCMAKE_PREFIX_PATH=/opt/conda/envs/py_3.12/lib/python3.12/site-packages/torch/share/cmake;/opt/rocm' -DRPU_BLAS=OpenBLAS -DBUILD_TEST=OFF -DRPU_USE_TORCH_BUFFERS=OFF -DCMAKE_BUILD_TYPE=Release 2>&1\ncmake --build build_gfx1100 -j16 2>&1\n"} +{"kind":"phase","ts":"2026-06-08T16:59:07Z","phase":"compile","seconds":158.420,"exit":0,"cmd":"bash -c \nexport HIP_VISIBLE_DEVICES=1\ncd /var/lib/jenkins/moat/agent_space/aihwkit-gfx1100-gpu1/new\ncmake -S . -B build_gfx1100 -GNinja -DUSE_HIP=ON -DUSE_CUDA=OFF -DRPU_CXX_STANDARD=20 -DCMAKE_HIP_ARCHITECTURES=gfx1100 -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ '-DCMAKE_PREFIX_PATH=/opt/conda/envs/py_3.12/lib/python3.12/site-packages/torch/share/cmake;/opt/rocm' -DRPU_BLAS=OpenBLAS -DBUILD_TEST=OFF -DRPU_USE_TORCH_BUFFERS=OFF -DCMAKE_BUILD_TYPE=Release 2>&1\ncmake --build build_gfx1100 -j16 2>&1\n"} +{"kind": "tokens", "ts": "2026-06-08T17:01:17Z", "tokens": 50822, "source": "validator"} +{"kind": "tokens", "ts": "2026-06-08T16:13:38Z", "tokens": 71037, "source": "validator"} +{"kind": "tokens", "ts": "2026-06-08T03:59:32Z", "tokens": 40723, "source": "validator"} +{"kind": "tokens", "ts": "2026-06-08T20:54:53Z", "tokens": 26777, "source": "revalidate-batch"} +{"kind":"phase","ts":"2026-06-19T20:16:45Z","phase":"compile","seconds":1.713,"exit":0,"cmd":"bash -c \nVENV=\"B:/develop/TheRock/external-builds/pytorch/.venv\"\nROCM_DEVEL=\"$VENV/Lib/site-packages/_rocm_sdk_devel\"\nROCM_LIBS=\"$VENV/Lib/site-packages/_rocm_sdk_libraries\"\nHOSTMATH=\"$ROCM_DEVEL/lib/host-math\"\nTORCHLIB=\"$VENV/Lib/site-packages/torch\"\nCLANG=\"$ROCM_DEVEL/lib/llvm/bin/clang++.exe\"\nTORCH_CMAKE=\"$TORCHLIB/share/cmake\"\n\nexport HIP_VISIBLE_DEVICES=1\nexport VIRTUAL_ENV=\"$VENV\"\nexport OPENBLAS_ROOT_DIR=\"$HOSTMATH\"\nexport PATH=\"$ROCM_DEVEL/lib/llvm/bin:$ROCM_DEVEL/bin:$PATH\"\n\ncmake -S \"B:/develop/moat/projects/aihwkit/src\" \\\n -B \"B:/develop/moat/projects/aihwkit/src/build_hip_gfx1101\" \\\n -G Ninja \\\n -DUSE_HIP=ON \\\n -DUSE_CUDA=OFF \\\n -DRPU_CXX_STANDARD=20 \\\n \"-DCMAKE_HIP_ARCHITECTURES=gfx1101\" \\\n \"-DCMAKE_HIP_COMPILER=$CLANG\" \\\n \"-DCMAKE_C_COMPILER=$CLANG\" \\\n \"-DCMAKE_CXX_COMPILER=$CLANG\" \\\n \"-DCMAKE_PREFIX_PATH=$TORCH_CMAKE;$ROCM_DEVEL\" \\\n -DRPU_BLAS=OpenBLAS \\\n \"-DOpenBLAS_INCLUDE_DIR=$HOSTMATH/include\" \\\n \"-DOpenBLAS_LIB=$HOSTMATH/lib/rocm-openblas.lib\" \\\n -DRPU_USE_TORCH_BUFFERS=OFF \\\n -DBUILD_TEST=OFF \\\n -DCMAKE_BUILD_TYPE=Release \\\n \"-DPYTHON_EXECUTABLE=$VENV/Scripts/python.exe\" 2>&1\necho \"CMake exit: $?\"\n"} +{"kind":"phase","ts":"2026-06-19T20:17:15Z","phase":"compile","seconds":6.206,"exit":0,"cmd":"bash -c \nVENV=\"B:/develop/TheRock/external-builds/pytorch/.venv\"\nROCM_DEVEL=\"$VENV/Lib/site-packages/_rocm_sdk_devel\"\nROCM_LIBS=\"$VENV/Lib/site-packages/_rocm_sdk_libraries\"\nHOSTMATH=\"$ROCM_DEVEL/lib/host-math\"\nTORCHLIB=\"$VENV/Lib/site-packages/torch\"\nCLANG=\"$ROCM_DEVEL/lib/llvm/bin/clang++.exe\"\nCLANGC=\"$ROCM_DEVEL/lib/llvm/bin/clang.exe\"\nTORCH_CMAKE=\"$TORCHLIB/share/cmake\"\n\nexport HIP_VISIBLE_DEVICES=1\nexport VIRTUAL_ENV=\"$VENV\"\nexport OPENBLAS_ROOT_DIR=\"$HOSTMATH\"\nexport PATH=\"$ROCM_DEVEL/lib/llvm/bin:$ROCM_DEVEL/bin:$PATH\"\n\ncmake -S \"B:/develop/moat/projects/aihwkit/src\" \\\n -B \"B:/develop/moat/projects/aihwkit/src/build_hip_gfx1101\" \\\n -G Ninja \\\n -DUSE_HIP=ON \\\n -DUSE_CUDA=OFF \\\n -DRPU_CXX_STANDARD=20 \\\n \"-DCMAKE_HIP_ARCHITECTURES=gfx1101\" \\\n \"-DCMAKE_HIP_COMPILER=$CLANG\" \\\n \"-DCMAKE_C_COMPILER=$CLANGC\" \\\n \"-DCMAKE_CXX_COMPILER=$CLANG\" \\\n \"-DCMAKE_PREFIX_PATH=$TORCH_CMAKE;$ROCM_DEVEL\" \\\n -DRPU_BLAS=OpenBLAS \\\n \"-DOpenBLAS_INCLUDE_DIR=$HOSTMATH/include\" \\\n \"-DOpenBLAS_LIB=$HOSTMATH/lib/rocm-openblas.lib\" \\\n -DRPU_USE_TORCH_BUFFERS=OFF \\\n -DBUILD_TEST=OFF \\\n -DCMAKE_BUILD_TYPE=Release \\\n \"-DPYTHON_EXECUTABLE=$VENV/Scripts/python.exe\" 2>&1\necho \"CMake exit: $?\"\n"} +{"kind":"phase","ts":"2026-06-19T20:17:50Z","phase":"compile","seconds":5.789,"exit":0,"cmd":"bash -c \nVENV=\"B:/develop/TheRock/external-builds/pytorch/.venv\"\nROCM_DEVEL=\"$VENV/Lib/site-packages/_rocm_sdk_devel\"\nROCM_LIBS=\"$VENV/Lib/site-packages/_rocm_sdk_libraries\"\nHOSTMATH=\"$ROCM_DEVEL/lib/host-math\"\nTORCHLIB=\"$VENV/Lib/site-packages/torch\"\nCLANG=\"$ROCM_DEVEL/lib/llvm/bin/clang++.exe\"\nCLANGC=\"$ROCM_DEVEL/lib/llvm/bin/clang.exe\"\nTORCH_CMAKE=\"$TORCHLIB/share/cmake\"\n\nexport HIP_VISIBLE_DEVICES=1\nexport VIRTUAL_ENV=\"$VENV\"\nexport OPENBLAS_ROOT_DIR=\"$HOSTMATH\"\nexport PATH=\"$ROCM_DEVEL/lib/llvm/bin:$ROCM_DEVEL/bin:$PATH\"\n\ncmake -S \"B:/develop/moat/projects/aihwkit/src\" \\\n -B \"B:/develop/moat/projects/aihwkit/src/build_hip_gfx1101\" \\\n -G Ninja \\\n -DUSE_HIP=ON \\\n -DUSE_CUDA=OFF \\\n -DRPU_CXX_STANDARD=20 \\\n \"-DCMAKE_HIP_ARCHITECTURES=gfx1101\" \\\n \"-DCMAKE_HIP_COMPILER=$CLANG\" \\\n \"-DCMAKE_C_COMPILER=$CLANGC\" \\\n \"-DCMAKE_CXX_COMPILER=$CLANG\" \\\n \"-DCMAKE_PREFIX_PATH=$TORCH_CMAKE;$ROCM_DEVEL;$ROCM_DEVEL/lib/cmake\" \\\n -DRPU_BLAS=OpenBLAS \\\n \"-DOpenBLAS_INCLUDE_DIR=$HOSTMATH/include\" \\\n \"-DOpenBLAS_LIB=$HOSTMATH/lib/rocm-openblas.lib\" \\\n -DRPU_USE_TORCH_BUFFERS=OFF \\\n -DBUILD_TEST=OFF \\\n -DCMAKE_BUILD_TYPE=Release \\\n \"-DPYTHON_EXECUTABLE=$VENV/Scripts/python.exe\" 2>&1\necho \"CMake exit: $?\"\n"} +{"kind":"phase","ts":"2026-06-19T20:26:09Z","phase":"compile","seconds":124.863,"exit":0,"cmd":"bash -c \nVENV=\"B:/develop/TheRock/external-builds/pytorch/.venv\"\nROCM_DEVEL=\"$VENV/Lib/site-packages/_rocm_sdk_devel\"\nexport PATH=\"$ROCM_DEVEL/lib/llvm/bin:$ROCM_DEVEL/bin:$PATH\"\nexport HIP_VISIBLE_DEVICES=1\n\ncd \"B:/develop/moat/projects/aihwkit/src/build_hip_gfx1101\"\nninja -j64 RPU_CPU.lib RPU_GPU.lib 2>&1 | tail -10\necho \"Build GPU+CPU exit: $?\"\n"} +{"kind":"phase","ts":"2026-06-19T20:29:41Z","phase":"test","seconds":0.426,"exit":4,"cmd":"B:/develop/TheRock/external-builds/pytorch/.venv/Scripts/python.exe -m pytest tests/test_specific_tiles.py -v --timeout=120"} +{"kind": "tokens", "ts": "2026-06-19T20:34:19Z", "tokens": 157196, "source": "validator (gfx1101)"} +{"kind": "tokens", "ts": "2026-06-08T15:01:20Z", "tokens": 100929, "source": "porter"} +{"kind": "tokens", "ts": "2026-06-23T22:08:43Z", "tokens": 38016, "source": "gfx90a audit cross-file fix (general-purpose subagent)"} diff --git a/projects/aihwkit/status.json b/projects/aihwkit/status.json new file mode 100644 index 00000000..ff4fdbe9 --- /dev/null +++ b/projects/aihwkit/status.json @@ -0,0 +1,139 @@ +{ + "schema_version": 3, + "name": "aihwkit", + "upstream_url": "https://github.com/IBM/aihwkit", + "fork_url": "https://github.com/AMD-Ecosystem/aihwkit", + "fork_default_branch": "master", + "priority": 0.0, + "ext_type": "cmake", + "adopted_at": "2026-06-04T01:19:43Z", + "updated_at": "2026-08-07T07:04:50Z", + "head_sha": "70577b578b7faa661611e905c0b9ec1a1d85bf33", + "depends_on": [], + "pr_url": "https://github.com/IBM/aihwkit/pull/770", + "pr_number": 770, + "pr_opened_at": "2026-06-08T22:00:49Z", + "porting": null, + "waivers": {}, + "pr_state": "open", + "license_spdx": "MIT", + "upstream_repo_id": 296369718, + "stage": "review-passed", + "platforms": { + "linux-gfx90a": { + "state": "completed", + "blocked": false, + "blocked_reason": null, + "validated_sha": "4bc09a0fa9a76d3038e6729ea03615fcc0beb76b", + "started_at": "2026-06-04T01:28:01Z", + "completed_at": "2026-06-08T16:12:07Z", + "updated_at": "2026-06-22T23:22:30Z", + "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": "4bc09a0fa9a76d3038e6729ea03615fcc0beb76b", + "method": "source-class", + "detail": "maintainer merged origin/master into PR #770; delta is setup.cfg mypy python_version, visualization.py plt.ylim style, and uv.lock only -- arch-independent, no C++/HIP/GPU code, behavior-preserving", + "at": "2026-06-22T23:22:30Z" + } + }, + "windows-gfx1201": { + "state": "completed", + "blocked": false, + "blocked_reason": null, + "validated_sha": "70577b578b7faa661611e905c0b9ec1a1d85bf33", + "started_at": null, + "completed_at": "2026-06-23T22:08:02Z", + "updated_at": "2026-06-24T17:03:22Z", + "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": "f415b1d4e58bf6352c90ec8993abeb821486971f", + "method": "source-class", + "detail": "CMake default-path change; explicit -D builds byte-identical", + "at": "2026-06-23T22:08:02Z" + } + }, + "windows-gfx1101": { + "state": "completed", + "blocked": false, + "blocked_reason": null, + "validated_sha": "70577b578b7faa661611e905c0b9ec1a1d85bf33", + "started_at": null, + "completed_at": "2026-06-23T22:08:02Z", + "updated_at": "2026-06-24T17:03:22Z", + "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": "f415b1d4e58bf6352c90ec8993abeb821486971f", + "method": "source-class", + "detail": "CMake default-path change; explicit -D builds byte-identical", + "at": "2026-06-23T22:08:02Z" + } + }, + "linux-gfx1100": { + "state": "completed", + "blocked": false, + "blocked_reason": null, + "validated_sha": "70577b578b7faa661611e905c0b9ec1a1d85bf33", + "started_at": null, + "completed_at": "2026-06-23T22:08:02Z", + "updated_at": "2026-06-24T17:03:22Z", + "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": "f415b1d4e58bf6352c90ec8993abeb821486971f", + "method": "source-class", + "detail": "CMake default-path change; explicit -D builds byte-identical", + "at": "2026-06-23T22:08:02Z" + } + } + } +} From 850b751cb7757ed2734016e5e95221f242850640 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Sun, 9 Aug 2026 01:03:00 +0000 Subject: [PATCH 2/4] aihwkit: revalidate linux-gfx90a at 70577b5 (binary-equiv carry-forward + GPU smoke) --- projects/aihwkit/notes.md | 77 ++++++++++++++++++++++++++++++++++++ projects/aihwkit/stats.jsonl | 6 +++ projects/aihwkit/status.json | 16 ++++---- 3 files changed, 91 insertions(+), 8 deletions(-) diff --git a/projects/aihwkit/notes.md b/projects/aihwkit/notes.md index 4ced631e..00e4b153 100644 --- a/projects/aihwkit/notes.md +++ b/projects/aihwkit/notes.md @@ -601,3 +601,80 @@ Test results (PYTHONPATH=src, HIP_VISIBLE_DEVICES=1): Total GPU-gated tests: 1521 passed, 383 skipped, 0 failed. Identical pass count to gfx90a. Verdict: PASS. Transitioning windows-gfx1101 to completed (validated_sha 42e0465). + +## Revalidation 2026-08-09 (linux-gfx90a) + +State: revalidate (validated_sha 4bc09a0 -> head_sha 70577b5). +GPU: AMD Instinct MI250X / MI250 (gfx90a, wave64). ROCm 7.2.1. HIP_VISIBLE_DEVICES=2. + +Delta 4bc09a0..70577b5 (1 commit, CMakeLists.txt only, +2/-3): drops the +`if(NOT DEFINED CMAKE_HIP_ARCHITECTURES) set(... "gfx90a" ...) endif()` pin, +leaving CMAKE_HIP_ARCHITECTURES to `enable_language(HIP)`'s own auto-detect +when a builder does not pass `-DCMAKE_HIP_ARCHITECTURES` explicitly. Our build +recipe always passes `-DCMAKE_HIP_ARCHITECTURES=gfx90a` explicitly, so this is +a default-path-only change here. `moatlib.py classify` returned `mixed` +(CMakeLists.txt token count differs), so did NOT auto-carry; ran the binary- +equivalence check per the mixed/unknown protocol. + +Binary equivalence check (both shas built from the same absolute checkout +path, fresh build dirs, explicit CMAKE_HIP_ARCHITECTURES=gfx90a): +``` +cd projects/aihwkit/src +git checkout 4bc09a0fa9a76d3038e6729ea03615fcc0beb76b +cmake -S . -B build_old -GNinja -DUSE_HIP=ON -DUSE_CUDA=OFF -DRPU_CXX_STANDARD=20 \ + -DCMAKE_HIP_ARCHITECTURES=gfx90a -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ \ + -DCMAKE_PREFIX_PATH="$TORCH_CMAKE;/opt/rocm" -DRPU_BLAS=OpenBLAS -DBUILD_TEST=OFF \ + -DRPU_USE_TORCH_BUFFERS=OFF +cmake --build build_old -j16 +git checkout 70577b578b7faa661611e905c0b9ec1a1d85bf33 +cmake -S . -B build_new -GNinja +cmake --build build_new -j16 +python3 utils/codeobj_diff.py projects/aihwkit/src/build_old projects/aihwkit/src/build_new +``` +Result: overall verdict `indeterminate`, but that came ONLY from CMake's own +compiler-ABI-detection scratch binaries (CMakeDetermineCompilerABI_*.bin, +CompilerIdC/CompilerIdCXX/CompilerIdHIP a.out, FindOpenMP ompver_*.bin) -- +plain host x86-64 ELF test programs with no HIP device code at all (confirmed +via `file` + `nm -D`: only libc/GCC weak symbols, identical between old and +new). These are not part of the aihwkit build output and never ship. The one +real artifact, `rpu_base.cpython-312-x86_64-linux-gnu.so`, reported +**identical** -- exported symbols + device ISA identical (53232 exports). + +GPU corroboration (HIP_VISIBLE_DEVICES=2, gfx90a wave64), rpu_base.so from the +head-sha build (build_new) copied into src/aihwkit/simulator/: +- Import smoke: `rpu_base.cuda.is_compiled()` True, `torch.cuda.get_device_name(0)` + = "AMD Instinct MI250X / MI250". +- tests/test_specific_tiles.py: **18/18 PASSED** (CRITICAL -- bit_line_maker + + pulsed-weight-update warp-size path). +- tests/test_simulator_tiles.py + tests/test_bindings_tiles.py `-k Cuda`: + **284 passed, 47 skipped, 0 failed**. +- tests/test_inference_tiles.py `-k Cuda`: **32 passed, 22 skipped, 0 failed**. + +CUDA no-regression gate: never previously recorded for this project at any +head_sha, so ran it here (compile-only, no NVIDIA GPU on host). Worked around +the two known environmental walls (ROCm-build torch/include missing +c10/cuda/impl/cuda_cmake_macros.h; torch/headeronly/util/complex.h duplicated- +token #if guard) by building against a scratch venv with pip-installed +`torch==2.11.0+cu128` (CUDA wheel). RPU_CUDA_ARCHITECTURES is already a fixed +cache-var list ("75;80;89", includes 80) in CMakeLists.txt -- not `native` +autodetect -- so no arch-pin patch was needed. +``` +python3 -m venv venv && source venv/bin/activate +pip install --index-url https://download.pytorch.org/whl/cu128 torch==2.11.0+cu128 +pip install numpy pybind11 scikit-build cmake +cd projects/aihwkit/src # at head_sha 70577b578b7faa661611e905c0b9ec1a1d85bf33 +cmake -S . -B build_cuda_head -GNinja -DUSE_CUDA=ON -DUSE_HIP=OFF -DRPU_CXX_STANDARD=17 \ + -DCMAKE_CUDA_COMPILER=/opt/conda/envs/cuda-12.8/bin/nvcc -DCMAKE_CUDA_ARCHITECTURES=80 \ + -DCMAKE_PREFIX_PATH="$TORCH_CMAKE" -DRPU_BLAS=OpenBLAS -DBUILD_TEST=OFF -DRPU_USE_TORCH_BUFFERS=OFF +cmake --build build_cuda_head -j16 +``` +Result: clean build, exit 0, produced rpu_base.cpython-312-x86_64-linux-gnu.so +(198 MB) with real CUDA sm_75/sm_80/sm_89 code objects for bit_line_maker.cu +and all RPU_GPU .cu sources. No `-fvisibility=hidden`-on-nvcc-style failure, +no `atomicAdd(double*)` overload failure (would have signaled a wrong/native +arch pin). CUDA path is a pure passthrough -- no regression. + +Jargon: `python3 utils/jargon.py --port aihwkit` -> clean. + +Verdict: PASS (binary-equivalence carry-forward, corroborated by real-GPU +smoke). Transitioning linux-gfx90a to completed (validated_sha 70577b5). diff --git a/projects/aihwkit/stats.jsonl b/projects/aihwkit/stats.jsonl index 64fb5b4d..f043b624 100644 --- a/projects/aihwkit/stats.jsonl +++ b/projects/aihwkit/stats.jsonl @@ -61,3 +61,9 @@ {"kind": "tokens", "ts": "2026-06-19T20:34:19Z", "tokens": 157196, "source": "validator (gfx1101)"} {"kind": "tokens", "ts": "2026-06-08T15:01:20Z", "tokens": 100929, "source": "porter"} {"kind": "tokens", "ts": "2026-06-23T22:08:43Z", "tokens": 38016, "source": "gfx90a audit cross-file fix (general-purpose subagent)"} +{"kind":"phase","ts":"2026-08-09T00:32:24Z","phase":"compile","seconds":135.757,"exit":0,"cmd":"cmake --build projects/aihwkit/src/build_old -j16"} +{"kind":"phase","ts":"2026-08-09T00:37:10Z","phase":"compile","seconds":136.898,"exit":0,"cmd":"cmake --build projects/aihwkit/src/build_new -j16"} +{"kind":"phase","ts":"2026-08-09T00:59:13Z","phase":"cuda-compile","seconds":660.559,"exit":0,"cmd":"cmake --build projects/aihwkit/src/build_cuda_head -j16"} +{"kind":"phase","ts":"2026-08-09T01:00:46Z","phase":"test","seconds":7.784,"exit":0,"cmd":"bash -c cd projects/aihwkit/src && HIP_VISIBLE_DEVICES=2 PYTHONPATH=src python -m pytest tests/test_specific_tiles.py -q 2>&1 | tail -40"} +{"kind":"phase","ts":"2026-08-09T01:01:46Z","phase":"test","seconds":55.814,"exit":0,"cmd":"bash -c cd projects/aihwkit/src && HIP_VISIBLE_DEVICES=2 PYTHONPATH=src python -m pytest tests/test_simulator_tiles.py tests/test_bindings_tiles.py -q -k Cuda 2>&1 | tail -20"} +{"kind":"phase","ts":"2026-08-09T01:02:01Z","phase":"test","seconds":9.320,"exit":0,"cmd":"bash -c cd projects/aihwkit/src && HIP_VISIBLE_DEVICES=2 PYTHONPATH=src python -m pytest tests/test_inference_tiles.py -q -k Cuda 2>&1 | tail -20"} diff --git a/projects/aihwkit/status.json b/projects/aihwkit/status.json index ff4fdbe9..b37ddc26 100644 --- a/projects/aihwkit/status.json +++ b/projects/aihwkit/status.json @@ -7,7 +7,7 @@ "priority": 0.0, "ext_type": "cmake", "adopted_at": "2026-06-04T01:19:43Z", - "updated_at": "2026-08-07T07:04:50Z", + "updated_at": "2026-08-09T01:02:52Z", "head_sha": "70577b578b7faa661611e905c0b9ec1a1d85bf33", "depends_on": [], "pr_url": "https://github.com/IBM/aihwkit/pull/770", @@ -24,10 +24,10 @@ "state": "completed", "blocked": false, "blocked_reason": null, - "validated_sha": "4bc09a0fa9a76d3038e6729ea03615fcc0beb76b", + "validated_sha": "70577b578b7faa661611e905c0b9ec1a1d85bf33", "started_at": "2026-06-04T01:28:01Z", - "completed_at": "2026-06-08T16:12:07Z", - "updated_at": "2026-06-22T23:22:30Z", + "completed_at": "2026-08-09T01:02:52Z", + "updated_at": "2026-08-09T01:02:52Z", "stats": { "tokens_total": 0, "tokens_approx": true, @@ -43,10 +43,10 @@ }, "last_agent": "validator", "carry_forward": { - "to": "4bc09a0fa9a76d3038e6729ea03615fcc0beb76b", - "method": "source-class", - "detail": "maintainer merged origin/master into PR #770; delta is setup.cfg mypy python_version, visualization.py plt.ylim style, and uv.lock only -- arch-independent, no C++/HIP/GPU code, behavior-preserving", - "at": "2026-06-22T23:22:30Z" + "to": "70577b578b7faa661611e905c0b9ec1a1d85bf33", + "method": "binary-equiv", + "detail": "CMakeLists.txt drops the gfx90a HIP_ARCHITECTURES default pin, letting enable_language(HIP) auto-detect when unset; our recipe always passes -DCMAKE_HIP_ARCHITECTURES=gfx90a explicitly so codegen is u", + "at": "2026-08-09T01:02:52Z" } }, "windows-gfx1201": { From bd93cea19453fb1582f1dffd9c3be01a1487377e Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Sun, 9 Aug 2026 21:41:19 +0000 Subject: [PATCH 3/4] aihwkit: commit the telemetry from this session --- projects/aihwkit/stats.jsonl | 1 + 1 file changed, 1 insertion(+) diff --git a/projects/aihwkit/stats.jsonl b/projects/aihwkit/stats.jsonl index f043b624..69ad58f9 100644 --- a/projects/aihwkit/stats.jsonl +++ b/projects/aihwkit/stats.jsonl @@ -67,3 +67,4 @@ {"kind":"phase","ts":"2026-08-09T01:00:46Z","phase":"test","seconds":7.784,"exit":0,"cmd":"bash -c cd projects/aihwkit/src && HIP_VISIBLE_DEVICES=2 PYTHONPATH=src python -m pytest tests/test_specific_tiles.py -q 2>&1 | tail -40"} {"kind":"phase","ts":"2026-08-09T01:01:46Z","phase":"test","seconds":55.814,"exit":0,"cmd":"bash -c cd projects/aihwkit/src && HIP_VISIBLE_DEVICES=2 PYTHONPATH=src python -m pytest tests/test_simulator_tiles.py tests/test_bindings_tiles.py -q -k Cuda 2>&1 | tail -20"} {"kind":"phase","ts":"2026-08-09T01:02:01Z","phase":"test","seconds":9.320,"exit":0,"cmd":"bash -c cd projects/aihwkit/src && HIP_VISIBLE_DEVICES=2 PYTHONPATH=src python -m pytest tests/test_inference_tiles.py -q -k Cuda 2>&1 | tail -20"} +{"kind": "tokens", "ts": "2026-08-09T14:30:54Z", "tokens": 114846, "source": "validator"} From 393f7403ea7510523c8992018c2315e73666bdf9 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Wed, 12 Aug 2026 22:11:02 +0000 Subject: [PATCH 4/4] aihwkit: published_sha backfilled -- the open PR shows 70577b578b7f (verified against the live PR head) --- projects/aihwkit/status.json | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/projects/aihwkit/status.json b/projects/aihwkit/status.json index b37ddc26..c6d0ca54 100644 --- a/projects/aihwkit/status.json +++ b/projects/aihwkit/status.json @@ -7,7 +7,7 @@ "priority": 0.0, "ext_type": "cmake", "adopted_at": "2026-06-04T01:19:43Z", - "updated_at": "2026-08-09T01:02:52Z", + "updated_at": "2026-08-12T22:11:02Z", "head_sha": "70577b578b7faa661611e905c0b9ec1a1d85bf33", "depends_on": [], "pr_url": "https://github.com/IBM/aihwkit/pull/770", @@ -135,5 +135,6 @@ "at": "2026-06-23T22:08:02Z" } } - } + }, + "published_sha": "70577b578b7faa661611e905c0b9ec1a1d85bf33" }