From 8f7fb01b13c3ee4352eed3d2ce61362a9672628a Mon Sep 17 00:00:00 2001 From: Ludovic Henry Date: Sun, 20 Sep 2026 20:52:21 +0000 Subject: [PATCH 1/2] warp-lang: Add version 1.16.0 Builds the CPU-only configuration of NVIDIA Warp for riscv64: the core runtime plus the bundled Clang/LLVM that JIT-compiles kernels to native code at run time. Upstream publishes no riscv64 wheel, and no prebuilt riscv64 Clang/LLVM SDK either, so the workflow compiles LLVM from source in the manylinux container the way upstream's own build-llvm-sdk.yml does for x86_64/aarch64. CUDA is off (upstream's --no-cuda, the configuration it ships macOS with): neither the CUDA toolkit nor libmathdx exists for riscv64, and the NVIDIA driver does not either, so a GPU build would have no device to run on. Four patches. 0001 and 0002 make the build work at all - the platform tables, and the rv64gc/lp64d ABI the Clang frontend does not default to at the cc1 level. 0003 and 0004 come from the first full run of upstream's test suite on the wheel, which accounted for every one of its 1552 errors: - 649 were one bug. RV64GC has no half-precision hardware, so Clang lowers builtin.h's _Float16 conversions to __extendhfsf2/__truncsfhf2, and wp_load_obj() resolves a JIT-compiled module's externals from a curated CRT table holding no compiler builtins - so every module with an fp16 kernel failed to materialize and took its non-fp16 kernels with it. x86_64 avoids this with +f16c and aarch64 with armv8-a's fcvt; RISC-V has no baseline equivalent (Zfh is optional and not in rv64gc), so 0003 converts in software, bit-for-bit what _Float16 yields elsewhere. - 903 were the test runner's hard-coded 3600s wall-clock limit. Warp compiles every kernel with Clang on the machine running the tests, and four riscv64 cores got through 137 of 180 suites inside the hour; the other 43 were marked crashed. 0004 reads the limit from WARP_TEST_SUITE_TIMEOUT. setup.py packages examples/assets and tests/assets, both git-lfs, so the wheel shipped pointer stubs and example_nvdb died on "NanoVDB signature not found". The blobs are fetched over the media CDN rather than by checking out with lfs: true, which would need a git-lfs client on the runner (gotcha 296). --- .github/workflows/build-warp-lang.yml | 229 ++++++++++++++++++ docs/packages/warp-lang.yaml | 14 ++ ...he-build-and-packaging-platform-tabl.patch | 138 +++++++++++ ...v64gc-lp64d-target-ABI-for-JIT-compi.patch | 63 +++++ ...software-on-RISC-V-targets-without-Z.patch | 167 +++++++++++++ ...vironment-raise-the-parallel-run-s-w.patch | 43 ++++ 6 files changed, 654 insertions(+) create mode 100644 .github/workflows/build-warp-lang.yml create mode 100644 docs/packages/warp-lang.yaml create mode 100644 patches/warp-lang/1.16.0/0001-Add-riscv64-to-the-build-and-packaging-platform-tabl.patch create mode 100644 patches/warp-lang/1.16.0/0002-clang-name-the-rv64gc-lp64d-target-ABI-for-JIT-compi.patch create mode 100644 patches/warp-lang/1.16.0/0003-Convert-fp16-in-software-on-RISC-V-targets-without-Z.patch create mode 100644 patches/warp-lang/1.16.0/0004-tests-let-the-environment-raise-the-parallel-run-s-w.patch diff --git a/.github/workflows/build-warp-lang.yml b/.github/workflows/build-warp-lang.yml new file mode 100644 index 00000000000..8e1fc2ce054 --- /dev/null +++ b/.github/workflows/build-warp-lang.yml @@ -0,0 +1,229 @@ +# SPDX-FileCopyrightText: 2026 The RISE Project +# SPDX-License-Identifier: MIT +--- +# Based on upstream's Linux wheel pipeline, which compiles the Clang/LLVM SDK from source in a +# manylinux container and then cuts one wheel per platform: +# https://github.com/NVIDIA/warp/blob/v1.16.0/.github/workflows/build-llvm-sdk.yml +# https://github.com/NVIDIA/warp/blob/v1.16.0/.gitlab-ci.yml +name: Build warp-lang wheels (riscv64) + +on: + workflow_dispatch: + inputs: + version: + description: 'Version glob to (re)build; empty builds every version of docs/packages/warp-lang.yaml not released yet' + required: false + default: '' + pull_request: + branches: [main] + paths: + - '.github/workflows/build-warp-lang.yml' + - 'docs/packages/warp-lang.yaml' + push: + branches: [main] + paths: + - '.github/workflows/build-warp-lang.yml' + - 'docs/packages/warp-lang.yaml' + +concurrency: + group: ${{ github.workflow }}-${{ github.head_ref || github.run_id }} + cancel-in-progress: true + +permissions: + contents: read # to fetch code (actions/checkout) + +env: + # Pinned explicitly so the gpl_sources job pulls the same image the statically linked + # libstdc++/libgcc in the wheel came from. + MANYLINUX_RISCV64_IMAGE: quay.io/pypa/manylinux_2_39_riscv64 + +jobs: + setup: + uses: $/.github/workflows/_setup.yml + with: + package: warp-lang + version: ${{ inputs.version }} + + build_wheels: + needs: [setup] + if: needs.setup.outputs.versions != '[]' + strategy: + fail-fast: false + matrix: + version: ${{ fromJSON(needs.setup.outputs.versions) }} + name: Build warp-lang ${{ matrix.version }} py3-none-manylinux_riscv64 + runs-on: ubuntu-24.04-riscv + timeout-minutes: 1440 # full Clang/LLVM compile from source + + env: + WARP_VERSION: ${{ matrix.version }} + + steps: + - name: Checkout warp v${{ env.WARP_VERSION }} + uses: actions/checkout@3d3c42e5aac5ba805825da76410c181273ba90b1 # v7.0.1 + with: + repository: NVIDIA/warp + ref: v${{ env.WARP_VERSION }} + persist-credentials: false + + - name: Checkout python-wheels + uses: actions/checkout@3d3c42e5aac5ba805825da76410c181273ba90b1 # v7.0.1 + with: + path: python-wheels + persist-credentials: false + + # setup.py packages examples/assets and tests/assets, and both are git-lfs; nothing + # guarantees a git-lfs client on the self-hosted runner, so fetch the blobs over the + # media CDN instead of checking out with lfs: true (gotcha 296). + - name: Fetch the git-lfs assets the wheel packages + run: | + for f in warp/examples/assets/* warp/tests/assets/*; do + grep -q '^version https://git-lfs' "$f" || continue + curl -sSfL -o "$f" \ + "https://media.githubusercontent.com/media/NVIDIA/warp/v${WARP_VERSION}/$f" + done + + - name: Apply patches + run: git apply python-wheels/patches/warp-lang/${{ env.WARP_VERSION }}/*.patch + + - name: Pull manylinux_riscv64 image + run: docker pull "${MANYLINUX_RISCV64_IMAGE}" + + - name: Write the JIT smoke test + run: | + cat > "${RUNNER_TEMP}/smoke.py" <<'EOF' + import numpy as np + + import warp as wp + + wp.init() + print(wp.get_devices()) + + + @wp.kernel + def saxpy(a: float, x: wp.array(dtype=float), y: wp.array(dtype=float)): + i = wp.tid() + y[i] = a * x[i] + wp.sqrt(y[i]) + + + x = wp.array(np.arange(16, dtype=np.float32), dtype=float, device="cpu") + y = wp.array(np.full(16, 4.0, dtype=np.float32), dtype=float, device="cpu") + wp.launch(saxpy, dim=16, inputs=[2.0, x, y], device="cpu") + np.testing.assert_allclose(y.numpy(), 2.0 * np.arange(16) + 2.0, rtol=1e-6) + print("JIT smoke test OK") + EOF + + # Upstream ships no prebuilt Clang/LLVM SDK for riscv64, so --build-llvm compiles it here + # instead of downloading one; --no-cuda is upstream's own CPU-only configuration. cmake and + # ninja are pinned as in upstream's tools/llvm/ci/build-linux.sh, ahead of the image's own + # cmake 4. libstdc++-static is CRB-only, and warp links -static-libstdc++/-static-libgcc. + - name: Build warp and the wheel + run: | + docker run --rm -v "$(pwd):/work" -v "${RUNNER_TEMP}/smoke.py:/smoke.py:ro" \ + -w /work -e WARP_VERSION "${MANYLINUX_RISCV64_IMAGE}" bash -c ' + set -euxo pipefail + dnf -y --enablerepo=crb install libstdc++-static + pybin=/opt/python/cp312-cp312/bin + "$pybin/pip" install -q --extra-index-url https://pypi.riseproject.dev/simple/ \ + numpy build setuptools wheel cmake==3.31.6 ninja==1.11.1.3 + export PATH="$pybin:$PATH" + "$pybin/python3" build_lib.py --no-cuda --build-llvm + file warp/bin/warp.so warp/bin/warp-clang.so + readelf -A warp/bin/warp-clang.so || true + ldd warp/bin/warp.so warp/bin/warp-clang.so + "$pybin/python3" -m build --wheel --no-isolation \ + -C--build-option=-Plinux-riscv64 -C--build-option=-Mmanylinux_2_39 + for tag in cp312-cp312 cp313-cp313 cp314-cp314 cp314-cp314t; do + "/opt/python/$tag/bin/python3" -m venv "/tmp/venv-$tag" + "/tmp/venv-$tag/bin/pip" install -q \ + --extra-index-url https://pypi.riseproject.dev/simple/ dist/*.whl + (cd /tmp && "/tmp/venv-$tag/bin/python3" /smoke.py) + done + ' + + - uses: actions/upload-artifact@043fb46d1a93c77aae656e7c1c64a875d1fc6a0a # v7.0.1 + with: + name: warp-lang-${{ env.WARP_VERSION }}-py3-none-manylinux_riscv64 + path: dist/*.whl + if-no-files-found: error + + test_wheels: + needs: [setup, build_wheels] + if: needs.setup.outputs.versions != '[]' + strategy: + fail-fast: false + matrix: + version: ${{ fromJSON(needs.setup.outputs.versions) }} + name: Test warp-lang ${{ matrix.version }} py3-none-manylinux_riscv64 + runs-on: ubuntu-24.04-riscv + timeout-minutes: 1440 # every kernel in the suite is JIT-compiled by Clang on the runner + + env: + WARP_VERSION: ${{ matrix.version }} + + steps: + - uses: actions/download-artifact@3e5f45b2cfb9172054b4087a40e8e0b5a5461e7c # v8.0.1 + with: + name: warp-lang-${{ env.WARP_VERSION }}-py3-none-manylinux_riscv64 + path: dist + + - name: Pull manylinux_riscv64 image + run: docker pull "${MANYLINUX_RISCV64_IMAGE}" + + # The runner's four cores compile every kernel in the suite with Clang, which does not fit + # unittest_parallel's 3600s default wall-clock limit (patch 0004 makes it settable). + - name: Run upstream's test suite against the wheel + run: | + docker run --rm -v "$(pwd)/dist:/dist:ro" -w /tmp -e WARP_TEST_SUITE_TIMEOUT=43200 \ + "${MANYLINUX_RISCV64_IMAGE}" bash -c ' + set -euxo pipefail + pybin=/opt/python/cp312-cp312/bin + "$pybin/pip" install -q --extra-index-url https://pypi.riseproject.dev/simple/ /dist/*.whl + "$pybin/python3" -m warp.tests --maxjobs 4 + ' + + gpl_sources: + needs: [setup] + if: needs.setup.outputs.versions != '[]' + strategy: + fail-fast: false + matrix: + version: ${{ fromJSON(needs.setup.outputs.versions) }} + name: Collect GPL sources (gcc) for warp-lang ${{ matrix.version }} + runs-on: ubuntu-24.04-riscv + + env: + WARP_VERSION: ${{ matrix.version }} + + steps: + - name: Collect gcc source RPM from manylinux_riscv64 + uses: riseproject-dev/python-wheels/actions/collect-gpl-sources@main + with: + image: ${{ env.MANYLINUX_RISCV64_IMAGE }} + packages: gcc + output: gpl-sources.tar + + - uses: actions/upload-artifact@043fb46d1a93c77aae656e7c1c64a875d1fc6a0a # v7.0.1 + with: + name: warp-lang-${{ env.WARP_VERSION }}-gpl-sources + path: gpl-sources.tar + if-no-files-found: error + + publish: + name: Publish warp-lang ${{ matrix.version }} + needs: [setup, build_wheels, test_wheels, gpl_sources] + if: needs.setup.outputs.versions != '[]' + strategy: + fail-fast: false + matrix: + version: ${{ fromJSON(needs.setup.outputs.versions) }} + permissions: + contents: write + pull-requests: write + uses: $/.github/workflows/_publish-wheel.yml + secrets: + app-private-key: ${{ secrets.RISEPROJECT_APP_PRIVATE_KEY }} + with: + artifact-pattern: warp-lang-${{ matrix.version }}-*-manylinux_riscv64 + gpl-sources-artifact: warp-lang-${{ matrix.version }}-gpl-sources + gpl-sources-description: gcc diff --git a/docs/packages/warp-lang.yaml b/docs/packages/warp-lang.yaml new file mode 100644 index 00000000000..f8c2e119872 --- /dev/null +++ b/docs/packages/warp-lang.yaml @@ -0,0 +1,14 @@ +package-name: warp-lang +source-code: https://github.com/NVIDIA/warp +license: Apache-2.0 +versions: +- version: 1.16.0 + patched: true + comment: >- + Built in upstream's own CPU-only configuration (``build_lib.py --no-cuda``), the one it + uses for macOS: kernels are JIT-compiled to native riscv64 code by the bundled + Clang/LLVM, and ``wp.get_devices()`` reports the CPU device only. The patches add riscv64 + to the build and packaging platform tables, name the rv64gc/lp64d ABI for the JIT (which + the Clang frontend does not default to), convert fp16 in software because rv64gc has no + half-precision hardware and the JIT resolves no compiler-rt builtins, and let the + environment raise the test runner's wall-clock limit. diff --git a/patches/warp-lang/1.16.0/0001-Add-riscv64-to-the-build-and-packaging-platform-tabl.patch b/patches/warp-lang/1.16.0/0001-Add-riscv64-to-the-build-and-packaging-platform-tabl.patch new file mode 100644 index 00000000000..c32cb636453 --- /dev/null +++ b/patches/warp-lang/1.16.0/0001-Add-riscv64-to-the-build-and-packaging-platform-tabl.patch @@ -0,0 +1,138 @@ +From 8b64e2e85baff0f1e96beb7b3ad4d25ac8cbd82e Mon Sep 17 00:00:00 2001 +From: Ludovic Henry +Date: Sun, 20 Sep 2026 12:00:00 +0000 +Subject: [PATCH 1/4] Add riscv64 to the build and packaging platform tables + +Upstream-Status: To upstream [no riscv64 CI runner or prebuilt Clang/LLVM SDK upstream yet] + +machine_architecture() in warp/_src/build_dll.py and in setup.py both raise +"Unrecognized machine architecture riscv64", which aborts build_lib.py before +anything is compiled. + +setup.py's `platforms` table is what WarpBDistWheel.get_tag() reads to stamp the +wheel's platform tag, so riscv64 needs an entry there as well. manylinux_2_39 is +the flavour of the image these wheels are built in +(quay.io/pypa/manylinux_2_39_riscv64, Rocky Linux 10, glibc 2.39). + +build_llvm.py picks LLVM_TARGETS_TO_BUILD from the same canonical architecture +string and otherwise falls back to X86, which would build an LLVM unable to emit +code for the machine it runs on; RISCV is the matching backend name. + +Nothing else in the tree is architecture-specific: warp/native carries no SIMD +intrinsics and the Linux compile flags are a generic -O3 -fPIC --std=c++17. +--- + build_llvm.py | 4 +++- + setup.py | 14 +++++++++----- + warp/_src/build_dll.py | 3 +++ + 3 files changed, 15 insertions(+), 6 deletions(-) + +diff --git a/build_llvm.py b/build_llvm.py +index 22a9b45..3d9e2cd 100644 +--- a/build_llvm.py ++++ b/build_llvm.py +@@ -150,7 +150,7 @@ def build_llvm_clang_from_source_for_arch(args, arch: str, llvm_source: str) -> + + Args: + args: Command line arguments +- arch: Architecture to build for ("aarch64" or "x86_64") ++ arch: Architecture to build for ("aarch64", "riscv64" or "x86_64") + llvm_source: Path to the LLVM source code + """ + +@@ -208,6 +208,8 @@ def build_llvm_clang_from_source_for_arch(args, arch: str, llvm_source: str) -> + + if arch == "aarch64": + target_backend = "AArch64" ++ elif arch == "riscv64": ++ target_backend = "RISCV" + else: + target_backend = "X86" + +diff --git a/setup.py b/setup.py +index b57e7ac..6562e89 100644 +--- a/setup.py ++++ b/setup.py +@@ -21,14 +21,14 @@ parser.add_argument( + "-P", + type=str, + default="", +- help="Wheel platform: windows-x86_64|linux-x86_64|linux-aarch64|macos-aarch64", ++ help="Wheel platform: windows-x86_64|linux-x86_64|linux-aarch64|linux-riscv64|macos-aarch64", + ) + parser.add_argument( + "--manylinux", + "-M", + type=str, + default="manylinux_2_28", +- help="Manylinux flavor for Linux wheels: manylinux_2_28|manylinux_2_34", ++ help="Manylinux flavor for Linux wheels: manylinux_2_28|manylinux_2_34|manylinux_2_39", + ) + args = parser.parse_known_args()[0] + +@@ -36,12 +36,15 @@ args = parser.parse_known_args()[0] + # returns a canonical machine architecture string + # - "x86_64" for x86-64, aka. AMD64, aka. x64 + # - "aarch64" for AArch64, aka. ARM64 ++# - "riscv64" for 64-bit RISC-V, aka. RV64 + def machine_architecture() -> str: + machine = platform.machine() + if machine == "x86_64" or machine == "AMD64": + return "x86_64" + if machine == "aarch64" or machine == "arm64": + return "aarch64" ++ if machine == "riscv64": ++ return "riscv64" + raise RuntimeError(f"Unrecognized machine architecture {machine}") + + +@@ -76,6 +79,7 @@ platforms = [ + Platform("windows", "x86_64", "Windows x86-64", ".dll", "win_amd64"), + Platform("linux", "x86_64", "Linux x86-64", ".so", "manylinux_2_28_x86_64"), + Platform("linux", "aarch64", "Linux AArch64", ".so", "manylinux_2_34_aarch64"), ++ Platform("linux", "riscv64", "Linux RISC-V 64", ".so", "manylinux_2_39_riscv64"), + Platform("macos", "aarch64", "macOS ARM64", ".dylib", "macosx_11_0_arm64"), + ] + +@@ -132,7 +136,7 @@ if args.command == "bdist_wheel": + if len(detected_platforms) > 1: + print("Libraries for multiple platforms were detected.") + print("Run `python -m build --wheel -C--build-option=-P` to select a specific one.") +- print("Available platforms: windows-x86_64, linux-x86_64, linux-aarch64, macos-aarch64") ++ print("Available platforms: windows-x86_64, linux-x86_64, linux-aarch64, linux-riscv64, macos-aarch64") + # Select the libraries corresponding with the this machine's platform + for p in platforms: + if p.os == machine_os() and p.arch == machine_architecture(): +@@ -154,8 +158,8 @@ class WarpBDistWheel(bdist_wheel): + # setuptools.Command can validate the command line options. + user_options: ClassVar[list[tuple[str, str, str]]] = [ + *bdist_wheel.user_options, +- ("platform=", "P", "Wheel platform: windows-x86_64|linux-x86_64|linux-aarch64|macos-aarch64"), +- ("manylinux=", "M", "Manylinux flavor for Linux wheels: manylinux_2_28|manylinux_2_34"), ++ ("platform=", "P", "Wheel platform: windows-x86_64|linux-x86_64|linux-aarch64|linux-riscv64|macos-aarch64"), ++ ("manylinux=", "M", "Manylinux flavor for Linux wheels: manylinux_2_28|manylinux_2_34|manylinux_2_39"), + ] + + def initialize_options(self): +diff --git a/warp/_src/build_dll.py b/warp/_src/build_dll.py +index f47f4a1..cdd61e7 100644 +--- a/warp/_src/build_dll.py ++++ b/warp/_src/build_dll.py +@@ -60,12 +60,15 @@ def machine_architecture() -> str: + """Return a canonical machine architecture string. + - "x86_64" for x86-64, aka. AMD64, aka. x64 + - "aarch64" for AArch64, aka. ARM64 ++ - "riscv64" for 64-bit RISC-V, aka. RV64 + """ + machine = platform.machine() + if machine == "x86_64" or machine == "AMD64": + return "x86_64" + if machine == "aarch64" or machine == "arm64": + return "aarch64" ++ if machine == "riscv64": ++ return "riscv64" + raise RuntimeError(f"Unrecognized machine architecture {machine}") + + +-- +2.43.0 + diff --git a/patches/warp-lang/1.16.0/0002-clang-name-the-rv64gc-lp64d-target-ABI-for-JIT-compi.patch b/patches/warp-lang/1.16.0/0002-clang-name-the-rv64gc-lp64d-target-ABI-for-JIT-compi.patch new file mode 100644 index 00000000000..c40f88ac060 --- /dev/null +++ b/patches/warp-lang/1.16.0/0002-clang-name-the-rv64gc-lp64d-target-ABI-for-JIT-compi.patch @@ -0,0 +1,63 @@ +From bb045ca51b3f98290f3194192d5a72fde6af1667 Mon Sep 17 00:00:00 2001 +From: Ludovic Henry +Date: Sun, 20 Sep 2026 12:05:00 +0000 +Subject: [PATCH 2/4] clang: name the rv64gc/lp64d target ABI for JIT-compiled + kernels + +Upstream-Status: To upstream [no riscv64 runner upstream to regression-test it on] + +warp-clang.so drives Clang through CompilerInvocation, i.e. at the cc1 level, +where none of the driver's per-target defaults are applied. On RV64 that leaves +clang::RISCVTargetInfo with a bare rv64i ISA (initFeatureMap sees no +-target-feature flags) and, because the invocation names no ABI, an ABIFLen of +zero -- the soft-float lp64 ABI. + +warp.so, crt.cpp and warp-clang.so itself are compiled by the host gcc for +rv64gc/lp64d, so every call a JIT-compiled kernel makes into them that passes or +returns a float or a double -- _wp_isfinite(double) and the other crt.cpp shims +-- would read the value from the wrong register class. The base ISA is wrong for +the same reason: without M/A/F/D the frontend emits integer-multiply and +soft-float libcalls for operations the hardware performs natively. + +Named explicitly rather than derived from the triple, because +LLVM_DEFAULT_TARGET_TRIPLE is whatever LLVM was configured with (build_llvm.py +sets riscv64-pc-linux) and cc1 infers neither a RISC-V ISA nor an ABI from the +triple's environment field. + +Mirrors the x86_64 (+f16c) and aarch64 (+reserve-x28) blocks beside it. +--- + warp/native/clang/clang.cpp | 18 ++++++++++++++++++ + 1 file changed, 18 insertions(+) + +diff --git a/warp/native/clang/clang.cpp b/warp/native/clang/clang.cpp +index d7734b0..f535c92 100644 +--- a/warp/native/clang/clang.cpp ++++ b/warp/native/clang/clang.cpp +@@ -231,6 +231,24 @@ static std::unique_ptr create_compiler( + args.push_back("+f16c"); + #endif + ++#if defined(__riscv) && __riscv_xlen == 64 ++ // cc1 defaults RV64 to a bare rv64i ISA and, with no ABI named, to soft-float lp64. ++ // warp.so and crt.cpp are compiled rv64gc/lp64d, so without this every kernel call ++ // into them (_wp_isfinite(double), ...) would pass floats in the wrong registers. ++ args.push_back("-target-abi"); ++ args.push_back("lp64d"); ++ args.push_back("-target-feature"); ++ args.push_back("+m"); ++ args.push_back("-target-feature"); ++ args.push_back("+a"); ++ args.push_back("-target-feature"); ++ args.push_back("+f"); ++ args.push_back("-target-feature"); ++ args.push_back("+d"); ++ args.push_back("-target-feature"); ++ args.push_back("+c"); ++#endif ++ + #if defined(__aarch64__) + if (tiles_in_stack_memory) { + // Static memory support is broken on AArch64 CPUs. As a workaround we reserve some stack memory on kernel +-- +2.43.0 + diff --git a/patches/warp-lang/1.16.0/0003-Convert-fp16-in-software-on-RISC-V-targets-without-Z.patch b/patches/warp-lang/1.16.0/0003-Convert-fp16-in-software-on-RISC-V-targets-without-Z.patch new file mode 100644 index 00000000000..6bea7fe0b8b --- /dev/null +++ b/patches/warp-lang/1.16.0/0003-Convert-fp16-in-software-on-RISC-V-targets-without-Z.patch @@ -0,0 +1,167 @@ +From 7bd5afd2c143e7db533bafc04c17ef4555c45a4b Mon Sep 17 00:00:00 2001 +From: Ludovic Henry +Date: Mon, 21 Sep 2026 10:08:00 +0000 +Subject: [PATCH 3/4] Convert fp16 in software on RISC-V targets without Zfh + +Upstream-Status: To upstream [no riscv64 runner upstream to regression-test it on] + +float_to_half()/half_to_float() convert through Clang's native _Float16 in +kernel code. RV64GC has no half-precision hardware, so Clang lowers both +conversions to the compiler-rt libcalls __extendhfsf2/__truncsfhf2, and +wp_load_obj() (warp/native/clang/clang.cpp) resolves a JIT-compiled module's +externals from a curated CRT table that carries no compiler builtins. Every +module holding an fp16 kernel therefore failed to materialize: + + JIT session error: Symbols not found: [ __extendhfsf2, __truncsfhf2 ] + Failed to lookup symbol: Failed to materialize symbols: ... + RuntimeError: Failed to find forward kernel '...' for device 'cpu' + +and took the module's non-fp16 kernels down with it -- 649 of the 1552 errors +in a full riscv64 run of warp.tests, spread over test_codegen, test_print, +test_spatial, test_transform, test_fabricarray and 40 more modules. + +x86_64 sidesteps this by adding +f16c to the cc1 invocation, which lowers both +conversions to vcvtph2ps/vcvtps2ph; aarch64 gets fcvt from armv8-a. RISC-V has +no such baseline extension -- Zfh is optional, is not part of rv64gc, and +asking for it would emit fcvt.h.s on hardware that need not implement it -- so +convert in software instead. + +The two routines are bit-for-bit what _Float16 yields where the hardware does +have half support: round to nearest with ties to even, signalling NaNs quieted, +NaN payloads preserved. Verified exhaustively against (_Float16) on x86_64 +(+f16c) over all 2^32 float bit patterns and all 2^16 half bit patterns, zero +differences, so a riscv64 kernel rounds exactly as an x86_64 or aarch64 one +does rather than picking up the ties-away rounding of the Giesen routine +warp.cpp uses for the host-side conversions. + +Guarded on !__riscv_zfh so a Zfh-enabled target keeps the native path. +--- + warp/native/builtin.h | 105 ++++++++++++++++++++++++++++++++++++++++++ + 1 file changed, 105 insertions(+) + +diff --git a/warp/native/builtin.h b/warp/native/builtin.h +index f0c23b5..b1ad4b7 100644 +--- a/warp/native/builtin.h ++++ b/warp/native/builtin.h +@@ -367,6 +367,109 @@ CUDA_CALLABLE inline float bfloat16_to_float(wp_bfloat16 x) { return wp_bfloat16 + + #elif defined(__clang__) + ++#if defined(__riscv) && !defined(__riscv_zfh) ++ ++// RISC-V without Zfh has no half-precision hardware, so Clang lowers every _Float16 conversion ++// to the compiler-rt libcalls __extendhfsf2/__truncsfhf2. A JIT-compiled module resolves its ++// externals from the curated CRT table in wp_load_obj() (warp/native/clang/clang.cpp), which ++// carries no compiler builtins, so a module holding any fp16 kernel fails to materialize with ++// "JIT session error: Symbols not found: [ __extendhfsf2, __truncsfhf2 ]". ++// ++// Convert in software instead. Both routines are bit-for-bit what _Float16 yields where the ++// hardware does have half support -- round to nearest with ties to even, signalling NaNs ++// quieted, payloads preserved -- verified exhaustively over all 2^32 floats and all 2^16 ++// halves, so kernels round exactly as they do on x86_64 (+f16c) and aarch64. ++CUDA_CALLABLE inline half float_to_half(float x) ++{ ++ unsigned int bits; ++ memcpy(&bits, &x, sizeof(bits)); ++ ++ const unsigned int sign = (bits >> 16) & 0x8000u; ++ const unsigned int magnitude = bits & 0x7fffffffu; ++ ++ unsigned int u; ++ if (magnitude >= 0x7f800000u) // Inf or NaN ++ { ++ const unsigned int mantissa = magnitude & 0x007fffffu; ++ u = sign | 0x7c00u | (mantissa ? ((mantissa >> 13) | 0x0200u) : 0u); // NaN quieted, Inf kept ++ } ++ else if (magnitude >= 0x47800000u) // 65536 and above overflows the half range ++ { ++ u = sign | 0x7c00u; ++ } ++ else if (magnitude >= 0x38800000u) // 2^-14 and above is a normal half ++ { ++ u = sign | (((magnitude >> 23) - 127 + 15) << 10) | ((magnitude >> 13) & 0x03ffu); ++ const unsigned int rest = magnitude & 0x1fffu; // the bits that do not fit ++ if (rest > 0x1000u || (rest == 0x1000u && (u & 1u))) ++ u += 1; // round to nearest, ties to even ++ } ++ else ++ { ++ const unsigned int exponent = magnitude >> 23; ++ const unsigned int shift = 126 - exponent; // a subnormal half's ulp is 2^-24 ++ if (exponent == 0 || shift > 24) ++ { ++ u = sign; // underflows to zero ++ } ++ else ++ { ++ const unsigned int mantissa = (magnitude & 0x007fffffu) | 0x00800000u; // implicit bit ++ const unsigned int rest = mantissa & ((1u << shift) - 1u); ++ const unsigned int tie = 1u << (shift - 1); ++ u = sign | (mantissa >> shift); ++ if (rest > tie || (rest == tie && (u & 1u))) ++ u += 1; // round to nearest, ties to even ++ } ++ } ++ ++ half h; ++ h.u = static_cast(u); ++ return h; ++} ++ ++CUDA_CALLABLE inline float half_to_float(half h) ++{ ++ const unsigned int sign = (static_cast(h.u) & 0x8000u) << 16; ++ const unsigned int exponent = (static_cast(h.u) >> 10) & 0x001fu; ++ const unsigned int mantissa = static_cast(h.u) & 0x03ffu; ++ ++ unsigned int bits; ++ if (exponent == 0x1fu) // Inf or NaN ++ { ++ bits = sign | 0x7f800000u | (mantissa ? ((mantissa << 13) | 0x00400000u) : 0u); ++ } ++ else if (exponent == 0u) // zero or subnormal ++ { ++ if (mantissa == 0u) ++ { ++ bits = sign; ++ } ++ else ++ { ++ // Renormalize: shift the mantissa up until its leading one leaves the field. ++ unsigned int m = mantissa; ++ unsigned int e = 127 - 15 + 1; ++ while ((m & 0x0400u) == 0u) ++ { ++ m <<= 1; ++ e -= 1; ++ } ++ bits = sign | (e << 23) | ((m & 0x03ffu) << 13); ++ } ++ } ++ else ++ { ++ bits = sign | ((exponent + 127 - 15) << 23) | (mantissa << 13); ++ } ++ ++ float val; ++ memcpy(&val, &bits, sizeof(val)); ++ return val; ++} ++ ++#else ++ + // _Float16 is Clang's native half-precision floating-point type + CUDA_CALLABLE inline half float_to_half(float x) + { +@@ -381,6 +484,8 @@ CUDA_CALLABLE inline float half_to_float(half h) + return static_cast(f16); + } + ++#endif // __riscv && !__riscv_zfh ++ + #ifndef WP_NO_BFLOAT16 + CUDA_CALLABLE inline wp_bfloat16 float_to_bfloat16(float x) + { +-- +2.43.0 + diff --git a/patches/warp-lang/1.16.0/0004-tests-let-the-environment-raise-the-parallel-run-s-w.patch b/patches/warp-lang/1.16.0/0004-tests-let-the-environment-raise-the-parallel-run-s-w.patch new file mode 100644 index 00000000000..cb3d8859a63 --- /dev/null +++ b/patches/warp-lang/1.16.0/0004-tests-let-the-environment-raise-the-parallel-run-s-w.patch @@ -0,0 +1,43 @@ +From 2436495b2ae8f03a9a72513aa9616f07fef456f6 Mon Sep 17 00:00:00 2001 +From: Ludovic Henry +Date: Mon, 21 Sep 2026 10:08:00 +0000 +Subject: [PATCH 4/4] tests: let the environment raise the parallel run's + wall-clock limit + +Upstream-Status: To upstream [no riscv64 runner upstream to regression-test it on] + +unittest_parallel caps the whole parallel run at a hard-coded 3600s with no +way to change it. Warp compiles every kernel with Clang on the machine running +the tests, so the suite's wall-clock cost scales with that machine rather than +with the suite: a full run on a 4-core riscv64 runner got through 137 of 180 +suites in the hour, and the other 43 were marked crashed -- 903 of the 1552 +errors in the run -- with nothing actually wrong. + +Read the limit from WARP_TEST_SUITE_TIMEOUT, keeping 3600 as the default so +nothing changes for anyone who does not set it. +--- + warp/_src/thirdparty/unittest_parallel.py | 8 +++++--- + 1 file changed, 5 insertions(+), 3 deletions(-) + +diff --git a/warp/_src/thirdparty/unittest_parallel.py b/warp/_src/thirdparty/unittest_parallel.py +index 4f81bd2..66ef92b 100644 +--- a/warp/_src/thirdparty/unittest_parallel.py ++++ b/warp/_src/thirdparty/unittest_parallel.py +@@ -38,9 +38,11 @@ except ImportError: + + # The following variables are NVIDIA Modifications + START_DIRECTORY = os.path.join(os.path.dirname(__file__), "..") # The directory to start test discovery +-_SUITE_TIMEOUT = ( +- 3600 # Timeout in seconds: total wall-clock limit for parallel execution, per-suite limit during isolated fallback +-) ++# Timeout in seconds: total wall-clock limit for parallel execution, per-suite limit during ++# isolated fallback. Overridable because the limit that fits a CI machine is a property of the ++# machine, not of the suite: every kernel is compiled by Clang on the host, so a slower or ++# narrower CPU spends proportionally longer here with nothing wrong. ++_SUITE_TIMEOUT = int(os.environ.get("WARP_TEST_SUITE_TIMEOUT", 3600)) + _WARP_CACHE_PATH_ENV = "WARP_CACHE_PATH" + + +-- +2.43.0 + From 13620cecb5fc6a07ce8daf22f9f7272ba5139d91 Mon Sep 17 00:00:00 2001 From: Ludovic Henry Date: Thu, 24 Sep 2026 11:28:20 +0000 Subject: [PATCH 2/2] warp-lang: raise the tile-OOB subprocess timeout on riscv64 The test job failed with subprocess.TimeoutExpired on all four TestTileOOB CPU cases: _run_in_subprocess() hard-codes 60s for a subprocess to import warp and JIT-compile a debug-mode kernel with Clang, and this runner's Clang JIT needs longer than that (other single-module compiles in the same run already take 120-130s). Adds patch 0005, making the timeout overridable via WARP_TEST_TILE_OOB_TIMEOUT the same way patch 0004 already does for unittest_parallel's whole-run cap, and sets it to 600s in the workflow. --- .github/workflows/build-warp-lang.yml | 6 ++- docs/packages/warp-lang.yaml | 3 +- ...vironment-raise-the-tile-OOB-subproc.patch | 53 +++++++++++++++++++ 3 files changed, 60 insertions(+), 2 deletions(-) create mode 100644 patches/warp-lang/1.16.0/0005-tests-let-the-environment-raise-the-tile-OOB-subproc.patch diff --git a/.github/workflows/build-warp-lang.yml b/.github/workflows/build-warp-lang.yml index 8e1fc2ce054..7d54739283a 100644 --- a/.github/workflows/build-warp-lang.yml +++ b/.github/workflows/build-warp-lang.yml @@ -172,9 +172,13 @@ jobs: # The runner's four cores compile every kernel in the suite with Clang, which does not fit # unittest_parallel's 3600s default wall-clock limit (patch 0004 makes it settable). + # TestTileOOB additionally forks a subprocess per case to import warp and JIT-compile a + # debug-mode kernel within a hard-coded 60s (patch 0005 makes that settable too); observed + # single-module Clang compiles elsewhere in the suite already run 120-130s on this runner. - name: Run upstream's test suite against the wheel run: | - docker run --rm -v "$(pwd)/dist:/dist:ro" -w /tmp -e WARP_TEST_SUITE_TIMEOUT=43200 \ + docker run --rm -v "$(pwd)/dist:/dist:ro" -w /tmp \ + -e WARP_TEST_SUITE_TIMEOUT=43200 -e WARP_TEST_TILE_OOB_TIMEOUT=600 \ "${MANYLINUX_RISCV64_IMAGE}" bash -c ' set -euxo pipefail pybin=/opt/python/cp312-cp312/bin diff --git a/docs/packages/warp-lang.yaml b/docs/packages/warp-lang.yaml index f8c2e119872..02113a47fd2 100644 --- a/docs/packages/warp-lang.yaml +++ b/docs/packages/warp-lang.yaml @@ -11,4 +11,5 @@ versions: to the build and packaging platform tables, name the rv64gc/lp64d ABI for the JIT (which the Clang frontend does not default to), convert fp16 in software because rv64gc has no half-precision hardware and the JIT resolves no compiler-rt builtins, and let the - environment raise the test runner's wall-clock limit. + environment raise both of the test runner's hard-coded wall-clock limits (the parallel + run's overall cap and the tile-OOB tests' per-subprocess cap). diff --git a/patches/warp-lang/1.16.0/0005-tests-let-the-environment-raise-the-tile-OOB-subproc.patch b/patches/warp-lang/1.16.0/0005-tests-let-the-environment-raise-the-tile-OOB-subproc.patch new file mode 100644 index 00000000000..d9f888323e3 --- /dev/null +++ b/patches/warp-lang/1.16.0/0005-tests-let-the-environment-raise-the-tile-OOB-subproc.patch @@ -0,0 +1,53 @@ +From a8546cb3bd1952d17723b11ae6af75de1081c5ca Mon Sep 17 00:00:00 2001 +From: Ludovic Henry +Date: Thu, 24 Sep 2026 12:00:00 +0000 +Subject: [PATCH 5/5] tests: let the environment raise the tile-OOB subprocess + timeout + +Upstream-Status: To upstream [no riscv64 runner upstream to regression-test it on] + +_run_in_subprocess() gives the child process a hard-coded 60s to import warp, +JIT-compile the debug-mode tile-bounds kernel with Clang, and abort -- and all +four CPU cases in TestTileOOB hit that ceiling on the riscv64 runner +(subprocess.TimeoutExpired at exactly 60s each), the same shape as +unittest_parallel's own hard-coded wall-clock cap that patch 0004 already +makes overridable: warp compiles every kernel on the machine running the +tests, so how long that takes is a property of the machine, not of this tiny +module's five one-line kernels. + +Read the limit from WARP_TEST_TILE_OOB_TIMEOUT, keeping 60 as the default so +nothing changes for anyone who does not set it. +--- + warp/tests/tile/test_tile_oob.py | 8 +++++++- + 1 file changed, 7 insertions(+), 1 deletion(-) + +diff --git a/warp/tests/tile/test_tile_oob.py b/warp/tests/tile/test_tile_oob.py +index 508971a..9e4b79e 100644 +--- a/warp/tests/tile/test_tile_oob.py ++++ b/warp/tests/tile/test_tile_oob.py +@@ -1,6 +1,7 @@ + # SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + # SPDX-License-Identifier: Apache-2.0 + ++import os + import signal + import subprocess + import sys +@@ -13,8 +14,13 @@ from warp.tests.unittest_utils import add_function_test, get_cuda_test_devices + # in debug. Subprocesses then amortize compilation through the normal kernel cache. + wp.set_module_options({"mode": "debug"}) + ++# Overridable because the 60s the subprocess gets to import warp and JIT-compile the debug ++# kernel is a property of the machine running it, not of the kernel: on a slower CPU the ++# Clang invocation alone can cost more than that with nothing wrong. ++_SUBPROCESS_TIMEOUT = int(os.environ.get("WARP_TEST_TILE_OOB_TIMEOUT", 60)) + +-def _run_in_subprocess(func_name: str, device, timeout: int = 60): ++ ++def _run_in_subprocess(func_name: str, device, timeout: int = _SUBPROCESS_TIMEOUT): + result = subprocess.run( + [ + sys.executable, +-- +2.43.0 +