Skip to content

[DUMP] sm_86: kernel 01 (RTX 3090, Ampere) #95

Description

@swesty

Target SM version

sm_86 (RTX 3090)

GPU model

NVIDIA GeForce RTX 3090

CUDA toolkit version

12.8 (nvcc V12.8.93)

NVIDIA driver version

580.159.03

Which kernel chapter

corpus/basics/01_vector_add

Additional context

Hi — I have an RTX 3090 (Ampere, sm_86) available and noticed sm_86 isn't yet covered anywhere in the corpus (only sm_89 and sm_120 are present). I'd like to contribute sm_86 dumps, starting with kernel 01.

Environment is ready and the pipeline is verified — I already produced a clean sm_86 dump of kernel 01 locally to confirm it builds:

  • Pop!_OS, driver 580.159.03
  • CUDA toolkit 12.8 (nvcc V12.8.93), gcc 11.4
  • cuobjdump + Nsight Compute (ncu 2025.1.1) available
  • source commit: 0790169

A few things I'd like you to confirm before I do the work, to keep dumps comparable:

  1. Compile flags. I used nvcc -arch=sm_86 <kernel>.cu (plus -ptx / -cubin). Do you want any specific flags (e.g. -O, -lineinfo, --generate-line-info, a fixed -arch/-code pairing) to match how the sm_120/sm_89 dumps were made?
  2. Toolkit version. Mine is 12.8; your reference env shows 13.2. sm_86 is fully supported on 12.8 — is that acceptable, or do you want a specific toolkit?
  3. Layout. CONTRIBUTING.md documents dumps/sm_XX.sass + sm_XX.meta.md, but the repo currently uses an sm_89/ subfolder with NN.ptx, NN.sass, env.txt. Which layout should I follow for sm_86?
  4. NCU. Reproducibility requirements list an NCU report, but the existing sm_89 dumps don't include one. Do you want .ncu-rep for sm_86 too? (I can run it on the physical 3090.)

Scope I can cover: the 12 kernels that already have sm_89 dumps — basics 01–08, warp_collectives 09–10, math_and_spills 11–12 — one dump per PR per your workflow. The tensor_cores chapters (13–25) target Blackwell features (fp8 QMMA, fp4, sparse MMA, stmatrix) that Ampere can't compile, so those are out of scope for a 3090.

I'll wait for your go-ahead on flags/layout before opening any PR.

Metadata

Metadata

Assignees

No one assigned

    Labels

    No labels
    No labels

    Projects

    No projects

    Milestone

    No milestone

    Relationships

    None yet

    Development

    No branches or pull requests

    Issue actions