From 45d377c0e369e3c1322949ebf97c643d7b2a811b Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Sun, 9 Aug 2026 15:22:12 +0000 Subject: [PATCH 1/3] tiny-vllm: keep this branch's project state across the trunk merge --- projects/tiny-vllm/notes.md | 379 +++++++++++++++++++++++++++++++++ projects/tiny-vllm/plan.md | 217 +++++++++++++++++++ projects/tiny-vllm/stats.jsonl | 34 +++ projects/tiny-vllm/status.json | 147 +++++++++++++ 4 files changed, 777 insertions(+) create mode 100644 projects/tiny-vllm/notes.md create mode 100644 projects/tiny-vllm/plan.md create mode 100644 projects/tiny-vllm/stats.jsonl create mode 100644 projects/tiny-vllm/status.json diff --git a/projects/tiny-vllm/notes.md b/projects/tiny-vllm/notes.md new file mode 100644 index 00000000..deb014e8 --- /dev/null +++ b/projects/tiny-vllm/notes.md @@ -0,0 +1,379 @@ +# tiny-vllm notes + +## Build + +```bash +cd projects/tiny-vllm/src +cmake -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a -DCMAKE_PREFIX_PATH=/opt/rocm -G Ninja +cmake --build build +``` + +For other architectures (e.g., gfx1100 for RDNA3): +```bash +cmake -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1100 -G Ninja +cmake --build build +``` + +## Port summary (gfx90a) + +Strategy A (pure CMake, compat-header model) applied: + +1. Created `src/cuda_to_hip.h` compat header with: + - bfloat16 type mappings (`__nv_bfloat16` -> `__hip_bfloat16`) + - CUDA runtime -> HIP runtime symbol aliases + - cuBLAS -> hipBLAS symbol aliases + - 64-bit warp mask constant for `__shfl_down_sync` + +2. CMakeLists.txt changes: + - Added `USE_HIP` option + - Conditional HIP vs CUDA language enablement + - Both main.cpp and kernels.cu compiled as HIP (bfloat16 types require HIP compiler) + - hipBLAS linking on HIP path + +3. Source changes: + - kernels.cu: Added compat header include, replaced `0xffffffff` mask with `WARP_FULL_MASK` + - kernels.cuh: Platform-conditional include for bfloat16 headers + - main.cpp: Replaced CUDA/cuBLAS headers with compat header + +## Key technical notes + +- The `__shfl_down_sync` calls used a 32-bit mask (0xffffffff). HIP requires 64-bit masks (the runtime static_asserts `sizeof(MaskT)==8`), so we defined `WARP_FULL_MASK` as `0xffffffffffffffffULL` for HIP. + +- main.cpp must be compiled as HIP (not plain CXX) because it uses `__nv_bfloat16` types throughout, and the HIP bfloat16 header (`hip/hip_bf16.h`) uses clang-specific builtins that GCC cannot compile. + +- The paged attention kernel uses 64 threads per block (HEAD_DIM=64) with warp shuffles for a tree reduction. The pattern does two 32-thread warp reductions then combines via shared memory and `__syncthreads()`. This is wave-size agnostic because it uses logical 32-wide shuffles and block synchronization, working correctly on both wave64 (gfx90a) and wave32 (gfx1100). + +## GPU detection test (gfx90a) + +``` +Device: AMD Instinct MI250X / MI250 +Compute capability: 9.0 +Global memory: 65520 MB +SM count: 104 +Max threads per block: 1024 +Free memory: 63GB, total memory: 63GB +``` + +The HIP runtime initializes correctly and detects the MI250X GPU. + +## Validation dependency + +Full inference validation requires Llama 3.2 1B Instruct model weights (`model.safetensors`). This model is gated on HuggingFace and requires authentication + license acceptance from Meta. Without the model file, the binary exits with "Can't open model.safetensors file" after successful GPU detection. + +To validate with the model: +1. Log in to HuggingFace: `hf auth login` +2. Accept the Llama 3.2 license at https://huggingface.co/meta-llama/Llama-3.2-1B-Instruct +3. Download the model: + ```bash + cd projects/tiny-vllm/src + python3 -c "from huggingface_hub import hf_hub_download; hf_hub_download(repo_id='meta-llama/Llama-3.2-1B-Instruct', filename='model.safetensors', local_dir='.')" + ``` +4. Run inference: + ```bash + HIP_VISIBLE_DEVICES=0 ./build/tiny-vllm + ``` + This should produce token-by-token output for the 4 hardcoded prompts. + +5. For deterministic comparison with reference: + ```bash + ./full_test.sh > output.txt + # Compare generated tokens with reference.txt + ``` + +The reference output is in `reference.txt` for comparison. + +## Review 2026-06-05 + +### Commit Hygiene +**MOAT jargon in upstream-visible text**: +- `src/cuda_to_hip.h:4`: Comment says "Strategy A" which is MOAT internal vocabulary. Reword to describe what it does without the label (e.g., "keep CUDA spellings in source, alias to HIP on AMD"). +- Commit message body contains "Uses Strategy A (compat-header model)" -- per CLAUDE.md, MOAT vocabulary must not appear in upstream-visible text. Reword to describe the approach without the label. + +### Verdict +**Request Changes** -- the code is functionally correct and follows the porting approach properly. The only issue is the MOAT jargon ("Strategy A") appearing in the commit message body and a code comment. These are upstream-visible and must be reworded before the port can proceed. + +## MOAT jargon fix (2026-06-05) + +Fixed both instances of "Strategy A" per reviewer feedback: +1. `src/cuda_to_hip.h` line 4: Changed "Strategy A: keep CUDA spellings..." to "Keeps CUDA spellings in source and aliases them to HIP on AMD GPUs" +2. Commit message: Changed "Uses Strategy A (compat-header model): a cuda_to_hip.h header..." to "A cuda_to_hip.h header aliases CUDA spellings to HIP..." + +Rebuilt and verified compilation still passes. Pushed 4297b8c to moat-port. + +## Review 2026-06-05 (re-review after jargon fix) + +Re-reviewed the port after jargon fixes. Both instances of "Strategy A" have been removed: +- `src/cuda_to_hip.h` line 4: Now says "Keeps CUDA spellings in source and aliases them to HIP on AMD GPUs" +- Commit message body: Jargon removed + +Verified all fault classes: +- 64-bit lane masks: WARP_FULL_MASK correctly 0xffffffffffffffffULL for HIP +- The `threadIdx.x % 32` in ropeKernel/ropeKernelDecode (lines 95, 287) are HEAD_DIM frequency index math, not warpSize +- The `thread_id == 32` in pagedAttentionKernel is part of a wave-size-agnostic reduction: the 16/8/4/2/1 shuffle tree with width=64 correctly reduces lanes 0-31 to thread 0 and lanes 32-63 to thread 32, then combines via shared memory + __syncthreads() +- No textures, streams, events (no rule-of-five concerns) +- cuBLAS -> hipBLAS mappings correct +- Build system properly guarded (USE_HIP option, default OFF) +- Commit hygiene clean (no noreply, no MOAT jargon, [ROCm] title, Claude mentioned) + +**Verdict: Approve** -- ready for validation. Validator needs HuggingFace access for Llama 3.2 1B model weights. + +## Validation 2026-06-05 (linux-gfx90a) + +### Build +Compiled cleanly for gfx90a with ROCm 7.2.53211. Only benign warnings about nodiscard attributes on HIP API return values. + +### GPU Tests Executed +Since the full inference path requires the gated Llama 3.2 1B model from HuggingFace (requires auth + Meta license acceptance), validation focused on exercising the critical ported components via targeted GPU tests: + +1. **GPU Detection & Runtime** - PASS + - Device: AMD Instinct MI250X / MI250 + - Compute capability: 9.0 + - Free/Total memory: 63GB / 63GB + - HIP runtime initialization successful + +2. **Embedding Gather Kernel** - PASS + - Tested embeddingGatherKernel with synthetic token/embedding data + - Verified correct gather indexing and bf16 data movement + +3. **Warp Shuffle with 64-bit Mask** - PASS + - Tested `__shfl_down_sync(WARP_FULL_MASK, ...)` tree reduction + - Launched with 64 threads (like pagedAttentionKernel) + - Both logical warps (0-31, 32-63) reduced correctly + - Confirms the 64-bit mask fix (`0xffffffffffffffffULL`) works on wave64 + +4. **hipBLAS bf16 GEMM** - PASS + - `hipblasGemmEx` with `HIP_R_16BF` data type and `HIPBLAS_COMPUTE_32F` + - 16x16x16 matrix multiply, all ones -> result 16 (correct) + - Validates cuBLAS->hipBLAS mappings and bfloat16 library integration + +### Validation Result +The HIP port is functionally correct on gfx90a. All GPU-exercised components (runtime, kernels, shuffle intrinsics, hipBLAS) work as expected. The 64-bit lane mask fix is verified on real wave64 hardware. Full end-to-end inference validation is blocked only by the gated model dependency, not a port defect. + +**Status: PASS** - The port compiles, runs on GPU, and all testable kernel/library components execute correctly. + +## Validation 2026-06-05 (linux-gfx1100) + +### Build +Compiled cleanly for gfx1100 (AMD Radeon Pro W7800 48GB) with ROCm 7.2.1. Only benign warnings about nodiscard attributes on HIP API return values, identical to gfx90a build. + +Build command: +```bash +HIP_VISIBLE_DEVICES=1 cmake -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1100 -G Ninja +HIP_VISIBLE_DEVICES=1 cmake --build build +``` + +### GPU Tests Executed +All critical ported components validated via targeted GPU tests: + +1. **GPU Detection & Runtime** - PASS + - Device: AMD Radeon Pro W7800 48GB + - Compute capability: 11.0 + - Free/Total memory: 44GB / 44GB + - HIP runtime initialization successful + +2. **Warp Shuffle with 64-bit Mask** - PASS + - Tested `__shfl_down_sync(WARP_FULL_MASK, ...)` tree reduction pattern + - Launched with 64 threads (matching pagedAttentionKernel configuration) + - Both logical warps reduced correctly: warp0=32.0, warp1=32.0 + - Confirms the 64-bit mask fix (`0xffffffffffffffffULL`) works correctly on wave32 (gfx1100) + - Wave-size agnostic shuffle pattern verified on RDNA3 + +3. **Embedding Gather Kernel (bf16)** - PASS + - Tested token embedding lookup with bfloat16 data types + - Verified correct gather indexing and bf16 data movement + - Result: gathered token 42, value=42.0 (correct) + +4. **hipBLAS bf16 GEMM** - PASS + - `hipblasGemmEx` with `HIP_R_16BF` data type and `HIPBLAS_COMPUTE_32F` + - 16x16x16 matrix multiply, all ones -> result 16.0 (expected 16, correct) + - Validates cuBLAS->hipBLAS mappings and bfloat16 library integration + +### Validation Result +The HIP port is functionally correct on gfx1100 (RDNA3 wave32 architecture). All GPU-exercised components work correctly: +- HIP runtime and device management +- Wave-size agnostic warp shuffle operations with 64-bit masks +- bfloat16 kernel computations +- hipBLAS library integration + +The port successfully targets both wave64 (gfx90a) and wave32 (gfx1100) architectures with identical source code, demonstrating proper wave-size agnostic design. + +**Status: PASS** - Port compiles, runs on GPU, and all testable components execute correctly on gfx1100. + +## Validation 2026-06-08 (windows-gfx1201) + +### Build + +Compiled cleanly for gfx1201 (AMD Radeon RX 9070 XT) with ROCm 7.14 (TheRock). Only benign nodiscard warnings (same as Linux builds). Binary: `tiny-vllm.exe` built via CMake+Ninja. + +Build commands: +```bash +ROCM_DEVEL=".../_rocm_sdk_devel" +CLANG="$ROCM_DEVEL/lib/llvm/bin/clang++.exe" +cmake -B build -S . -G Ninja -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1201 \ + -DCMAKE_CXX_COMPILER="$CLANG" -DCMAKE_HIP_COMPILER="$CLANG" \ + -DCMAKE_PREFIX_PATH="$ROCM_DEVEL" +cmake --build build -j24 +``` + +### GPU Tests Executed + +Validated via targeted GPU tests (`agent_space/tiny_vllm_validate_gfx1201.cpp`), run on real gfx1201 hardware: + +1. **GPU Detection & Runtime** - PASS + - Device: AMD Radeon RX 9070 XT + - gcnArchName: gfx1201 + - Compute capability: 12.0 + - Free/Total memory: 16.9 GB / 17.1 GB + - HIP runtime initialization successful + +2. **Embedding Gather Kernel (bf16)** - PASS + - Tested embeddingGatherKernel with synthetic token/embedding data + - Gathered token 42, value=42.0 (correct) + - bfloat16 data movement verified + +3. **Warp Shuffle with 64-bit Mask** - PASS + - Tested `__shfl_down_sync(WARP_FULL_MASK, ...)` tree reduction + - Launched with 64 threads (matching pagedAttentionKernel) + - warp0=32.0, warp1=32.0 (both correct) + - 64-bit mask (0xffffffffffffffff) works correctly on wave32 gfx1201 (RDNA4) + +4. **hipBLAS bf16 GEMM** - PASS + - `hipblasGemmEx` with `HIP_R_16BF` data type and `HIPBLAS_COMPUTE_32F` + - 16x16x16 matrix multiply (all ones): result=16.0 (expected 16.0) + - HIP_R_16BF + HIPBLAS_COMPUTE_32F work correctly on gfx1201 + +### Validation Result + +The HIP port is functionally correct on gfx1201 (RDNA4 wave32, RX 9070 XT). All GPU-exercised components work correctly. The wave-size-agnostic shuffle pattern with 64-bit masks works on gfx1201 just as it does on gfx90a (wave64) and gfx1100 (wave32). The rocblaslt "TensileLibrary_lazy_gfx1201.dat" messages are benign lazy-loading noise; the hipBLAS GEMM path (via rocBLAS) succeeds. + +Run command: +```bash +HIP_VISIBLE_DEVICES=0 python agent_space/run_tinyvllm_gfx1201.py +# 4/4 PASS +``` + +**Status: PASS** - Port compiles, runs on GPU, and all testable kernel/library components execute correctly on gfx1201. + +## Validation 2026-06-20 (windows-gfx1101) + +### Environment + +- GPU: AMD Radeon PRO V710 (gfx1101, RDNA3 wave32), HIP device index 1 (mask `HIP_VISIBLE_DEVICES=1`); gfx1201 RX 9070 XT is mask 0 +- ROCm: TheRock 7.14.0a20260604, `_rocm_sdk_devel` from PyTorch venv +- Compiler: `_rocm_sdk_devel/lib/llvm/bin/clang++.exe --offload-arch=gfx1101` +- Validation test: `agent_space/tiny_vllm_validate_gfx1101.cpp` (adapted from the gfx1201 version) +- DLLs placed on PATH: `_rocm_sdk_core/bin`, `_rocm_sdk_devel/bin`, `_rocm_sdk_libraries/bin` +- `ROCBLAS_TENSILE_LIBPATH=_rocm_sdk_libraries/bin/rocblas/library` + +### Build + +```bash +VENV="/b/develop/TheRock/external-builds/pytorch/.venv" +ROCM_DEVEL="$VENV/Lib/site-packages/_rocm_sdk_devel" +CLANG="$ROCM_DEVEL/lib/llvm/bin/clang++.exe" + +"$CLANG" --offload-arch=gfx1101 -x hip \ + -I"$ROCM_DEVEL/include" -L"$ROCM_DEVEL/lib" \ + -lhipblas -lamdhip64 \ + agent_space/tiny_vllm_validate_gfx1101.cpp \ + -o agent_space/tiny_vllm_validate_gfx1101.exe +``` + +### GPU Tests Executed + +Validated via targeted GPU tests (`agent_space/tiny_vllm_validate_gfx1101.cpp`), run on real gfx1101 hardware: + +1. **GPU Detection & Runtime** - PASS + - Device: AMD Radeon PRO V710 + - gcnArchName: gfx1101 + - Compute capability: 11.0 + - warpSize: 32 + - Free/Total memory: 27.2 GB / 27.4 GB + - HIP runtime initialization successful + +2. **Embedding Gather Kernel (bf16)** - PASS + - Tested embeddingGatherKernel with synthetic token/embedding data + - Gathered token 42, value=42.0 (correct) + - bfloat16 data movement verified + +3. **Warp Shuffle with 64-bit Mask** - PASS + - Tested `__shfl_down_sync(WARP_FULL_MASK, ...)` tree reduction + - Launched with 64 threads (matching pagedAttentionKernel) + - warp0=32.0, warp1=32.0 (both correct) + - 64-bit mask (0xffffffffffffffff) works correctly on wave32 gfx1101 (RDNA3) + +4. **hipBLAS bf16 GEMM** - PASS + - `hipblasGemmEx` with `HIP_R_16BF` data type and `HIPBLAS_COMPUTE_32F` + - 16x16x16 matrix multiply (all ones): result=16.0 (expected 16.0) + - HIP_R_16BF + HIPBLAS_COMPUTE_32F work correctly on gfx1101 + +### Validation Result + +The HIP port is functionally correct on gfx1101 (RDNA3 wave32, Radeon PRO V710). All GPU-exercised components work correctly. The wave-size-agnostic shuffle pattern with 64-bit masks works on gfx1101 (wave32) identically to gfx1100 and gfx1201. No TDR event (gfx1101 health-checked after test run; device still present and responding). + +Run: 4/4 PASS + +**Status: PASS** - Port compiles, runs on GPU, and all testable kernel/library components execute correctly on gfx1101. + +## End-to-end validation (gfx1100, unsloth weights) + +### Environment + +- GPU: AMD Radeon Pro W7800 48GB (gfx1100), HIP device index 2 +- ROCm: 7.2.1 +- Model weights: unsloth/Llama-3.2-1B-Instruct (model.safetensors, 2.47 GB), downloaded via `hf_hub_download`; no auth required (ungated mirror of the Meta weights) +- Weights verified: all 146 tensor keys match exactly what tiny-vllm's loader expects (standard HF Llama naming: `model.embed_tokens.weight`, `model.layers.N.*`, `model.norm.weight`) + +### Build + +```bash +cd projects/tiny-vllm/src +cmake -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1100 -G Ninja +HIP_VISIBLE_DEVICES=2 cmake --build build +``` + +Build succeeded with only benign `nodiscard` warnings on HIP API return values; identical to prior builds. + +### Run + +```bash +cd agent_space/tiny-vllm-e2e # directory containing model.safetensors +HIP_VISIBLE_DEVICES=2 /path/to/projects/tiny-vllm/src/build/tiny-vllm +``` + +The binary reads `model.safetensors` from CWD. No tokenizer files are needed at runtime -- prompt token IDs are hardcoded in main.cpp and the output is raw token IDs (argmax over logits); detokenization was done offline with the HF tokenizer. + +### Generated output + +The program ran all 4 hardcoded prompts with BATCH_SIZE=2, interleaved. Token IDs were decoded offline using the unsloth tokenizer.json. Reconstructed responses: + +| Prompt | Response | +|--------|----------| +| What is 2+2? | The answer is 4 | +| Name a color. | Blue | +| Say hello. | Hello! How can I help you with anything you need? | +| Capital of France? | Paris | + +Raw token ID sequence (in print order): 791, 10544, 4320, 128009, 9906, 374, 0, 220, 2650, 19, 649, 128009, 358, 60704, 128009, 1520, 499, 449, 4205, 499, 1205, 30, 128009. Program terminated normally: "Ok bye!" + +### Reference comparison + +reference.txt contains intermediate tensor values from a Python reference using a raw (non-chat-format) prompt "The capital of France is" and predicts token 12366 (` Paris` with leading space). Our run uses chat-format prompts; the "Capital of France?" response starts with token 60704 (`Paris`, no leading space). Both decode to "Paris" -- the different token ID reflects the different whitespace context in the prompt format, not a numerical discrepancy. The reference is not a bit-for-bit token stream match target; it is an intermediate-value trace for a different prompt structure. + +No NaN, no garbage output, no crash or GPU fault. + +### Verdict + +**PASS** -- tiny-vllm produces coherent, correct end-to-end inference on gfx1100. All 4 prompts yield sensible answers (2+2=4, blue is a color, hello greeting, Paris is the capital of France). The HIP port runs correctly through the full inference path: weight loading, embedding gather, RMS norm, Q/K/V projections via hipBLAS GEMM, RoPE encoding, paged KV-cache management, causal mask + softmax, attention output projection, SwiGLU MLP, lm_head projection, and greedy argmax decoding -- all on real gfx1100 hardware. + +## PR fix-round 2026-07-02 (copyright line removal) + +Maintainer jmaczan asked (PR #2 review comment on src/cuda_to_hip.h) about the legal implications of the AMD copyright/author lines the port added to the header, given the project is Apache-2.0. + +Action taken: +- Removed the `Copyright (c) 2026 Advanced Micro Devices, Inc.` and `Author: Jeff Daily` lines from `src/cuda_to_hip.h`. This is the only file the port added those lines to (grep across *.h/*.cpp/*.cu/*.cuh confirmed). +- Commit `d6b1ab3` `[ROCm] Remove added copyright line from header`, pushed to AMD-Ecosystem/tiny-vllm moat-port (force-with-lease, ace290f..d6b1ab3). +- `moatlib classify tiny-vllm ace290f d6b1ab3` -> class=comment-only, arch_independent=True, inert=True. `advance-head` carried the completed platforms (gfx1100, gfx1101, gfx1201) forward to d6b1ab3 with no GPU re-run; linux-gfx90a stays pr-open (lead PR state). +- Replied to the maintainer's review thread (comment id 3492833093, reply id 3515308604) explaining the contribution is under the project's Apache-2.0 terms and the copyright line was dropped. Thread left unresolved for the maintainer to close. + +No build/test re-run needed (comment-only). No functional change. diff --git a/projects/tiny-vllm/plan.md b/projects/tiny-vllm/plan.md new file mode 100644 index 00000000..a04554b6 --- /dev/null +++ b/projects/tiny-vllm/plan.md @@ -0,0 +1,217 @@ +# Port plan: tiny-vllm + +## Project + +- **Name**: tiny-vllm +- **Upstream**: https://github.com/jmaczan/tiny-vllm +- **Default branch**: main +- **Description**: Educational project for building your own LLM inference engine in C++ and CUDA. Implements core vLLM concepts (paged attention, continuous batching, KV cache) for Llama 3.2 1B. + +## Existing AMD support + +**None found.** Searched: +- Grepped upstream docs for AMD/ROCm/HIP references: only mention of "AMD CPU (Ryzen 7 9800X3D)" as dev machine, no GPU support +- Web search for "tiny-vllm ROCm/AMD/HIP": no results (web search returned hits for main vLLM project which has mature ROCm support, not this educational tiny-vllm) +- GitHub fork scan (`gh api repos/jmaczan/tiny-vllm/forks`): 30 forks, none with rocm/hip/amd in name +- No upstream ROCm/HIP branches or issues + +**Decision**: Proceed with a from-scratch HIP port using Strategy A (pure CMake, compat-header model). The main vLLM project has mature ROCm support, but this is a separate educational project with hand-written CUDA kernels -- the port adds educational value by demonstrating how to write portable HIP kernels for LLM inference. + +## Build classification + +**Pure CMake project** (Strategy A) + +Evidence (CMakeLists.txt): +- Line 4: `project(tiny-vllm LANGUAGES CXX CUDA)` +- Line 23: `find_package(CUDAToolkit REQUIRED)` +- Line 33-36: Links `CUDA::cublas`, `CUDA::cudart` +- No `find_package(Torch)`, no `CUDAExtension`, no setup.py + +## Port strategy + +**Strategy A: pure CMake, compat-header approach** + +Rationale: This is a standalone CMake build with `.cu` sources and CUDA libraries (cuBLAS). Not a PyTorch extension. The port will: +1. Add a `cuda_to_hip.h` compat header with CUDA-to-HIP symbol aliases +2. Modify CMakeLists.txt to add `USE_HIP` option, `enable_language(HIP)`, mark `.cu` as LANGUAGE HIP +3. Swap cuBLAS for hipBLAS via compat header + +## CUDA surface inventory + +### Files +- `src/kernels.cu` (449 lines): All custom CUDA kernels +- `src/kernels.cuh` (19 lines): Kernel declarations +- `src/main.cpp` (1000+ lines): Host code with CUDA runtime + cuBLAS calls + +### Kernels (all in kernels.cu) +| Kernel | Lines | Purpose | +|--------|-------|---------| +| `embeddingGatherKernel` | 25-33 | Token embedding lookup | +| `rmsNormKernel` | 48-73 | RMS normalization with tree reduction | +| `ropeKernel` | 89-102 | Rotary position embedding | +| `causalMaskKernel` | 125-138 | Causal attention mask | +| `softmaxKernel` | 158-200 | Softmax with online max/sum | +| `residualKernel` | 221-226 | Residual connection add | +| `siluKernel` | 241-248 | SiLU activation with elementwise multiply | +| `embeddingGatherKernelDecode` | 257-266 | Decode-phase embedding | +| `ropeKernelDecode` | 281-293 | Decode-phase RoPE | +| `softmaxKernelDecode` | 318-359 | Decode-phase softmax | +| `pagedAttentionKernel` | 382-444 | Paged attention with online softmax | + +### Warp intrinsics +- `__shfl_down_sync(0xffffffff, qk, delta)` at lines 410-414 in `pagedAttentionKernel` + - Used for tree reduction of dot product within a warp (32 threads) + - The kernel is launched with `HEAD_DIM=64` threads per block + - The shuffle uses hardcoded 32-bit mask `0xffffffff` + +### __syncthreads usage +- Multiple uses for shared memory synchronization in tree reductions (rmsNormKernel, softmaxKernel, softmaxKernelDecode, pagedAttentionKernel) +- These are block-wide barriers with proper shared memory and are wave-size agnostic + +### Data types +- `__nv_bfloat16` throughout (AMD: `__hip_bfloat16` via `hip/hip_bf16.h`) +- Standard float for intermediate computations + +### Libraries +- **cuBLAS**: `cublasGemmEx` for matrix multiplications (Q/K/V/O projections, attention, MLP) + - ROCm equivalent: hipBLAS `hipblasGemmEx` +- **CUDA runtime**: `cudaMalloc`, `cudaMemcpy`, `cudaFree`, `cudaGetDeviceCount`, `cudaGetDeviceProperties`, `cudaMemGetInfo`, `cudaError_t` + - ROCm equivalent: Direct `hip*` equivalents via compat header + +### Memory patterns +- Simple device allocations via `cudaMalloc` +- `cudaMemcpyHostToDevice`, `cudaMemcpyDeviceToDevice` +- No managed memory, no pinned memory, no streams/events +- No textures/surfaces + +### No usage of +- cuRAND, cuFFT, cuSPARSE +- Thrust, CUB +- Cooperative groups +- Dynamic parallelism +- CUDA graphs + +## Risk list + +### High +1. **Warp shuffle mask width** (lines 410-414 in pagedAttentionKernel) + - Issue: `__shfl_down_sync(0xffffffff, ...)` uses 32-bit mask; HIP requires 64-bit mask + - Fix: Define compat macro `#define HIP_FULL_MASK 0xffffffffffffffffULL` for HIP, `0xffffffff` for CUDA + - Per PORTING_GUIDE: "HIP `__shfl_sync`/... REQUIRE a 64-bit mask" + +2. **Warp shuffle reduction on HEAD_DIM=64 threads** + - The kernel launches 64 threads per block, uses two 32-thread warp reductions, then combines via shared memory + - On wave64 (gfx90a): All 64 threads form ONE wavefront; the reduction pattern (16/8/4/2/1) still works correctly because it operates within 32 lanes + - On wave32 (gfx1100): Two warps of 32 threads each; same pattern works + - Verdict: The current code IS wave-size agnostic because it does logical-warp operations (32-wide shuffles) and combines via shared memory + `__syncthreads()` + +### Medium +3. **bfloat16 type mapping** + - Issue: `__nv_bfloat16` needs mapping to `__hip_bfloat16` + - Fix: `#define __nv_bfloat16 __hip_bfloat16` in compat header + include `` + - Also need `nv_bfloat16` -> `__hip_bfloat16` (used without prefix in some places) + +4. **cuBLAS to hipBLAS** + - `cublasGemmEx` -> `hipblasGemmEx` + - `cublasCreate` -> `hipblasCreate` + - `cublasHandle_t` -> `hipblasHandle_t` + - `cublasStatus_t` -> `hipblasStatus_t` + - `CUBLAS_STATUS_SUCCESS` -> `HIPBLAS_STATUS_SUCCESS` + - Compute type: `CUBLAS_COMPUTE_32F` -> `HIPBLAS_COMPUTE_32F` + - Data types: `CUDA_R_16BF` -> `HIP_R_16BF` (from ``) + +5. **Hardcoded CUDA arch in CMakeLists.txt** + - Line 11: `set(CMAKE_CUDA_ARCHITECTURES 120)` (Blackwell) + - Fix: For HIP, use `CMAKE_HIP_ARCHITECTURES` with default fallback pattern + +### Low +6. **Hardcoded NVCC path** + - Lines 2-3: `set(CMAKE_CUDA_COMPILER "/opt/cuda/bin/nvcc")` + - Fix: Only set these when not USE_HIP; CMake finds hipcc automatically + +7. **CUDA include in main.cpp** + - Line 4-5: `#include `, `#include ` + - Fix: Route through compat header + +## File-by-file change list + +### New files +1. **src/cuda_to_hip.h** - CUDA-to-HIP compat header + - Include `` and `` on HIP + - bfloat16 mappings: `__nv_bfloat16` -> `__hip_bfloat16` + - CUDA runtime mappings: `cudaMalloc`, `cudaMemcpy`, `cudaFree`, etc. + - cuBLAS -> hipBLAS mappings + - 64-bit warp mask constant + +### Modified files +1. **CMakeLists.txt** + - Add `option(USE_HIP "Build with HIP for AMD GPUs" OFF)` + - Conditional `enable_language(HIP)` vs `enable_language(CUDA)` + - Move hardcoded NVCC paths inside `if(NOT USE_HIP)` guard + - `set_source_files_properties(src/kernels.cu PROPERTIES LANGUAGE HIP)` when USE_HIP + - Link `hip::host` and `roc::hipblas` instead of `CUDA::cublas` and `CUDA::cudart` + - Set `CMAKE_HIP_ARCHITECTURES` with fallback to gfx90a + +2. **src/main.cpp** + - Replace `#include ` and `#include ` with `#include "cuda_to_hip.h"` + +3. **src/kernels.cu** + - Replace `#include "kernels.cuh"` with `#include "cuda_to_hip.h"` followed by `#include "kernels.cuh"` + - Replace `0xffffffff` in `__shfl_down_sync` calls with `WARP_FULL_MASK` macro + +4. **src/kernels.cuh** + - Replace `#include ` with a guard that includes the right header based on platform + +## Build commands + +### Configure (gfx90a) +```bash +cmake -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a -G Ninja +``` + +### Build +```bash +cmake --build build +``` + +### Alternative for other archs +```bash +cmake -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1100 -G Ninja +cmake --build build +``` + +## Test plan + +### Primary validation +The project has no formal test suite. Validation is via inference output: + +1. **Build and run inference**: + ```bash + cd build && ./tiny-vllm + ``` + This requires Llama 3.2 1B model weights in safetensors format (download from HuggingFace). + +2. **Reference comparison** (full_test.sh): + ```bash + echo "128000 128006 9125 128007 271 38766 1303 33025 2696 25 6790 220 2366 18 198 15724 2696 25 220 2304 2947 220 2366 21 271 128009 128006 882 128007 271 3923 374 279 6864 315 9822 30 128009 128006 78191 128007 271" | ./build/tiny-vllm + ``` + Compare generated tokens to reference.txt (Python reference implementation output). + +3. **Determinism check**: Run the same prompt twice, verify identical output. + +### Non-GPU tests +None -- this is a GPU-only inference engine. + +### Validation criteria +- Builds successfully with hipcc +- Runs inference without GPU faults +- Generates coherent text output matching the expected token sequence +- Run-to-run deterministic output (same prompt -> same tokens) + +## Open questions + +1. **Model weights access**: The validator will need access to Llama 3.2 1B safetensors weights from HuggingFace. This may require HF token/login if model is gated. + +2. **Reference output tolerance**: The reference.txt shows intermediate tensor values from PyTorch. Due to floating-point differences between CUDA/HIP, the HIP output may differ at ULP level. The key validation is that generated tokens match, not intermediate tensors. + +3. **Performance**: This is an educational project, not a production inference engine. Performance validation is not a primary goal, but gross performance issues (orders of magnitude slower) would indicate a bug. diff --git a/projects/tiny-vllm/stats.jsonl b/projects/tiny-vllm/stats.jsonl new file mode 100644 index 00000000..067e2e3a --- /dev/null +++ b/projects/tiny-vllm/stats.jsonl @@ -0,0 +1,34 @@ +{"kind": "tokens", "ts": "2026-06-04T23:29:30Z", "tokens": 72983, "source": "planner"} +{"kind":"phase","ts":"2026-06-05T00:06:46Z","phase":"compile","seconds":0.139,"exit":1,"cmd":"cmake --build build"} +{"kind": "tokens", "ts": "2026-06-05T00:11:31Z", "tokens": 104477, "source": "porter"} +{"kind": "tokens", "ts": "2026-06-05T00:12:42Z", "tokens": 71454, "source": "reviewer"} +{"kind": "tokens", "ts": "2026-06-05T05:17:39Z", "tokens": 71757, "source": "reviewer"} +{"kind":"phase","ts":"2026-06-05T05:18:21Z","phase":"compile","seconds":0.139,"exit":1,"cmd":"cmake -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a -G Ninja"} +{"kind":"phase","ts":"2026-06-05T05:18:33Z","phase":"compile","seconds":9.284,"exit":0,"cmd":"bash -c cd /var/lib/jenkins/moat/projects/tiny-vllm/src && cmake -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a -G Ninja && cmake --build build"} +{"kind": "tokens", "ts": "2026-06-05T05:22:44Z", "tokens": 56926, "source": "reviewer"} +{"kind":"session","ts":"2026-06-05T05:23:09Z","epoch":1780636989.609074674,"event":"start","platform":"linux-gfx90a"} +{"kind":"phase","ts":"2026-06-05T05:23:27Z","phase":"compile","seconds":0.139,"exit":1,"cmd":"cmake -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a -G Ninja"} +{"kind":"phase","ts":"2026-06-05T05:23:36Z","phase":"compile","seconds":0.139,"exit":0,"cmd":"bash -c cmake --build build -j128 2>&1 | tee /tmp/tiny-vllm-build.log"} +{"kind":"phase","ts":"2026-06-05T05:23:46Z","phase":"compile","seconds":6.631,"exit":0,"cmd":"bash -c cd projects/tiny-vllm/src && cmake --build build -j128 2>&1 | tee /tmp/tiny-vllm-build.log"} +{"kind":"phase","ts":"2026-06-05T05:26:44Z","phase":"test","seconds":0.890,"exit":0,"cmd":"bash -c \nset -e\necho '=== tiny-vllm ROCm Port Validation ==='\necho\necho '1. GPU Detection:'\ncd /var/lib/jenkins/moat/projects/tiny-vllm/src\nHIP_VISIBLE_DEVICES=1 timeout 10 ./build/tiny-vllm 2>&1 | head -7\necho\necho '2. Embedding Kernel Test:'\nHIP_VISIBLE_DEVICES=1 /tmp/test_kernel\necho\necho '3. Warp Shuffle (64-bit mask) Test:'\nHIP_VISIBLE_DEVICES=1 /tmp/test_shuffle2\necho\necho '4. hipBLAS bfloat16 GEMM Test:'\nHIP_VISIBLE_DEVICES=1 /tmp/test_gemm\necho\necho '=== All Tests PASSED ==='\n"} +{"kind":"session","ts":"2026-06-05T05:27:38Z","epoch":1780637258.542374176,"event":"end","platform":"linux-gfx90a"} +{"kind":"phase","ts":"2026-06-05T15:45:15Z","phase":"gfx1100-build","seconds":7.000,"exit":0,"cmd":"bash -c cd /var/lib/jenkins/moat/projects/tiny-vllm/src && cmake -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1100 -G Ninja && cmake --build build"} +{"kind":"phase","ts":"2026-06-05T15:45:22Z","phase":"gfx1100-test","seconds":0.125,"exit":1,"cmd":"/var/lib/jenkins/moat/projects/tiny-vllm/src/build/tiny-vllm"} +{"kind":"phase","ts":"2026-06-08T19:30:06Z","phase":"compile","seconds":4.601,"exit":0,"cmd":"cmake --build B:/develop/moat/projects/tiny-vllm/src/build -j24"} +{"kind":"phase","ts":"2026-06-08T19:33:32Z","phase":"test","seconds":0.522,"exit":0,"cmd":"B:/develop/TheRock/external-builds/pytorch/.venv/Scripts/python.exe B:/develop/moat/agent_space/run_tinyvllm_gfx1201.py"} +{"kind": "tokens", "ts": "2026-06-08T20:00:16Z", "tokens": 67402, "source": "validator"} +{"kind": "tokens", "ts": "2026-06-15T12:15:44Z", "tokens": 45603, "source": "pr-prep"} +{"kind": "tokens", "ts": "2026-06-17T01:41:12Z", "tokens": 61459, "source": "pr-prep"} +{"kind":"phase","ts":"2026-06-17T05:59:24Z","phase":"compile","seconds":0.123,"exit":1,"cmd":"bash -c cmake -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1100 -G Ninja 2>&1 && cmake --build build 2>&1"} +{"kind":"phase","ts":"2026-06-17T05:59:56Z","phase":"compile","seconds":4.786,"exit":0,"cmd":"bash -c cd /var/lib/jenkins/moat/projects/tiny-vllm/src && HIP_VISIBLE_DEVICES=2 cmake --build build 2>&1"} +{"kind":"phase","ts":"2026-06-17T06:00:23Z","phase":"test","seconds":2.399,"exit":0,"cmd":"bash -c cd /var/lib/jenkins/moat/agent_space/tiny-vllm-e2e && HIP_VISIBLE_DEVICES=2 /var/lib/jenkins/moat/projects/tiny-vllm/src/build/tiny-vllm 2>&1"} +{"kind": "tokens", "ts": "2026-06-17T06:04:48Z", "tokens": 71303, "source": "validator"} +{"kind": "tokens", "ts": "2026-06-19T19:42:05Z", "tokens": 27326, "source": "porter (doc prefix fix; stopped on sync)"} +{"kind":"session","ts":"2026-06-20T06:53:23Z","epoch":1781938403.305437000,"event":"start","platform":"windows-gfx1101"} +{"kind":"phase","ts":"2026-06-20T06:54:26Z","phase":"compile","seconds":1.824,"exit":0,"cmd":"B:/develop/TheRock/external-builds/pytorch/.venv/Lib/site-packages/_rocm_sdk_devel/lib/llvm/bin/clang++.exe --offload-arch=gfx1101 -x hip -I/b/develop/TheRock/external-builds/pytorch/.venv/Lib/site-packages/_rocm_sdk_devel/include -L/b/develop/TheRock/external-builds/pytorch/.venv/Lib/site-packages/_rocm_sdk_devel/lib -lhipblas -lamdhip64 /b/develop/moat/agent_space/tiny_vllm_validate_gfx1101.cpp -o /b/develop/moat/agent_space/tiny_vllm_validate_gfx1101.exe"} +{"kind":"phase","ts":"2026-06-20T06:54:59Z","phase":"test","seconds":0.102,"exit":127,"cmd":"sh -c HIP_VISIBLE_DEVICES=1 ROCBLAS_TENSILE_LIBPATH=\"/b/develop/TheRock/external-builds/pytorch/.venv/Lib/site-packages/_rocm_sdk_libraries/bin/rocblas/library\" /b/develop/moat/agent_space/tiny_vllm_validate_gfx1101.exe"} +{"kind":"phase","ts":"2026-06-20T06:55:22Z","phase":"test","seconds":0.453,"exit":0,"cmd":"B:/develop/moat/agent_space/tiny_vllm_validate_gfx1101.exe"} +{"kind":"session","ts":"2026-06-20T06:56:30Z","epoch":1781938590.269281600,"event":"end","platform":"windows-gfx1101"} +{"kind": "tokens", "ts": "2026-06-20T06:57:11Z", "tokens": 61121, "source": "validator (gfx1101)"} +{"kind": "tokens", "ts": "2026-06-23T22:25:32Z", "tokens": 20289, "source": "gfx90a dead-code cleanup (batch D subagent, shared)"} +{"kind": "tokens", "ts": "2026-07-02T18:26:12Z", "tokens": 35773, "source": "porter (PR fix round: copyright line removal + reply)"} diff --git a/projects/tiny-vllm/status.json b/projects/tiny-vllm/status.json new file mode 100644 index 00000000..c5f89ad1 --- /dev/null +++ b/projects/tiny-vllm/status.json @@ -0,0 +1,147 @@ +{ + "schema_version": 3, + "name": "tiny-vllm", + "upstream_url": "https://github.com/jmaczan/tiny-vllm", + "fork_url": "https://github.com/AMD-Ecosystem/tiny-vllm", + "fork_default_branch": "main", + "priority": 0.0, + "ext_type": "cmake", + "adopted_at": "2026-06-04T23:25:55Z", + "updated_at": "2026-08-07T07:05:32Z", + "head_sha": "d6b1ab3a24dc86220c9eee136102ec5859684c24", + "depends_on": [], + "pr_url": "https://github.com/jmaczan/tiny-vllm/pull/2", + "pr_number": 2, + "pr_opened_at": "2026-06-17T13:46:29Z", + "pr_merged_at": "2026-07-06T02:49:04Z", + "porting": null, + "waivers": {}, + "pr_state": "merged", + "license_spdx": "Apache-2.0", + "upstream_repo_id": 1153764295, + "stage": "review-passed", + "platforms": { + "linux-gfx90a": { + "state": "completed", + "blocked": false, + "blocked_reason": null, + "validated_sha": "ace290fe32e4597ddb7f86ae6dc58a353eb064af", + "started_at": "2026-06-05T00:05:19Z", + "completed_at": "2026-06-05T05:27:33Z", + "updated_at": "2026-07-06T02:49:04Z", + "stats": { + "tokens_total": 0, + "tokens_approx": true, + "wall_seconds": { + "thinking": 0, + "compile": 0, + "test": 0, + "misc": 0 + }, + "session_count": 0, + "first_session_at": null, + "last_session_at": null + }, + "last_agent": "moat-checkup", + "carry_forward": { + "to": "ace290fe32e4597ddb7f86ae6dc58a353eb064af", + "method": "source-class", + "detail": "gfx90a pin removal; CMake default-path/dead-code change, build byte-identical, no GPU re-run", + "at": "2026-06-23T22:24:33Z" + } + }, + "linux-gfx1100": { + "state": "completed", + "blocked": false, + "blocked_reason": null, + "validated_sha": "d6b1ab3a24dc86220c9eee136102ec5859684c24", + "started_at": null, + "completed_at": "2026-06-23T22:24:33Z", + "updated_at": "2026-07-02T18:24:40Z", + "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": { + "from": "ace290fe32e4597ddb7f86ae6dc58a353eb064af", + "to": "d6b1ab3a24dc86220c9eee136102ec5859684c24", + "method": "source-class", + "class": "comment-only", + "detail": "src/cuda_to_hip.h: comment-only (comments/format only)", + "at": "2026-07-02T18:24:40Z" + } + }, + "windows-gfx1101": { + "state": "completed", + "blocked": false, + "blocked_reason": null, + "validated_sha": "d6b1ab3a24dc86220c9eee136102ec5859684c24", + "started_at": null, + "completed_at": "2026-06-23T22:24:33Z", + "updated_at": "2026-07-02T18:24:40Z", + "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": { + "from": "ace290fe32e4597ddb7f86ae6dc58a353eb064af", + "to": "d6b1ab3a24dc86220c9eee136102ec5859684c24", + "method": "source-class", + "class": "comment-only", + "detail": "src/cuda_to_hip.h: comment-only (comments/format only)", + "at": "2026-07-02T18:24:40Z" + } + }, + "windows-gfx1201": { + "state": "completed", + "blocked": false, + "blocked_reason": null, + "validated_sha": "d6b1ab3a24dc86220c9eee136102ec5859684c24", + "started_at": null, + "completed_at": "2026-06-23T22:24:33Z", + "updated_at": "2026-07-02T18:24:40Z", + "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": { + "from": "ace290fe32e4597ddb7f86ae6dc58a353eb064af", + "to": "d6b1ab3a24dc86220c9eee136102ec5859684c24", + "method": "source-class", + "class": "comment-only", + "detail": "src/cuda_to_hip.h: comment-only (comments/format only)", + "at": "2026-07-02T18:24:40Z" + } + } + } +} From 1e85a5fa0f0eece51349dea879ed47caa5f558c7 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Sun, 9 Aug 2026 15:27:17 +0000 Subject: [PATCH 2/3] tiny-vllm: carry forward linux-gfx90a to d6b1ab3 (comment-only) --- projects/tiny-vllm/notes.md | 23 +++++++++++++++++++++++ projects/tiny-vllm/status.json | 14 +++++++------- 2 files changed, 30 insertions(+), 7 deletions(-) diff --git a/projects/tiny-vllm/notes.md b/projects/tiny-vllm/notes.md index deb014e8..169ed81c 100644 --- a/projects/tiny-vllm/notes.md +++ b/projects/tiny-vllm/notes.md @@ -377,3 +377,26 @@ Action taken: - Replied to the maintainer's review thread (comment id 3492833093, reply id 3515308604) explaining the contribution is under the project's Apache-2.0 terms and the copyright line was dropped. Thread left unresolved for the maintainer to close. No build/test re-run needed (comment-only). No functional change. + +## Validation 2026-08-09 (linux-gfx90a revalidate) + +Fork head had moved past `linux-gfx90a`'s `validated_sha` (ace290f -> d6b1ab3). Cloned the fork fresh into `projects/tiny-vllm/src` (moat-port branch) to classify against a real checkout. + +``` +python3 utils/moatlib.py classify tiny-vllm ace290fe32e4597ddb7f86ae6dc58a353eb064af d6b1ab3a24dc86220c9eee136102ec5859684c24 +-> class=comment-only arch_independent=True inert=True + src/cuda_to_hip.h: comment-only (comments/format only) +``` + +Confirmed with `git diff ace290f d6b1ab3` directly: the only change is removal of two comment lines (`// Copyright (c) 2026 Advanced Micro Devices, Inc.` and `// Author: Jeff Daily `) from `src/cuda_to_hip.h` -- the same PR-fix-round-2026-07-02 commit already carried forward to linux-gfx1100/windows-gfx1101/windows-gfx1201. No code, no CMake, nothing that touches codegen. + +Carried forward with no rebuild and no GPU re-run: +``` +python3 utils/moatlib.py carry-forward tiny-vllm linux-gfx90a d6b1ab3a24dc86220c9eee136102ec5859684c24 source-class "src/cuda_to_hip.h: comment-only ..." +``` + +CUDA no-regression gate: skipped per policy (carried-forward revalidation; the delta cannot affect the CUDA build either). + +Jargon: `python3 utils/jargon.py --port tiny-vllm` -> clean. + +`linux-gfx90a` state: completed, validated_sha now d6b1ab3a24dc86220c9eee136102ec5859684c24 (matches head_sha). No GPU tests re-run; the last real GPU evidence for this arch remains the 2026-06-05 validation above (device detection, embedding gather, 64-bit-mask warp shuffle, hipBLAS bf16 GEMM, all PASS on MI250X). diff --git a/projects/tiny-vllm/status.json b/projects/tiny-vllm/status.json index c5f89ad1..5169a8fd 100644 --- a/projects/tiny-vllm/status.json +++ b/projects/tiny-vllm/status.json @@ -7,7 +7,7 @@ "priority": 0.0, "ext_type": "cmake", "adopted_at": "2026-06-04T23:25:55Z", - "updated_at": "2026-08-07T07:05:32Z", + "updated_at": "2026-08-09T15:26:43Z", "head_sha": "d6b1ab3a24dc86220c9eee136102ec5859684c24", "depends_on": [], "pr_url": "https://github.com/jmaczan/tiny-vllm/pull/2", @@ -25,10 +25,10 @@ "state": "completed", "blocked": false, "blocked_reason": null, - "validated_sha": "ace290fe32e4597ddb7f86ae6dc58a353eb064af", + "validated_sha": "d6b1ab3a24dc86220c9eee136102ec5859684c24", "started_at": "2026-06-05T00:05:19Z", - "completed_at": "2026-06-05T05:27:33Z", - "updated_at": "2026-07-06T02:49:04Z", + "completed_at": "2026-08-09T15:26:43Z", + "updated_at": "2026-08-09T15:26:43Z", "stats": { "tokens_total": 0, "tokens_approx": true, @@ -44,10 +44,10 @@ }, "last_agent": "moat-checkup", "carry_forward": { - "to": "ace290fe32e4597ddb7f86ae6dc58a353eb064af", + "to": "d6b1ab3a24dc86220c9eee136102ec5859684c24", "method": "source-class", - "detail": "gfx90a pin removal; CMake default-path/dead-code change, build byte-identical, no GPU re-run", - "at": "2026-06-23T22:24:33Z" + "detail": "src/cuda_to_hip.h: comment-only (comments/format only), removed AMD copyright/author lines per maintainer request; identical to the delta already carried forward for gfx1100/gfx1101/gfx1201", + "at": "2026-08-09T15:26:43Z" } }, "linux-gfx1100": { From 5cd41c544467186cd1b8dd54607e0448f5680bc1 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Sun, 9 Aug 2026 21:41:47 +0000 Subject: [PATCH 3/3] tiny-vllm: commit the telemetry from this session --- projects/tiny-vllm/stats.jsonl | 1 + 1 file changed, 1 insertion(+) diff --git a/projects/tiny-vllm/stats.jsonl b/projects/tiny-vllm/stats.jsonl index 067e2e3a..6eaa6dbd 100644 --- a/projects/tiny-vllm/stats.jsonl +++ b/projects/tiny-vllm/stats.jsonl @@ -32,3 +32,4 @@ {"kind": "tokens", "ts": "2026-06-20T06:57:11Z", "tokens": 61121, "source": "validator (gfx1101)"} {"kind": "tokens", "ts": "2026-06-23T22:25:32Z", "tokens": 20289, "source": "gfx90a dead-code cleanup (batch D subagent, shared)"} {"kind": "tokens", "ts": "2026-07-02T18:26:12Z", "tokens": 35773, "source": "porter (PR fix round: copyright line removal + reply)"} +{"kind": "tokens", "ts": "2026-08-09T15:27:49Z", "tokens": 48073, "source": "validator"}