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