Team Ai
Datasetpublic

williamhtan/kernelsight

KernelSight v4 Per-timestep workload labels for GPU execution traces. KernelSight pairs every GPU workload trace with a dense, per-timestep workload labeling. Each snapshot is a [24, 512] counter image — 24 hardware-counter channels sampled across 512 equal-width time bins — paired with per-bin labels drawn from a two-level hierarchy of 12 coarse (L1) and 73 fine (L2) workload classes. The goal is to label what a kernel is doing at each instant (matmul, attention, reduction… See the full description on the dataset page: https://huggingface.co/datasets/williamhtan/kernelsight.

sourceHugging Faceapache-2.0updated 4mo agoView on Hugging Face
0likes20downloads
Dataset Card

KernelSight v4

Per-timestep workload labels for GPU execution traces.

KernelSight pairs every GPU workload trace with a dense, per-timestep workload labeling. Each snapshot is a [24, 512] counter image — 24 hardware-counter channels sampled across 512 equal-width time bins — paired with per-bin labels drawn from a two-level hierarchy of 12 coarse (L1) and 73 fine (L2) workload classes. The goal is to label what a kernel is doing at each instant (matmul, attention, reduction, memory movement, …) rather than the single coarse bottleneck label a profiler assigns per launch.

All counters are collected on a single NVIDIA H100 80GB HBM3 (Hopper, sm_90a). No data augmentation is applied: each snapshot is the trace as measured, and class imbalance is handled at training time rather than by resampling.

Snapshots1,444
Tensor shape[24, 512] (24 channels × 512 time bins)
Label vocab12 L1 classes · 73 L2 classes
Labeled segments45,860
Overlap ground truth472 snapshots (has_overlap=1)
Standard splittrain 1,124 / val 160 / test 160
HardwareNVIDIA H100 80GB HBM3 (sm_90a)
On-disk size~68 MB (4,332 .npz files)
LicenseApache-2.0

Quick start

The dataset ships as raw NumPy .npz files in a fixed directory layout, indexed by JSON split files. It is not loadable through datasets.load_dataset(); download the folder and read the .npz files directly with NumPy.

python
from huggingface_hub import snapshot_download
import numpy as np, json, os

root = snapshot_download(repo_id="williamhtan/kernelsight", repo_type="dataset")

# Load a split (paths inside are relative to `root`)
split = json.load(open(os.path.join(root, "splits/train.json")))
rec = split["traces"][0]

# Input tensor
t = np.load(os.path.join(root, rec["path"]), allow_pickle=True)
X = t["data"]                      # (24, 512) float32

# Labels live alongside the input
lpath = rec["path"].replace("/input/", "/labels/").replace("tensor_input.npz", "labels.npz")
l = np.load(os.path.join(root, lpath), allow_pickle=True)
y_l1  = l["workload_l1"]           # (512,) int32, -1 where unlabeled
y_l2  = l["workload_l2"]           # (512,) int32
mask  = l["mask_labeled"]          # (512,) uint8
y_mh  = l["workload_l1_multihot"]  # (512, 12) uint8, multi-label overlap track

A reference PyTorch Dataset is bundled at tools/sass_dataloader_stub.py, and the label vocabularies live in tools/workload_taxonomy.py (single source of truth). Full schema documentation is in `data_info.md`.


Directory layout

kernelsight_dataset_v4/
├── README.md                 # this card
├── MANIFEST_v4.md            # release notes / per-class counts
├── data_info.md              # full on-disk schema reference
├── splits/                   # 7 split JSONs (see Splits)
├── tools/
│   ├── workload_taxonomy.py  # 12 L1 + 73 L2 + 8 flags + 5 spatial (source of truth)
│   └── sass_dataloader_stub.py
└── kernels/<motif>/_out/<variant>/
    ├── input/tensor_input.npz       # [24, 512] profiler heatmap + metadata
    ├── labels/labels.npz            # per-bin + per-segment L1/L2 labels, vocabs
    └── fingerprint/fingerprint.npz  # 32-D instruction-mix fingerprint

Each <variant> is one snapshot (e.g. parameter-swept geometry like cutlass_gemm/_out/m8192_n1024_k4096_.../). Split JSON path fields are relative to the dataset root and point at .../input/tensor_input.npz.


Input tensor — tensor_input.npz

  • —data — [24, 512] float32: 24 counter channels × 512 time bins.
  • —time_edges_ns — [513] int64: bin-boundary timestamps (bins are equal-width per trace, ~0.5 ms for fast matmul up to ~30 ms for long scatter; the window is clipped to the kernel-active span).
  • —counter_names — [24]: channel names. kernels, kernel_names, kernel_function_index — per-launch identity metadata.

Channels divide into five semantic groups. Each channel is divided by a fixed physical-scale divisor (pipe/throughput ÷100, warp-stall ÷64, coalescing ÷8) so values land in ~[0, 1] while preserving magnitude differences (no per-channel min/max rescale).

RowsGroupSourceChannels
0–6Pipe signatureCUPTItensor_op_hmma, xu, fma, alu, lsu, cbu, tma
7–8Memory accessCUPTI/ncuhit: l2, atom: lts_atomic_input_pct
9–12Discriminatorsncu/NVBitshort_scoreboard, barrier, pred_on_per_inst_ratio, gmem_coalesce_ratio
13–16System bandwidthNsight SystemsSMs Active %, DRAM Read %, DRAM Write %, Tensor Active %
17–23SASS modalityNVBitcompute_fma, compute_tensor, memory_global, memory_shared, memory_tma, control, misc

Labels — labels.npz

Per-bin arrays (length T = 512):

KeydtypeMeaning
workload_l1int32L1 class id per bin (-1 if unlabeled)
workload_l2int32L2 class id per bin (-1 if unlabeled)
workload_l1_multihotuint8 [T,12]Multi-hot per-bin L1 (overlap track)
workload_l2_multihotuint8 [T,73]Multi-hot per-bin L2
multihot_n_activeuint8 [T]# active L1 classes per bin
multihot_has_overlapuint8 []1 if any bin asserts ≥2 classes
segment_idint32 [T]0-based segment ordinal per bin (-1 if none)
mask_any_kerneluint8 [T]1 if a kernel interval overlaps this bin
mask_labeleduint8 [T]1 if workload_l1 >= 0
time_edges_nsint64 [T+1]Bin boundaries

Per-segment arrays (length S, varies by motif): segment_starts, segment_ends, segment_label_l1, segment_label_l2, segment_kernel_names, segment_predecessor_l1/l2, segment_position, attribute_flags [S,8].

Vocabularies (carried in every labels.npz): vocab_l1[12], vocab_l2[73], attribute_flag_names[8], spatial_state_vocab[5], l2_parent_l1[73].

The single-label fields are always a subset of the multi-hot tracks. On the sequential corpus the multi-hot is effectively one-hot; genuine overlap comes from 472 cutlass_ws_overlap snapshots whose producer (TMA load → memory_movement) and consumer (WGMMA → matmul) phases co-occur, derived from device %globaltimer markers (independent of the 24 counter channels — no label leakage).


Label taxonomy

L1 (12): matmul, conv, activation, normalization, softmax, pooling, reduction, attention, loss, elementwise, memory_movement, other.

L2 (73), nested under L1 parents:

L1L2 classes
matmulbmm, gemm, matvec
convconv1d_standard, conv2d_depthwise, conv2d_pointwise, conv2d_standard, conv3d_standard, convtranspose1d, convtranspose2d, convtranspose3d
activationelu, gelu, hardsigmoid, hardswish, hardtanh, leaky_relu, mish, other, relu, selu, sigmoid, softplus, softsign, swish, tanh
normalizationbatchnorm, frobeniusnorm, groupnorm, instancenorm, l1norm, l2norm, layernorm, rmsnorm
softmaxlog_softmax, logsumexp, softmax
poolingavg_pool, global_avg_pool, max_pool
reductionargmax, argmin, cumprod, cumsum, max, mean, min, prod, sum
attentionscaled_dot_product
losscross_entropy, hinge, huber, kldiv, mse, triplet_margin
elementwiseadd, bias_add, cast, clamp, div, mul, residual_add, scalar_multiplication, scaling, sub
memory_movementcopy, embedding, gather, scatter, transpose
otherdropout, misc

Attribute flags (8, multi-label per segment): sparse, tma, cluster, masked, persistent, vectorized_store, atomic_accum, ldgsts.

Spatial-state vocab (5, exposed for the model side): uniform, wavefront_transition, tail_effect, load_imbalanced, hotspot.


Corpus composition

SourceMotifSnapshotsNotes
Microbenchvector_add20coalesced BW-bound elementwise
Microbenchgather17random-indexed memory movement
Microbenchreduction16tree + atomic reductions
Microbenchscatter31atomic histogram scatter
Microbenchwgmma1tensor-core GEMM baseline
KernelBenchkernelbench480PyTorch L1 + L2 ops (11 populated L1 classes)
CUTLASScutlass_gemm278ex48 TF32 warp-specialized GEMM
CUTLASScutlass_fmha85ex88 FlashAttention-3
CUTLASScutlass_ws_overlap472ex48 + device %globaltimer markers
CUTLASScutlass_fp8_gemm14ex54 FP8 WS-GEMM
CUTLASScutlass_sparse_gemm18ex62 2:4 structured sparsity
CUTLASScutlass_grouped_gemm12ex57 grouped GEMM

Per-L1 distribution (snapshots containing each class / labeled segments): matmul 849 / 3,257 · activation 147 / 11,654 · reduction 125 / 7,155 · conv 98 / 6,418 · attention 92 / 295 · elementwise 86 / 2,887 · normalization 79 / 6,546 · pooling 62 / 2,019 · memory_movement 48 / 48 · loss 42 / 4,860 · softmax 28 /

  1. 1.The corpus is heavily imbalanced (matmul dominates bin count via long CUTLASS traces), which motivates a class-balanced objective and macro-F1.

Splits

Each split JSON is {split, seed, n, traces, notes}; every traces[i] records path, motif, n_kernels, n_unique_kernels, T, l1_labels, l2_labels, dominant_l1, dominant_l2.

Standard disjoint partition (L2-stratified, trace-level leak-free):

Splitn
train.json1,124
val.json160
test.json160

This measures within-kernel generalization: most test traces share kernel identity with training and differ in geometry/precision/sweep parameters.

Overlapping analysis tags (views over the same corpus, not a partition):

TagnSelects
iid.json433random IID sample
param_ood.json956parameter-sweep variants (fixed op, unseen geometry)
composed.json1,124multi-kernel / multi-segment traces
length_ood.json0reserved (empty in v4)

Collection methodology

Workloads come from three sources — hand-written CUDA microbenchmarks isolating canonical GPU behaviors, the KernelBench Level-1 / Level-2 problem suite, and CUTLASS Hopper examples (the single largest contributor, ~61% of the corpus) spanning six warp-specialized datapaths (TF32, FP8, 2:4-sparse, grouped GEMM, FlashAttention-3, and WS-GEMM with device markers).

Each workload is profiled by three complementary collectors and fused onto one time grid:

  1. 1.NVBit — SASS-level dynamic binary instrumentation: per-PC instruction mix and coalescing statistics.
  2. 2.CUPTI Range Profiler — replays each kernel for a 19-metric warp-stall taxonomy (stall reasons, pipe utilizations, occupancy).
  3. 3.Nsight Systems — samples system throughput at ~10 kHz alongside the CUDA/NVTX timeline; the only natively time-resolved source, so it defines the time grid.

Labels come from NVTX markers + kernel boundaries, with kernel-name + SASS pattern matching resolving the L2 class. The full collector fork and per-motif run.sh reproduction harness are part of the KernelSight project (not bundled in this dataset distribution, which ships the rendered tensors, labels, and splits).


Changes from v3.1

  • —Dropped megakernel (1 PoC snapshot) and tiled_gemm_poc (590 hand-written PoC snapshots).
  • —Added three CUTLASS Hopper datapaths: FP8 (ex54), 2:4 sparse (ex62), grouped (ex57).
  • —Selective KernelBench expansion (activation, normalization, pooling, reduction, elementwise) and geometry sweeps over microbenchmarks and CUTLASS GEMM/FMHA.
  • —Corpus 262 → 1,444 snapshots; overlap ground truth 29 → 472 snapshots.
  • —CI: 26,996 assertions passed, 0 failed.

See `MANIFEST_v4.md` for full release notes.


Limitations

  • —Counters are from a single H100 (sm_90a); cross-architecture transfer is out of scope.
  • —Overlap timing is coarse: device-marker spans resolve producer/consumer envelopes (≈ whole launch), so overlap is annotated at launch granularity.
  • —All 24 channels carry real signal, but many rows are legitimately zero where the hardware is inactive for a given motif.
  • —The spatial_state vocab is exposed for the model side; per-bin spatial-state derivation is not provided.

License & provenance

Released under Apache-2.0. Derived workloads retain their upstream licenses:

  • —KernelBench problems — MIT (Scaling Intelligence Lab, Stanford University).
  • —CUTLASS examples — BSD-3-Clause (NVIDIA Corporation).

The profiler tooling builds on the Intra-Kernel Profiler (NVBit / CUPTI / Nsight Systems). This release contains only derived, aggregated counter tensors and labels — no third-party source code.

Citation

bibtex
@misc{tan2026kernelsight,
  title  = {KernelSight: Per-Timestep Workload Labeling of GPU Execution Traces},
  author = {Tan, William},
  year   = {2026},
  note   = {CS231N project, Stanford University},
  howpublished = {\url{https://huggingface.co/datasets/williamhtan/kernelsight}}
}