Scaffold hippihx HIP op zoo base

.gitignore

@@ -0,0 +1,36 @@
+# Build trees
+/build/
+/build-*/
+/cmake-build-*/
+/out/
+/fatbin/
+
+# HIP / clang artifacts
+*.o
+*.a
+*.so
+*.dylib
+*.hsaco
+*.hipfb
+*.co
+*.bc
+compile_commands.json
+
+# Python
+__pycache__/
+*.py[cod]
+*.egg-info/
+.eggs/
+dist/
+.venv/
+.pytest_cache/
+.mypy_cache/
+.ruff_cache/
+.coverage
+htmlcov/
+
+# Editors / OS
+.idea/
+.vscode/
+*.swp
+.DS_Store

CMakeLists.txt

@@ -0,0 +1,130 @@
+cmake_minimum_required(VERSION 3.21)
+
+# Validate the fatbin slot before enabling languages so a broken host
+# compiler cannot mask a multi-arch misconfigure.
+include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/HippihxArch.cmake")
+
+# Languages are decided after we know whether hipcc is present. CXX is
+# always enabled so this tree stays configure-able in CI without ROCm.
+project(hippihx VERSION 0.0.0 LANGUAGES CXX)
+
+include(CheckLanguage)
+check_language(HIP)
+
+option(HIPPIHX_FORCE_HOST_STUB
+       "Compile .hip sources as host C++ (no hipcc). For layout CI only." OFF)
+
+set(HIPPIHX_HAS_HIP OFF)
+if(NOT HIPPIHX_FORCE_HOST_STUB AND CMAKE_HIP_COMPILER)
+  enable_language(HIP)
+  set(HIPPIHX_HAS_HIP ON)
+  message(STATUS "hippihx: HIP compiler ${CMAKE_HIP_COMPILER}")
+else()
+  if(HIPPIHX_FORCE_HOST_STUB)
+    message(STATUS "hippihx: host stub forced (HIPPIHX_FORCE_HOST_STUB=ON)")
+  else()
+    message(WARNING
+      "hippihx: no HIP compiler — building the host compile stub. "
+      "On V620 use ROCm ${HIPPIHX_ROCM_PIN} hipcc with -DHIPPIHX_ARCH=gfx1030.")
+  endif()
+endif()
+
+if(HIPPIHX_WAVE_UNVERIFIED)
+  message(STATUS "hippihx: fatbin slot ${HIPPIHX_ARCH} (RDNA2 Skillfish; wave unverified)")
+elseif(DEFINED HIPPIHX_WAVE_SIZE)
+  message(STATUS "hippihx: fatbin slot ${HIPPIHX_ARCH} (wave${HIPPIHX_WAVE_SIZE})")
+else()
+  message(STATUS "hippihx: fatbin slot ${HIPPIHX_ARCH}")
+endif()
+message(STATUS "hippihx: ROCm pin for V620 is ${HIPPIHX_ROCM_PIN}")
+
+add_library(hippihx STATIC)
+add_library(hippihx::hippihx ALIAS hippihx)
+
+set_target_properties(hippihx PROPERTIES
+  OUTPUT_NAME "hippihx_${HIPPIHX_ARCH}"
+  ARCHIVE_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/fatbin/${HIPPIHX_ARCH}"
+  POSITION_INDEPENDENT_CODE ON
+  CXX_STANDARD 17
+  CXX_STANDARD_REQUIRED ON
+)
+
+if(HIPPIHX_HAS_HIP)
+  set_target_properties(hippihx PROPERTIES
+    HIP_STANDARD 17
+    HIP_STANDARD_REQUIRED ON
+  )
+endif()
+
+target_include_directories(hippihx PUBLIC
+  "${CMAKE_CURRENT_SOURCE_DIR}/include"
+)
+
+# Shared DOT source: FA / EXL3 / AWQ-W4A16 / moe.shared.
+# Same .hip for every DOT slot (gfx1030, gfx110x, gfx1151, gfx103x, gfx1013).
+# gfx1013 builds; wave is VERIFY — never force -mwavefrontsize32.
+# gfx900 does not load these. Never #ifdef WMMA. Never HSA_OVERRIDE.
+set(HIPPIHX_DOT_SOURCES
+  tiles/attn/fa_fdot2/kernel.hip
+  tiles/gemm/w4a16_fdot2/kernel.hip
+  tiles/gemm/exl3_3inst/kernel.hip
+  tiles/moe/shared/kernel.hip
+)
+
+# Non-DOT contracts compile for every built fatbin (stubs).
+set(HIPPIHX_COMMON_SOURCES
+  tiles/smoke.hip
+  tiles/attn/gdn_scan/kernel.hip
+  tiles/attn/kda_scan/kernel.hip
+  tiles/attn/qsa_indexer/kernel.hip
+  tiles/attn/dsa_nope/kernel.hip
+  tiles/moe/routed/kernel.hip
+  tiles/moe/leftover_bf16/kernel.hip
+  tiles/sequence/causal_conv/kernel.hip
+  tiles/comm/pcie/kernel.hip
+)
+
+set(HIPPIHX_TILE_SOURCES ${HIPPIHX_COMMON_SOURCES})
+if(HIPPIHX_IS_DOT_SLOT)
+  list(APPEND HIPPIHX_TILE_SOURCES ${HIPPIHX_DOT_SOURCES})
+  if(HIPPIHX_WAVE_UNVERIFIED)
+    message(STATUS
+      "hippihx: DOT shared source ON for ${HIPPIHX_ARCH} "
+      "(portable/unoptimized; wave VERIFY — do not force -mwavefrontsize32; "
+      "never HSA_OVERRIDE; bind on arch + wave)")
+  elseif(HIPPIHX_DOT_UNOPTIMIZED)
+    message(STATUS
+      "hippihx: DOT shared source ON for ${HIPPIHX_ARCH} "
+      "(portable/unoptimized; own --offload-arch; never HSA_OVERRIDE)")
+  else()
+    message(STATUS "hippihx: DOT shared source ON for ${HIPPIHX_ARCH} (wave32, no WMMA gate)")
+  endif()
+else()
+  message(STATUS "hippihx: non-DOT slot ${HIPPIHX_ARCH} — no FA/EXL3 DOT (mad_mix / pk_fma)")
+endif()
+
+if(NOT HIPPIHX_HAS_HIP)
+  foreach(_src IN LISTS HIPPIHX_TILE_SOURCES)
+    set_source_files_properties("${_src}" PROPERTIES LANGUAGE CXX)
+  endforeach()
+  target_compile_definitions(hippihx PRIVATE HIPPIHX_HOST_STUB=1)
+endif()
+
+target_sources(hippihx PRIVATE ${HIPPIHX_TILE_SOURCES})
+hippihx_apply_arch_flags(hippihx)
+
+# Host-callable smoke symbol so a later Python extension can link a real .a.
+add_executable(hippihx_smoke_host tiles/smoke_host.cpp)
+target_link_libraries(hippihx_smoke_host PRIVATE hippihx::hippihx)
+target_include_directories(hippihx_smoke_host PRIVATE
+  "${CMAKE_CURRENT_SOURCE_DIR}/include"
+)
+set_target_properties(hippihx_smoke_host PROPERTIES
+  RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/fatbin/${HIPPIHX_ARCH}"
+)
+
+# Convenience: document slots without mixing them into this .a.
+add_custom_target(hippihx_list_fatbin_slots
+  COMMAND ${CMAKE_COMMAND} -E echo
+          "DOT: gfx1030 gfx110x gfx1151 gfx103x gfx1013(BC-250 wave VERIFY); stub gfx900; Later gfx906(Vega20)"
+)

