diff --git a/.github/workflows/spirv-ci-linux.yml b/.github/workflows/spirv-ci-linux.yml index 0e64947bc..72c92bb4d 100644 --- a/.github/workflows/spirv-ci-linux.yml +++ b/.github/workflows/spirv-ci-linux.yml @@ -56,26 +56,46 @@ jobs: # ---- Build LLVM + Clang + LLD + in-tree SPIRV translator -------------- - name: Configure LLVM + # CMAKE_INSTALL_LIBDIR=lib because the manylinux container's + # GNUInstallDirs defaults to lib64, which trips up downstream + # CMake projects (rocm-examples uses CMAKE_HIP_COMPILER_ROCM_ROOT + # which expects /lib). + # compiler-rt is included with builtins only (all other features + # off): hip-lang-config.cmake hard-requires libclang_rt.builtins.a + # for any downstream project that does `enable_language(HIP)`. run: | cmake -G Ninja -S llvm-project/llvm -B build \ -DCMAKE_BUILD_TYPE=Release \ + -DCMAKE_INSTALL_LIBDIR=lib \ -DLLVM_ENABLE_ASSERTIONS=ON \ -DLLVM_ENABLE_PROJECTS="clang;lld" \ + -DLLVM_ENABLE_RUNTIMES="compiler-rt" \ -DLLVM_TARGETS_TO_BUILD="AMDGPU;X86;SPIRV" \ -DLLVM_INCLUDE_TESTS=ON \ -DLLVM_INSTALL_GTEST=ON \ - -DLLVM_LIT_ARGS="-sv --no-progress-bar" + -DLLVM_LIT_ARGS="-sv --no-progress-bar" \ + -DCOMPILER_RT_BUILD_BUILTINS=ON \ + -DCOMPILER_RT_BUILD_LIBFUZZER=OFF \ + -DCOMPILER_RT_BUILD_MEMPROF=OFF \ + -DCOMPILER_RT_BUILD_PROFILE=OFF \ + -DCOMPILER_RT_BUILD_SANITIZERS=OFF \ + -DCOMPILER_RT_BUILD_XRAY=OFF \ + -DCOMPILER_RT_BUILD_CTX_PROFILE=OFF \ + -DCOMPILER_RT_BUILD_ORC=OFF \ + -DCOMPILER_RT_BUILD_GWP_ASAN=OFF - name: Build LLVM + Clang + amd-llvm-spirv + test deps # *-test-depends pull in all tools needed for lit (FileCheck, not, # llc, llvm-*, clang, opt, etc.) and stay current with upstream. - run: ninja -C build llvm-test-depends clang-test-depends amd-llvm-spirv + # `builtins` builds libclang_rt.builtins.a (see Configure comment). + run: ninja -C build llvm-test-depends clang-test-depends amd-llvm-spirv builtins # ---- Build device-libs (standalone, against built LLVM) -------------- - name: Configure device-libs run: | cmake -G Ninja -S llvm-project/amd/device-libs -B build-device-libs \ -DCMAKE_BUILD_TYPE=Release \ + -DCMAKE_INSTALL_LIBDIR=lib \ -DCMAKE_PREFIX_PATH=$PWD/build \ -DLLVM_DIR=$PWD/build/lib/cmake/llvm @@ -87,9 +107,13 @@ jobs: # LLVM_EXTERNAL_SPIRV_LLVM_TRANSLATOR_SOURCE_DIR points Comgr at # the in-tree translator headers so COMGR_SPIRV_TRANSLATOR_AVAILABLE # turns ON (otherwise translator-dependent lit tests are UNSUPPORTED). + # CMAKE_INSTALL_LIBDIR=lib so the install lands in staging/lib (not + # lib64) and the runtime resolver finds libamd_comgr.so.3 there + # instead of falling through to /opt/rocm on the test runner. run: | cmake -G Ninja -S llvm-project/amd/comgr -B build-comgr \ -DCMAKE_BUILD_TYPE=Release \ + -DCMAKE_INSTALL_LIBDIR=lib \ -DCMAKE_PREFIX_PATH="$PWD/build;$PWD/build-device-libs" \ -DLLVM_DIR=$PWD/build/lib/cmake/llvm \ -DLLVM_EXTERNAL_SPIRV_LLVM_TRANSLATOR_SOURCE_DIR=$PWD/llvm-project/llvm/projects/SPIRV-LLVM-Translator \ @@ -98,23 +122,125 @@ jobs: - name: Build Comgr run: ninja -C build-comgr amd_comgr + # ---- In-pipeline HIP runtime build ---------------------------------- + # Builds ROCR-Runtime (libhsa-runtime64.so) and CLR (libamdhip64.so) + # against the just-built clang, plus installs comgr + device-libs into + # the same staging tree. Result: a self-contained ROCm install layout + # under `staging/` that PR clang can target via `--rocm-path` to + # compile and run HIP for `--offload-arch=amdgcnspirv`. + # + # Why in-pipeline (vs pulling a prebuilt nightly tarball): TheRock's + # public `therock-nightly-tarball` bucket only publishes consumer + # (dgpu) families on a daily cadence; gfx94X / gfx950 (dcgpu) tarballs + # stopped Oct 2025, and TheRock's own multi-arch CI for those families + # builds the runtime in-pipeline for the same reason. + - name: Install build deps for runtime build + run: | + dnf install -y --setopt=install_weak_deps=False \ + numactl-devel elfutils-libelf-devel libdrm-devel || true + + - name: Install comgr + device-libs into staging + # CMake `--install` lays them out under /include/amd_comgr, + # /lib/cmake/amd_comgr, /amdgcn/bitcode/ — the + # installed-tree layout that clang's `--rocm-path` and CLR's + # find_package(amd_comgr) both expect. Build-dir layouts won't do. + run: | + cmake --install build-comgr --prefix $PWD/staging + cmake --install build-device-libs --prefix $PWD/staging + + - name: Checkout rocm-systems (pinned) + uses: actions/checkout@de0fac2e4500dabe0009e67214ff5f5447ce83dd # v6.0.2 + with: + repository: ROCm/rocm-systems + ref: 8bb3b73c117e5630106540447268ccad771906a4 # develop, 2026-05-15 + path: rocm-systems + fetch-depth: 1 + persist-credentials: false + + - name: Checkout rocm-cmake (pinned) + # Provides ROCM cmake modules (ROCMConfig) that CLR's build needs. + uses: actions/checkout@de0fac2e4500dabe0009e67214ff5f5447ce83dd # v6.0.2 + with: + repository: ROCm/rocm-cmake + ref: 4d391d7bda62a0f476dbadb383e99f4508c7415d # develop, 2026-05-12 + path: rocm-cmake + fetch-depth: 1 + persist-credentials: false + + - name: Configure ROCR-Runtime + # LLVM_DIR/Clang_DIR are explicit because rocr-runtime's + # trap_handler/CMakeLists.txt does find_package(Clang) and will + # otherwise grab a stale system-installed clang config. + # CMAKE_INSTALL_LIBDIR=lib because the manylinux container defaults + # to lib64 and clang's --rocm-path looks at /lib. + run: | + cmake -G Ninja -S rocm-systems/projects/rocr-runtime -B build-rocr \ + -DCMAKE_BUILD_TYPE=Release \ + -DCMAKE_C_COMPILER=$PWD/build/bin/clang \ + -DCMAKE_CXX_COMPILER=$PWD/build/bin/clang++ \ + -DBUILD_SHARED_LIBS=ON \ + -DCMAKE_INSTALL_PREFIX=$PWD/staging \ + -DCMAKE_INSTALL_LIBDIR=lib \ + -DLLVM_DIR=$PWD/build/lib/cmake/llvm \ + -DClang_DIR=$PWD/build/lib/cmake/clang + + - name: Build ROCR-Runtime + # ROCM_PATH points clang at our staging tree when ROCR's image lib + # compiles OpenCL blit bitcode (no /opt/rocm in the CI container). + env: + ROCM_PATH: ${{ github.workspace }}/staging + run: ninja -C build-rocr install + + - name: Configure CLR (HIP runtime) + # ROCM_KPACK_ENABLED, HIP_ENABLE_ROCPROFILER_REGISTER, HIPCC_BIN_DIR= + # are explicit-disable knobs for optional deps we don't need; without + # them the configure errors out with REQUIRED-not-found. + run: | + cmake -G Ninja -S rocm-systems/projects/clr -B build-clr \ + -DCMAKE_BUILD_TYPE=Release \ + -DCMAKE_C_COMPILER=$PWD/build/bin/clang \ + -DCMAKE_CXX_COMPILER=$PWD/build/bin/clang++ \ + -DBUILD_SHARED_LIBS=ON \ + -DCMAKE_INSTALL_PREFIX=$PWD/staging \ + -DCMAKE_INSTALL_LIBDIR=lib \ + -DCMAKE_PREFIX_PATH="$PWD/staging;$PWD/build;$PWD/rocm-cmake/share/rocm" \ + -DHIP_PLATFORM=amd \ + -DHIP_COMMON_DIR=$PWD/rocm-systems/projects/hip \ + -DCLR_BUILD_HIP=ON \ + -DROCM_KPACK_ENABLED=OFF \ + -DHIP_ENABLE_ROCPROFILER_REGISTER=OFF \ + -DHIPCC_BIN_DIR= + + - name: Build CLR + run: ninja -C build-clr install + + - name: Stage translator binary for runtime SPIRV JIT + # HIP runtime calls Comgr at kernel-launch to translate the embedded + # SPIR-V to AMDGPU code. Comgr looks for the translator at + # /../bin/amd-llvm-spirv. + run: | + mkdir -p staging/bin + cp -a build/bin/amd-llvm-spirv staging/bin/ + # ---- Strip + upload artifact ----------------------------------------- # Strip binaries to keep the artifact under GHA's 10GB cap and shorten # upload/download time. Tests don't need debug symbols. `--strip-unneeded` # preserves dynamic symbols needed at link/load time. - name: Strip binaries run: | - find build build-comgr build-device-libs \ + find build build-comgr build-device-libs build-rocr build-clr staging \ -type f \( -executable -o -name '*.so*' -o -name '*.a' \) \ -exec strip --strip-unneeded {} + 2>/dev/null || true - # Tar before upload: actions/upload-artifact@v4 strips +x bits and + # Tar before upload: actions/upload-artifact@v7 strips +x bits and # excludes hidden files (loses FetchContent .git dirs). - name: Tar build trees - run: tar -cf linux-build-tree.tar build build-comgr build-device-libs + run: | + tar -cf linux-build-tree.tar \ + build build-comgr build-device-libs build-rocr build-clr staging - name: Upload build tree artifact - uses: actions/upload-artifact@v4 + uses: actions/upload-artifact@v7 with: name: linux-build-tree path: linux-build-tree.tar @@ -157,7 +283,7 @@ jobs: persist-credentials: false - name: Download build tree artifact - uses: actions/download-artifact@v4 + uses: actions/download-artifact@v8 with: name: linux-build-tree @@ -260,7 +386,7 @@ jobs: persist-credentials: false - name: Download build tree artifact - uses: actions/download-artifact@v4 + uses: actions/download-artifact@v8 with: name: linux-build-tree @@ -303,7 +429,7 @@ jobs: persist-credentials: false - name: Download build tree artifact - uses: actions/download-artifact@v4 + uses: actions/download-artifact@v8 with: name: linux-build-tree @@ -331,3 +457,193 @@ jobs: env: AMD_COMGR_REDIRECT_LOGS: stderr run: ctest --test-dir build-comgr --output-on-failure --rerun-failed + + # ===================================================================== + # Test - rocm-examples (compile + run hello-world via amdgcnspirv) + # ===================================================================== + # Drives PR clang through the SPIRV translator on a real HIP example + # and JIT-runs the SPIR-V on a gfx942 GPU. Catches translator/codegen + # bugs that pass lit but fail at runtime kernel-load. Informational + # initially; promote to required once stable. + test_rocm_examples: + name: Test rocm-examples + needs: build + runs-on: linux-gfx942-1gpu-ossci-rocm + timeout-minutes: 45 + # Run inside the same manylinux container as the Build job (glibc + # consistency with the built libs) with GPU passthrough — keeps the + # job isolated from whatever ROCm install or Comgr cache might exist + # on the persistent host runner. + container: + image: ghcr.io/rocm/therock_build_manylinux_x86_64@sha256:702a5133851e6d1daf1207d2c9fbb01c2667914a5b6dc5a01faeb3ce66ea6421 + options: --device=/dev/kfd --device=/dev/dri --group-add video + + steps: + - name: Checkout rocm-examples (pinned) + uses: actions/checkout@de0fac2e4500dabe0009e67214ff5f5447ce83dd # v6.0.2 + with: + repository: ROCm/rocm-examples + ref: e260595e2ec06f8a8714bd67bb1bc93d6c27e589 # amd-staging, 2026-05-28 + path: rocm-examples + fetch-depth: 1 + persist-credentials: false + + - name: Download build tree artifact + uses: actions/download-artifact@v8 + with: + name: linux-build-tree + + - name: Untar build trees + run: tar -xmf linux-build-tree.tar + + - name: Install runtime deps in container + run: dnf install -y numactl-libs + + # ---- HIP-Basic: build via rocm-examples' own CMake ------------------- + - name: Build HIP-Basic via rocm-examples CMake + # Uses rocm-examples' upstream CMakeLists, which already auto-skips + # SPIR-V-incompatible examples (assembly_to_executable, + # llvm_ir_to_executable, opengl/vulkan_interop, hello_world_cuda, + # hipify, sobel_filter) via find_package / GPU_RUNTIME gates. + # The HIP_PLATFORM=amd + ROCM_PATH=$STAGING combo bypasses the + # hipconfig probe that hip-config.cmake does (hipconfig isn't + # staged). CMAKE_TRY_COMPILE_TARGET_TYPE=STATIC_LIBRARY avoids + # try_compile linking against an unfound HIP runtime during + # find_package. + run: | + STAGING=$PWD/staging + LLVM_PATH=$PWD/build + cmake -S rocm-examples/HIP-Basic -B examples-build \ + -DCMAKE_BUILD_TYPE=Release \ + -DCMAKE_CXX_COMPILER=$LLVM_PATH/bin/clang++ \ + -DCMAKE_HIP_COMPILER=$LLVM_PATH/bin/clang++ \ + -DCMAKE_HIP_COMPILER_ROCM_ROOT=$STAGING \ + -DCMAKE_HIP_ARCHITECTURES=amdgcnspirv \ + -DCMAKE_TRY_COMPILE_TARGET_TYPE=STATIC_LIBRARY \ + -DCMAKE_DISABLE_FIND_PACKAGE_Perl=TRUE \ + -DHIP_PLATFORM=amd \ + -DROCM_PATH=$STAGING \ + -DCMAKE_HIP_FLAGS="--rocm-path=$STAGING --rtlib=libgcc -unwindlib=libgcc" \ + -DCMAKE_HIP_LINK_FLAGS="--rtlib=libgcc -unwindlib=libgcc -L$STAGING/lib -Wl,-rpath,$STAGING/lib -Wl,--unresolved-symbols=ignore-in-shared-libs" \ + -DCMAKE_CXX_FLAGS="--rtlib=libgcc -unwindlib=libgcc" \ + -DCMAKE_EXE_LINKER_FLAGS="--rtlib=libgcc -unwindlib=libgcc -L$STAGING/lib -Wl,-rpath,$STAGING/lib -Wl,--unresolved-symbols=ignore-in-shared-libs" \ + -DCMAKE_SHARED_LINKER_FLAGS="--rtlib=libgcc -unwindlib=libgcc" + cmake --build examples-build -j + + # ---- Applications + Tutorials: bash loop (their CMake pulls in -------- + # hipcub/rocrand/rocfft find_packages we don't have staged) ----------- + - name: Compile rocm-examples Applications + Tutorials + # Builds the example sets ROCm/RocmCIForSPIRV's + # spirv_external_llvm.sh covers that aren't in HIP-Basic: + # - APPLICATIONS: Applications/ + External/ include (6) + # - REDUCTION: Tutorials/reduction v1..v9, c++20 (9) + # HIP-Basic itself is built via CMake above. + run: | + set +e + STAGING=$PWD/staging + CLANG=$PWD/build/bin/clang++ + + APPLICATIONS=( + bitonic_sort convolution fdtd floyd_warshall histogram prefix_sum + ) + REDUCTION_VERSIONS=(v1 v2 v3 v4 v5 v6 v7 v8 v9) + + BASE_FLAGS=( + -x hip --offload-arch=amdgcnspirv + -D__HIP_PLATFORM_AMD__ + --rocm-path=$STAGING + --rocm-device-lib-path=$STAGING/amdgcn/bitcode + -L $STAGING/lib -Wl,-rpath,$STAGING/lib + -Wl,--unresolved-symbols=ignore-in-shared-libs + -lamdhip64 -lhsa-runtime64 + ) + + BUILD_FAILED=() + BUILD_OK=() + + build_one() { + local name="$1" src="$2" extra_inc="$3" cxxstd="${4:-c++17}" + if [ ! -f "$src" ]; then + echo "SKIP $name (no source at $src)" + return 0 + fi + if $CLANG -std=$cxxstd "${BASE_FLAGS[@]}" \ + -I rocm-examples/Common $extra_inc \ + -o "$name" "$src" 2>/tmp/build_err.txt; then + BUILD_OK+=("$name") + echo "BUILD OK: $name" + else + BUILD_FAILED+=("$name") + echo "BUILD FAIL: $name" + head -8 /tmp/build_err.txt + fi + } + + for ex in "${APPLICATIONS[@]}"; do + build_one "$ex" "rocm-examples/Applications/$ex/main.hip" "-I rocm-examples/External" + done + for v in "${REDUCTION_VERSIONS[@]}"; do + build_one "reduction_$v" "rocm-examples/Tutorials/reduction/example/$v.hip" \ + "-I rocm-examples/Tutorials/reduction/include" "c++20" + done + + echo "" + echo "=== Compile summary (Applications + Tutorials) ===" + echo "Built: ${#BUILD_OK[@]}" + echo "Failed: ${#BUILD_FAILED[@]}" + for f in "${BUILD_FAILED[@]}"; do echo " - $f"; done + + printf '%s\n' "${BUILD_OK[@]}" > built_examples.txt + + [ ${#BUILD_FAILED[@]} -eq 0 ] + + - name: Run rocm-examples + # Runs everything built: HIP-Basic binaries from examples-build/bin + # plus the Applications + Tutorials binaries from the bash-loop + # compile. Gates on exit code 0 (60s timeout each). + # + # hip_cooperative_groups is currently a known failure + # ("Cannot select intrinsic %llvm.amdgcn.s.wait.asynccnt") and is + # intentionally left in the gating set — the whole point of this + # job is to keep surfacing this kind of break until the AMDGPU + # backend gets a lowering pattern. Remove this note once fixed. + run: | + set +e + STAGING=$PWD/staging + export LD_LIBRARY_PATH=$STAGING/lib + + RUN_OK=() + RUN_FAILED=() + + run_one() { + local exe="$1" label="$2" + if timeout 60 "$exe" > "/tmp/$(basename $exe).out" 2>&1; then + RUN_OK+=("$label") + echo "RUN OK: $label" + else + local rc=$? + RUN_FAILED+=("$label") + echo "RUN FAIL: $label (exit $rc)" + head -10 "/tmp/$(basename $exe).out" + fi + } + + # HIP-Basic (CMake build → examples-build/bin/) + for exe in examples-build/bin/*; do + [ -x "$exe" ] || continue + run_one "$exe" "$(basename $exe)" + done + + # Applications + Tutorials (bash-loop build → cwd) + while IFS= read -r exe; do + [ -z "$exe" ] && continue + run_one "./$exe" "$exe" + done < built_examples.txt + + echo "" + echo "=== Run summary ===" + echo "Passed: ${#RUN_OK[@]}" + echo "Failed: ${#RUN_FAILED[@]}" + for f in "${RUN_FAILED[@]}"; do echo " - $f"; done + + [ ${#RUN_FAILED[@]} -eq 0 ]