Compute numerical error between *output* and *reference*. Returns ``(correctness, exceeds)`` where *correctness* is a :class:`Correctness` ca

Write solution.json and execute the build_ext template in *cwd* with ext.load mocked. Returns the ``ext.load`` mock so callers can inspect how it

Copy source data into the next pool offset and return views. Returns inputs followed by zero-filled outputs (for DPS kernels). Each c

Benchmark GPU time using CUPTI activity tracing for precise kernel timing. CUPTI (CUDA Profiling Tools Interface) provides hardware-level pr

Benchmark the runtime of the provided function. Derived from triton.testing.do_bench (MIT licence), with fixes from sol-bench: explicit synch

Generate a tensor using heuristics based on the input name. Returns None if no heuristic matches, falling back to _rand_tensor.

Time the execution of a callable using CUDA events. Creates a :class:`ShiftingMemoryPoolAllocator` from *inputs* and *outputs* so each timed

Verify current GPU/DRAM clocks match expected frequencies via nvidia-smi. Queries all GPUs and checks that current frequencies are within *tolera

Write one Trace as strictly valid JSON to the real stdout. Uses allow_nan=False so that any unexpected NaN/Inf in the trace dict raises Value

Write all staging files and run eval_driver.py in a subprocess. Args: tmp_path: Temporary directory to use as the staging directory.

Stage compilation files and return (command, artifact_path). Writes build_ext.py, solution.json, and C++/CUDA source files to output_

Lock GPU and DRAM clocks for the given device. Looks up the device in the preset table, then calls ``sudo nvidia-smi -lgc`` and ``sudo nvidia

Resolve a relative path against blob roots, handling partial overlap. Tries ``root / rel`` first, then progressively strips leading components

Stage execution files and return the command to run. Writes eval_driver.py, definition.json, workload.jsonl, solution.json to output_

Get concrete axis values from variable axis values. Resolves all expressions. Parameters ---------- var_axes_values : dict[st

Set random seeds for reproducibility across Python, PyTorch CPU and CUDA.

Summarize timing statistics based on return mode.

Parse JSONL stdout from eval_driver.py into Trace objects. Each line starting with '{' is parsed as a Trace JSON object. Non-JSON lin

Get the clock preset for a given GPU device name. Returns None if the device is not in the preset table. Parameters ---------- devic

Capture ``id()`` of named functions from a namespace. Call this **before** user code is imported. Pass the returned dict to :func:`check_eva

Return a combined patch context that mocks verify_clocks → True and time.sleep → noop.

Verify that critical eval-driver functions have not been replaced. Compares the current ``id()`` of each snapshotted name against the value c

Detect lazy/proxy tensors in the user output. Uses strict ``type()`` equality — not ``isinstance`` — so any subclass (including ``FakeTensor`

Generate input tensors in definition order. Returns a list of input values (tensors or scalars) in the same order as definition.inputs.

Load safetensors inputs for a workload. Safetensors blobs are resolved from the staging directory (passed via ``blob_roots``). The first roo

Clear the cache buffer by zeroing it.

Mirrors the _make_eval wrapper from eval_driver.py.

Run a reward-hack check; emit REWARD_HACK trace and return True if detected.

Check whether clocks were locked at startup. Reads the ``SOL_EXECBENCH_CLOCKS_LOCKED`` environment variable set by the Docker entrypoint or s

Return the eval_driver.py template source.

Detect if user code spawned background threads. Capture ``threading.active_count()`` before and after the user call and pass both values here

Quantize a tensor to FP4 E2M1 and pack into uint8 (2 FP4 values per byte). Args: x: Input tensor of shape (..., cols) with values in rang

Return the SM version of the current GPU (e.g. 90, 100), or 0 if unavailable.

Run a command in the given directory and return the result.

Clone tensor arguments to prevent cross-iteration data contamination. Returns fresh copies of all tensor arguments so each benchmark iteration

Probe whether GPU clock locking is available via ``sudo nvidia-smi``. Runs ``sudo -n nvidia-smi -lgc 1`` and immediately resets the lock. Return

Reset GPU and DRAM clocks. Best-effort — errors are logged but not raised.

Call fn in the correct calling convention and return a list of output tensors. DPS=True (destination-passing style): fn(*inputs, *pre_all

Create a 256 MB buffer for clearing L2 cache before benchmark runs.

Get concrete tensor shapes given variable axis values. Parameters ---------- tensors : Iterable[TensorSpec] List

Infer destination-passing style by checking the ``run()`` signature. If the last parameter of ``run()`` matches the last output name in the d

Create a mock CUDA Event class for testing do_bench without GPU.

Run a command in the given directory and return the result.

Allocate output tensors based on definition and resolved axis values. Parameters ---------- definition : Definition The kernel de

Detect if torch.cuda.Event.elapsed_time has been patched. Compares the current function identity against the address captured at module load

Helper function to check common tensor properties

Get concrete input shapes given variable axis values. Parameters ---------- var_values : Optional[dict[str, int]], default=No

Inspect traces for correctness. Returns a summary dict with pass/fail counts and per-workload latencies.

Check if dtype is an integer or boolean type.

Check if the benchmark execution was successful. Returns ------- bool True if this is a regular trace with succes

Assemble a definition.json dict from a dataset row.

Compute a deterministic hash of the solution content.

Get dtype string to Python type mapping (cached).

Lazily build dtype string to torch dtype mapping.

Get frozenset of integer and boolean dtypes (cached).

Detect the local GPU's SM version via nvidia-smi.

Auto-inject -gencode flags when no explicit arch flag is set. Blackwell targets get sm_100a (required for tcgen05/TMEM instructions).

Load evil solution with kernel content inlined from the sample directory.

Load definition, solution, and workloads from an example directory.

Load definition dict + workload dicts from staging directory.

Load definition, solution, and workloads from a self-contained sample directory.

Create a minimal valid solution JSON string with given compile_options.

Print a rich table summarizing evaluation traces.