feature: Add pseudo-GPU device compilation harness for pass verification

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] <source.cpp> [-- <extra clang args>]
+#
+# Options:
+#   --llvm <prefix>   LLVM install prefix (default: /opt/homebrew/opt/llvm@19;
+#                     the pass plugin is picked from build-pseudogpu-llvm<ver>)
+#   --build <dir>     Build dir containing src/pass/libProteusPass.dylib
+#   --raw-ir <file>   Also emit pre-pass device IR (no plugin) for shape study
+#   --log <file>      Copy the pass trace here (default: <source>.proteus-pass.log)
+#   -O<n>             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

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 <cstddef>
+#include <cstdint>
+
+// 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 <typename T>
+static inline hipError_t hipMallocManaged(T **Ptr, size_t Size,
+                                          unsigned Flags = 1) {
+  return hipMallocManaged(reinterpret_cast<void **>(Ptr), Size, Flags);
+}
+
+template <typename T>
+static inline hipError_t hipMalloc(T **Ptr, size_t Size) {
+  return hipMalloc(reinterpret_cast<void **>(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