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 +