morluto/rea

Fork 0

mirror of https://github.com/morluto/rea.git synced 2026-11-05 06:45:22 +00:00

[GH-ISSUE #6594] [Bug] PTX 1.3 body parameters disappear while kernel coverage is complete #3981

New issue

Closed

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

bsakaguchi commented

2026-10-31 00:45:45 +00:00

Owner

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

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

The static GPU CLI returns complete inspection for PTX 1.3 whose kernel body declares .param .u32 input, but reports parameter_count: 0, parameter_bytes: 0, and parameters: []. The input contains one explicit four-byte kernel parameter. A PTX 1.4 control that moves the same declaration into the entry's parameter list returns the expected one parameter and four bytes.

Reproduction

Executed once per input. REA_CHECKOUT denotes a built source checkout at the revision below; run these commands from a disposable directory with Node.js v24.22.1 on PATH.

cat > legacy-body-params.ptx <<'PTX'
.version 1.3
.target sm_10
.entry legacy {
.param .u32 input;
.reg .u32 r;
ld.param.u32 r, [input];
exit;
}
PTX
node "$REA_CHECKOUT/scripts/rea.mjs" inspect-gpu-code-objects legacy-body-params.ptx --analysis kernels --detail complete --json

The executed control was:

.version 1.4
.target sm_10
.entry modern(.param .u32 input) {
.reg .u32 r;
ld.param.u32 r, [input];
exit;
}

It was saved as modern-params-control.ptx and passed to the same CLI command with that filename.

Expected behavior

Read the body-declared kernel parameter for PTX 1.0–1.3, or explicitly mark its parameter facts unknown and the inspection partial/unsupported. Returning an authoritative zero count with complete coverage discards a declaration in the selected artifact. NVIDIA's .entry specification documents that versions 1.0–1.3 declare kernel parameters in the body, while 1.4 and later put them in the entry parameter list. REA documents PTX parameter declarations without a minimum PTX version.

Impact and scope

Parameter signature evidence is silently wrong for the demonstrated legacy PTX input, while the object and payload coverage claim completeness and carry no limitation about ignored body parameters. One legacy input and one modern control were observed. The scope established here is static inspection of a synthetic PTX 1.3 artifact; no CUDA compiler, driver, execution, or embedded artifact was tested. This matters when inspecting historical PTX, even on a host without a legacy GPU runtime.

REA version

6.7.0; source revision 642f6571ea9d333ea50114c25e648dbb5b8a3c46. The source launcher's --version printed 6.7.0.

Runtime and platform

Node.js v24.22.1; Omarchy 4.0.6, Linux 7.2.8-9-omarchy, x86_64, glibc 2.44. Static reader only; no CUDA assembler, GPU runtime, Hopper, or Ghidra was invoked.

Target and artifact context

Synthetic standalone PTX 1.3, target sm_10, one body-declared .u32 parameter; 108 bytes; SHA-256 a39d54e1b86b75472a554db23cd96d838ed3f7813f74fb80bcc0258fa47df90c.

Interface

CLI

Earliest failure stage

Native analysis or reconstruction

Relevant output or evidence

Fields selected from the actual CLI normalized_result:

{
  "format": "ptx",
  "coverage": { "discovery": "complete", "payloads": "complete" },
  "objects": [{ "kind": "ptx", "target": "sm_10", "inspection": "complete", "limitations": [] }],
  "kernels": [{
    "name": "legacy",
    "parameter_count": 0,
    "parameter_bytes": 0,
    "parameters": [],
    "registers": { "virtual": [{ "type": ".u32", "count": 1 }] }
  }]
}

Both commands exited 0 with empty stderr. The 1.4 control returned parameter_count: 1, parameter_bytes: 4 and parameters: [{"ordinal":0,"name":"input","offset":0,"size":4,"type":".u32","hidden":null}], also with complete coverage. The legacy report retains the full PTX text, so the ignored declaration remains available as source evidence despite the incorrect parameter facts.

Expected source: NVIDIA's version-specific .entry parameter rule above; the expected size follows the explicit .u32 declaration.
Failure source: observed public CLI result; source-derived diagnosis below.
Evidence provenance: observed.
Local verification: reproduced.
Reproduction completeness: complete; fixture validity is grounded in the published grammar, not an assembler run.

Typed outcome or recovery expected

Return the one parameter, or retain null/unknown parameter facts with an explicit limitation and partial coverage when legacy parameter decoding is unavailable.

Analysis or suspected cause

Source-derived: parameter collection only runs for a parenthesized entry parameter list. The body loop skips every declaration other than .reg and .shared, including the legacy kernel .param. The emitted kernel consequently publishes the initial empty list and zero byte total. Header recognition admits this PTX version, and the inventory consumer preserves the parser's complete status.

Additional context

Open issue bodies and all-state searches for PTX, PTX 1.3, and body/parameter combinations found no matching report. Related closed #6493 concerns rejecting an entry with an omitted optional parameter list; this report concerns silently dropping a parameter from the body after that spelling has been accepted. A fresh pre-submission search also found no duplicate.

NOT_ENOUGH_INFO: first affected release and prevalence of legacy PTX among users were not investigated.

Triage assessment

Impact level: P2
Assessment status: supported
Impact basis: An explicitly declared kernel parameter disappears from normalized facts while inspection is reported complete for a valid bounded legacy PTX artifact.
Workaround status: available
Workaround basis: The emitted complete Evidence retains the original PTX text, where the parameter can be read manually. The executed 1.4 translation control also yields the correct declaration, but changes artifact bytes and version and is only suitable for a disposable analysis copy.

Originally created by @coygeek on GitHub (Oct 23, 2026). Original GitHub issue: https://github.com/morluto/rea/issues/6594 ### 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 The static GPU CLI returns complete inspection for PTX 1.3 whose kernel body declares `.param .u32 input`, but reports `parameter_count: 0`, `parameter_bytes: 0`, and `parameters: []`. The input contains one explicit four-byte kernel parameter. A PTX 1.4 control that moves the same declaration into the entry's parameter list returns the expected one parameter and four bytes. ### Reproduction Executed once per input. `REA_CHECKOUT` denotes a built source checkout at the revision below; run these commands from a disposable directory with Node.js v24.22.1 on `PATH`. ```sh cat > legacy-body-params.ptx <<'PTX' .version 1.3 .target sm_10 .entry legacy { .param .u32 input; .reg .u32 r; ld.param.u32 r, [input]; exit; } PTX node "$REA_CHECKOUT/scripts/rea.mjs" inspect-gpu-code-objects legacy-body-params.ptx --analysis kernels --detail complete --json ``` The executed control was: ```ptx .version 1.4 .target sm_10 .entry modern(.param .u32 input) { .reg .u32 r; ld.param.u32 r, [input]; exit; } ``` It was saved as `modern-params-control.ptx` and passed to the same CLI command with that filename. ### Expected behavior Read the body-declared kernel parameter for PTX 1.0–1.3, or explicitly mark its parameter facts unknown and the inspection partial/unsupported. Returning an authoritative zero count with complete coverage discards a declaration in the selected artifact. NVIDIA's [`.entry` specification](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#kernel-and-function-directives-entry) documents that versions 1.0–1.3 declare kernel parameters in the body, while 1.4 and later put them in the entry parameter list. REA documents [PTX parameter declarations](https://github.com/morluto/rea/blob/642f6571ea9d333ea50114c25e648dbb5b8a3c46/docs/gpu-code-objects.md) without a minimum PTX version. ### Impact and scope Parameter signature evidence is silently wrong for the demonstrated legacy PTX input, while the object and payload coverage claim completeness and carry no limitation about ignored body parameters. One legacy input and one modern control were observed. The scope established here is static inspection of a synthetic PTX 1.3 artifact; no CUDA compiler, driver, execution, or embedded artifact was tested. This matters when inspecting historical PTX, even on a host without a legacy GPU runtime. ### REA version 6.7.0; source revision `642f6571ea9d333ea50114c25e648dbb5b8a3c46`. The source launcher's `--version` printed `6.7.0`. ### Runtime and platform Node.js v24.22.1; Omarchy 4.0.6, Linux 7.2.8-9-omarchy, x86_64, glibc 2.44. Static reader only; no CUDA assembler, GPU runtime, Hopper, or Ghidra was invoked. ### Target and artifact context Synthetic standalone PTX 1.3, target `sm_10`, one body-declared `.u32` parameter; 108 bytes; SHA-256 `a39d54e1b86b75472a554db23cd96d838ed3f7813f74fb80bcc0258fa47df90c`. ### Interface CLI ### Earliest failure stage Native analysis or reconstruction ### Relevant output or evidence Fields selected from the actual CLI `normalized_result`: ```json { "format": "ptx", "coverage": { "discovery": "complete", "payloads": "complete" }, "objects": [{ "kind": "ptx", "target": "sm_10", "inspection": "complete", "limitations": [] }], "kernels": [{ "name": "legacy", "parameter_count": 0, "parameter_bytes": 0, "parameters": [], "registers": { "virtual": [{ "type": ".u32", "count": 1 }] } }] } ``` Both commands exited 0 with empty stderr. The 1.4 control returned `parameter_count: 1`, `parameter_bytes: 4` and `parameters: [{"ordinal":0,"name":"input","offset":0,"size":4,"type":".u32","hidden":null}]`, also with complete coverage. The legacy report retains the full PTX text, so the ignored declaration remains available as source evidence despite the incorrect parameter facts. - Expected source: NVIDIA's version-specific `.entry` parameter rule above; the expected size follows the explicit `.u32` declaration. - Failure source: observed public CLI result; source-derived diagnosis below. - Evidence provenance: observed. - Local verification: reproduced. - Reproduction completeness: complete; fixture validity is grounded in the published grammar, not an assembler run. ### Typed outcome or recovery expected Return the one parameter, or retain null/unknown parameter facts with an explicit limitation and partial coverage when legacy parameter decoding is unavailable. ### Analysis or suspected cause Source-derived: [parameter collection](https://github.com/morluto/rea/blob/642f6571ea9d333ea50114c25e648dbb5b8a3c46/src/gpu/PtxKernels.ts#L216-L259) only runs for a parenthesized entry parameter list. [The body loop](https://github.com/morluto/rea/blob/642f6571ea9d333ea50114c25e648dbb5b8a3c46/src/gpu/PtxKernels.ts#L272-L297) skips every declaration other than `.reg` and `.shared`, including the legacy kernel `.param`. [The emitted kernel](https://github.com/morluto/rea/blob/642f6571ea9d333ea50114c25e648dbb5b8a3c46/src/gpu/PtxKernels.ts#L320-L328) consequently publishes the initial empty list and zero byte total. [Header recognition](https://github.com/morluto/rea/blob/642f6571ea9d333ea50114c25e648dbb5b8a3c46/src/gpu/PtxKernels.ts#L404-L415) admits this PTX version, and the inventory consumer preserves the parser's complete status. ### Additional context Open issue bodies and all-state searches for `PTX`, `PTX 1.3`, and body/parameter combinations found no matching report. Related closed [#6493](https://github.com/morluto/rea/issues/6493) concerns rejecting an entry with an omitted optional parameter list; this report concerns silently dropping a parameter from the body after that spelling has been accepted. A fresh pre-submission search also found no duplicate. NOT_ENOUGH_INFO: first affected release and prevalence of legacy PTX among users were not investigated. ## Triage assessment - Impact level: P2 - Assessment status: supported - Impact basis: An explicitly declared kernel parameter disappears from normalized facts while inspection is reported complete for a valid bounded legacy PTX artifact. - Workaround status: available - Workaround basis: The emitted complete Evidence retains the original PTX text, where the parameter can be read manually. The executed 1.4 translation control also yields the correct declaration, but changes artifact bytes and version and is only suitable for a disposable analysis copy.

bsakaguchi closed this issue

2026-10-31 00:45:45 +00:00

bsakaguchi referenced this issue

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

[PR #4101] [CLOSED] fix(ghidra): compare DOS relocation windows with the final loaded memory #6056

bsakaguchi referenced this issue

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

[PR #4156] [CLOSED] fix(ghidra): compare DOS relocation windows with final loaded bytes #6094

bsakaguchi referenced this issue

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

[PR #4155] [CLOSED] fix(pe): retain empty resources at the section's virtual end #6095