How to submit

Download a challenge package, replace submission.py, and check it on your own GPUs with the package's self-test:

python harness.py submission.py

Then send it to the judge. KERNELS_TOKEN is your own token from the organizers, which puts your name on your submissions. --run is the name they are grouped under, and a run belongs to whoever used its name first:

export KERNELS_TOKEN=...
python submit.py --run my-team

submit.py is in every package. It uploads the code, waits for the judge, and prints every shape's speedup, the score and the rank. Its first line links to the submission's page, which shows the same result and refreshes while the evaluation runs. Stopping submit.py does not stop the evaluation:

python submit.py --status 42

A submission with several files is a directory with submission.py at its root, and a build.sh there runs first:

python submit.py path/to/dir --run my-team

The same content is judged once per challenge, so submitting it again reports the earlier result.

What the judge does

An evaluation runs on the judge's own node, in two processes.

The first process imports the submission, runs prepare, and captures kernel_fn to a CUDA graph for every shape. The inputs it sees are made up. It records the graphs, the cubins and the contents of memory, and then it ends.

The second process rebuilds the graphs and replays them on hidden inputs. None of the submission's host code runs in it. It alternates replays of the baseline and of the submission, 30 times per shape, on cold caches. Each replay gets new hidden inputs, and the outputs of every replay are checked against the reference.

Rules beyond TASK.md

• prepare runs once per GPU, and its state serves every shape.
• kernel_fn must capture to a graph of kernels, memsets and device-to-device copies. Host callbacks, conditional nodes, child graphs and copies from host memory are refused.
• A submission can't capture graphs of its own.
• Values in __device__ and __constant__ variables don't reach the replay, so keep such data in allocated memory.
• Multicast memory must come from driver.symmetric(comm.workspace.ranks, size), which a replay rebuilds.

The score

A shape's speedup is the median of 30 paired ratios of baseline time to submission time. A challenge's score is the geometric mean of its shapes' speedups. A submission that fails any test or shape isn't ranked. A shape that runs faster than its floor is flagged and still counts.

Each evaluation's page shows the node, the GPUs, the driver, every shape's times and confidence interval, and the failure reason if there is one.