Skip to content

Keep ptxas resource counts scoped to the first kernel report - #12

Merged
jinhongyii merged 1 commit into
mlc-ai:mainfrom
adenzhou1350:codex/ptxas-first-kernel-resources-20260930
Sep 30, 2026
Merged

jinhongyii merged 1 commit into
mlc-ai:mainfrom
adenzhou1350:codex/ptxas-first-kernel-resources-20260930

Conversation

@adenzhou1350

Copy link
Copy Markdown
Contributor

When a cubin contains multiple kernels, parse_ptxas can combine the first kernel's register count with static shared memory or spills from a later kernel. For example, a real CUDA 13.0 SM103a compile reports 10 registers and no static SMEM for plain, then 128 bytes SMEM for shared; parsing the combined log incorrectly assigns 128 bytes to plain.

Scope matching to the first entry-function report, as the API already documents. Logs without compile headers use the first function-properties block; abbreviated single-report logs remain supported. Preserve absent fields rather than filling them from another function.

Validation:

  • Focused tests with pytest -q -n 2 -p no:cacheprovider test_dump_kernel.py against isolated source-package overlays: base 2 failed, 3 passed, patched 5 passed. Both new tests fail without this change.
  • Real nvcc -cubin -arch=sm_103a -Xptxas=-v and cuobjdump --dump-sass on a B300 environment, CUDA 13.0.88. Compiled the plain/shared pair in both declaration orders: the original parser mixes fields when the first report lacks SMEM; the patched parser equals the isolated first report in both orders.
  • This is a resource-report correctness fix. It does not change generated kernels or claim a kernel speedup. The full NumSim/GPU suite was not run for this text-parser change.

The native reproducer uses these kernels (put shared before plain to obtain the plain report first with this toolkit):

extern "C" __global__ void shared(float* y, const float* x) {
    __shared__ float s[32];
    s[threadIdx.x] = x[threadIdx.x];
    __syncthreads();
    y[threadIdx.x] = s[(threadIdx.x + 1) % 32] * 3.0f;
}
extern "C" __global__ void plain(float* y, const float* x) {
    y[threadIdx.x] = x[threadIdx.x] * 2.0f;
}

@jinhongyii jinhongyii left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM

@jinhongyii
jinhongyii merged commit 8ab46f4 into mlc-ai:main Sep 30, 2026
3 checks passed
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.

2 participants