From 8f7fb01b13c3ee4352eed3d2ce61362a9672628a Mon Sep 17 00:00:00 2001 From: Ludovic Henry Date: Sun, 20 Sep 2026 20:52:21 +0000 Subject: [PATCH] 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 0000000000..8e1fc2ce05 --- /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 0000000000..f8c2e11987 --- /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 0000000000..c32cb63645 --- /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 0000000000..c40f88ac06 --- /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 0000000000..6bea7fe0b8 --- /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 0000000000..cb3d8859a6 --- /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 +