diff --git a/AGENTS.md b/AGENTS.md index fdf579f..eeda3e2 100644 --- a/AGENTS.md +++ b/AGENTS.md @@ -38,7 +38,7 @@ CMake. **Tools:** - `tools/warpforth-translate/warpforth-translate.cpp` - Translation tool entry point - `tools/warpforth-opt/warpforth-opt.cpp` - Optimization tool entry point -- `tools/warpforth-runner/warpforth-runner.cpp` - PTX execution tool for GPU kernels +- `gpu_test/warpforth_runner.py` - PTX execution tool for GPU kernels ## Tools Usage @@ -55,8 +55,8 @@ CMake. ./build/bin/warpforth-opt --warpforth-pipeline | \ ./build/bin/warpforth-translate --mlir-to-ptx > kernel.ptx -# Execute PTX on GPU -./warpforth-runner kernel.ptx --param i64[]:1,2,3 --param i64:42 --output-param 0 --output-count 3 +# Execute PTX on GPU from a JSON request +uv run python gpu_test/warpforth_runner.py kernel.ptx request.json result.json ``` ## Adding New Operations diff --git a/README.md b/README.md index 846b739..514eaa0 100644 --- a/README.md +++ b/README.md @@ -10,6 +10,7 @@ See the [documentation](https://tetsuo-cpp.github.io/warpforth/). - CMake - C++17 compiler - CUDA toolkit (for GPU execution) +- `cuda-python>=12.8,<13` and an NVIDIA driver (for GPU execution) - [uv](https://github.com/astral-sh/uv) (for Python test tooling) ## Building @@ -56,17 +57,38 @@ Compile to PTX: ./build/bin/warpforthc matmul.forth -o matmul.ptx --arch sm_80 ``` -Test on a GPU (A is 2x4 row-major, B is 4x3 row-major, C is 2x3 output): +Test on a GPU (A is 2x4 row-major, B is 4x3 row-major, C is 2x3 output). +Save the launch options in `request.json`: + +```json +{ + "kernel": "main", + "grid": [6, 1, 1], + "block": [1, 1, 1], + "params": [ + {"type": "i64[]", "values": [1,2,3,4,5,6,7,8]}, + {"type": "i64[]", "values": [1,2,3,4,5,6,7,8,9,10,11,12]}, + {"type": "i64[]", "values": [0,0,0,0,0,0]} + ], + "outputs": [{"param": 2, "count": 6}] +} +``` ```bash -./build/bin/warpforth-runner matmul.ptx \ - --param 'i64[]:1,2,3,4,5,6,7,8' \ - --param 'i64[]:1,2,3,4,5,6,7,8,9,10,11,12' \ - --param 'i64[]:0,0,0,0,0,0' \ - --grid 6,1,1 --block 1,1,1 \ - --output-param 2 --output-count 6 +uv run python gpu_test/warpforth_runner.py matmul.ptx request.json result.json ``` +`kernel` is required. `grid` and `block` default to +`[1,1,1]`; `outputs` defaults to `[]`. Each parameter has type `i64` or `f64` +with a scalar `value`, or type `i64[]` or `f64[]` with a nonempty array `values`. +Each output selects an array parameter by zero-based `param` index and may +specify `count` (default: the full array; zero is allowed). Multiple outputs +are returned in request order. Values must be signed 64-bit integers or finite +floats; nonfinite GPU results produce an error rather than invalid JSON. +The runner writes `result.json` with `{"status":"ok","outputs":[{"param":2,"type":"i64[]","values":[70,80,90,158,184,210]}]}`. +On failure it writes `{"status":"error","error":"..."}` to that file and exits +with status 1. No request or result data is sent through stdin or stdout. + ## Toolchain | Tool | Description | @@ -74,7 +96,7 @@ Test on a GPU (A is 2x4 row-major, B is 4x3 row-major, C is 2x3 output): | `warpforthc` | Compiles Forth source to PTX | | `warpforth-translate` | Translates from Forth source to MLIR and MLIR to PTX assembly | | `warpforth-opt` | Runs individual MLIR passes or entire pipeline | -| `warpforth-runner` | Executes PTX kernels on a GPU for testing | +| `gpu_test/warpforth_runner.py` | Executes PTX kernels on a GPU for testing | These tools can be composed for debugging or inspecting intermediate stages: diff --git a/gpu_test/conftest.py b/gpu_test/conftest.py index 8578f5b..0b5179c 100644 --- a/gpu_test/conftest.py +++ b/gpu_test/conftest.py @@ -3,6 +3,7 @@ from __future__ import annotations import atexit +import json import logging import os import subprocess @@ -18,6 +19,8 @@ from requests import RequestException from vastai import VastAI +from gpu_test.warpforth_runner import validate_request + if TYPE_CHECKING: from collections.abc import Generator from typing import Self @@ -26,7 +29,7 @@ PROJECT_ROOT = Path(__file__).resolve().parent.parent WARPFORTHC = PROJECT_ROOT / "build" / "bin" / "warpforthc" -RUNNER_SRC = PROJECT_ROOT / "tools" / "warpforth-runner" / "warpforth-runner.cpp" +RUNNER_SRC = PROJECT_ROOT / "gpu_test" / "warpforth_runner.py" MAX_COST_PER_HOUR = 0.50 POLL_INTERVAL_S = 10 @@ -329,7 +332,7 @@ def _wait_for_ssh(self) -> None: ) self._attach_ssh_key() self._wait_for_sshd() - self._compile_runner() + self._install_runner() return time.sleep(POLL_INTERVAL_S) @@ -356,14 +359,15 @@ def _wait_for_sshd(self) -> None: ) raise TimeoutError(msg) - def _compile_runner(self) -> None: - """Upload warpforth-runner.cpp and compile it on the remote host.""" - self.scp_upload(RUNNER_SRC, f"{REMOTE_TMP}/warpforth-runner.cpp") - nvcc_cmd = ( - f"nvcc -o {REMOTE_TMP}/warpforth-runner" - f" {REMOTE_TMP}/warpforth-runner.cpp -lcuda -std=c++17" + def _install_runner(self) -> None: + """Upload the Python runner and install its CUDA bindings without nvcc.""" + self.scp_upload(RUNNER_SRC, f"{REMOTE_TMP}/warpforth_runner.py") + self.ssh_run( + "apt-get update && apt-get install -y python3 python3-venv && " + f"python3 -m venv {REMOTE_TMP}/warpforth-venv && " + f"{REMOTE_TMP}/warpforth-venv/bin/pip install 'cuda-python>=12.8,<13'", + timeout=300, ) - self.ssh_run(nvcc_cmd, timeout=60) def _record_labeled_instance_ids( self, @@ -494,8 +498,9 @@ def ssh_run(self, cmd: str, *, timeout: int = 120) -> str: check=False, ) if result.returncode != 0: - msg = f"SSH command failed (rc={result.returncode}):\n{result.stderr}" - raise RuntimeError(msg) + raise subprocess.CalledProcessError( + result.returncode, cmd, output=result.stdout, stderr=result.stderr + ) return result.stdout def scp_upload(self, local_path: str | Path, remote_path: str) -> None: @@ -515,6 +520,23 @@ def scp_upload(self, local_path: str | Path, remote_path: str) -> None: check=True, ) + def scp_download(self, remote_path: str, local_path: str | Path) -> None: + """Download a file from the remote instance via SCP.""" + subprocess.run( + [ + "scp", + *self._ssh_options(), + "-P", + str(self.ssh_port), + f"root@{self.ssh_host}:{remote_path}", + str(local_path), + ], + capture_output=True, + text=True, + timeout=60, + check=True, + ) + def _parse_array_type(type_spec: str) -> tuple[str, int]: """Parse 'i64[256]' or 'f64[256]' into (base_type, size).""" @@ -603,7 +625,8 @@ def run( block: tuple[int, int, int] = (1, 1, 1), output_param: int = 0, output_count: int | None = None, - ) -> list[int] | list[float]: + outputs: list[dict[str, int]] | None = None, + ) -> list[int] | list[float] | list[list[int] | list[float]]: """Compile Forth source locally, execute on remote GPU, return output values. Param buffer sizes are derived from the Forth source's 'param' declarations. @@ -611,6 +634,8 @@ def run( - Array params: list of int or float (padded with zeros to declared size) - Scalar params: int or float Params not in the dict are zero-initialized. + With outputs=[{"param": index, "count": optional_count}, ...], returns + one values list per output. Otherwise returns the single output's values. """ # Parse kernel name and param declarations kernel_name = _parse_kernel_name(forth_source) @@ -621,36 +646,10 @@ def run( params = params or {} - # Validate output_param - if output_param < 0 or output_param >= len(decls): - msg = f"output_param {output_param} out of range (have {len(decls)} params)" - raise ValueError(msg) - if not decls[output_param].is_array: - name = decls[output_param].name - msg = f"output_param {output_param} ('{name}') is a scalar and cannot be read back" - raise ValueError(msg) - # Compile locally ptx = self.compiler.compile_source(forth_source) - # Write PTX to temp file and upload - with tempfile.NamedTemporaryFile(mode="w", suffix=".ptx", delete=False) as f: - f.write(ptx) - ptx_path = Path(f.name) - - try: - self.session.scp_upload(ptx_path, f"{REMOTE_TMP}/kernel.ptx") - finally: - ptx_path.unlink() - - # Build remote command - cmd_parts = [ - f"{REMOTE_TMP}/warpforth-runner", - f"{REMOTE_TMP}/kernel.ptx", - "--kernel", - kernel_name, - ] - + wire_params = [] for decl in decls: if decl.is_array: values = params.get(decl.name, []) @@ -661,35 +660,56 @@ def run( buf = [zero] * decl.size for i, v in enumerate(values): buf[i] = v - cmd_parts.extend(["--param", f"{decl.base_type}[]:{','.join(str(v) for v in buf)}"]) + wire_params.append({"type": f"{decl.base_type}[]", "values": buf}) else: value = params.get(decl.name, 0.0 if decl.base_type == "f64" else 0) if isinstance(value, list): msg = f"Scalar param '{decl.name}' expects a scalar, got list" raise TypeError(msg) - cmd_parts.extend(["--param", f"{decl.base_type}:{value}"]) - - cmd_parts.extend( - [ - "--grid", - f"{grid[0]},{grid[1]},{grid[2]}", - "--block", - f"{block[0]},{block[1]},{block[2]}", - "--output-param", - str(output_param), - ] - ) + wire_params.append({"type": decl.base_type, "value": value}) + single_output = {"param": output_param} if output_count is not None: - cmd_parts.extend(["--output-count", str(output_count)]) - - cmd = " ".join(cmd_parts) - stdout = self.session.ssh_run(cmd, timeout=120) - - # Parse CSV output — type depends on the output param - out_type = decls[output_param].base_type - parse = float if out_type == "f64" else int - return [parse(v) for v in stdout.strip().split(",")] + single_output["count"] = output_count + request = { + "kernel": kernel_name, + "grid": list(grid), + "block": list(block), + "params": wire_params, + "outputs": outputs if outputs is not None else [single_output], + } + validate_request(request) + response = self._execute(ptx, request) + values = [output["values"] for output in response["outputs"]] + return values if outputs is not None else values[0] + + def _execute(self, ptx: str, request: dict) -> dict: + """Upload input files and download the result, including runner errors.""" + with tempfile.TemporaryDirectory() as directory: + ptx_path = Path(directory) / "kernel.ptx" + request_path = Path(directory) / "request.json" + result_path = Path(directory) / "result.json" + ptx_path.write_text(ptx) + request_path.write_text(json.dumps(request, allow_nan=False)) + self.session.scp_upload(ptx_path, f"{REMOTE_TMP}/kernel.ptx") + self.session.scp_upload(request_path, f"{REMOTE_TMP}/request.json") + # A failed invocation must never read a previous invocation's result. + remote_result = f"{REMOTE_TMP}/result-{uuid4().hex}.json" + try: + self.session.ssh_run( + f"{REMOTE_TMP}/warpforth-venv/bin/python {REMOTE_TMP}/warpforth_runner.py" + f" {REMOTE_TMP}/kernel.ptx {REMOTE_TMP}/request.json {remote_result}", + timeout=120, + ) + except subprocess.CalledProcessError as error: + # Exit 1 writes a runner error file; SSH failures remain transport errors. + if error.returncode != 1: + raise + self.session.scp_download(remote_result, result_path) + response = json.loads(result_path.read_text()) + if response["status"] != "ok": + raise RuntimeError(response["error"]) + return response # --- Fixtures --- diff --git a/gpu_test/test_kernels.py b/gpu_test/test_kernels.py index 96a08f3..f2e974c 100644 --- a/gpu_test/test_kernels.py +++ b/gpu_test/test_kernels.py @@ -14,6 +14,17 @@ pytestmark = pytest.mark.gpu +def test_multiple_outputs(kernel_runner: KernelRunner) -> None: + result = kernel_runner.run( + forth_source=( + "\\! kernel main\n\\! param INTS i64[2]\n\\! param FLOATS f64[2]\n" + "42 INTS ! 3.25 FLOATS F!" + ), + outputs=[{"param": 1}, {"param": 0, "count": 1}], + ) + assert result == [[3.25, 0.0], [42]] + + # --- Arithmetic --- diff --git a/gpu_test/test_runner.py b/gpu_test/test_runner.py new file mode 100644 index 0000000..6d13d4f --- /dev/null +++ b/gpu_test/test_runner.py @@ -0,0 +1,326 @@ +"""GPU-independent tests of the JSON runner and its SSH integration.""" + +from __future__ import annotations + +import ctypes +import json +import subprocess +from pathlib import Path +from unittest.mock import Mock + +import cuda.bindings +import pytest +from cuda.bindings import driver + +from gpu_test import warpforth_runner +from gpu_test.conftest import KernelRunner, VastSession + + +@pytest.fixture +def ptx_file(tmp_path: Path, request: pytest.FixtureRequest) -> Path: + path = tmp_path / "kernel.ptx" + path.write_bytes(b"// PTX\n" + getattr(request, "param", b"")) + return path + + +@pytest.fixture +def request_data() -> dict: + return { + "kernel": "main", + "grid": [2, 3, 4], + "block": [5, 6, 7], + "params": [ + {"type": "i64[]", "values": [-(2**63), 2**63 - 1]}, + {"type": "f64", "value": 3.14}, + {"type": "f64[]", "values": [1.25, -2.5]}, + {"type": "i64", "value": -42}, + ], + "outputs": [{"param": 2}, {"param": 0, "count": 1}, {"param": 0, "count": 0}], + } + + +@pytest.fixture +def json_files(tmp_path: Path, request_data: dict) -> tuple[Path, Path]: + request = tmp_path / "request.json" + request.write_text(json.dumps(request_data)) + return request, tmp_path / "result.json" + + +@pytest.fixture +def cuda_driver(monkeypatch: pytest.MonkeyPatch) -> Mock: + fake = Mock(spec=driver) + fake.CUresult = driver.CUresult + for name in ( + "cuInit", + "cuMemcpyHtoD", + "cuMemcpyDtoH", + "cuLaunchKernel", + "cuCtxSynchronize", + "cuMemFree", + "cuModuleUnload", + "cuCtxDestroy", + ): + getattr(fake, name).return_value = (driver.CUresult.CUDA_SUCCESS,) + fake.cuDeviceGet.return_value = (0, 0) + fake.cuCtxCreate.return_value = (0, 10) + fake.cuModuleLoadData.return_value = (0, 20) + fake.cuModuleGetFunction.return_value = (0, 30) + fake.cuGetErrorName.return_value = (0, b"CUDA_ERROR_INVALID_PTX") + fake.cuGetErrorString.return_value = (0, b"the provided PTX was invalid") + allocations = [] + + def allocate(size: int) -> tuple: + storage = ctypes.create_string_buffer(size) + allocations.append(storage) + return (0, ctypes.addressof(storage)) + + def copy(destination: int, source: int, size: int) -> tuple: + ctypes.memmove(destination, source, size) + return (0,) + + fake.cuMemAlloc.side_effect = allocate + fake.cuMemcpyHtoD.side_effect = copy + fake.cuMemcpyDtoH.side_effect = copy + monkeypatch.setattr(cuda.bindings, "driver", fake) + return fake + + +@pytest.mark.parametrize("ptx_file", [b"", b"\0"], indirect=True) +def test_launch_and_multiple_outputs(cuda_driver: Mock, request_data: dict, ptx_file: Path) -> None: + def launch(*args: object) -> tuple: + assert args[:9] == (30, 2, 3, 4, 5, 6, 7, 0, 0) + pointers = (ctypes.c_void_p * 4).from_address(args[9]) + integer_address = ctypes.c_uint64.from_address(pointers[0]).value + integers = (ctypes.c_int64 * 2).from_address(integer_address) + assert list(integers) == [-(2**63), 2**63 - 1] + assert ctypes.c_double.from_address(pointers[1]).value == 3.14 + float_address = ctypes.c_uint64.from_address(pointers[2]).value + floats = (ctypes.c_double * 2).from_address(float_address) + assert list(floats) == [1.25, -2.5] + assert ctypes.c_int64.from_address(pointers[3]).value == -42 + integers[0] = 42 + floats[0] = 2.75 + return (0,) + + cuda_driver.cuLaunchKernel.side_effect = launch + assert warpforth_runner.execute(ptx_file, request_data) == { + "status": "ok", + "outputs": [ + {"param": 2, "type": "f64[]", "values": [2.75, -2.5]}, + {"param": 0, "type": "i64[]", "values": [42]}, + {"param": 0, "type": "i64[]", "values": []}, + ], + } + cuda_driver.cuModuleLoadData.assert_called_once_with(b"// PTX\n\0") + cuda_driver.cuModuleGetFunction.assert_called_once_with(20, b"main") + cuda_driver.cuCtxSynchronize.assert_called_once() + assert cuda_driver.cuMemcpyDtoH.call_count == 2 + assert cuda_driver.cuMemFree.call_count == 2 + cuda_driver.cuModuleUnload.assert_called_once_with(20) + cuda_driver.cuCtxDestroy.assert_called_once_with(10) + + +def test_no_parameters_or_outputs(cuda_driver: Mock, ptx_file: Path) -> None: + request = {"kernel": "main"} + assert warpforth_runner.execute(ptx_file, request) == {"status": "ok", "outputs": []} + cuda_driver.cuLaunchKernel.assert_called_once_with(30, 1, 1, 1, 1, 1, 1, 0, 0, 0, 0) + + +@pytest.mark.parametrize( + ("stage", "allocations", "modules"), + [ + ("cuModuleLoadData", 0, 0), + ("cuMemcpyHtoD", 1, 1), + ("cuLaunchKernel", 2, 1), + ("cuCtxSynchronize", 2, 1), + ("cuMemcpyDtoH", 2, 1), + ], +) +def test_cuda_errors_cleanup( # noqa: PLR0913 + cuda_driver: Mock, + request_data: dict, + ptx_file: Path, + stage: str, + allocations: int, + modules: int, +) -> None: + call = getattr(cuda_driver, stage) + call.side_effect = None + call.return_value = (driver.CUresult.CUDA_ERROR_INVALID_PTX,) + with pytest.raises(RuntimeError, match="CUDA_ERROR_INVALID_PTX: the provided PTX was invalid"): + warpforth_runner.execute(ptx_file, request_data) + assert cuda_driver.cuMemFree.call_count == allocations + assert cuda_driver.cuModuleUnload.call_count == modules + cuda_driver.cuCtxDestroy.assert_called_once() + + +@pytest.mark.parametrize( + "update", + [ + {"kernel": ""}, + {"grid": [1, 2]}, + {"block": [0, 1, 1]}, + {"grid": [True, 1, 1]}, + {"params": [{"type": "i32", "value": 1}]}, + {"params": [{"type": "i64", "value": 2**63}]}, + {"params": [{"type": "i64", "value": 1.5}]}, + {"params": [{"type": "i64[]", "values": []}]}, + {"params": [{"type": "f64", "value": float("inf")}]}, + {"outputs": [{"param": -1}]}, + {"outputs": [{"param": 4}]}, + {"outputs": [{"param": 1}]}, + {"outputs": [{"param": 0, "count": -1}]}, + {"outputs": [{"param": 0, "count": 3}]}, + ], +) +def test_invalid_request( + cuda_driver: Mock, request_data: dict, update: dict, ptx_file: Path +) -> None: + request_data.update(update) + with pytest.raises((ValueError, TypeError)): + warpforth_runner.execute(ptx_file, request_data) + cuda_driver.cuInit.assert_not_called() + + +@pytest.mark.parametrize("text", [None, "not json", "null", "{}"]) +def test_cli_invalid_input( + text: str | None, json_files: tuple[Path, Path], capsys: pytest.CaptureFixture, ptx_file: Path +) -> None: + request, result = json_files + if text is None: + request.unlink() + else: + request.write_text(text) + assert warpforth_runner.main([str(ptx_file), str(request), str(result)]) == 1 + captured = capsys.readouterr() + assert captured.out == captured.err == "" + assert json.loads(result.read_text())["status"] == "error" + + +@pytest.mark.parametrize("contents", [None, b"", b"// PTX\0truncated"]) +def test_cli_ptx_file_errors( + contents: bytes | None, + ptx_file: Path, + json_files: tuple[Path, Path], +) -> None: + if contents is None: + ptx_file.unlink() + else: + ptx_file.write_bytes(contents) + request, result = json_files + assert warpforth_runner.main([str(ptx_file), str(request), str(result)]) == 1 + response = json.loads(result.read_text()) + assert response["status"] == "error" + assert response["error"] + + +def test_cli_success( + cuda_driver: Mock, + ptx_file: Path, + json_files: tuple[Path, Path], + capsys: pytest.CaptureFixture, +) -> None: + request, result = json_files + result.write_text("stale result") + assert warpforth_runner.main([str(ptx_file), str(request), str(result)]) == 0 + captured = capsys.readouterr() + assert captured.out == captured.err == "" + assert json.loads(result.read_text())["status"] == "ok" + cuda_driver.cuLaunchKernel.assert_called_once() + + +def test_cli_cuda_error( + cuda_driver: Mock, + ptx_file: Path, + json_files: tuple[Path, Path], + capsys: pytest.CaptureFixture, +) -> None: + cuda_driver.cuModuleLoadData.return_value = (driver.CUresult.CUDA_ERROR_INVALID_PTX,) + request, result = json_files + assert warpforth_runner.main([str(ptx_file), str(request), str(result)]) == 1 + captured = capsys.readouterr() + assert captured.out == captured.err == "" + assert json.loads(result.read_text()) == { + "status": "error", + "error": "CUDA_ERROR_INVALID_PTX: the provided PTX was invalid", + } + + +FORTH = "\\! kernel main\n\\! param a i64[3]\n\\! param b f64[2]\n\\! param c f64\n" + + +@pytest.mark.parametrize("outputs", [None, [{"param": 1}, {"param": 0, "count": 1}]]) +def test_harness_json(outputs: list | None) -> None: + session, compiler = Mock(), Mock() + compiler.compile_source.return_value = "// PTX\n" + uploaded = [] + session.scp_upload.side_effect = lambda path, _remote: uploaded.append(Path(path).read_bytes()) + session.ssh_run.return_value = "ignored stdout" + session.scp_download.side_effect = lambda _remote, path: path.write_text( + json.dumps({"status": "ok", "outputs": [{"values": [1]}, {"values": [2]}]}) + ) + result = KernelRunner(session, compiler).run( + FORTH, params={"a": [2**62], "c": 1.5}, outputs=outputs + ) + assert result == ([1] if outputs is None else [[1], [2]]) + request = json.loads(uploaded[1]) + assert set(request) == {"kernel", "grid", "block", "params", "outputs"} + assert uploaded[0] == b"// PTX\n" + command = session.ssh_run.call_args.args[0] + for call in session.scp_upload.call_args_list: + local_path, remote_path = call.args + assert not Path(local_path).exists() + assert f" {remote_path}" in command + remote_result, local_result = session.scp_download.call_args.args + assert command.endswith(f" {remote_result}") + assert not local_result.exists() + assert "stdin" not in session.ssh_run.call_args.kwargs + assert request["params"] == [ + {"type": "i64[]", "values": [2**62, 0, 0]}, + {"type": "f64[]", "values": [0.0, 0.0]}, + {"type": "f64", "value": 1.5}, + ] + assert request["outputs"] == ([{"param": 0}] if outputs is None else outputs) + + +@pytest.mark.parametrize("returncode", [1, 255]) +def test_harness_errors(returncode: int) -> None: + session, compiler = Mock(), Mock() + compiler.compile_source.return_value = "// PTX" + session.ssh_run.side_effect = subprocess.CalledProcessError(returncode, "runner") + session.scp_download.side_effect = lambda _remote, path: path.write_text( + '{"status":"error","error":"CUDA_ERROR_INVALID_PTX"}' + ) + expected = RuntimeError if returncode == 1 else subprocess.CalledProcessError + with pytest.raises(expected): + KernelRunner(session, compiler).run(FORTH) + assert session.scp_download.call_count == (1 if returncode == 1 else 0) + + +def test_scp_download(monkeypatch: pytest.MonkeyPatch, tmp_path: Path) -> None: + session = object.__new__(VastSession) + session.ssh_host = "host" + session.ssh_port = 22 + monkeypatch.setattr(session, "_ssh_options", list) + run = Mock() + monkeypatch.setattr(subprocess, "run", run) + destination = tmp_path / "result.json" + session.scp_download("/tmp/result.json", destination) # noqa: S108 + assert run.call_args.args[0] == [ + "scp", + "-P", + "22", + "root@host:/tmp/result.json", + str(destination), + ] + assert run.call_args.kwargs["check"] is True + + +def test_install_runner() -> None: + session = Mock() + VastSession._install_runner(session) # noqa: SLF001 + assert session.scp_upload.call_args.args[0].name == "warpforth_runner.py" + command = session.ssh_run.call_args.args[0] + assert "pip install 'cuda-python>=12.8,<13'" in command + assert "python3 -m venv" in command + assert "nvcc" not in command diff --git a/gpu_test/warpforth_runner.py b/gpu_test/warpforth_runner.py new file mode 100644 index 0000000..2972027 --- /dev/null +++ b/gpu_test/warpforth_runner.py @@ -0,0 +1,167 @@ +"""Execute a PTX file using JSON request and result files.""" + +from __future__ import annotations + +import argparse +import ctypes +import json +import math +import sys +from contextlib import ExitStack +from pathlib import Path + + +def validate_param(param: dict) -> None: + kind = param["type"] + if kind not in ("i64", "f64", "i64[]", "f64[]"): + msg = f"Unsupported parameter type: {kind}" + raise ValueError(msg) + values = param["values"] if kind.endswith("[]") else [param["value"]] + if not isinstance(values, list) or not values: + msg = "Array parameters must contain a nonempty values list" + raise ValueError(msg) + for value in values: + if kind.startswith("i64"): + valid = type(value) is int and -(2**63) <= value < 2**63 + else: + valid = type(value) in (int, float) and math.isfinite(value) + if not valid: + msg = f"Invalid {kind} value: {value}" + raise ValueError(msg) + + +def validate_request(request: dict) -> None: + """Validate the wire protocol before making any CUDA calls.""" + kernel = request["kernel"] + if not isinstance(kernel, str) or not kernel or "\0" in kernel: + msg = "kernel must be a nonempty name without NUL bytes" + raise ValueError(msg) + for name in ("grid", "block"): + dims = request.get(name, [1, 1, 1]) + if ( + not isinstance(dims, list) + or len(dims) != 3 # noqa: PLR2004 - CUDA uses three-dimensional launches + or any(type(d) is not int or not 0 < d < 2**32 for d in dims) + ): + msg = f"{name} must contain three positive uint32 dimensions" + raise ValueError(msg) + params = request.get("params", []) + if not isinstance(params, list): + msg = "params must be a list" + raise TypeError(msg) + for param in params: + validate_param(param) + validate_outputs(request.get("outputs", []), params) + + +def validate_outputs(outputs: list, params: list) -> None: + if not isinstance(outputs, list): + msg = "outputs must be a list" + raise TypeError(msg) + for output in outputs: + index = output["param"] + if type(index) is not int or not 0 <= index < len(params): + msg = f"Output parameter index out of range: {index}" + raise ValueError(msg) + param = params[index] + if not param["type"].endswith("[]"): + msg = f"Output parameter {index} is a scalar" + raise ValueError(msg) + count = output.get("count", len(param["values"])) + if type(count) is not int or not 0 <= count <= len(param["values"]): + msg = f"Output count out of range for parameter {index}: {count}" + raise ValueError(msg) + + +def execute(ptx_path: Path, request: dict) -> dict: + validate_request(request) + # warpforthc emits NUL-terminated PTX; accept both terminated and plain text. + ptx = ptx_path.read_bytes().rstrip(b"\0") + if not ptx or b"\0" in ptx: + msg = "PTX file must contain nonempty PTX without embedded NUL bytes" + raise ValueError(msg) + # Import here so missing bindings/driver libraries also produce JSON errors. + from cuda.bindings import driver # noqa: PLC0415 + + def check(result: tuple) -> object: + error, *values = result + if error != driver.CUresult.CUDA_SUCCESS: + _, name = driver.cuGetErrorName(error) + _, description = driver.cuGetErrorString(error) + msg = f"{name.decode()}: {description.decode()}" + raise RuntimeError(msg) + return values[0] if values else None + + check(driver.cuInit(0)) + device = check(driver.cuDeviceGet(0)) + with ExitStack() as resources: + context = check(driver.cuCtxCreate(0, device)) + resources.callback(driver.cuCtxDestroy, context) + module = check(driver.cuModuleLoadData(ptx + b"\0")) + resources.callback(driver.cuModuleUnload, module) + kernel = check(driver.cuModuleGetFunction(module, request["kernel"].encode())) + buffers = {} + arguments = [] + for index, param in enumerate(request.get("params", [])): + scalar_type = ctypes.c_int64 if param["type"].startswith("i64") else ctypes.c_double + if param["type"].endswith("[]"): + host = (scalar_type * len(param["values"]))(*param["values"]) + pointer = check(driver.cuMemAlloc(ctypes.sizeof(host))) + resources.callback(driver.cuMemFree, pointer) + check(driver.cuMemcpyHtoD(pointer, ctypes.addressof(host), ctypes.sizeof(host))) + buffers[index] = (host, pointer) + arguments.append(ctypes.c_uint64(int(pointer))) + else: + arguments.append(scalar_type(param["value"])) + # Keep argument storage alive until the synchronous launch completes. + pointers = (ctypes.c_void_p * len(arguments))(*(ctypes.addressof(a) for a in arguments)) + check( + driver.cuLaunchKernel( + kernel, + *request.get("grid", [1, 1, 1]), + *request.get("block", [1, 1, 1]), + 0, + 0, + ctypes.addressof(pointers) if arguments else 0, + 0, + ) + ) + check(driver.cuCtxSynchronize()) + outputs = [] + for output in request.get("outputs", []): + index = output["param"] + host, pointer = buffers[index] + count = output.get("count", len(host)) + if count: + check(driver.cuMemcpyDtoH(ctypes.addressof(host), pointer, count * 8)) + outputs.append( + { + "param": index, + "type": request["params"][index]["type"], + "values": list(host)[:count], + } + ) + return {"status": "ok", "outputs": outputs} + + +def main(argv: list[str] | None = None) -> int: + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("ptx", type=Path, help="PTX file to execute") + parser.add_argument("request", type=Path, help="JSON launch options file") + parser.add_argument("result", type=Path, help="JSON result file to write") + args = parser.parse_args(argv) + try: + response = execute(args.ptx, json.loads(args.request.read_text())) + # Reject nonfinite GPU results rather than emitting nonstandard JSON. + text = json.dumps(response, allow_nan=False) + except Exception as error: # noqa: BLE001 - all failures belong in the wire response + text = json.dumps({"status": "error", "error": str(error)}) + status = 1 + else: + status = 0 + args.result.write_text(text + "\n") + return status + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/pyproject.toml b/pyproject.toml index b1c1386..c14e774 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -3,6 +3,7 @@ name = "warpforth" version = "0.1.0" requires-python = ">=3.11" dependencies = [ + "cuda-python>=12.8,<13", "lit>=18.1.0", "numpy", "pytest", diff --git a/tools/CMakeLists.txt b/tools/CMakeLists.txt index 7d4740f..f9e8857 100644 --- a/tools/CMakeLists.txt +++ b/tools/CMakeLists.txt @@ -1,4 +1,3 @@ add_subdirectory(warpforth-translate) add_subdirectory(warpforth-opt) add_subdirectory(warpforthc) -add_subdirectory(warpforth-runner) diff --git a/tools/warpforth-runner/CMakeLists.txt b/tools/warpforth-runner/CMakeLists.txt deleted file mode 100644 index a2786f2..0000000 --- a/tools/warpforth-runner/CMakeLists.txt +++ /dev/null @@ -1,9 +0,0 @@ -find_package(CUDAToolkit QUIET) - -if(CUDAToolkit_FOUND) - add_executable(warpforth-runner warpforth-runner.cpp) - target_link_libraries(warpforth-runner PRIVATE CUDA::cuda_driver) - target_compile_features(warpforth-runner PRIVATE cxx_std_17) -else() - message(STATUS "CUDAToolkit not found — skipping warpforth-runner") -endif() diff --git a/tools/warpforth-runner/warpforth-runner.cpp b/tools/warpforth-runner/warpforth-runner.cpp deleted file mode 100644 index 67e1d3f..0000000 --- a/tools/warpforth-runner/warpforth-runner.cpp +++ /dev/null @@ -1,347 +0,0 @@ -/// warpforth-runner: Execute PTX kernels via the CUDA Driver API. -/// -/// Single-file C++ program designed to be uploaded and compiled on a remote -/// GPU host with `nvcc -o warpforth-runner warpforth-runner.cpp -lcuda -/// -std=c++17`. -/// -/// Usage: -/// warpforth-runner kernel.ptx --param i64[]:1,2,3 --param f64:3.14 \ -/// --grid 4,1,1 --block 64,1,1 --kernel main \ -/// --output-param 0 --output-count 3 - -#include - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#define CHECK_CU(call) \ - do { \ - CUresult err = (call); \ - if (err != CUDA_SUCCESS) { \ - const char *errStr = nullptr; \ - cuGetErrorString(err, &errStr); \ - std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__ << ": " \ - << (errStr ? errStr : "unknown") << "\n"; \ - exit(1); \ - } \ - } while (0) - -template struct ArrayParam { - std::vector values; - CUdeviceptr devicePtr = 0; -}; - -template struct ScalarParam { - T value; -}; - -using Param = std::variant, ArrayParam, - ScalarParam, ScalarParam>; - -template static void allocDevice(ArrayParam &arr) { - size_t bytes = arr.values.size() * sizeof(T); - CHECK_CU(cuMemAlloc(&arr.devicePtr, bytes)); - CHECK_CU(cuMemcpyHtoD(arr.devicePtr, arr.values.data(), bytes)); -} - -template -static void printOutput(ArrayParam &arr, size_t count) { - std::vector output(arr.values.size()); - CHECK_CU(cuMemcpyDtoH(output.data(), arr.devicePtr, - arr.values.size() * sizeof(T))); - for (size_t i = 0; i < count; ++i) { - if (i > 0) - std::cout << ","; - if constexpr (std::is_floating_point_v) - std::cout << std::setprecision(17) << output[i]; - else - std::cout << output[i]; - } - std::cout << "\n"; -} - -static void *kernelArgPtr(Param &p) { - if (auto *a = std::get_if>(&p)) - return &a->devicePtr; - if (auto *a = std::get_if>(&p)) - return &a->devicePtr; - if (auto *s = std::get_if>(&p)) - return &s->value; - return &std::get>(p).value; -} - -static bool isScalar(const Param &p) { - return std::holds_alternative>(p) || - std::holds_alternative>(p); -} - -struct Dims { - unsigned x = 1, y = 1, z = 1; -}; - -static int parseIntArg(std::string_view s, std::string_view optName) { - int value = 0; - auto [ptr, ec] = std::from_chars(s.data(), s.data() + s.size(), value); - if (ec != std::errc{} || ptr != s.data() + s.size()) { - std::cerr << "Error: " << optName << " expects an integer, got: " << s - << "\n"; - exit(1); - } - return value; -} - -static Dims parseDims(std::string_view s) { - Dims d; - const char *p = s.data(); - const char *end = s.data() + s.size(); - - auto dimsErr = [&]() { - std::cerr << "Error: expected 3 comma-separated values, got: " << s << "\n"; - exit(1); - }; - - auto [p1, ec1] = std::from_chars(p, end, d.x); - if (ec1 != std::errc{} || p1 == end || *p1 != ',') - dimsErr(); - - auto [p2, ec2] = std::from_chars(p1 + 1, end, d.y); - if (ec2 != std::errc{} || p2 == end || *p2 != ',') - dimsErr(); - - auto [p3, ec3] = std::from_chars(p2 + 1, end, d.z); - if (ec3 != std::errc{} || p3 != end) - dimsErr(); - - return d; -} - -static Param parseParam(std::string_view s) { - std::string input(s); - - auto colonPos = input.find(':'); - if (colonPos == std::string::npos) { - std::cerr << "Error: --param requires type prefix (e.g. i64:42 or " - "f64[]:1.0,2.0), got: " - << s << "\n"; - exit(1); - } - - std::string typePrefix = input.substr(0, colonPos); - std::string valueStr = input.substr(colonPos + 1); - - if (valueStr.empty()) { - std::cerr << "Error: --param requires at least one value, got: " << s - << "\n"; - exit(1); - } - - // Parse comma-separated values into a typed vector - auto parseValues = [&](auto convert) { - using T = decltype(convert(std::string{})); - std::vector vals; - std::istringstream iss(valueStr); - std::string token; - while (std::getline(iss, token, ',')) - vals.push_back(convert(token)); - return vals; - }; - - auto toI64 = [&](const std::string &tok) -> int64_t { - try { - return std::stoll(tok); - } catch (const std::exception &) { - std::cerr << "Error: invalid integer value '" << tok << "' in --param " - << s << "\n"; - exit(1); - } - }; - auto toF64 = [&](const std::string &tok) -> double { - try { - return std::stod(tok); - } catch (const std::exception &) { - std::cerr << "Error: invalid float value '" << tok << "' in --param " << s - << "\n"; - exit(1); - } - }; - - if (typePrefix == "i64[]") - return Param{ArrayParam{parseValues(toI64)}}; - if (typePrefix == "f64[]") - return Param{ArrayParam{parseValues(toF64)}}; - - // Scalars — must be exactly one value - if (valueStr.find(',') != std::string::npos) { - std::cerr << "Error: scalar param expects exactly one value, got: " << s - << "\n"; - exit(1); - } - - if (typePrefix == "i64") - return Param{ScalarParam{toI64(valueStr)}}; - if (typePrefix == "f64") - return Param{ScalarParam{toF64(valueStr)}}; - - std::cerr << "Error: unsupported param type '" << typePrefix - << "' (expected i64, i64[], f64, or f64[]), got: " << s << "\n"; - exit(1); -} - -static std::string readFile(std::string_view path) { - std::ifstream f(std::string(path), std::ios::binary); - if (!f) { - std::cerr << "Error: cannot open " << path << "\n"; - exit(1); - } - std::ostringstream ss; - ss << f.rdbuf(); - return ss.str(); -} - -int main(int argc, char **argv) { - const char *ptxFile = nullptr; - const char *kernelName = nullptr; - std::vector params; - Dims grid, block; - int outputParam = 0; - int outputCount = -1; // -1 = all - - // Parse arguments - for (int i = 1; i < argc; ++i) { - std::string_view arg = argv[i]; - auto needsValue = [&](std::string_view opt) { - if (++i >= argc) { - std::cerr << "Error: " << opt << " requires a value\n"; - exit(1); - } - }; - if (arg == "--param") { - needsValue("--param"); - params.push_back(parseParam(argv[i])); - } else if (arg == "--grid") { - needsValue("--grid"); - grid = parseDims(argv[i]); - } else if (arg == "--block") { - needsValue("--block"); - block = parseDims(argv[i]); - } else if (arg == "--output-param") { - needsValue("--output-param"); - outputParam = parseIntArg(argv[i], "--output-param"); - } else if (arg == "--output-count") { - needsValue("--output-count"); - outputCount = parseIntArg(argv[i], "--output-count"); - } else if (arg == "--kernel") { - needsValue("--kernel"); - kernelName = argv[i]; - } else if (arg[0] == '-') { - std::cerr << "Error: unknown option " << arg << "\n"; - exit(1); - } else { - ptxFile = argv[i]; - } - } - - if (!ptxFile) { - std::cerr << "Usage: warpforth-runner kernel.ptx --kernel NAME " - "[--param i64[]:V,...] [--param f64[]:V,...] " - "[--param i64:V] [--param f64:V] [--grid X,Y,Z] " - "[--block X,Y,Z] [--output-param N] [--output-count N]\n"; - return 1; - } - - if (!kernelName) { - std::cerr << "Error: --kernel NAME is required\n"; - return 1; - } - - if (params.empty()) { - std::cerr << "Error: at least one --param is required\n"; - return 1; - } - - if (outputParam < 0 || outputParam >= static_cast(params.size())) { - std::cerr << "Error: output-param " << outputParam << " out of range (have " - << params.size() << " params)\n"; - return 1; - } - - if (isScalar(params[outputParam])) { - std::cerr << "Error: output-param " << outputParam - << " is a scalar (cannot read back)\n"; - return 1; - } - - // Read PTX - std::string ptx = readFile(ptxFile); - - // Initialize CUDA - CHECK_CU(cuInit(0)); - - CUdevice device; - CHECK_CU(cuDeviceGet(&device, 0)); - - CUcontext ctx; - CHECK_CU(cuCtxCreate(&ctx, 0, device)); - - // Load PTX module - CUmodule module; - CHECK_CU(cuModuleLoadData(&module, ptx.c_str())); - - CUfunction func; - CHECK_CU(cuModuleGetFunction(&func, module, kernelName)); - - // Allocate device buffers for array params - for (auto &p : params) { - if (auto *a = std::get_if>(&p)) - allocDevice(*a); - else if (auto *a = std::get_if>(&p)) - allocDevice(*a); - } - - // Set up kernel parameters — Driver API expects array of pointers to args - std::vector kernelArgs(params.size()); - for (size_t i = 0; i < params.size(); ++i) - kernelArgs[i] = kernelArgPtr(params[i]); - - // Launch kernel - CHECK_CU(cuLaunchKernel(func, grid.x, grid.y, grid.z, block.x, block.y, - block.z, 0, nullptr, kernelArgs.data(), nullptr)); - - CHECK_CU(cuCtxSynchronize()); - - // Copy back and print output param - size_t count = outputCount >= 0 ? static_cast(outputCount) : 0; - if (auto *iArr = std::get_if>(¶ms[outputParam])) { - if (outputCount < 0) - count = iArr->values.size(); - printOutput(*iArr, count); - } else { - auto &fArr = std::get>(params[outputParam]); - if (outputCount < 0) - count = fArr.values.size(); - printOutput(fArr, count); - } - - // Cleanup — only free device memory for array params - for (auto &p : params) { - if (auto *a = std::get_if>(&p)) - cuMemFree(a->devicePtr); - else if (auto *a = std::get_if>(&p)) - cuMemFree(a->devicePtr); - } - cuModuleUnload(module); - cuCtxDestroy(ctx); - - return 0; -} diff --git a/uv.lock b/uv.lock index f2bb18e..69df46d 100644 --- a/uv.lock +++ b/uv.lock @@ -469,6 +469,53 @@ wheels = [ { url = "https://files.pythonhosted.org/packages/aa/50/a9caea39ad19c431c1a3f8a31114df65b260cdfe67786b6c7e7c040c4c44/cryptography-49.0.0-pp311-pypy311_pp73-win_amd64.whl", hash = "sha256:be9fcb48a55f023493482827d4f459bd263cc20efde64f204b97c123201850c6", size = 3783731, upload-time = "2026-06-12T20:02:43.319Z" }, ] +[[package]] +name = "cuda-bindings" +version = "12.9.7" +source = { registry = "https://pypi.org/simple" } +dependencies = [ + { name = "cuda-pathfinder" }, +] +wheels = [ + { url = "https://files.pythonhosted.org/packages/40/f3/f9d1095f90d2a4df24cfcafe7487fd9444c6dacb94e3722be6fedd8ac26c/cuda_bindings-12.9.7-cp311-cp311-manylinux_2_24_aarch64.manylinux_2_28_aarch64.whl", hash = "sha256:16043ef5b15ab88fe9954c5c2061b1d8007591b27f2c916331056de0ebc6187e", size = 7114834, upload-time = "2026-05-27T18:44:07.746Z" }, + { url = "https://files.pythonhosted.org/packages/3a/8a/1251e1794b69865aacd5629936006b18ea0816a495de4ecea9a825556eb3/cuda_bindings-12.9.7-cp311-cp311-manylinux_2_24_x86_64.manylinux_2_28_x86_64.whl", hash = "sha256:c6496a88d84b1209d6651b0370c19c26319e157c22f6d018bf9a358cd8049041", size = 7647147, upload-time = "2026-05-27T18:44:09.4Z" }, + { url = "https://files.pythonhosted.org/packages/1e/39/158392f6572e6e0def70ca39029c46b75e02ea4a43c63ff7320b3d180a29/cuda_bindings-12.9.7-cp311-cp311-win_amd64.whl", hash = "sha256:c392ffa5010ef4073bfd9dfff4d1ae56032094ed52d3d732014f8e41a73e6b59", size = 7218081, upload-time = "2026-05-27T18:44:11.104Z" }, + { url = "https://files.pythonhosted.org/packages/32/45/557d4ed1fa54f0c7db8aee083229f624990d69f7d00f55477eed5c7e169a/cuda_bindings-12.9.7-cp312-cp312-manylinux_2_24_aarch64.manylinux_2_28_aarch64.whl", hash = "sha256:0666d3c082ef8f4b2d670950589373550e9f3bf564d635dd883f24a0b40402ff", size = 7071026, upload-time = "2026-05-27T18:44:13.356Z" }, + { url = "https://files.pythonhosted.org/packages/91/97/e3c6e58ece26a053419ba0a18444b5443cfc64451bbf37f84e8143b8bdca/cuda_bindings-12.9.7-cp312-cp312-manylinux_2_24_x86_64.manylinux_2_28_x86_64.whl", hash = "sha256:4c7ef48c5e13ae90f3b2ecfb72f8e99ac43c8f4c43e67e1325b8aae331453687", size = 7611059, upload-time = "2026-05-27T18:44:15.252Z" }, + { url = "https://files.pythonhosted.org/packages/6d/39/afaa3de4d491a55af8961081e0b69c08d51bfbe471c359a7bddb4a28ca41/cuda_bindings-12.9.7-cp312-cp312-win_amd64.whl", hash = "sha256:3c089aaf4f5f570ec50244c68f5a2b00a2c9a8e01e04219fd2e36e340be0d88b", size = 7400841, upload-time = "2026-05-27T18:44:17.164Z" }, + { url = "https://files.pythonhosted.org/packages/eb/7b/f1575e41e1a17dc2f2a408b2e8e864c9324e41e3e23f6401e5efc54c152a/cuda_bindings-12.9.7-cp313-cp313-manylinux_2_24_aarch64.manylinux_2_28_aarch64.whl", hash = "sha256:266379e4942051f544a8e7ea1a30ead8d7e8199b6b30fcdc8917cae2bf614e61", size = 6978549, upload-time = "2026-05-27T18:44:18.839Z" }, + { url = "https://files.pythonhosted.org/packages/9d/dc/62d62eb4f91eb721bcf46da51b13e9872ccd8fa7e60eb8ba7b7baeac72c6/cuda_bindings-12.9.7-cp313-cp313-manylinux_2_24_x86_64.manylinux_2_28_x86_64.whl", hash = "sha256:59cf4a37b0d662ba15037c9ceebe1a306ebf2c01a8235a09be13cd07094fdb74", size = 7457675, upload-time = "2026-05-27T18:44:20.637Z" }, + { url = "https://files.pythonhosted.org/packages/43/b2/753fe88151001d0dc23f56a8e119fe06b991b0d1a885fa02f9852b12f523/cuda_bindings-12.9.7-cp313-cp313-win_amd64.whl", hash = "sha256:5bd89dcb78475a6d8a4620ea94b74edf0cbbeacee6d1622d8f94452c1e8d3f15", size = 7360097, upload-time = "2026-05-27T18:44:22.405Z" }, + { url = "https://files.pythonhosted.org/packages/f9/77/94d9b85f26add6fe9c9cb7c4ec3b96bc598f7ea5cfbd7490cc0a36adf5be/cuda_bindings-12.9.7-cp313-cp313t-manylinux_2_24_aarch64.manylinux_2_28_aarch64.whl", hash = "sha256:2dbcd4801954eb3508f4dc2fa0d0c8eb93eb3f45326fd61be2731418c371e7a0", size = 6870886, upload-time = "2026-05-27T18:44:24.164Z" }, + { url = "https://files.pythonhosted.org/packages/04/dd/3ec34b569e1b990b11276feba306bf8f446656cc38e8ed0f49b5facfeffa/cuda_bindings-12.9.7-cp313-cp313t-manylinux_2_24_x86_64.manylinux_2_28_x86_64.whl", hash = "sha256:3747ea132642416786a8e31bf229032df3a7856911ae5426a7be53d032df183d", size = 7345663, upload-time = "2026-05-27T18:44:26.333Z" }, + { url = "https://files.pythonhosted.org/packages/b8/c8/d79a20ba396e7ab2dfdd4b72b62356972b25b88aee2ded49a70c797ddea1/cuda_bindings-12.9.7-cp313-cp313t-win_amd64.whl", hash = "sha256:64f7ade7a7a3b69001489753acc21706d9dbda32db8deb68a767a0a0aab30b68", size = 7780136, upload-time = "2026-05-27T18:44:28.121Z" }, + { url = "https://files.pythonhosted.org/packages/68/e4/075052d42872cf8162da53f14447a4b8abc004c3750e4b724ee502428da0/cuda_bindings-12.9.7-cp314-cp314-manylinux_2_24_aarch64.manylinux_2_28_aarch64.whl", hash = "sha256:775960ac9e530717f3b48e165cc6f68684fa9a4141764fd923e4c1a9820acc73", size = 7060090, upload-time = "2026-05-27T18:44:30.281Z" }, + { url = "https://files.pythonhosted.org/packages/ec/cd/3289c810a4d45e5364a3387a74b4c9b6f6f57ee96ae0e5b537cc61dec242/cuda_bindings-12.9.7-cp314-cp314-manylinux_2_24_x86_64.manylinux_2_28_x86_64.whl", hash = "sha256:3c47ec1a7a441d91aab32339951df7a1be53451121a12c094bba51467717a35a", size = 7504419, upload-time = "2026-05-27T18:44:31.992Z" }, + { url = "https://files.pythonhosted.org/packages/5b/a0/c429fdcfa5aae181415504c5085ea5944f782b417dd16a7f2a14be0da80d/cuda_bindings-12.9.7-cp314-cp314-win_amd64.whl", hash = "sha256:1e2a4f2ec5b67408c04bb4fbed45d214b66de1f00ee2e972865cacb8708d4e1e", size = 7493876, upload-time = "2026-05-27T18:44:33.618Z" }, + { url = "https://files.pythonhosted.org/packages/11/43/472a6281c3d94e71687e27c657a8f60718d3579b4d94c41deea503165f8a/cuda_bindings-12.9.7-cp314-cp314t-manylinux_2_24_aarch64.manylinux_2_28_aarch64.whl", hash = "sha256:00a833d399b31071fab4cf3de2929840ae462dc4848116eeff033d09219e7116", size = 6899146, upload-time = "2026-05-27T18:44:35.556Z" }, + { url = "https://files.pythonhosted.org/packages/2b/13/10c1d0b32a9da65142d213e0733d748457fb3fd066aee4317335266f15c6/cuda_bindings-12.9.7-cp314-cp314t-manylinux_2_24_x86_64.manylinux_2_28_x86_64.whl", hash = "sha256:11aeafa2b33995f890086b3fb0f062075176d956e9b6a6fe1a699dddc413f6ad", size = 7369087, upload-time = "2026-05-27T18:44:37.359Z" }, + { url = "https://files.pythonhosted.org/packages/33/10/c71a07cd2a1d4db119bada1848b4752a874ccfe4927d419bfdd05f250920/cuda_bindings-12.9.7-cp314-cp314t-win_amd64.whl", hash = "sha256:ece8dfbc22e6de96a26940ab9887eb3cfe1fc1bc3966169391cdb866bb82bb64", size = 8208198, upload-time = "2026-05-27T18:44:39.053Z" }, +] + +[[package]] +name = "cuda-pathfinder" +version = "1.8.1" +source = { registry = "https://pypi.org/simple" } +wheels = [ + { url = "https://files.pythonhosted.org/packages/9d/e6/22df83f82f9bc26cb1c42265cf14d34d4908dba2a0f261bd7b28244acb00/cuda_pathfinder-1.8.1-py3-none-any.whl", hash = "sha256:ae0137ff9e56ea97499bcbf54f5f2778ec25f3266715ac86da192a795af982a8", size = 62552, upload-time = "2026-09-02T16:55:28.64Z" }, +] + +[[package]] +name = "cuda-python" +version = "12.9.7" +source = { registry = "https://pypi.org/simple" } +dependencies = [ + { name = "cuda-bindings" }, +] +wheels = [ + { url = "https://files.pythonhosted.org/packages/c1/9d/05e753afbaac3f92691059b3ba875589c98a425d69e5808cec32b31b580c/cuda_python-12.9.7-py3-none-any.whl", hash = "sha256:23a1fc406d491eef7a7e985095725cb7b20a04a7bd9b7a66400e5c86e082e0aa", size = 7597, upload-time = "2026-05-27T19:50:32.605Z" }, +] + [[package]] name = "curlify" version = "3.0.0" @@ -1895,6 +1942,7 @@ name = "warpforth" version = "0.1.0" source = { virtual = "." } dependencies = [ + { name = "cuda-python" }, { name = "lit" }, { name = "numpy" }, { name = "pytest" }, @@ -1916,6 +1964,7 @@ docs = [ [package.metadata] requires-dist = [ + { name = "cuda-python", specifier = ">=12.8,<13" }, { name = "lit", specifier = ">=18.1.0" }, { name = "numpy" }, { name = "pytest" },