From 10135adaeea9bdb41da0a0d9860a6e8bc511cfe7 Mon Sep 17 00:00:00 2001 From: john bowen Date: Mon, 12 Oct 2026 16:06:47 -0700 Subject: [PATCH 1/7] Add pseudo-GPU device compilation harness for pass verification Drive the real HIP device pipeline (clang -x hip targeting amdgcn) on machines without ROCm or a GPU. A minimal hip/hip_runtime.h stub plus -nogpuinc -nogpulib --offload-device-only lets unmodified GPU tests compile with ProteusPass loaded, so lambda kernel-arg and clobbering analysis decisions are verifiable ahead of time from the PROTEUS_DEBUG_OUTPUT pass trace. Verified against Homebrew LLVM 19 and 22: lambda_clobber_shapes reproduces all 12 expected decisions. Known quirk handled by the driver: upstream clang registers -fpass-plugin pipeline callbacks twice for HIP device compiles (ROCm clang does not), so the second ProteusPass run dies on the duplicate _jit_bitcode global after analysis completes; the script treats exactly that fatal as benign. Co-Authored-By: Claude Fable 5 --- tools/pseudo-gpu/compile-device.sh | 131 +++++++++++++++++++++ tools/pseudo-gpu/include/hip/hip_runtime.h | 124 +++++++++++++++++++ 2 files changed, 255 insertions(+) create mode 100755 tools/pseudo-gpu/compile-device.sh create mode 100644 tools/pseudo-gpu/include/hip/hip_runtime.h diff --git a/tools/pseudo-gpu/compile-device.sh b/tools/pseudo-gpu/compile-device.sh new file mode 100755 index 00000000..1f5cb470 --- /dev/null +++ b/tools/pseudo-gpu/compile-device.sh @@ -0,0 +1,131 @@ +#!/usr/bin/env bash +# Pseudo-GPU device compilation harness. +# +# Runs the real HIP *device* compilation pipeline (clang -x hip, amdgcn +# target, ProteusPass as a pass plugin) on a machine without ROCm or a GPU, +# so lambda kernel-arg / clobbering analysis changes can be verified at +# compile time from the pass trace (PROTEUS_DEBUG_OUTPUT). +# +# Usage: +# tools/pseudo-gpu/compile-device.sh [options] [-- ] +# +# Options: +# --llvm LLVM install prefix (default: /opt/homebrew/opt/llvm@19; +# the pass plugin is picked from build-pseudogpu-llvm) +# --build Build dir containing src/pass/libProteusPass.dylib +# --raw-ir Also emit pre-pass device IR (no plugin) for shape study +# --log Copy the pass trace here (default: .proteus-pass.log) +# -O Optimization level (default -O1) +# +# Exit status: 0 if compilation reached the end of the Proteus analysis. +# NOTE: upstream clang (unlike ROCm clang) registers -fpass-plugin pipeline +# callbacks twice for HIP device compiles, so ProteusPass runs twice and the +# second run hits the (expected-unique) _jit_bitcode global and calls +# reportFatalError. That fatal happens after the analysis and manifest are +# complete, so this script treats exactly that error as benign tail noise. +# Every analysis result appears twice in the trace for the same reason. + +set -u + +REPO_ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")/../.." && pwd)" +LLVM_PREFIX=/opt/homebrew/opt/llvm@19 +BUILD_DIR= +RAW_IR= +LOG_OUT= +OPT_LEVEL=-O1 +SOURCE= +EXTRA_ARGS=() + +while [[ $# -gt 0 ]]; do + case "$1" in + --llvm) LLVM_PREFIX="$2"; shift 2 ;; + --build) BUILD_DIR="$2"; shift 2 ;; + --raw-ir) RAW_IR="$2"; shift 2 ;; + --log) LOG_OUT="$2"; shift 2 ;; + -O*) OPT_LEVEL="$1"; shift ;; + --) shift; EXTRA_ARGS=("$@"); break ;; + *) SOURCE="$1"; shift ;; + esac +done + +if [[ -z "$SOURCE" ]]; then + echo "error: no source file given" >&2 + exit 2 +fi + +# The plugin compile runs inside a temp dir (to contain .proteus-logs), so +# user-supplied paths must survive the cd. +case "$SOURCE" in /*) ;; *) SOURCE="$PWD/$SOURCE" ;; esac +if [[ -n "$RAW_IR" ]]; then case "$RAW_IR" in /*) ;; *) RAW_IR="$PWD/$RAW_IR" ;; esac; fi +if [[ -n "$LOG_OUT" ]]; then case "$LOG_OUT" in /*) ;; *) LOG_OUT="$PWD/$LOG_OUT" ;; esac; fi + +if [[ -z "$BUILD_DIR" ]]; then + VER=$("$LLVM_PREFIX/bin/llvm-config" --version | cut -d. -f1) + BUILD_DIR="$REPO_ROOT/build-pseudogpu-llvm$VER" +fi +PLUGIN="$BUILD_DIR/src/pass/libProteusPass.dylib" +if [[ ! -f "$PLUGIN" ]]; then + echo "error: $PLUGIN not found; configure with:" >&2 + echo " cmake -S $REPO_ROOT -B $BUILD_DIR -DLLVM_INSTALL_DIR=$LLVM_PREFIX \\" >&2 + echo " -DCMAKE_C_COMPILER=$LLVM_PREFIX/bin/clang -DCMAKE_CXX_COMPILER=$LLVM_PREFIX/bin/clang++ \\" >&2 + echo " -DCMAKE_CXX_FLAGS=-DPROTEUS_ENABLE_HIP=1" >&2 + echo " cmake --build $BUILD_DIR --target ProteusPass -j 8" >&2 + exit 2 +fi + +COMMON_ARGS=( + -x hip --offload-arch=gfx90a --offload-device-only + -nogpulib -nogpuinc -std=c++17 "$OPT_LEVEL" + -DPROTEUS_ENABLE_HIP=1 + -I "$REPO_ROOT/tools/pseudo-gpu/include" + -I "$REPO_ROOT/include" +) + +WORK_DIR=$(mktemp -d /tmp/pseudo-gpu.XXXXXX) +trap 'rm -rf "$WORK_DIR"' EXIT + +if [[ -n "$RAW_IR" ]]; then + "$LLVM_PREFIX/bin/clang++" "${COMMON_ARGS[@]}" ${EXTRA_ARGS[@]+"${EXTRA_ARGS[@]}"} \ + -Xclang -disable-llvm-passes -S -emit-llvm -o "$RAW_IR" "$SOURCE" || + { echo "error: raw-IR compile failed" >&2; exit 1; } + echo "raw device IR: $RAW_IR" +fi + +STDERR_FILE="$WORK_DIR/stderr.txt" +(cd "$WORK_DIR" && PROTEUS_DEBUG_OUTPUT=1 "$LLVM_PREFIX/bin/clang++" \ + "${COMMON_ARGS[@]}" ${EXTRA_ARGS[@]+"${EXTRA_ARGS[@]}"} \ + -fpass-plugin="$PLUGIN" \ + -S -emit-llvm -o "$WORK_DIR/device.post.ll" "$SOURCE") 2> "$STDERR_FILE" +STATUS=$? + +LOG_FILE=$(ls "$WORK_DIR"/.proteus-logs/proteus-pass.*.log 2>/dev/null | head -1) +if [[ -z "$LOG_FILE" ]]; then + echo "error: no proteus-pass trace produced; compile failed before the pass:" >&2 + cat "$STDERR_FILE" >&2 + exit 1 +fi + +if [[ -z "$LOG_OUT" ]]; then + LOG_OUT="$(basename "${SOURCE%.cpp}").proteus-pass.log" +fi +cp "$LOG_FILE" "$LOG_OUT" + +if [[ $STATUS -ne 0 ]]; then + if ! grep -q "Expected unique name for jit module global variable" "$STDERR_FILE"; then + echo "error: compile failed with an unexpected error:" >&2 + cat "$STDERR_FILE" >&2 + exit 1 + fi + echo "note: known benign double-run fatal (_jit_bitcode) ignored;" \ + "results below are from the completed first pass run" +fi + +echo "pass trace: $LOG_OUT" +echo +echo "== Lambda kernel-arg analysis decisions (first run) ==" +# The trace repeats once per pass run; report each unique line once. +grep -E "Found termination case from function|Kernel arg analysis failed for functor" \ + "$LOG_OUT" | sort -u | sed \ + -e 's/.*Found termination case from function /SPECIALIZE kernel=/' \ + -e 's/.*Kernel arg analysis failed for functor beginning at */DECLINE /' +exit 0 diff --git a/tools/pseudo-gpu/include/hip/hip_runtime.h b/tools/pseudo-gpu/include/hip/hip_runtime.h new file mode 100644 index 00000000..26049f53 --- /dev/null +++ b/tools/pseudo-gpu/include/hip/hip_runtime.h @@ -0,0 +1,124 @@ +#ifndef PROTEUS_PSEUDO_GPU_HIP_RUNTIME_H +#define PROTEUS_PSEUDO_GPU_HIP_RUNTIME_H + +// Minimal HIP API surface for "pseudo-GPU" compilation on machines without +// ROCm (e.g. macOS dev boxes). This supports compile-time verification of +// ProteusPass analyses by driving the real HIP device pipeline: +// +// clang++ -x hip --offload-arch=gfx90a --offload-device-only \ +// -nogpulib -nogpuinc -I tools/pseudo-gpu/include ... \ +// -fpass-plugin=libProteusPass.dylib -S -emit-llvm +// +// The pass runs on the amdgcn device module exactly as in a production HIP +// build; PROTEUS_DEBUG_OUTPUT=1 prints the analysis trace at compile time. +// Nothing here is intended for execution or linking. + +#include +#include + +// Target attribute macros normally supplied by HIP's clang wrapper headers +// (skipped by -nogpuinc). +#ifndef __global__ +#define __host__ __attribute__((host)) +#define __device__ __attribute__((device)) +#define __global__ __attribute__((global)) +#define __shared__ __attribute__((shared)) +#define __constant__ __attribute__((constant)) +#define __managed__ __attribute__((managed)) +#endif + +#ifndef __forceinline__ +#define __forceinline__ inline __attribute__((always_inline)) +#endif +#define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__))) + +struct dim3 { + unsigned x, y, z; + constexpr __host__ __device__ dim3(unsigned X = 1, unsigned Y = 1, + unsigned Z = 1) + : x(X), y(Y), z(Z) {} +}; + +typedef enum hipError_t { + hipSuccess = 0, + hipErrorUnknown = 999, +} hipError_t; + +typedef struct ihipStream_t *hipStream_t; + +typedef enum hipMemcpyKind { + hipMemcpyHostToHost = 0, + hipMemcpyHostToDevice = 1, + hipMemcpyDeviceToHost = 2, + hipMemcpyDeviceToDevice = 3, + hipMemcpyDefault = 4, +} hipMemcpyKind; + +extern "C" { +hipError_t hipLaunchKernel(const void *Func, dim3 Grid, dim3 Block, + void **Args, size_t SharedMem, hipStream_t Stream); +hipError_t hipDeviceSynchronize(void); +const char *hipGetErrorString(hipError_t Err); +hipError_t hipMalloc(void **Ptr, size_t Size); +hipError_t hipMallocManaged(void **Ptr, size_t Size, unsigned Flags); +hipError_t hipFree(void *Ptr); +hipError_t hipMemcpy(void *Dst, const void *Src, size_t Size, + hipMemcpyKind Kind); +hipError_t hipMemcpyFromSymbol(void *Dst, const void *Symbol, size_t Size, + size_t Offset, hipMemcpyKind Kind); +hipError_t hipGetSymbolAddress(void **DevPtr, const void *Symbol); +hipError_t hipStreamCreate(hipStream_t *Stream); +hipError_t hipStreamSynchronize(hipStream_t Stream); +hipError_t hipStreamDestroy(hipStream_t Stream); + +// Triple-chevron <<<...>>> launch support expected by CodeGen. +hipError_t __hipPushCallConfiguration(dim3 Grid, dim3 Block, + size_t SharedMem = 0, + hipStream_t Stream = nullptr); +hipError_t __hipPopCallConfiguration(dim3 *Grid, dim3 *Block, + size_t *SharedMem, hipStream_t *Stream); +} // extern "C" + +template +static inline hipError_t hipMallocManaged(T **Ptr, size_t Size, + unsigned Flags = 1) { + return hipMallocManaged(reinterpret_cast(Ptr), Size, Flags); +} + +template +static inline hipError_t hipMalloc(T **Ptr, size_t Size) { + return hipMalloc(reinterpret_cast(Ptr), Size); +} + +static inline hipError_t hipMemcpyFromSymbol(void *Dst, const void *Symbol, + size_t Size) { + return hipMemcpyFromSymbol(Dst, Symbol, Size, 0, hipMemcpyDeviceToHost); +} + +#if defined(__HIP_DEVICE_COMPILE__) +// Declaration only; AMDGPU lowers printf during codegen, which the pseudo-GPU +// pipeline never reaches (-emit-llvm). +extern "C" __device__ int printf(const char *Fmt, ...); + +// Work-item builtins give the same IR shapes as the real HIP headers. +#define threadIdx \ + (dim3{(unsigned)__builtin_amdgcn_workitem_id_x(), \ + (unsigned)__builtin_amdgcn_workitem_id_y(), \ + (unsigned)__builtin_amdgcn_workitem_id_z()}) +#define blockIdx \ + (dim3{(unsigned)__builtin_amdgcn_workgroup_id_x(), \ + (unsigned)__builtin_amdgcn_workgroup_id_y(), \ + (unsigned)__builtin_amdgcn_workgroup_id_z()}) +// Declaration-only accessors: resolve to extern device functions in IR, which +// is sufficient for pass-level analysis (never linked or executed). +extern "C" __device__ unsigned __pseudo_gpu_block_dim(int Axis); +extern "C" __device__ unsigned __pseudo_gpu_grid_dim(int Axis); +#define blockDim \ + (dim3{__pseudo_gpu_block_dim(0), __pseudo_gpu_block_dim(1), \ + __pseudo_gpu_block_dim(2)}) +#define gridDim \ + (dim3{__pseudo_gpu_grid_dim(0), __pseudo_gpu_grid_dim(1), \ + __pseudo_gpu_grid_dim(2)}) +#endif // __HIP_DEVICE_COMPILE__ + +#endif // PROTEUS_PSEUDO_GPU_HIP_RUNTIME_H From ad23c206329185106839a08d2ad8c4767b6dd226 Mon Sep 17 00:00:00 2001 From: john bowen Date: Tue, 13 Oct 2026 10:54:03 -0700 Subject: [PATCH 2/7] Add PROTEUS_PSEUDO_GPU build mode with per-test device-compile targets cmake -DPROTEUS_PSEUDO_GPU=ON (GPU toolchains off) adds a .HIP target per GPU test that drives the real HIP device pipeline -- clang -x hip targeting amdgcn with ProteusPass loaded -- and emits the post-pass LLVM module to pseudo-gpu/.device.ll, requiring no ROCm or GPU. The pass objects are built HIP-flavored and with PROTEUS_PSEUDO_GPU, which downgrades the duplicate _jit_bitcode fatal to a debug-logged skip: upstream clang registers pass-plugin pipeline callbacks twice for HIP device compiles, so the pass legitimately re-runs there; production builds keep the fatal. Includes the stub HIP header and standalone driver script (tools/pseudo-gpu) previously developed on feature/bowen/pseudo-gpu-harness, now with forced stub inclusion and a hipLaunchKernelGGL macro so all registered lambda tests compile. --- CMakeLists.txt | 20 ++++++++ src/pass/ProteusPass.cpp | 14 +++++- tools/pseudo-gpu/PseudoGPUTargets.cmake | 56 ++++++++++++++++++++++ tools/pseudo-gpu/compile-device.sh | 3 ++ tools/pseudo-gpu/include/hip/hip_runtime.h | 3 ++ 5 files changed, 95 insertions(+), 1 deletion(-) create mode 100644 tools/pseudo-gpu/PseudoGPUTargets.cmake diff --git a/CMakeLists.txt b/CMakeLists.txt index b0562a47..c14d1427 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -23,6 +23,10 @@ option(ENABLE_DEVELOPER_COMPILER_FLAGS "Enable developer compiler flags: -Wall - option(PROTEUS_INSTALL_IMPL_HEADERS "Install implementation headers" OFF) option(PROTEUS_ENABLE_PYTHON "Build Python bindings" OFF) option(PROTEUS_PYTHON_WHEEL "Build wheel-oriented Python package layout" OFF) +option(PROTEUS_PSEUDO_GPU + "Add per-test pseudo-GPU device-compile targets (.HIP) that run the \ +real HIP device pipeline through ProteusPass and emit LLVM IR, without \ +ROCm or a GPU" OFF) set(PROTEUS_PYTHON_PACKAGE_DIR "proteus" CACHE STRING "Install-relative path for the Proteus Python package") @@ -111,8 +115,24 @@ include(cmake/ProteusFunctions.cmake) include(GNUInstallDirs) include(CMakePackageConfigHelpers) +if(PROTEUS_PSEUDO_GPU AND (PROTEUS_ENABLE_HIP OR PROTEUS_ENABLE_CUDA)) + message(FATAL_ERROR "PROTEUS_PSEUDO_GPU is for builds without a GPU " + "toolchain; disable PROTEUS_ENABLE_HIP/CUDA") +endif() + add_subdirectory(src) +if(PROTEUS_PSEUDO_GPU) + # Build the pass HIP-flavored so isDeviceKernel recognizes amdgpu_kernel + # functions without requiring the HIP toolchain, and let it tolerate + # upstream clang re-running it on the device module (see emitModuleDevice); + # production builds keep the unique-name fatal. Scoped to the pass objects + # so the runtime library is unaffected. + target_compile_definitions(ProteusPassObjects PRIVATE PROTEUS_ENABLE_HIP=1 + PROTEUS_PSEUDO_GPU=1) + include(tools/pseudo-gpu/PseudoGPUTargets.cmake) +endif() + if(ENABLE_TESTS) enable_testing() endif() diff --git a/src/pass/ProteusPass.cpp b/src/pass/ProteusPass.cpp index 57af8f1e..87711f33 100644 --- a/src/pass/ProteusPass.cpp +++ b/src/pass/ProteusPass.cpp @@ -976,9 +976,21 @@ class ProteusPassImpl { Constant *JitModule = ConstantDataArray::get( M.getContext(), ArrayRef((const uint8_t *)Bitcode.data(), Bitcode.size())); - if (M.getNamedGlobal(GVName)) + if (M.getNamedGlobal(GVName)) { +#ifdef PROTEUS_PSEUDO_GPU + // Upstream clang's HIP device pipeline registers pass-plugin pipeline + // callbacks twice (ROCm clang runs the pass once), so pseudo-GPU + // compiles re-run this pass on an already-processed module. The first + // embedding is authoritative; skip the duplicate so the device compile + // can emit its module. + DEBUG(Logger::logs("proteus-pass") + << "Skipping duplicate jit module global " << GVName << "\n"); + return; +#else reportFatalError("Expected unique name for jit module global variable " + GVName); +#endif + } auto *GV = new GlobalVariable(M, JitModule->getType(), /* isConstant */ true, GlobalValue::ExternalLinkage, JitModule, GVName); diff --git a/tools/pseudo-gpu/PseudoGPUTargets.cmake b/tools/pseudo-gpu/PseudoGPUTargets.cmake new file mode 100644 index 00000000..74cee43c --- /dev/null +++ b/tools/pseudo-gpu/PseudoGPUTargets.cmake @@ -0,0 +1,56 @@ +# Pseudo-GPU device-compile targets. +# +# For every GPU test, adds a `.HIP` target that runs the real HIP +# device pipeline (clang -x hip targeting amdgcn, ProteusPass loaded) and +# emits the post-pass LLVM module to pseudo-gpu/.device.ll -- the +# equivalent of -S -emit-llvm, with no ROCm or GPU required. Executables and +# runtime behavior are out of scope; use these targets to inspect pass +# decisions (with PROTEUS_DEBUG_OUTPUT=1) and IR shapes ahead of time. +# +# cmake -B build -DPROTEUS_PSEUDO_GPU=ON -DLLVM_INSTALL_DIR=... \ +# -DCMAKE_C_COMPILER=.../clang -DCMAKE_CXX_COMPILER=.../clang++ +# make -C build lambda_clobber_shapes.HIP +# make -C build pseudo-gpu-all + +set(PSEUDO_GPU_CLANG "${LLVM_INSTALL_DIR}/bin/clang++") +if(NOT EXISTS "${PSEUDO_GPU_CLANG}") + message(FATAL_ERROR "PROTEUS_PSEUDO_GPU requires clang++ in LLVM_INSTALL_DIR" + " (${PSEUDO_GPU_CLANG} not found)") +endif() + +set(PSEUDO_GPU_OUT_DIR "${CMAKE_BINARY_DIR}/pseudo-gpu") +file(MAKE_DIRECTORY "${PSEUDO_GPU_OUT_DIR}") + +set(PSEUDO_GPU_STUB_HEADER + "${PROJECT_SOURCE_DIR}/tools/pseudo-gpu/include/hip/hip_runtime.h") + +set(PSEUDO_GPU_FLAGS + -x hip --offload-arch=gfx90a --offload-device-only + -nogpulib -nogpuinc -std=c++17 -O1 + -DPROTEUS_ENABLE_HIP=1 + -I "${PROJECT_SOURCE_DIR}/tools/pseudo-gpu/include" + -I "${PROJECT_SOURCE_DIR}/include" + # Real HIP force-includes clang's runtime wrapper before user code; mirror + # that so target attribute macros exist regardless of include order. + -include hip/hip_runtime.h) + +file(GLOB PSEUDO_GPU_TEST_SOURCES "${PROJECT_SOURCE_DIR}/tests/gpu/*.cpp") + +set(PSEUDO_GPU_ALL_OUTPUTS) +foreach(PseudoSource IN LISTS PSEUDO_GPU_TEST_SOURCES) + get_filename_component(PseudoTestName "${PseudoSource}" NAME_WE) + set(PseudoOutput "${PSEUDO_GPU_OUT_DIR}/${PseudoTestName}.device.ll") + add_custom_command( + OUTPUT "${PseudoOutput}" + COMMAND "${PSEUDO_GPU_CLANG}" ${PSEUDO_GPU_FLAGS} + -fpass-plugin=$ + -S -emit-llvm -o "${PseudoOutput}" "${PseudoSource}" + DEPENDS "${PseudoSource}" ProteusPass "${PSEUDO_GPU_STUB_HEADER}" + WORKING_DIRECTORY "${PSEUDO_GPU_OUT_DIR}" + COMMENT "Pseudo-GPU device compile: ${PseudoTestName}" + VERBATIM) + add_custom_target("${PseudoTestName}.HIP" DEPENDS "${PseudoOutput}") + list(APPEND PSEUDO_GPU_ALL_OUTPUTS "${PseudoOutput}") +endforeach() + +add_custom_target(pseudo-gpu-all DEPENDS ${PSEUDO_GPU_ALL_OUTPUTS}) diff --git a/tools/pseudo-gpu/compile-device.sh b/tools/pseudo-gpu/compile-device.sh index 1f5cb470..1f522eab 100755 --- a/tools/pseudo-gpu/compile-device.sh +++ b/tools/pseudo-gpu/compile-device.sh @@ -79,6 +79,9 @@ COMMON_ARGS=( -DPROTEUS_ENABLE_HIP=1 -I "$REPO_ROOT/tools/pseudo-gpu/include" -I "$REPO_ROOT/include" + # Real HIP force-includes clang's runtime wrapper before user code; mirror + # that so target attribute macros exist regardless of include order. + -include hip/hip_runtime.h ) WORK_DIR=$(mktemp -d /tmp/pseudo-gpu.XXXXXX) diff --git a/tools/pseudo-gpu/include/hip/hip_runtime.h b/tools/pseudo-gpu/include/hip/hip_runtime.h index 26049f53..98830d00 100644 --- a/tools/pseudo-gpu/include/hip/hip_runtime.h +++ b/tools/pseudo-gpu/include/hip/hip_runtime.h @@ -79,6 +79,9 @@ hipError_t __hipPopCallConfiguration(dim3 *Grid, dim3 *Block, size_t *SharedMem, hipStream_t *Stream); } // extern "C" +#define hipLaunchKernelGGL(Kernel, Grid, Block, SharedMem, Stream, ...) \ + (Kernel)<<<(Grid), (Block), (SharedMem), (Stream)>>>(__VA_ARGS__) + template static inline hipError_t hipMallocManaged(T **Ptr, size_t Size, unsigned Flags = 1) { From 9539e8a297fbaca803cbe4fb2c87f8b638e931ec Mon Sep 17 00:00:00 2001 From: john bowen Date: Thu, 15 Oct 2026 11:29:01 -0700 Subject: [PATCH 3/7] Add specialization ground truth for the lambda GPU suite tests/gpu/lambda-spec-groundtruth.json records per-test specialize/decline decisions parsed from the ProteusPass debug trace; tools/lambda-spec-ci/lambda_spec_ci.py captures, checks, and renders it. The trace lines are identical for ROCm and pseudo-GPU runs, so one truth file serves both; CI can fail on any decision drift. Initial capture is from the pseudo-GPU pipeline and is marked as requiring ROCm re-capture before being treated as hardware ground truth. --- docs/dev/lambda-spec-groundtruth.md | 86 ++++++++ tests/gpu/lambda-spec-groundtruth.json | 292 +++++++++++++++++++++++++ tools/lambda-spec-ci/lambda_spec_ci.py | 179 +++++++++++++++ 3 files changed, 557 insertions(+) create mode 100644 docs/dev/lambda-spec-groundtruth.md create mode 100644 tests/gpu/lambda-spec-groundtruth.json create mode 100644 tools/lambda-spec-ci/lambda_spec_ci.py diff --git a/docs/dev/lambda-spec-groundtruth.md b/docs/dev/lambda-spec-groundtruth.md new file mode 100644 index 00000000..7507cef3 --- /dev/null +++ b/docs/dev/lambda-spec-groundtruth.md @@ -0,0 +1,86 @@ +# Lambda GPU suite: specialization ground truth + +`tests/gpu/lambda-spec-groundtruth.json` records, for every lambda GPU test, +which kernels the analysis specializes and which functors it conservatively +declines. It is the reviewable source of truth for analysis behavior: any +change to the pass that shifts a decision must either be a bug fix (update +the JSON in the same commit, explaining the flip) or is a regression (CI +fails). + +## Running the check + +Logs come from the ProteusPass debug trace, which is identical for real ROCm +builds and the pseudo-GPU pipeline: + + # 1. produce one log per test (any build that runs the pass), e.g.: + PROTEUS_DEBUG_OUTPUT=1 ctest -R 'lambda_.*HIP' ... # ROCm + # or compile each test --offload-device-only with the pass plugin + # (pseudo-GPU), then collect .proteus-logs/proteus-pass.*.log into + # /.log + + # 2. diff against ground truth: + python3 tools/lambda-spec-ci/lambda_spec_ci.py check \ + --logs --truth tests/gpu/lambda-spec-groundtruth.json + + # re-capture (after an intended behavior change): + python3 tools/lambda-spec-ci/lambda_spec_ci.py capture \ + --logs --truth tests/gpu/lambda-spec-groundtruth.json \ + --pipeline --note "" + +## Provenance caveat + +The current JSON was captured from the pseudo-GPU pipeline (LLVM 19, -O1, +gfx90a) at the branch tip. Compile-time decisions have matched ROCm in spot +checks, but this file MUST be re-captured from a ROCm run before being +treated as hardware ground truth -- pseudo-GPU "green" has disagreed with +ROCm before. Tests with no entry produce no analysis decisions in this +pipeline: lambda_dominating_store (ill-formed, unregistered in CMake) and +lambda_large_warning (needs __syncthreads in the pseudo-GPU stub). + +Known expected-fail pinned in the suite: lambda_interior_slot_overwrite +currently SPECIALIZES (type-erased interior write; soundness delegated to +the frontend contract -- see docs/dev/design/lambda-closure-provenance.md). +Its runtime FileCheck fails by design until the AST diagnostic lands. + +## Current decisions + +| Test | Specializes | Declines | +| --- | --- | --- | +| lambda | kernel | — | +| lambda-two | kernel | — | +| lambda_aggregate_store_reconciliation | — | makeBody | +| lambda_call_shapes | kernelSelectSame | makeIndirectBody, makeReturnBody, makeSameReturnBody, makeSelectSameBody | +| lambda_clobber_query_range | kernelUnrelatedPrefixOverwrite | — | +| lambda_clobber_query_suffix | kernelUnrelatedSuffixOverwrite | — | +| lambda_clobber_shapes | kernelCallOverwrite, kernelLocalSupersedesCall, kernelNestedCallOverwrite, kernelNonoverlapOverwrite, kernelPredecessorOverwrite | makeAmbiguousLoopBody, makeDifferentBranchBody, makeDifferentCallBranchBody, makeSameBranchBody, makeSameCallBranchBody, makeSameLoopBody, makeUnrelatedCallBody | +| lambda_declaration_clobber | — | makeBody | +| lambda_def | kernel | — | +| lambda_factory | kernel | — | +| lambda_factory_forwarded | kernel | — | +| lambda_factory_multi | — | — | +| lambda_factory_so_runner | — | — | +| lambda_factory_tu1 | _GLOBAL__N_1 | — | +| lambda_factory_tu2 | _GLOBAL__N_1 | — | +| lambda_host_device | kernel | — | +| lambda_index_first_wrapper | kernel | — | +| lambda_insertvalue_offsets | kernelDoubleShift, kernelPadded, kernelShift | — | +| lambda_inst_visitor_call_clobber | kernelInstVisitorCallClobber, kernelInstVisitorCallDoesntClobber, kernelInstVisitorOrderedMemcpys | makeUntraceableBody | +| lambda_interior_slot_overwrite | — | makeInteriorBody | +| lambda_memory_intrinsic_shapes | kernelNonoverlapMemset | makeDynamicBody, makeMoveBody | +| lambda_mfem_style_launch | MockRajaInterface | — | +| lambda_mixed_field_overwrite | — | makeBody | +| lambda_multiple | kernel | — | +| lambda_multiple_callers | — | makeBody | +| lambda_pathological | kernel_case1, kernel_case2_a, kernel_case2_b, kernel_case4, kernel_case5 | — | +| lambda_phi_analysis | kernel_phi, kernel_phi_arg1, kernel_phi_changed_rc | — | +| lambda_pointer_field_load | kernelPointerFieldLoad | — | +| lambda_portability_aggregate_shapes | kernelAggregateFieldCall, kernelCallContext, kernelNestedReturn, kernelReturnedAlias, kernelReturnedFieldAlias | — | +| lambda_provenance_shapes | kernelArrayMember, kernelMemcpySource, kernelNegativeGep, kernelOffsetMemcpy, kernelPartialMemcpy, kernelReturnedInteriorPointer | — | +| lambda_reconciliation_adversarial | — | makeBytePair, makeConditionalPair, makeDuplicatePair, makeHeterogeneousPair, makePermutedPair, makeTriple | +| lambda_size_test | kernelEnclosedFunctor, kernelEnvelopeSuffix | — | +| lambda_slot_selection_supported | kernelCallAfterInvoke, kernelPhiSlot, kernelSelectedSlot, kernelTrailingReadOnlyCall | — | +| lambda_spec_test | kernel | — | +| lambda_store_order | kernelCaptureFromArg, kernelLatestStore, kernelStoreOrder | makeAmbiguousBody, makeMutatedBody | +| lambda_thread_privatizer | kernel_thread_privatizer | — | +| lambda_unknown_clobber | — | makeAggregateHiddenBody, makeAtomicBody, makeCompareExchangeBody, makeIndirectCallBody, makeOpaqueCallBody, makeSelectedStoreBody | +| lambda_unknown_clobber_decl | — | — | diff --git a/tests/gpu/lambda-spec-groundtruth.json b/tests/gpu/lambda-spec-groundtruth.json new file mode 100644 index 00000000..45ffa37c --- /dev/null +++ b/tests/gpu/lambda-spec-groundtruth.json @@ -0,0 +1,292 @@ +{ + "meta": { + "generated": "2026-10-15", + "commit": "34051bce", + "pipeline": "pseudo-gpu-llvm19", + "note": "Baseline at branch tip (pre affine-delta fix). Captured from the pseudo-GPU device pipeline; MUST be re-captured from a ROCm run before being treated as hardware ground truth." + }, + "tests": { + "lambda-two": { + "specialize": [ + "kernel" + ], + "decline": [] + }, + "lambda": { + "specialize": [ + "kernel" + ], + "decline": [] + }, + "lambda_aggregate_store_reconciliation": { + "specialize": [], + "decline": [ + "makeBody" + ] + }, + "lambda_call_shapes": { + "specialize": [ + "kernelSelectSame" + ], + "decline": [ + "makeIndirectBody", + "makeReturnBody", + "makeSameReturnBody", + "makeSelectSameBody" + ] + }, + "lambda_clobber_query_range": { + "specialize": [ + "kernelUnrelatedPrefixOverwrite" + ], + "decline": [] + }, + "lambda_clobber_query_suffix": { + "specialize": [ + "kernelUnrelatedSuffixOverwrite" + ], + "decline": [] + }, + "lambda_clobber_shapes": { + "specialize": [ + "kernelCallOverwrite", + "kernelLocalSupersedesCall", + "kernelNestedCallOverwrite", + "kernelNonoverlapOverwrite", + "kernelPredecessorOverwrite" + ], + "decline": [ + "makeAmbiguousLoopBody", + "makeDifferentBranchBody", + "makeDifferentCallBranchBody", + "makeSameBranchBody", + "makeSameCallBranchBody", + "makeSameLoopBody", + "makeUnrelatedCallBody" + ] + }, + "lambda_declaration_clobber": { + "specialize": [], + "decline": [ + "makeBody" + ] + }, + "lambda_def": { + "specialize": [ + "kernel" + ], + "decline": [] + }, + "lambda_factory": { + "specialize": [ + "kernel" + ], + "decline": [] + }, + "lambda_factory_forwarded": { + "specialize": [ + "kernel" + ], + "decline": [] + }, + "lambda_factory_multi": { + "specialize": [], + "decline": [] + }, + "lambda_factory_so_runner": { + "specialize": [], + "decline": [] + }, + "lambda_factory_tu1": { + "specialize": [ + "_GLOBAL__N_1" + ], + "decline": [] + }, + "lambda_factory_tu2": { + "specialize": [ + "_GLOBAL__N_1" + ], + "decline": [] + }, + "lambda_host_device": { + "specialize": [ + "kernel" + ], + "decline": [] + }, + "lambda_index_first_wrapper": { + "specialize": [ + "kernel" + ], + "decline": [] + }, + "lambda_insertvalue_offsets": { + "specialize": [ + "kernelDoubleShift", + "kernelPadded", + "kernelShift" + ], + "decline": [] + }, + "lambda_inst_visitor_call_clobber": { + "specialize": [ + "kernelInstVisitorCallClobber", + "kernelInstVisitorCallDoesntClobber", + "kernelInstVisitorOrderedMemcpys" + ], + "decline": [ + "makeUntraceableBody" + ] + }, + "lambda_interior_slot_overwrite": { + "specialize": [], + "decline": [ + "makeInteriorBody" + ] + }, + "lambda_memory_intrinsic_shapes": { + "specialize": [ + "kernelNonoverlapMemset" + ], + "decline": [ + "makeDynamicBody", + "makeMoveBody" + ] + }, + "lambda_mfem_style_launch": { + "specialize": [ + "MockRajaInterface" + ], + "decline": [] + }, + "lambda_mixed_field_overwrite": { + "specialize": [], + "decline": [ + "makeBody" + ] + }, + "lambda_multiple": { + "specialize": [ + "kernel" + ], + "decline": [] + }, + "lambda_multiple_callers": { + "specialize": [], + "decline": [ + "makeBody" + ] + }, + "lambda_pathological": { + "specialize": [ + "kernel_case1", + "kernel_case2_a", + "kernel_case2_b", + "kernel_case4", + "kernel_case5" + ], + "decline": [] + }, + "lambda_phi_analysis": { + "specialize": [ + "kernel_phi", + "kernel_phi_arg1", + "kernel_phi_changed_rc" + ], + "decline": [] + }, + "lambda_pointer_field_load": { + "specialize": [ + "kernelPointerFieldLoad" + ], + "decline": [] + }, + "lambda_portability_aggregate_shapes": { + "specialize": [ + "kernelAggregateFieldCall", + "kernelCallContext", + "kernelNestedReturn", + "kernelReturnedAlias", + "kernelReturnedFieldAlias" + ], + "decline": [] + }, + "lambda_provenance_shapes": { + "specialize": [ + "kernelArrayMember", + "kernelMemcpySource", + "kernelNegativeGep", + "kernelOffsetMemcpy", + "kernelPartialMemcpy", + "kernelReturnedInteriorPointer" + ], + "decline": [] + }, + "lambda_reconciliation_adversarial": { + "specialize": [], + "decline": [ + "makeBytePair", + "makeConditionalPair", + "makeDuplicatePair", + "makeHeterogeneousPair", + "makePermutedPair", + "makeTriple" + ] + }, + "lambda_size_test": { + "specialize": [ + "kernelEnclosedFunctor", + "kernelEnvelopeSuffix" + ], + "decline": [] + }, + "lambda_slot_selection_supported": { + "specialize": [ + "kernelCallAfterInvoke", + "kernelPhiSlot", + "kernelSelectedSlot", + "kernelTrailingReadOnlyCall" + ], + "decline": [] + }, + "lambda_spec_test": { + "specialize": [ + "kernel" + ], + "decline": [] + }, + "lambda_store_order": { + "specialize": [ + "kernelCaptureFromArg", + "kernelLatestStore", + "kernelStoreOrder" + ], + "decline": [ + "makeAmbiguousBody", + "makeMutatedBody" + ] + }, + "lambda_thread_privatizer": { + "specialize": [ + "kernel_thread_privatizer" + ], + "decline": [] + }, + "lambda_unknown_clobber": { + "specialize": [], + "decline": [ + "makeAggregateHiddenBody", + "makeAtomicBody", + "makeCompareExchangeBody", + "makeIndirectCallBody", + "makeOpaqueCallBody", + "makeSelectedStoreBody" + ] + }, + "lambda_unknown_clobber_decl": { + "specialize": [], + "decline": [] + } + } +} diff --git a/tools/lambda-spec-ci/lambda_spec_ci.py b/tools/lambda-spec-ci/lambda_spec_ci.py new file mode 100644 index 00000000..6083d182 --- /dev/null +++ b/tools/lambda-spec-ci/lambda_spec_ci.py @@ -0,0 +1,179 @@ +#!/usr/bin/env python3 +"""Ground-truth tracking for lambda kernel-arg specialization decisions. + +The ProteusPass trace (PROTEUS_DEBUG_OUTPUT=1, .proteus-logs/proteus-pass.*.log) +records one line per analysis outcome: + + "Found termination case from function " -> SPECIALIZE + "[WARNING]: Kernel arg analysis failed for functor ..." -> DECLINE + +These lines are identical whether the log came from a real ROCm build (e.g. a +ctest run with PROTEUS_DEBUG_OUTPUT=1) or from the pseudo-GPU device pipeline, +so one ground-truth file serves both. + +Modes: + capture parse per-test log files into a ground-truth JSON + check parse per-test log files and diff against a ground-truth JSON + table render the JSON as a Markdown table + +Log layout for capture/check: a directory containing .log files, one per +test (concatenating that test's proteus-pass logs is fine; duplicate decisions +from repeated pass runs are deduplicated). +""" + +import argparse +import datetime +import json +import re +import subprocess +import sys +from pathlib import Path + +SPECIALIZE_RE = re.compile(r"Found termination case from function (\S+)") +DECLINE_RE = re.compile( + r"\[WARNING\]: Kernel arg analysis failed for functor beginning at\s+" + r"call \S+ @(\S+?)\(") + +# Itanium length-prefixed identifiers, e.g. 26kernelPredecessorOverwrite. +LENGTH_PREFIXED_RE = re.compile(r"(\d+)([A-Za-z_])") +PLUMBING_NAMES = {"proteus", "detail", "LambdaFunctorWrapper"} + + +def extract_identifiers(mangled: str): + """Yield the length-prefixed identifiers embedded in a mangled name.""" + pos = 0 + while True: + match = LENGTH_PREFIXED_RE.search(mangled, pos) + if not match: + return + length = int(match.group(1)) + start = match.end(1) + ident = mangled[start:start + length] + if len(ident) == length and re.fullmatch(r"[A-Za-z0-9_]+", ident): + yield ident + pos = start + length + else: + pos = match.end(1) + + +def friendly_name(mangled: str) -> str: + """Stable human-readable key for a kernel or wrapper-call mangled name. + + The first non-plumbing length-prefixed identifier is the symbol's own + name (the kernel for specializations; for wrapper-operator declines the + plumbing filter skips to the registered-lambda factory). Later + identifiers are template-argument noise and Itanium substitution refs. + """ + for ident in extract_identifiers(mangled): + if ident not in PLUMBING_NAMES and len(ident) >= 2: + return ident + return mangled + + +def parse_log(path: Path): + text = path.read_text(errors="replace") + specialize = sorted({friendly_name(M) for M in SPECIALIZE_RE.findall(text)}) + decline = sorted({friendly_name(M) for M in DECLINE_RE.findall(text)}) + return {"specialize": specialize, "decline": decline} + + +def collect(logs_dir: Path): + tests = {} + for log in sorted(logs_dir.glob("*.log")): + tests[log.stem] = parse_log(log) + if not tests: + sys.exit(f"error: no .log files under {logs_dir}") + return tests + + +def git_commit() -> str: + try: + return subprocess.run(["git", "rev-parse", "--short", "HEAD"], + capture_output=True, text=True, + check=True).stdout.strip() + except Exception: + return "unknown" + + +def cmd_capture(args): + truth = { + "meta": { + "generated": datetime.date.today().isoformat(), + "commit": git_commit(), + "pipeline": args.pipeline, + "note": args.note, + }, + "tests": collect(Path(args.logs)), + } + Path(args.truth).write_text(json.dumps(truth, indent=2) + "\n") + total = sum(len(T["specialize"]) + len(T["decline"]) + for T in truth["tests"].values()) + print(f"captured {len(truth['tests'])} tests, {total} decisions " + f"-> {args.truth}") + + +def cmd_check(args): + truth = json.loads(Path(args.truth).read_text()) + actual = collect(Path(args.logs)) + failures = [] + for test, expected in truth["tests"].items(): + if test not in actual: + failures.append(f"{test}: MISSING (no log captured)") + continue + for kind in ("specialize", "decline"): + want, got = set(expected[kind]), set(actual[test][kind]) + for name in sorted(want - got): + failures.append(f"{test}: {name} expected {kind.upper()}, " + f"decision absent") + for name in sorted(got - want): + failures.append(f"{test}: {name} unexpectedly {kind.upper()}") + for test in sorted(set(actual) - set(truth["tests"])): + print(f"note: {test} not in ground truth (new test?); " + f"re-capture to adopt it") + if failures: + print(f"LAMBDA-SPEC-CI: FAIL ({len(failures)} drift(s) vs " + f"{args.truth} @ {truth['meta'].get('commit')})") + for F in failures: + print(f" {F}") + sys.exit(1) + print(f"LAMBDA-SPEC-CI: OK ({len(truth['tests'])} tests match " + f"{args.truth} @ {truth['meta'].get('commit')})") + + +def cmd_table(args): + truth = json.loads(Path(args.truth).read_text()) + print("| Test | Specializes | Declines |") + print("| --- | --- | --- |") + for test, T in sorted(truth["tests"].items()): + spec = ", ".join(T["specialize"]) or "—" + decl = ", ".join(T["decline"]) or "—" + print(f"| {test} | {spec} | {decl} |") + + +def main(): + parser = argparse.ArgumentParser(description=__doc__) + sub = parser.add_subparsers(dest="mode", required=True) + + cap = sub.add_parser("capture") + cap.add_argument("--logs", required=True) + cap.add_argument("--truth", required=True) + cap.add_argument("--pipeline", default="unspecified", + help="e.g. rocm-6.4.3-llvm19 or pseudo-gpu-llvm19") + cap.add_argument("--note", default="") + cap.set_defaults(func=cmd_capture) + + chk = sub.add_parser("check") + chk.add_argument("--logs", required=True) + chk.add_argument("--truth", required=True) + chk.set_defaults(func=cmd_check) + + tab = sub.add_parser("table") + tab.add_argument("--truth", required=True) + tab.set_defaults(func=cmd_table) + + args = parser.parse_args() + args.func(args) + + +if __name__ == "__main__": + main() From 3e90e6c4766624e883e3436dfd3214717a8b72f1 Mon Sep 17 00:00:00 2001 From: john bowen Date: Thu, 15 Oct 2026 12:48:53 -0700 Subject: [PATCH 4/7] Sync spec ground truth to latest capture (37 tests) Adopts the newer decision map and doc from the analysis branch: lambda_interior_slot_overwrite is gone from the GPU suite there (its shape becomes a frontend error), and the metadata records the removal. The JSON remains pseudo-GPU provenance and must be re-captured -- against main for this branch's standalone PR, against ROCm before being treated as hardware truth. --- docs/dev/lambda-spec-groundtruth.md | 9 ++++----- tests/gpu/lambda-spec-groundtruth.json | 8 +------- 2 files changed, 5 insertions(+), 12 deletions(-) diff --git a/docs/dev/lambda-spec-groundtruth.md b/docs/dev/lambda-spec-groundtruth.md index 7507cef3..7dd8d97b 100644 --- a/docs/dev/lambda-spec-groundtruth.md +++ b/docs/dev/lambda-spec-groundtruth.md @@ -37,10 +37,10 @@ ROCm before. Tests with no entry produce no analysis decisions in this pipeline: lambda_dominating_store (ill-formed, unregistered in CMake) and lambda_large_warning (needs __syncthreads in the pseudo-GPU stub). -Known expected-fail pinned in the suite: lambda_interior_slot_overwrite -currently SPECIALIZES (type-erased interior write; soundness delegated to -the frontend contract -- see docs/dev/design/lambda-closure-provenance.md). -Its runtime FileCheck fails by design until the AST diagnostic lands. +The former expected-fail lambda_interior_slot_overwrite has been removed +from the GPU suite: its type-erased interior write is now rejected at +compile time by the AST checker (tests/ast/lambda_wrapper_storage_invalid.cpp, +per the frontend contract in docs/dev/design/lambda-closure-provenance.md). ## Current decisions @@ -65,7 +65,6 @@ Its runtime FileCheck fails by design until the AST diagnostic lands. | lambda_index_first_wrapper | kernel | — | | lambda_insertvalue_offsets | kernelDoubleShift, kernelPadded, kernelShift | — | | lambda_inst_visitor_call_clobber | kernelInstVisitorCallClobber, kernelInstVisitorCallDoesntClobber, kernelInstVisitorOrderedMemcpys | makeUntraceableBody | -| lambda_interior_slot_overwrite | — | makeInteriorBody | | lambda_memory_intrinsic_shapes | kernelNonoverlapMemset | makeDynamicBody, makeMoveBody | | lambda_mfem_style_launch | MockRajaInterface | — | | lambda_mixed_field_overwrite | — | makeBody | diff --git a/tests/gpu/lambda-spec-groundtruth.json b/tests/gpu/lambda-spec-groundtruth.json index 45ffa37c..ad1b8db5 100644 --- a/tests/gpu/lambda-spec-groundtruth.json +++ b/tests/gpu/lambda-spec-groundtruth.json @@ -3,7 +3,7 @@ "generated": "2026-10-15", "commit": "34051bce", "pipeline": "pseudo-gpu-llvm19", - "note": "Baseline at branch tip (pre affine-delta fix). Captured from the pseudo-GPU device pipeline; MUST be re-captured from a ROCm run before being treated as hardware ground truth." + "note": "Baseline at branch tip (pre affine-delta fix). Captured from the pseudo-GPU device pipeline; MUST be re-captured from a ROCm run before being treated as hardware ground truth. lambda_interior_slot_overwrite removed: its shape is now rejected by the AST checker (tests/ast/lambda_wrapper_storage_invalid.cpp)." }, "tests": { "lambda-two": { @@ -139,12 +139,6 @@ "makeUntraceableBody" ] }, - "lambda_interior_slot_overwrite": { - "specialize": [], - "decline": [ - "makeInteriorBody" - ] - }, "lambda_memory_intrinsic_shapes": { "specialize": [ "kernelNonoverlapMemset" From 630c72798ab8a6d7c2426434dca94201cf12262a Mon Sep 17 00:00:00 2001 From: john bowen Date: Thu, 15 Oct 2026 12:50:38 -0700 Subject: [PATCH 5/7] Pseudo-GPU glob: CONFIGURE_DEPENDS so new tests get targets without re-cmake --- tools/pseudo-gpu/PseudoGPUTargets.cmake | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/tools/pseudo-gpu/PseudoGPUTargets.cmake b/tools/pseudo-gpu/PseudoGPUTargets.cmake index 74cee43c..5b193a5a 100644 --- a/tools/pseudo-gpu/PseudoGPUTargets.cmake +++ b/tools/pseudo-gpu/PseudoGPUTargets.cmake @@ -34,7 +34,10 @@ set(PSEUDO_GPU_FLAGS # that so target attribute macros exist regardless of include order. -include hip/hip_runtime.h) -file(GLOB PSEUDO_GPU_TEST_SOURCES "${PROJECT_SOURCE_DIR}/tests/gpu/*.cpp") +# CONFIGURE_DEPENDS: recheck the listing on every build so newly added test +# sources get targets without a manual re-run of cmake. +file(GLOB PSEUDO_GPU_TEST_SOURCES CONFIGURE_DEPENDS + "${PROJECT_SOURCE_DIR}/tests/gpu/*.cpp") set(PSEUDO_GPU_ALL_OUTPUTS) foreach(PseudoSource IN LISTS PSEUDO_GPU_TEST_SOURCES) From 0996942909b2055aaee83c90a94ec1f5d8c6396b Mon Sep 17 00:00:00 2001 From: john bowen Date: Thu, 15 Oct 2026 13:01:08 -0700 Subject: [PATCH 6/7] Re-capture ground truth as the main-analysis baseline The JSON now records main's decisions over main's lambda suite (23 tests, 44 decisions) rather than a snapshot from the analysis feature branch, so this tooling PR is self-consistent and the feature branch's PR can show its behavioral delta as a diff against this file. --- docs/dev/lambda-spec-groundtruth.md | 33 ++----- tests/gpu/lambda-spec-groundtruth.json | 128 ++----------------------- 2 files changed, 13 insertions(+), 148 deletions(-) diff --git a/docs/dev/lambda-spec-groundtruth.md b/docs/dev/lambda-spec-groundtruth.md index 7dd8d97b..86b15697 100644 --- a/docs/dev/lambda-spec-groundtruth.md +++ b/docs/dev/lambda-spec-groundtruth.md @@ -29,31 +29,21 @@ builds and the pseudo-GPU pipeline: ## Provenance caveat -The current JSON was captured from the pseudo-GPU pipeline (LLVM 19, -O1, -gfx90a) at the branch tip. Compile-time decisions have matched ROCm in spot -checks, but this file MUST be re-captured from a ROCm run before being -treated as hardware ground truth -- pseudo-GPU "green" has disagreed with -ROCm before. Tests with no entry produce no analysis decisions in this +The current JSON is the baseline of MAIN's analysis over main's lambda +suite, captured from the pseudo-GPU pipeline (LLVM 19, -O1, gfx90a). +Compile-time decisions have matched ROCm in spot checks, but this file +should be re-captured from a ROCm run before being treated as hardware +ground truth. Tests with no entry produce no analysis decisions in this pipeline: lambda_dominating_store (ill-formed, unregistered in CMake) and lambda_large_warning (needs __syncthreads in the pseudo-GPU stub). -The former expected-fail lambda_interior_slot_overwrite has been removed -from the GPU suite: its type-erased interior write is now rejected at -compile time by the AST checker (tests/ast/lambda_wrapper_storage_invalid.cpp, -per the frontend contract in docs/dev/design/lambda-closure-provenance.md). - ## Current decisions | Test | Specializes | Declines | | --- | --- | --- | | lambda | kernel | — | | lambda-two | kernel | — | -| lambda_aggregate_store_reconciliation | — | makeBody | -| lambda_call_shapes | kernelSelectSame | makeIndirectBody, makeReturnBody, makeSameReturnBody, makeSelectSameBody | -| lambda_clobber_query_range | kernelUnrelatedPrefixOverwrite | — | -| lambda_clobber_query_suffix | kernelUnrelatedSuffixOverwrite | — | -| lambda_clobber_shapes | kernelCallOverwrite, kernelLocalSupersedesCall, kernelNestedCallOverwrite, kernelNonoverlapOverwrite, kernelPredecessorOverwrite | makeAmbiguousLoopBody, makeDifferentBranchBody, makeDifferentCallBranchBody, makeSameBranchBody, makeSameCallBranchBody, makeSameLoopBody, makeUnrelatedCallBody | -| lambda_declaration_clobber | — | makeBody | +| lambda_call_shapes | kernelSameReturns, kernelSelectSame | makeIndirectBody, makeReturnBody, makeSelectSameBody | | lambda_def | kernel | — | | lambda_factory | kernel | — | | lambda_factory_forwarded | kernel | — | @@ -64,22 +54,13 @@ per the frontend contract in docs/dev/design/lambda-closure-provenance.md). | lambda_host_device | kernel | — | | lambda_index_first_wrapper | kernel | — | | lambda_insertvalue_offsets | kernelDoubleShift, kernelPadded, kernelShift | — | -| lambda_inst_visitor_call_clobber | kernelInstVisitorCallClobber, kernelInstVisitorCallDoesntClobber, kernelInstVisitorOrderedMemcpys | makeUntraceableBody | -| lambda_memory_intrinsic_shapes | kernelNonoverlapMemset | makeDynamicBody, makeMoveBody | +| lambda_memory_intrinsic_shapes | kernelNonoverlapMemset, kernelOffsetMemmove | makeDynamicBody | | lambda_mfem_style_launch | MockRajaInterface | — | -| lambda_mixed_field_overwrite | — | makeBody | | lambda_multiple | kernel | — | -| lambda_multiple_callers | — | makeBody | | lambda_pathological | kernel_case1, kernel_case2_a, kernel_case2_b, kernel_case4, kernel_case5 | — | | lambda_phi_analysis | kernel_phi, kernel_phi_arg1, kernel_phi_changed_rc | — | | lambda_pointer_field_load | kernelPointerFieldLoad | — | -| lambda_portability_aggregate_shapes | kernelAggregateFieldCall, kernelCallContext, kernelNestedReturn, kernelReturnedAlias, kernelReturnedFieldAlias | — | | lambda_provenance_shapes | kernelArrayMember, kernelMemcpySource, kernelNegativeGep, kernelOffsetMemcpy, kernelPartialMemcpy, kernelReturnedInteriorPointer | — | -| lambda_reconciliation_adversarial | — | makeBytePair, makeConditionalPair, makeDuplicatePair, makeHeterogeneousPair, makePermutedPair, makeTriple | -| lambda_size_test | kernelEnclosedFunctor, kernelEnvelopeSuffix | — | -| lambda_slot_selection_supported | kernelCallAfterInvoke, kernelPhiSlot, kernelSelectedSlot, kernelTrailingReadOnlyCall | — | | lambda_spec_test | kernel | — | | lambda_store_order | kernelCaptureFromArg, kernelLatestStore, kernelStoreOrder | makeAmbiguousBody, makeMutatedBody | | lambda_thread_privatizer | kernel_thread_privatizer | — | -| lambda_unknown_clobber | — | makeAggregateHiddenBody, makeAtomicBody, makeCompareExchangeBody, makeIndirectCallBody, makeOpaqueCallBody, makeSelectedStoreBody | -| lambda_unknown_clobber_decl | — | — | diff --git a/tests/gpu/lambda-spec-groundtruth.json b/tests/gpu/lambda-spec-groundtruth.json index ad1b8db5..45fa64d4 100644 --- a/tests/gpu/lambda-spec-groundtruth.json +++ b/tests/gpu/lambda-spec-groundtruth.json @@ -1,9 +1,9 @@ { "meta": { "generated": "2026-10-15", - "commit": "34051bce", + "commit": "9cffdd4a", "pipeline": "pseudo-gpu-llvm19", - "note": "Baseline at branch tip (pre affine-delta fix). Captured from the pseudo-GPU device pipeline; MUST be re-captured from a ROCm run before being treated as hardware ground truth. lambda_interior_slot_overwrite removed: its shape is now rejected by the AST checker (tests/ast/lambda_wrapper_storage_invalid.cpp)." + "note": "Baseline of MAIN's analysis over main's lambda suite (branch base 2a8f200a). Captured from the pseudo-GPU device pipeline; re-capture from a ROCm run before treating as hardware ground truth. lambda_dominating_store (ill-formed, unregistered) and lambda_large_warning (stub lacks __syncthreads) produce no decisions." }, "tests": { "lambda-two": { @@ -18,59 +18,17 @@ ], "decline": [] }, - "lambda_aggregate_store_reconciliation": { - "specialize": [], - "decline": [ - "makeBody" - ] - }, "lambda_call_shapes": { "specialize": [ + "kernelSameReturns", "kernelSelectSame" ], "decline": [ "makeIndirectBody", "makeReturnBody", - "makeSameReturnBody", "makeSelectSameBody" ] }, - "lambda_clobber_query_range": { - "specialize": [ - "kernelUnrelatedPrefixOverwrite" - ], - "decline": [] - }, - "lambda_clobber_query_suffix": { - "specialize": [ - "kernelUnrelatedSuffixOverwrite" - ], - "decline": [] - }, - "lambda_clobber_shapes": { - "specialize": [ - "kernelCallOverwrite", - "kernelLocalSupersedesCall", - "kernelNestedCallOverwrite", - "kernelNonoverlapOverwrite", - "kernelPredecessorOverwrite" - ], - "decline": [ - "makeAmbiguousLoopBody", - "makeDifferentBranchBody", - "makeDifferentCallBranchBody", - "makeSameBranchBody", - "makeSameCallBranchBody", - "makeSameLoopBody", - "makeUnrelatedCallBody" - ] - }, - "lambda_declaration_clobber": { - "specialize": [], - "decline": [ - "makeBody" - ] - }, "lambda_def": { "specialize": [ "kernel" @@ -129,23 +87,13 @@ ], "decline": [] }, - "lambda_inst_visitor_call_clobber": { - "specialize": [ - "kernelInstVisitorCallClobber", - "kernelInstVisitorCallDoesntClobber", - "kernelInstVisitorOrderedMemcpys" - ], - "decline": [ - "makeUntraceableBody" - ] - }, "lambda_memory_intrinsic_shapes": { "specialize": [ - "kernelNonoverlapMemset" + "kernelNonoverlapMemset", + "kernelOffsetMemmove" ], "decline": [ - "makeDynamicBody", - "makeMoveBody" + "makeDynamicBody" ] }, "lambda_mfem_style_launch": { @@ -154,24 +102,12 @@ ], "decline": [] }, - "lambda_mixed_field_overwrite": { - "specialize": [], - "decline": [ - "makeBody" - ] - }, "lambda_multiple": { "specialize": [ "kernel" ], "decline": [] }, - "lambda_multiple_callers": { - "specialize": [], - "decline": [ - "makeBody" - ] - }, "lambda_pathological": { "specialize": [ "kernel_case1", @@ -196,16 +132,6 @@ ], "decline": [] }, - "lambda_portability_aggregate_shapes": { - "specialize": [ - "kernelAggregateFieldCall", - "kernelCallContext", - "kernelNestedReturn", - "kernelReturnedAlias", - "kernelReturnedFieldAlias" - ], - "decline": [] - }, "lambda_provenance_shapes": { "specialize": [ "kernelArrayMember", @@ -217,33 +143,6 @@ ], "decline": [] }, - "lambda_reconciliation_adversarial": { - "specialize": [], - "decline": [ - "makeBytePair", - "makeConditionalPair", - "makeDuplicatePair", - "makeHeterogeneousPair", - "makePermutedPair", - "makeTriple" - ] - }, - "lambda_size_test": { - "specialize": [ - "kernelEnclosedFunctor", - "kernelEnvelopeSuffix" - ], - "decline": [] - }, - "lambda_slot_selection_supported": { - "specialize": [ - "kernelCallAfterInvoke", - "kernelPhiSlot", - "kernelSelectedSlot", - "kernelTrailingReadOnlyCall" - ], - "decline": [] - }, "lambda_spec_test": { "specialize": [ "kernel" @@ -266,21 +165,6 @@ "kernel_thread_privatizer" ], "decline": [] - }, - "lambda_unknown_clobber": { - "specialize": [], - "decline": [ - "makeAggregateHiddenBody", - "makeAtomicBody", - "makeCompareExchangeBody", - "makeIndirectCallBody", - "makeOpaqueCallBody", - "makeSelectedStoreBody" - ] - }, - "lambda_unknown_clobber_decl": { - "specialize": [], - "decline": [] } } } From 8e4faeeef6358b03a1c2360793d60cabb143df92 Mon Sep 17 00:00:00 2001 From: john bowen Date: Thu, 15 Oct 2026 13:02:48 -0700 Subject: [PATCH 7/7] Add CPU-runner CI for lambda specialization decisions tools/lambda-spec-ci/run_sweep.sh drives the pseudo-GPU device pipeline over every lambda GPU test and checks the resulting specialize/decline decisions against the committed ground truth; the new ci-lambda-spec.yml workflow runs it in a plain ubuntu:24.04 container with apt.llvm.org clang -- no GPU, ROCm, or conda -- failing on decision drift and uploading the pass traces on failure. Tests absent from the ground truth (expected not to compile in this pipeline) are skipped with a note; a traceless test that IS in the ground truth fails the job. --- .github/workflows/ci-lambda-spec.yml | 79 ++++++++++++++++++++++++++++ tools/lambda-spec-ci/run_sweep.sh | 59 +++++++++++++++++++++ 2 files changed, 138 insertions(+) create mode 100644 .github/workflows/ci-lambda-spec.yml create mode 100755 tools/lambda-spec-ci/run_sweep.sh diff --git a/.github/workflows/ci-lambda-spec.yml b/.github/workflows/ci-lambda-spec.yml new file mode 100644 index 00000000..c602e762 --- /dev/null +++ b/.github/workflows/ci-lambda-spec.yml @@ -0,0 +1,79 @@ +name: Lambda specialization decisions + +# Pseudo-GPU verification of the lambda kernel-arg analysis: compile every +# lambda GPU test through the real HIP device pipeline (clang -x hip +# targeting amdgcn) with ProteusPass loaded, on a plain CPU container -- no +# GPU or ROCm -- and fail on any drift of specialize/decline decisions +# against tests/gpu/lambda-spec-groundtruth.json. + +on: + pull_request: + paths: + - ".github/workflows/ci-lambda-spec.yml" + - "CMakeLists.txt" + - "include/**" + - "src/**" + - "tests/gpu/**" + - "tools/lambda-spec-ci/**" + - "tools/pseudo-gpu/**" + push: + branches: + - main + paths: + - ".github/workflows/ci-lambda-spec.yml" + - "CMakeLists.txt" + - "include/**" + - "src/**" + - "tests/gpu/**" + - "tools/lambda-spec-ci/**" + - "tools/pseudo-gpu/**" + +jobs: + lambda-spec: + name: pseudo-GPU decision check (LLVM ${{ matrix.llvm }}) + runs-on: ubuntu-latest + container: ubuntu:24.04 + strategy: + matrix: + llvm: [19] + steps: + - name: Install toolchain + run: | + apt-get update + DEBIAN_FRONTEND=noninteractive apt-get install -y \ + wget gnupg software-properties-common cmake ninja-build \ + python3 git zlib1g-dev + wget -qO- https://apt.llvm.org/llvm-snapshot.gpg.key \ + | tee /etc/apt/trusted.gpg.d/apt.llvm.org.asc + add-apt-repository -y \ + "deb http://apt.llvm.org/noble/ llvm-toolchain-noble-${{ matrix.llvm }} main" + apt-get update + apt-get install -y clang-${{ matrix.llvm }} llvm-${{ matrix.llvm }}-dev \ + libclang-${{ matrix.llvm }}-dev libzstd-dev libedit-dev libcurl4-openssl-dev + + - uses: actions/checkout@v4 + + - name: Build ProteusPass (pseudo-GPU mode) + run: | + cmake -S . -B build -G Ninja \ + -DCMAKE_BUILD_TYPE=Release \ + -DCMAKE_C_COMPILER=clang-${{ matrix.llvm }} \ + -DCMAKE_CXX_COMPILER=clang++-${{ matrix.llvm }} \ + -DLLVM_INSTALL_DIR=/usr/lib/llvm-${{ matrix.llvm }} \ + -DPROTEUS_PSEUDO_GPU=ON + cmake --build build --target ProteusPass -j "$(nproc)" + + - name: Sweep lambda suite and check decisions + run: | + bash tools/lambda-spec-ci/run_sweep.sh \ + clang++-${{ matrix.llvm }} \ + "$(pwd)/build/src/pass/libProteusPass.so" \ + "$(pwd)/build/lambda-spec-logs" + + - name: Upload traces on failure + if: failure() + uses: actions/upload-artifact@v7 + with: + name: lambda-spec-logs-llvm${{ matrix.llvm }} + path: build/lambda-spec-logs + retention-days: 7 diff --git a/tools/lambda-spec-ci/run_sweep.sh b/tools/lambda-spec-ci/run_sweep.sh new file mode 100755 index 00000000..51d823d9 --- /dev/null +++ b/tools/lambda-spec-ci/run_sweep.sh @@ -0,0 +1,59 @@ +#!/usr/bin/env bash +# Run the pseudo-GPU device pipeline over every lambda GPU test, collecting +# one ProteusPass trace per test, then check the decisions against the +# committed ground truth. Needs no GPU or ROCm: any host clang++ with the +# AMDGPU backend (stock LLVM releases include it) can run it. +# +# Usage: tools/lambda-spec-ci/run_sweep.sh [logs-dir] +# Exits nonzero on compile failure of a test that has ground-truth entries, +# or on any decision drift. + +set -u + +REPO_ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")/../.." && pwd)" +CLANGXX="${1:?usage: run_sweep.sh [logs-dir]}" +PLUGIN="${2:?usage: run_sweep.sh [logs-dir]}" +LOGS_DIR="${3:-$REPO_ROOT/build/lambda-spec-logs}" +TRUTH="$REPO_ROOT/tests/gpu/lambda-spec-groundtruth.json" + +rm -rf "$LOGS_DIR" +mkdir -p "$LOGS_DIR" + +status=0 +for SOURCE in "$REPO_ROOT"/tests/gpu/lambda*.cpp; do + TEST=$(basename "$SOURCE" .cpp) + WORK_DIR=$(mktemp -d) + (cd "$WORK_DIR" && PROTEUS_DEBUG_OUTPUT=1 "$CLANGXX" \ + -x hip --offload-arch=gfx90a --offload-device-only \ + -nogpulib -nogpuinc -std=c++17 -O1 \ + -DPROTEUS_ENABLE_HIP=1 \ + -I "$REPO_ROOT/tools/pseudo-gpu/include" \ + -I "$REPO_ROOT/include" \ + -include hip/hip_runtime.h \ + -fpass-plugin="$PLUGIN" \ + -S -emit-llvm -o /dev/null "$SOURCE") \ + > "$WORK_DIR/stdout.txt" 2> "$WORK_DIR/stderr.txt" + if compgen -G "$WORK_DIR/.proteus-logs/proteus-pass.*.log" > /dev/null; then + cat "$WORK_DIR"/.proteus-logs/proteus-pass.*.log > "$LOGS_DIR/$TEST.log" + else + # No trace at all: fail loudly only for tests the ground truth knows, + # since some sources are expected not to compile in this pipeline. + if python3 - "$TRUTH" "$TEST" <<'EOF' +import json, sys +truth = json.load(open(sys.argv[1])) +sys.exit(0 if sys.argv[2] in truth["tests"] else 1) +EOF + then + echo "ERROR: $TEST produced no pass trace; compile output:" >&2 + tail -20 "$WORK_DIR/stderr.txt" >&2 + status=1 + else + echo "note: $TEST produced no trace (not in ground truth); skipping" + fi + fi + rm -rf "$WORK_DIR" +done + +[ $status -ne 0 ] && exit $status +exec python3 "$REPO_ROOT/tools/lambda-spec-ci/lambda_spec_ci.py" check \ + --logs "$LOGS_DIR" --truth "$TRUTH"