CONTRIBUTING.md

@@ -0,0 +1,78 @@
+# Contributing
+
+## Where work belongs
+
+| Change | Land in |
+|---|---|
+| Tile contract, HIP ISA, LDS / `__launch_bounds__` lock | **this repo** (`tiles/…`) |
+| `plan` / `bind` / `run`, scratch specs | **this repo** (`hippihx/`) |
+| Fatbin / CMake / ROCm pin | **this repo** |
+| vLLM registration, flags, model hooks | `opengfx1030/vllm-rdna` `rdna_extras` (not from this tree) |
+| AWQ / `3inst` produce | packer tools — **not** a hippihx directory |
+| Upstream vLLM | **do not open PRs** |
+
+Do not copy CUDA, CuTe, CE, or NVFP4 objects from
+`local-inference-lab/b12x`. Copy the **plan / bind / run** verbs only.
+
+## ROCm / arch
+
+- V620 = **gfx1030**, **wave32**, **ROCm 7.14**.
+- **gfx1100/1101/1102** are first-class DOT consumers of the **same** FA /
+  EXL3 / AWQ / `moe.shared` source. One `--offload-arch` per fatbin.
+- gfx900 is a Vega stub (`mad_mix` / `pk_fma`). It does **not** load DOT
+  tiles.
+- **BC-250 is gfx1013** (Cyan Skillfish): **RDNA2**, same generation as
+  V620/Deck, built portable/unoptimized. Different GFX than Navi21 or
+  Van Gogh — not Steam Deck. RADV/llama.cpp report Skillfish warp size
+  64 — do not force `-mwavefrontsize32` or `HSA_OVERRIDE`. Serve bind
+  keys on **arch + wave size**. **gfx906 is Vega20/MI50**, Later non-DOT,
+  **not** BC-250.
+- gfx1151 and Deck gfx103x (including Steam Deck **gfx1033 / wave32**)
+  are built portable DOT fatbins (same stubs, can run, not dest-tuned).
+- CMake must refuse multi-arch lists and gfx906, and never imply
+  `HSA_OVERRIDE` / foreign ISA load.
+- Shared DOT tiles: **wave32 only**, **no `fdot2.bf16`**, **never
+  `#ifdef WMMA`**. WMMA is a gfx110x-only Later overlay.
+
+## Landing a tile
+
+1. Directory already exists under `tiles/<class>/…`. Do not name it after a
+   product (no `qwen38/`, no `glm53/`).
+2. Fill the tile README **locks** (LDS bytes, `__launch_bounds__`, wave,
+   DOT unit) *before* claiming the kernel is real.
+3. One HIP entry that will become one `torch.ops` / V1 symbol. No second
+   Triton path in this library.
+4. Scratch sized in `plan`. Serve zeros it for page-commit. `bind` views
+   only. `run` is capture-safe (no D2H).
+5. Host tests for the protocol stay torch-free until a HIP extension exists.
+6. If the tile is DOT, include `hippihx/dot.hpp` (not a WMMA header) and
+   mark `make_op(..., dot=True)`. Do not add a per-SKU copy of the file.
+   gfx1151 / gfx103x / gfx1013 compile the same header (portable, not
+   dest-tuned). Deck is wave32; gfx1013 does not force wave32. Never
+   HSA_OVERRIDE a dest object onto those chips.
+
+## Build checks
+
+```bash
+# Layout / host stub (no ROCm). Prefer g++ if clang cannot find libstdc++.
+cmake -S . -B build -DHIPPIHX_FORCE_HOST_STUB=ON -DCMAKE_CXX_COMPILER=g++
+cmake --build build
+./build/fatbin/gfx1030/hippihx_smoke_host
+
+# Silicon (V620)
+cmake -S . -B build -DHIPPIHX_ARCH=gfx1030
+cmake --build build
+
+pip install -e ".[dev]"
+pytest
+```
+
+Dump `hipcc -Rpass-analysis=kernel-resource-usage` (VGPR / LDS /
+`waves_per_eu`, `v_dot2c`) in the tile README when ISA lands. Spill or a
+bad ISel is a drop, not a “tune later”.
+
+## Python
+
+Keep `hippihx` importable without torch and without a built fatbin. Optional
+HIP extensions can load lazily later. `list_ops()` must stay in lockstep
+with `tiles/` class directories.