[GH-ISSUE #6771] [Bug] PTX numeric percent register names abort kernel metadata decoding #4126

New issue

Closed

opened 2026-10-31 00:47:00 +00:00 by bsakaguchi · 0 comments

bsakaguchi commented

2026-10-31 00:47:00 +00:00

Owner

Originally created by @coygeek on GitHub (Oct 23, 2026).
Original GitHub issue: https://github.com/morluto/rea/issues/6771

Before submitting

I searched the open and closed issues for a related report.
I removed credentials, tokens, private paths, and sensitive target data.

Area

Native analysis or decompilation

Observed behavior

inspect-gpu-code-objects recognizes a standalone PTX file containing .reg .b32 %0;, but marks the object malformed with PTX register declaration has no type or name and returns no kernels. An otherwise identical control using %r0 returns the kernel and one virtual .b32 register.

Reproduction

Executed once for each fixture through the built public CLI. From a disposable directory, with REA_CHECKOUT pointing to the source checkout below:

cat > numeric.ptx <<'PTX'
.version 7.0
.target sm_80
.address_size 64
.visible .entry kernel() {
.reg .b32 %0;
mov.u32 %0, 0;
ret;
}
PTX
node "$REA_CHECKOUT/scripts/rea.mjs" inspect-gpu-code-objects numeric.ptx --analysis kernels --format json

The executed control used the same text with both %0 occurrences changed to %r0, then ran the same command against that control file. Both commands exited 0 with empty stderr.

Expected behavior

Retain kernel and its one scalar virtual register. NVIDIA's PTX identifier grammar permits a leading percent character followed by one or more letters, digits, underscores, or dollar signs; %0 satisfies that grammar. REA's GPU documentation says scalar register declarations count each bare or percent-prefixed name.

Impact and scope

A valid 107-byte PTX fixture loses its entire kernel declaration because of a register identifier. One failure and one passing control were observed. Embedded PTX and CUDA assembler/runtime behavior were not tested. Renaming registers in a copy changes the artifact bytes and digest; no option preserving the original input was demonstrated.

REA version

Source revision 99d2d7ad22bcbd780919760008aba6a88e7d1077, package version 6.7.0; latest published release checked as rea-agents-6.7.0.

Runtime and platform

Node.js v24.22.1; Linux x86_64. Static REA reader; no GPU, CUDA assembler, or native decompiler was invoked.

Target and artifact context

Generated standalone UTF-8 PTX 7.0, target sm_80, address size 64. Failing input SHA-256: 549398d9de9408f88d8d5f8c2be2c023b0cf65dbac057b1888e6e1ef6c4d95ee.

Interface

CLI

Earliest failure stage

Native analysis or reconstruction

Relevant output or evidence

Selected fields from the actual failing CLI normalized_result:

{
  "format": "ptx",
  "objects": [{
    "kind": "ptx",
    "target": "sm_80",
    "inspection": "malformed",
    "limitations": ["PTX register declaration has no type or name"]
  }],
  "kernels": [],
  "coverage": { "discovery": "complete", "payloads": "partial" }
}

The control returned inspection: "complete", coverage.payloads: "complete", and kernel kernel with registers.virtual: [{"type":".b32","count":1}].

Observed failure and control evidence came from actual public CLI runs. Fixture validity is grounded in the published identifier grammar, not an assembler run. Source-derived diagnosis follows below.

Typed outcome or recovery expected

Return the supported kernel declaration and one virtual scalar register without misclassifying the input as malformed.

Analysis or suspected cause

Source-derived: the identifier expression and register parser allow an optional % only when followed by a letter, underscore, or dollar sign. %0 therefore never becomes a name; the parser consumes it as part of the type and rejects at the semicolon.

Additional context

Related closed #6592 concerned the entry name %kernel. This fixture uses the ordinary entry name kernel and exercises a numeric register identifier in the body; the current identifier guard still excludes this valid spelling. Searches for PTX/register/identifier reports found no report of this remaining trigger. No first affected version or real-producer prevalence was established.

Triage assessment

Impact level: P2
Assessment status: supported
Impact basis: A bounded valid PTX input loses all kernel facts with a false malformed diagnosis; changing only the register spelling restores decoding in the observed control.
Workaround status: unverified
Workaround basis: The observed renamed copy changes the selected artifact and digest. No workflow preserving the original input was demonstrated.

Originally created by @coygeek on GitHub (Oct 23, 2026). Original GitHub issue: https://github.com/morluto/rea/issues/6771 ## Before submitting - [x] I searched the open and closed issues for a related report. - [x] I removed credentials, tokens, private paths, and sensitive target data. ## Area Native analysis or decompilation ## Observed behavior `inspect-gpu-code-objects` recognizes a standalone PTX file containing `.reg .b32 %0;`, but marks the object `malformed` with `PTX register declaration has no type or name` and returns no kernels. An otherwise identical control using `%r0` returns the kernel and one virtual `.b32` register. ## Reproduction Executed once for each fixture through the built public CLI. From a disposable directory, with `REA_CHECKOUT` pointing to the source checkout below: ```sh cat > numeric.ptx <<'PTX' .version 7.0 .target sm_80 .address_size 64 .visible .entry kernel() { .reg .b32 %0; mov.u32 %0, 0; ret; } PTX node "$REA_CHECKOUT/scripts/rea.mjs" inspect-gpu-code-objects numeric.ptx --analysis kernels --format json ``` The executed control used the same text with both `%0` occurrences changed to `%r0`, then ran the same command against that control file. Both commands exited 0 with empty stderr. ## Expected behavior Retain `kernel` and its one scalar virtual register. NVIDIA's [PTX identifier grammar](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#identifiers) permits a leading percent character followed by one or more letters, digits, underscores, or dollar signs; `%0` satisfies that grammar. REA's [GPU documentation](https://github.com/morluto/rea/blob/99d2d7ad22bcbd780919760008aba6a88e7d1077/docs/gpu-code-objects.md#L109) says scalar register declarations count each bare or percent-prefixed name. ## Impact and scope A valid 107-byte PTX fixture loses its entire kernel declaration because of a register identifier. One failure and one passing control were observed. Embedded PTX and CUDA assembler/runtime behavior were not tested. Renaming registers in a copy changes the artifact bytes and digest; no option preserving the original input was demonstrated. ## REA version Source revision `99d2d7ad22bcbd780919760008aba6a88e7d1077`, package version `6.7.0`; latest published release checked as `rea-agents-6.7.0`. ## Runtime and platform Node.js v24.22.1; Linux x86_64. Static REA reader; no GPU, CUDA assembler, or native decompiler was invoked. ## Target and artifact context Generated standalone UTF-8 PTX 7.0, target `sm_80`, address size 64. Failing input SHA-256: `549398d9de9408f88d8d5f8c2be2c023b0cf65dbac057b1888e6e1ef6c4d95ee`. ## Interface CLI ## Earliest failure stage Native analysis or reconstruction ## Relevant output or evidence Selected fields from the actual failing CLI `normalized_result`: ```json { "format": "ptx", "objects": [{ "kind": "ptx", "target": "sm_80", "inspection": "malformed", "limitations": ["PTX register declaration has no type or name"] }], "kernels": [], "coverage": { "discovery": "complete", "payloads": "partial" } } ``` The control returned `inspection: "complete"`, `coverage.payloads: "complete"`, and kernel `kernel` with `registers.virtual: [{"type":".b32","count":1}]`. Observed failure and control evidence came from actual public CLI runs. Fixture validity is grounded in the published identifier grammar, not an assembler run. Source-derived diagnosis follows below. ## Typed outcome or recovery expected Return the supported kernel declaration and one virtual scalar register without misclassifying the input as malformed. ## Analysis or suspected cause Source-derived: the [identifier expression and register parser](https://github.com/morluto/rea/blob/99d2d7ad22bcbd780919760008aba6a88e7d1077/src/gpu/PtxKernels.ts#L90) allow an optional `%` only when followed by a letter, underscore, or dollar sign. `%0` therefore never becomes a name; the parser consumes it as part of the type and rejects at the semicolon. ## Additional context Related closed [#6592](https://github.com/morluto/rea/issues/6592) concerned the entry name `%kernel`. This fixture uses the ordinary entry name `kernel` and exercises a numeric register identifier in the body; the current identifier guard still excludes this valid spelling. Searches for PTX/register/identifier reports found no report of this remaining trigger. No first affected version or real-producer prevalence was established. ## Triage assessment - Impact level: P2 - Assessment status: supported - Impact basis: A bounded valid PTX input loses all kernel facts with a false malformed diagnosis; changing only the register spelling restores decoding in the observed control. - Workaround status: unverified - Workaround basis: The observed renamed copy changes the selected artifact and digest. No workflow preserving the original input was demonstrated.

bsakaguchi closed this issue

2026-10-31 00:47:00 +00:00

bsakaguchi referenced this issue

2026-10-31 01:11:39 +00:00

[PR #4126] [CLOSED] fix(artifacts): interpret ZIP current-directory prefixes #6074