Repository navigation
chore: prepare shared host and Tile crates for v0.5.0 #2763
Workflow file for this run
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| --- | |
| name: examples-compile | |
| # Path-filtered SIMT lane (same roots as ci.yml). | |
| on: | |
| push: | |
| branches: [main] | |
| paths: | |
| - 'cuda-oxide/**' | |
| - 'cuda-bindings/**' | |
| - 'cuda-core/**' | |
| - 'cuda-core-derive/**' | |
| - 'cuda-async/**' | |
| - 'Cargo.toml' | |
| - 'Cargo.lock' | |
| - '.github/workflows/**' | |
| pull_request: | |
| branches: [main] | |
| paths: | |
| - 'cuda-oxide/**' | |
| - 'cuda-bindings/**' | |
| - 'cuda-core/**' | |
| - 'cuda-core-derive/**' | |
| - 'cuda-async/**' | |
| - 'Cargo.toml' | |
| - 'Cargo.lock' | |
| - '.github/workflows/**' | |
| # Each commit on main gets its own group, so it starts at once. Any other ref | |
| # still shares one group and cancels the older run. | |
| concurrency: | |
| group: ${{ github.workflow }}-${{ github.ref_name }}${{ github.ref_name == 'main' && format('-{0}', github.sha) || '' }} | |
| cancel-in-progress: ${{ github.ref_name != 'main' }} | |
| permissions: | |
| contents: read | |
| jobs: | |
| examples-compile: | |
| name: compile all examples (device pipeline, GPU-less) | |
| runs-on: ubuntu-latest | |
| # Every example crate: backend build + shared dep tree + per-example | |
| # device codegen. Typically well under an hour with the shared | |
| # CARGO_TARGET_DIR; the timeout leaves headroom for cold runners. | |
| # Leave job-level headroom around the 50-minute full compile and the | |
| # focused 10-minute materialization gate, plus toolkit/LLVM setup. | |
| timeout-minutes: 75 | |
| steps: | |
| - name: Clone repository | |
| uses: actions/checkout@v4 | |
| # Every example workspace must pin the same dependency versions. The | |
| # shared CARGO_TARGET_DIR below keys build caches on the resolved | |
| # dependency subtree, so one drifted Cargo.lock silently costs a full | |
| # dependency-tree rebuild (bindgen over cuda.h included) for that one | |
| # example -- a 15-70s spike in an otherwise ~2s-per-example run. | |
| # Fails fast, before any toolkit install. | |
| - name: Example lockfiles resolve identical dependency versions | |
| run: bash cuda-oxide/scripts/sync-example-locks.sh --check | |
| # Needed for cuda.h and curand.h: bindgen runs in the shared | |
| # cuda-bindings build script. Same setup as unit-tests.yml / clippy.yml. | |
| - name: Install CUDA toolkit | |
| uses: Jimver/cuda-toolkit@3d45d157f327c09c04b50ee6ccdea2d9d017ec76 # v0.2.35 | |
| id: cuda-toolkit | |
| with: | |
| cuda: "13.1.1" | |
| method: network | |
| sub-packages: '["nvcc", "cudart-dev"]' | |
| # nvJitLink is packaged separately from cuda-nvcc on Linux. Keep it | |
| # in this already-toolkit-backed job so the materialization gate | |
| # below exercises the real public command without adding work to | |
| # the latency-sensitive fmt job. | |
| non-cuda-sub-packages: '["libcurand-dev", "libnvjitlink-dev"]' | |
| # The device pipeline shells out to LLVM's llc/opt for PTX | |
| # generation. This pin is the FLOOR GATE: if someone lands IR that | |
| # needs a newer llc, this job is where the floor bump becomes | |
| # visible instead of on contributors' machines. LLVM 23 is the true | |
| # floor: the bf16x2 to packed-FP8 conversions (PTX 9.1 forms, used | |
| # by generated_intrinsics_blackwell) need `+ptx91`, which llc-22 | |
| # does not know. It ignored the unknown feature and emitted | |
| # `.version 8.8` PTX that failed only at JIT; the pipeline now | |
| # refuses a PTX 9.1 request below LLVM 23 with a named error, so an | |
| # under-floor llc fails loudly here. The previous floor, LLVM 22 for | |
| # the ldmatrix/stmatrix family (llc-21 leaked unknown intrinsics | |
| # into the PTX as `.extern .func`), is covered by the same gate. | |
| # If a future ubuntu-latest codename has no llvm-toolchain-*-23 | |
| # suite, apt-get update fails loudly; recover by pinning runs-on to | |
| # the previous image or bumping the LLVM major here. | |
| - name: Install LLVM 23 (llc/opt for the NVPTX backend) | |
| run: | | |
| set -euo pipefail | |
| source /etc/os-release | |
| sudo install -d -m 0755 /etc/apt/keyrings | |
| wget -qO- https://apt.llvm.org/llvm-snapshot.gpg.key \ | |
| | sudo tee /etc/apt/keyrings/apt-llvm-org.asc >/dev/null | |
| echo "deb [signed-by=/etc/apt/keyrings/apt-llvm-org.asc] https://apt.llvm.org/${VERSION_CODENAME}/ llvm-toolchain-${VERSION_CODENAME}-23 main" \ | |
| | sudo tee /etc/apt/sources.list.d/llvm-23.list | |
| sudo apt-get update | |
| sudo apt-get install -y --no-install-recommends llvm-23 | |
| llc-23 --version | head -2 | |
| opt-23 --version | head -2 | |
| - name: Compile every example (smoketest --compile-only) | |
| # Step-level timeout, below the job's 75: a job-level timeout | |
| # CANCELS the job and `if: failure()` never fires, losing the | |
| # logs. A step timeout fails the step, so the upload still runs. | |
| timeout-minutes: 50 | |
| env: | |
| CUDA_TOOLKIT_PATH: ${{ steps.cuda-toolkit.outputs.CUDA_PATH }} | |
| CUDA_OXIDE_LLC: /usr/bin/llc-23 | |
| # Each example is its own cargo workspace; without a shared | |
| # target dir the dep tree rebuilds ~85 times and blows the | |
| # runner's ~14 GB disk (same trick as clippy.yml). smoketest.sh | |
| # honours an externally-set CARGO_TARGET_DIR and still builds | |
| # the backend .so at its fixed discovery path. | |
| CARGO_TARGET_DIR: ${{ runner.temp }}/examples-shared-target | |
| SMOKETEST_LOG_DIR: ${{ runner.temp }}/smoketest-logs | |
| # The runner's ~14 GB disk also holds the CUDA toolkit, LLVM, | |
| # cargo-oxide, and ~200 example workspaces; rustc incremental dirs | |
| # under rustc-codegen-cuda/target/ are what used to fill it and fail | |
| # later examples with ENOSPC / ld bus errors. The cutile-rs source | |
| # tree itself is a few MB and must stay: cutile_inter_kernel | |
| # path-depends on cutile-rs/cutile and cutile-rs/cutile-compiler. | |
| CARGO_INCREMENTAL: "0" | |
| run: | | |
| set -euo pipefail | |
| sysroot="$(rustc --print sysroot)" | |
| export LIBRARY_PATH="${sysroot}/lib:${LIBRARY_PATH:-}" | |
| export LD_LIBRARY_PATH="${sysroot}/lib:${LD_LIBRARY_PATH:-}" | |
| cuda-oxide/scripts/smoketest.sh --compile-only --no-color | |
| - name: Artifact modes keep shared Cargo dependencies fresh | |
| timeout-minutes: 10 | |
| working-directory: cuda-oxide | |
| env: | |
| CUDA_TOOLKIT_PATH: ${{ steps.cuda-toolkit.outputs.CUDA_PATH }} | |
| CUDA_OXIDE_LLC: /usr/bin/llc-23 | |
| CARGO_TARGET_DIR: ${{ runner.temp }}/examples-shared-target | |
| CARGO_INCREMENTAL: "0" | |
| run: | | |
| set -euo pipefail | |
| sysroot="$(rustc --print sysroot)" | |
| export LIBRARY_PATH="${sysroot}/lib:${LIBRARY_PATH:-}" | |
| export LD_LIBRARY_PATH="${CUDA_TOOLKIT_PATH}/lib64:${sysroot}/lib:${LD_LIBRARY_PATH:-}" | |
| assert_embedded_artifact() { | |
| local binary="$1" | |
| local expected_version="$2" | |
| local expected_kind="$3" | |
| local expected_target="$4" | |
| local expected_strings=("${@:5}") | |
| local section_file="${RUNNER_TEMP}/$(basename "${binary}")-${expected_kind}.oxart" | |
| objcopy --dump-section ".oxart=${section_file}" "${binary}" | |
| local version name_len target_len target record_offset kind data_offset data_len magic | |
| version="$(od -An -tu2 -j8 -N2 "${section_file}" | tr -d '[:space:]')" | |
| name_len="$(od -An -tu2 -j16 -N2 "${section_file}" | tr -d '[:space:]')" | |
| target_len="$(od -An -tu2 -j18 -N2 "${section_file}" | tr -d '[:space:]')" | |
| target="$(od -An -tc -j$((32 + name_len)) -N"${target_len}" "${section_file}" | tr -d '[:space:]')" | |
| record_offset=$((32 + name_len + target_len)) | |
| kind="$(od -An -tu2 -j"${record_offset}" -N2 "${section_file}" | tr -d '[:space:]')" | |
| data_offset="$(od -An -tu4 -j$((record_offset + 4)) -N4 "${section_file}" | tr -d '[:space:]')" | |
| data_len="$(od -An -tu4 -j$((record_offset + 8)) -N4 "${section_file}" | tr -d '[:space:]')" | |
| test "${version}" = "${expected_version}" | |
| test "${target}" = "${expected_target}" | |
| test "${kind}" = "${expected_kind}" | |
| if [ "${expected_kind}" = "512" ]; then | |
| magic="$(od -An -tx1 -j"${data_offset}" -N4 "${section_file}" | tr -d '[:space:]')" | |
| test "${magic}" = "7f454c46" | |
| if [ "${#expected_strings[@]}" -gt 0 ]; then | |
| local payload_file="${section_file}.cubin" | |
| dd if="${section_file}" of="${payload_file}" bs=1 \ | |
| skip="${data_offset}" count="${data_len}" status=none | |
| local expected_string | |
| for expected_string in "${expected_strings[@]}"; do | |
| grep -Fqx "${expected_string}" < <(strings "${payload_file}") | |
| done | |
| fi | |
| fi | |
| } | |
| assert_only_expected_rebuilt() { | |
| local log="$1" | |
| local allowed="$2" | |
| shift 2 | |
| local unexpected | |
| unexpected="$( | |
| sed -nE 's/^[[:space:]]*Compiling ([^[:space:]]+).*/\1/p' "${log}" \ | |
| | sort -u \ | |
| | grep -Ev "${allowed}" \ | |
| || true | |
| )" | |
| if [ -n "${unexpected}" ]; then | |
| echo "Codegen mode change rebuilt shared dependencies:" >&2 | |
| echo "${unexpected}" >&2 | |
| return 1 | |
| fi | |
| local package | |
| for package in "$@"; do | |
| if ! grep -Eq "^[[:space:]]*Compiling ${package}([[:space:]]|$)" "${log}"; then | |
| echo "Expected ${package} to rebuild, but it stayed fresh" >&2 | |
| return 1 | |
| fi | |
| done | |
| } | |
| cargo oxide build vecadd --arch sm_86 2>&1 \ | |
| | tee "${RUNNER_TEMP}/vecadd-ptx-build.log" | |
| assert_only_expected_rebuilt \ | |
| "${RUNNER_TEMP}/vecadd-ptx-build.log" '^vecadd$' vecadd | |
| assert_embedded_artifact "${CARGO_TARGET_DIR}/release/vecadd" 1 256 sm_86 | |
| cargo oxide build vecadd --materialize-cubin --arch sm_86 2>&1 \ | |
| | tee "${RUNNER_TEMP}/vecadd-cubin-build.log" | |
| assert_only_expected_rebuilt \ | |
| "${RUNNER_TEMP}/vecadd-cubin-build.log" '^vecadd$' vecadd | |
| # Default compile policy remains wire-compatible artifact v1; the | |
| # payload kind is what must change from PTX (256) to cubin (512). | |
| assert_embedded_artifact "${CARGO_TARGET_DIR}/release/vecadd" 1 512 sm_86 | |
| # IKET metadata is consumed by the profiler rather than device code. | |
| # Guard the complete libNVVM + nvJitLink path: an IR-only assertion | |
| # would miss internalization dropping these globals from the cubin. | |
| cargo oxide build iket_trace --materialize-cubin --arch sm_120 2>&1 \ | |
| | tee "${RUNNER_TEMP}/iket-trace-cubin-build.log" | |
| assert_only_expected_rebuilt \ | |
| "${RUNNER_TEMP}/iket-trace-cubin-build.log" '^iket_trace$' iket_trace | |
| assert_embedded_artifact \ | |
| "${CARGO_TARGET_DIR}/release/iket_trace" 1 512 sm_120 \ | |
| __iket_meta_info \ | |
| vecadd.completed.with.a.name.longer.than.thirty.one.bytes | |
| # Switching back must invalidate vecadd's device artifact without | |
| # throwing away the shared host dependency cache. | |
| cargo oxide build vecadd --arch sm_86 2>&1 \ | |
| | tee "${RUNNER_TEMP}/vecadd-ptx-return-build.log" | |
| assert_only_expected_rebuilt \ | |
| "${RUNNER_TEMP}/vecadd-ptx-return-build.log" '^vecadd$' vecadd | |
| assert_embedded_artifact "${CARGO_TARGET_DIR}/release/vecadd" 1 256 sm_86 | |
| # A kernel-producing path dependency must be invalidated too, while | |
| # ordinary CUDA/host dependencies remain reusable. | |
| cargo oxide build cross_crate_kernel --arch sm_86 2>&1 \ | |
| | tee "${RUNNER_TEMP}/cross-crate-sm86-build.log" | |
| assert_only_expected_rebuilt \ | |
| "${RUNNER_TEMP}/cross-crate-sm86-build.log" \ | |
| '^(kernel-lib|cross_crate_kernel)$' kernel-lib cross_crate_kernel | |
| cargo oxide build cross_crate_kernel --arch sm_90 2>&1 \ | |
| | tee "${RUNNER_TEMP}/cross-crate-sm90-build.log" | |
| assert_only_expected_rebuilt \ | |
| "${RUNNER_TEMP}/cross-crate-sm90-build.log" \ | |
| '^(kernel-lib|cross_crate_kernel)$' kernel-lib cross_crate_kernel | |
| - name: Upload failure logs | |
| if: failure() | |
| uses: actions/upload-artifact@v4 | |
| with: | |
| name: smoketest-compile-logs | |
| # One glob per log the artifact-modes step writes. `iket-trace-*` | |
| # matched neither of the other two, so the one assertion that | |
| # exercises the full libNVVM + nvJitLink path -- and the only one | |
| # that can fail on internalization dropping `__iket_meta_info` from | |
| # the cubin -- was the one whose log never reached the artifact. | |
| # `if-no-files-found: ignore` below made that silent. | |
| path: | | |
| ${{ runner.temp }}/smoketest-logs/ | |
| ${{ runner.temp }}/vecadd-*-build.log | |
| ${{ runner.temp }}/iket-trace-*-build.log | |
| ${{ runner.temp }}/cross-crate-*-build.log | |
| if-no-files-found: ignore | |
| retention-days: 7 |