Skip to content

Add an optional workflow that runs the GPU tests on a simulated H100 - #16

Open
saqibkh wants to merge 1 commit into
KernelTuner:mainfrom
saqibkh:add-gpu-sim-ci
Open

saqibkh wants to merge 1 commit into
KernelTuner:mainfrom
saqibkh:add-gpu-sim-ci

Conversation

@saqibkh

@saqibkh saqibkh commented Nov 9, 2026

Copy link
Copy Markdown

The CI here builds the Catch2 tests but runs only ~[GPU], because the hosted runners have no GPU, so the GPU-tagged tests never run. This adds a separate, optional workflow .github/workflows/gpu-sim.yml that builds the tests and runs them on a simulated NVIDIA H100, using PantheonSim (via the pantheongpu/setup-pantheonsim action, pinned by SHA). I am the author of PantheonSim, so please weigh this accordingly.

It does not touch the existing workflows and is not meant to be a required check. A simulator cannot reproduce hardware-specific behavior or timing, so it complements real-GPU testing and does not replace it.

What it does:

  • Installs CUDA 13.0 and the simulated H100, configures with CMAKE_CUDA_ARCHITECTURES=90-virtual and CMAKE_CUDA_RUNTIME_LIBRARY=Shared (the simulator runs PTX only and needs a shared CUDA runtime), and builds with -j3.
  • Runs kernel_float_tests through vgpu run, including the GPU-tagged tests, then runs the pi example.
  • Triggers: workflow_dispatch and pull_request to main, limited to include/, tests/, examples/, CMakeLists.txt and the workflow itself. A concurrency group cancels superseded runs, with permissions: contents: read.

What is left out, with the reason in comments in the file:

  • 4 of the 181 test cases are skipped: binary operators and constant eq tests for __nv_bfloat16 (the simulator compares .bf16 values by raw bits in setp, a known simulator bug), ternary operators for __half (the reference x * y + z compiles to mul.f16 and add.f16 without a rounding modifier; on an RTX 3060 they are fused into an fma, the simulator rounds twice, so the last bit can differ) and ternary operators for __nv_bfloat16 (this one also fails on an RTX 3060, sm_86, so it is not a simulator difference).
  • The vector_add and vector_add_tiling examples are not run: they index empty std::vectors (output_expected[i] and output_result are never sized) and segfault on an RTX 3060 as well. I did not change them here.
  • Only one architecture (H100), one CUDA version (13.0) and a Release build are covered.

Side observation, not touched by this PR: on an RTX 3060 (sm_86, CUDA 13.0, Release) the cache modifier tests for all five types fail with error 717 (illegal memory address) when run on their own, while they pass on the simulator.

Result on a fork: the run passed. 177 of 181 test cases ran and all passed (All tests passed (5264 assertions in 177 test cases)), and the pi example ran. The job took about 13.7 minutes on a cold run, of which about 12.3 minutes was the action setting up the simulator on its first run. The build took about 1 minute and the tests and the example a few seconds.

Run: https://github.com/saqibkh/kernel_float/actions/runs/42659400821

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant