morluto/rea

Fork 0

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

[GH-ISSUE #6626] [Bug] Call prototypes erase known PTX virtual register totals #4010

New issue

Closed

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

bsakaguchi commented

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

Owner

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

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

Adding Fproto: .callprototype _ (.reg .b32 _); to a kernel with the supported declaration .reg .b32 %r<4>; changes its virtual .b32 count from 4 to null and payload coverage from complete to partial. Both the parser and the real shared application service/provider produced this result. The prototype describes a callee signature; it does not declare another kernel register or make the four actual declarations unknown. The retained PTX still exactly matches the input.

Reproduction

Prerequisites: the source checkout identified under REA version, its locked dependencies already installed, and the Node runtime below. No build, GPU, NVIDIA driver, CUDA toolkit, Hopper or Ghidra is required for this source-boundary probe. Run from that checkout. Save this script as /tmp/rea-repro/probe.mts after creating /tmp/rea-repro. It is the executed probe, with its work and checkout supplied as arguments rather than private paths.

import { pathToFileURL } from 'node:url';
import { join, resolve } from 'node:path';
import { writeFile } from 'node:fs/promises';
const repo=resolve(process.argv[2]), work=resolve(process.argv[3]);
const load=(p:string)=>import(pathToFileURL(join(repo,p)).href);
const {readPtxKernels}=await load('src/gpu/PtxKernels.ts');
const {GpuCodeObjectsService}=await load('src/application/gpu/GpuCodeObjectsService.ts');
const {GpuCodeObjectsProvider}=await load('src/gpu/GpuCodeObjectsProvider.ts');
const service=new GpuCodeObjectsService(new GpuCodeObjectsProvider());
console.log(JSON.stringify({node:process.version,platform:process.platform,arch:process.arch}));
try {
 for(const [name,prototype] of [['control',''],['prototype','Fproto: .callprototype _ (.reg .b32 _);']] as const){
  const text='.version 8.8\n.target sm_80\n.address_size 64\n.visible .entry k() {\n.reg .b32 %r<4>;\n'+prototype+'\nret;\n}\n';
  const path=join(work,name+'.ptx'); await writeFile(path,text);
  const module=[...readPtxKernels(text,'object')].map(({kernel,coverage}:any)=>({coverage,name:kernel.name,virtual:kernel.registers.virtual,limitations:kernel.limitations}));
  const result=await service.inspect({path,analysis:'kernels',detail:'complete'});
  if(!result.ok)throw result.error;
  const r=result.value.normalized_result;
  console.log(JSON.stringify({case:name,module,service:{coverage:r.coverage,objects:r.objects.map(({inspection,limitations}:any)=>({inspection,limitations})),kernels:r.kernels.map(({name,registers,limitations}:any)=>({name,virtual:registers.virtual,limitations})),retained_text_matches:r.ptx_artifacts[0].text===text}}));
 }
}finally{await service.close();}

Execute with isolated home/configuration directories:

REPO="$PWD"
WORK=/tmp/rea-repro
mkdir -p "$WORK"/home "$WORK"/config "$WORK"/cache "$WORK"/data "$WORK"/state "$WORK"/run "$WORK"/temp
env PATH="$PATH" HOME="$WORK/home" USERPROFILE="$WORK/home" XDG_CONFIG_HOME="$WORK/config" XDG_CACHE_HOME="$WORK/cache" XDG_DATA_HOME="$WORK/data" XDG_STATE_HOME="$WORK/state" XDG_RUNTIME_DIR="$WORK/run" TMPDIR="$WORK/temp" node --import "$REPO/node_modules/tsx/dist/loader.mjs" "$WORK/probe.mts" "$REPO" "$WORK"

Each of the two inputs was executed once through each boundary before drafting, with HOME/USERPROFILE/XDG/TMPDIR redirected beneath the evidence directory. The affected input is ordinary standalone PTX:

.version 8.8
.target sm_80
.address_size 64
.visible .entry k() {
.reg .b32 %r<4>;
Fproto: .callprototype _ (.reg .b32 _);
ret;
}

Expected behavior

Retain the known .b32 virtual count of 4 and complete decoding for this otherwise supported kernel. REA's PTX contract counts scalar register names/ranges and preserves known facts. NVIDIA's PTX ISA 8.8 .callprototype specification permits base-type formal parameters in register state space and the sink name _; this directive describes an indirect-call prototype, rather than declaring caller-owned register variables. It was introduced in PTX 2.1 and requires sm_20 or later, so the fixture's declared version and target satisfy those requirements. Validity is grounded in this specification; the fixture was not independently assembled.

Impact and scope

A supported kernel's known virtual register total becomes unavailable solely because an ordinary indirect-call prototype appears in its body. The public-facing shared service retains the wrong null total and partial coverage in Evidence. This is bounded metadata loss; parameter facts and original PTX survive. Prevalence among real compiler outputs and effects on other register types were not measured. Reading/counting the four names from the retained declaration provides a manual workaround; no source file is changed by inspection. No security-boundary crossing or target execution is involved.

REA version

Package 6.7.0; source revision 00a9e95de91a4bd4d2a41f8aee704a53fa575750. This is a source-boundary finding; no installed-package or earlier-release comparison was run.

Runtime and platform

Node.js v26.12.1; Linux x64, kernel 7.2.8-9-omarchy; locked tsx 4.23.15. The probe used the already prepared source checkout and source loader. No install, build or test suite ran.

Target and artifact context

Two small UTF-8 PTX 8.8 files targeting sm_80 with 64-bit addresses, one zero-argument entry, and four .b32 virtual registers. The affected input adds one unused .callprototype with a register-space formal parameter. The control differs only by omission of that directive.

Interface

Other

Earliest failure stage

Native analysis or reconstruction

Relevant output or evidence

Observed process exit: 0. Stderr: empty. Selected fields from the saved JSON output:

[
  {"case":"control","module":{"coverage":"complete","virtual":[{"type":".b32","count":4}]},"service":{"payloads":"complete","inspection":"complete","virtual":[{"type":".b32","count":4}],"retained_text_matches":true}},
  {"case":"prototype","module":{"coverage":"partial","virtual":[{"type":".b32","count":null}]},"service":{"payloads":"partial","inspection":"partial","virtual":[{"type":".b32","count":null}],"retained_text_matches":true}}
]

The affected kernel adds exactly this limitation: Unsupported PTX register declarations or overflowing totals leave the affected virtual register count unknown. Discovery remains complete in both service results.

Expected source: the pinned REA contract and NVIDIA PTX specification linked above.
Failure source: the executed parser and shared-service comparison shown above; original script, both exact PTX inputs, stdout and empty stderr were preserved for the research handoff.
Evidence provenance: observed runtime failure and control; expected syntax and reachability are attributed source evidence.
Local verification: reproduced.
Reproduction completeness: complete for the parser and production shared application-service/provider boundary; CLI/MCP transport and independent producer validation were not executed.

Typed outcome or recovery expected

Return the existing .b32 count 4, preserving the retained text, parameter facts and complete coverage for both inputs. A proposed restoration check compares these same two files at the shared service and then through a public adapter. Keep genuinely unsupported register declarations unknown with partial coverage; restoring the count must not discard that existing contract.

Analysis or suspected cause

Source-derived: the body scanner sends every .reg token encountered in an entry body to the resource-declaration parser, including formal types nested inside .callprototype. The declaration parser encounters ) after _, treats the construct as unsupported, consumes through the prototype's semicolon, and merges a null into the actual register total. A format-aware distinction between resource declarations and prototype signatures would preserve known facts without requiring physical-register or call analysis. This is a restoration direction, not an implemented fix.

Additional context

Source-derived public reachability: CLI and MCP both call the exercised shared service. Standalone recognition and payload extraction admit the valid header, and kernel-mode decoding invokes this parser. analysis: "kernels" is required; inventory-only requests do not parse entries. Equivalent public CLI and MCP checks remain proposed, not observed. Existing register unit coverage exercises actual declarations, unsupported totals and malformed ranges, but does not distinguish prototype formal types. No real NVIDIA producer/assembler, GPU runtime, Hopper or Ghidra lane was run. The exercised provider is REA's actual static reader, not a mocked provider. No source edits, issue filing or repair verification was performed.

Triage assessment

Impact level: P2
Assessment status: supported
Impact basis: The executed production shared service loses a known virtual-register count and downgrades payload coverage for the demonstrated ordinary PTX prototype; the impact is bounded metadata loss with retained source evidence.
Workaround status: available
Workaround basis: Both executions retained PTX text exactly, including the supported four-register declaration, so the affected fact can be recovered manually without changing or executing the target; a correct structured total still requires a parser fix.

Originally created by @coygeek on GitHub (Oct 23, 2026). Original GitHub issue: https://github.com/morluto/rea/issues/6626 ### Before submitting - [x] 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 Adding `Fproto: .callprototype _ (.reg .b32 _);` to a kernel with the supported declaration `.reg .b32 %r<4>;` changes its virtual `.b32` count from `4` to `null` and payload coverage from `complete` to `partial`. Both the parser and the real shared application service/provider produced this result. The prototype describes a callee signature; it does not declare another kernel register or make the four actual declarations unknown. The retained PTX still exactly matches the input. ### Reproduction Prerequisites: the source checkout identified under REA version, its locked dependencies already installed, and the Node runtime below. No build, GPU, NVIDIA driver, CUDA toolkit, Hopper or Ghidra is required for this source-boundary probe. Run from that checkout. Save this script as `/tmp/rea-repro/probe.mts` after creating `/tmp/rea-repro`. It is the executed probe, with its work and checkout supplied as arguments rather than private paths. ```ts import { pathToFileURL } from 'node:url'; import { join, resolve } from 'node:path'; import { writeFile } from 'node:fs/promises'; const repo=resolve(process.argv[2]), work=resolve(process.argv[3]); const load=(p:string)=>import(pathToFileURL(join(repo,p)).href); const {readPtxKernels}=await load('src/gpu/PtxKernels.ts'); const {GpuCodeObjectsService}=await load('src/application/gpu/GpuCodeObjectsService.ts'); const {GpuCodeObjectsProvider}=await load('src/gpu/GpuCodeObjectsProvider.ts'); const service=new GpuCodeObjectsService(new GpuCodeObjectsProvider()); console.log(JSON.stringify({node:process.version,platform:process.platform,arch:process.arch})); try { for(const [name,prototype] of [['control',''],['prototype','Fproto: .callprototype _ (.reg .b32 _);']] as const){ const text='.version 8.8\n.target sm_80\n.address_size 64\n.visible .entry k() {\n.reg .b32 %r<4>;\n'+prototype+'\nret;\n}\n'; const path=join(work,name+'.ptx'); await writeFile(path,text); const module=[...readPtxKernels(text,'object')].map(({kernel,coverage}:any)=>({coverage,name:kernel.name,virtual:kernel.registers.virtual,limitations:kernel.limitations})); const result=await service.inspect({path,analysis:'kernels',detail:'complete'}); if(!result.ok)throw result.error; const r=result.value.normalized_result; console.log(JSON.stringify({case:name,module,service:{coverage:r.coverage,objects:r.objects.map(({inspection,limitations}:any)=>({inspection,limitations})),kernels:r.kernels.map(({name,registers,limitations}:any)=>({name,virtual:registers.virtual,limitations})),retained_text_matches:r.ptx_artifacts[0].text===text}})); } }finally{await service.close();} ``` Execute with isolated home/configuration directories: ```sh REPO="$PWD" WORK=/tmp/rea-repro mkdir -p "$WORK"/home "$WORK"/config "$WORK"/cache "$WORK"/data "$WORK"/state "$WORK"/run "$WORK"/temp env PATH="$PATH" HOME="$WORK/home" USERPROFILE="$WORK/home" XDG_CONFIG_HOME="$WORK/config" XDG_CACHE_HOME="$WORK/cache" XDG_DATA_HOME="$WORK/data" XDG_STATE_HOME="$WORK/state" XDG_RUNTIME_DIR="$WORK/run" TMPDIR="$WORK/temp" node --import "$REPO/node_modules/tsx/dist/loader.mjs" "$WORK/probe.mts" "$REPO" "$WORK" ``` Each of the two inputs was executed once through each boundary before drafting, with HOME/USERPROFILE/XDG/TMPDIR redirected beneath the evidence directory. The affected input is ordinary standalone PTX: ```ptx .version 8.8 .target sm_80 .address_size 64 .visible .entry k() { .reg .b32 %r<4>; Fproto: .callprototype _ (.reg .b32 _); ret; } ``` ### Expected behavior Retain the known `.b32` virtual count of `4` and complete decoding for this otherwise supported kernel. [REA's PTX contract](https://github.com/morluto/rea/blob/00a9e95de91a4bd4d2a41f8aee704a53fa575750/docs/gpu-code-objects.md#L79-L90) counts scalar register names/ranges and preserves known facts. NVIDIA's [PTX ISA 8.8 `.callprototype` specification](https://docs.nvidia.com/cuda/archive/12.9.1/parallel-thread-execution/index.html#control-flow-directives-callprototype) permits base-type formal parameters in register state space and the sink name `_`; this directive describes an indirect-call prototype, rather than declaring caller-owned register variables. It was introduced in PTX 2.1 and requires sm_20 or later, so the fixture's declared version and target satisfy those requirements. Validity is grounded in this specification; the fixture was not independently assembled. ### Impact and scope A supported kernel's known virtual register total becomes unavailable solely because an ordinary indirect-call prototype appears in its body. The public-facing shared service retains the wrong null total and partial coverage in Evidence. This is bounded metadata loss; parameter facts and original PTX survive. Prevalence among real compiler outputs and effects on other register types were not measured. Reading/counting the four names from the retained declaration provides a manual workaround; no source file is changed by inspection. No security-boundary crossing or target execution is involved. ### REA version Package 6.7.0; source revision `00a9e95de91a4bd4d2a41f8aee704a53fa575750`. This is a source-boundary finding; no installed-package or earlier-release comparison was run. ### Runtime and platform Node.js v26.12.1; Linux x64, kernel 7.2.8-9-omarchy; locked tsx 4.23.15. The probe used the already prepared source checkout and source loader. No install, build or test suite ran. ### Target and artifact context Two small UTF-8 PTX 8.8 files targeting sm_80 with 64-bit addresses, one zero-argument entry, and four `.b32` virtual registers. The affected input adds one unused `.callprototype` with a register-space formal parameter. The control differs only by omission of that directive. ### Interface Other ### Earliest failure stage Native analysis or reconstruction ### Relevant output or evidence Observed process exit: 0. Stderr: empty. Selected fields from the saved JSON output: ```json [ {"case":"control","module":{"coverage":"complete","virtual":[{"type":".b32","count":4}]},"service":{"payloads":"complete","inspection":"complete","virtual":[{"type":".b32","count":4}],"retained_text_matches":true}}, {"case":"prototype","module":{"coverage":"partial","virtual":[{"type":".b32","count":null}]},"service":{"payloads":"partial","inspection":"partial","virtual":[{"type":".b32","count":null}],"retained_text_matches":true}} ] ``` The affected kernel adds exactly this limitation: `Unsupported PTX register declarations or overflowing totals leave the affected virtual register count unknown.` Discovery remains complete in both service results. - Expected source: the pinned REA contract and NVIDIA PTX specification linked above. - Failure source: the executed parser and shared-service comparison shown above; original script, both exact PTX inputs, stdout and empty stderr were preserved for the research handoff. - Evidence provenance: observed runtime failure and control; expected syntax and reachability are attributed source evidence. - Local verification: reproduced. - Reproduction completeness: complete for the parser and production shared application-service/provider boundary; CLI/MCP transport and independent producer validation were not executed. ### Typed outcome or recovery expected Return the existing `.b32` count `4`, preserving the retained text, parameter facts and complete coverage for both inputs. A proposed restoration check compares these same two files at the shared service and then through a public adapter. Keep genuinely unsupported register declarations unknown with partial coverage; restoring the count must not discard that existing contract. ### Analysis or suspected cause Source-derived: [the body scanner](https://github.com/morluto/rea/blob/00a9e95de91a4bd4d2a41f8aee704a53fa575750/src/gpu/PtxKernels.ts#L329-L343) sends every `.reg` token encountered in an entry body to the resource-declaration parser, including formal types nested inside `.callprototype`. [The declaration parser](https://github.com/morluto/rea/blob/00a9e95de91a4bd4d2a41f8aee704a53fa575750/src/gpu/PtxKernels.ts#L90-L137) encounters `)` after `_`, treats the construct as unsupported, consumes through the prototype's semicolon, and merges a null into the actual register total. A format-aware distinction between resource declarations and prototype signatures would preserve known facts without requiring physical-register or call analysis. This is a restoration direction, not an implemented fix. ### Additional context Source-derived public reachability: [CLI](https://github.com/morluto/rea/blob/00a9e95de91a4bd4d2a41f8aee704a53fa575750/src/cli/gpuCommands.ts#L90-L117) and [MCP](https://github.com/morluto/rea/blob/00a9e95de91a4bd4d2a41f8aee704a53fa575750/src/server/registerGpuCodeObjectsTools.ts#L20-L30) both call the exercised shared service. [Standalone recognition and payload extraction](https://github.com/morluto/rea/blob/00a9e95de91a4bd4d2a41f8aee704a53fa575750/src/gpu/GpuInventory.ts#L228-L252) admit the valid header, and [kernel-mode decoding](https://github.com/morluto/rea/blob/00a9e95de91a4bd4d2a41f8aee704a53fa575750/src/gpu/GpuInventory.ts#L344-L376) invokes this parser. `analysis: "kernels"` is required; inventory-only requests do not parse entries. Equivalent public CLI and MCP checks remain proposed, not observed. Existing [register unit coverage](https://github.com/morluto/rea/blob/00a9e95de91a4bd4d2a41f8aee704a53fa575750/src/gpu/PtxKernels.test.ts#L51-L91) exercises actual declarations, unsupported totals and malformed ranges, but does not distinguish prototype formal types. No real NVIDIA producer/assembler, GPU runtime, Hopper or Ghidra lane was run. The exercised provider is REA's actual static reader, not a mocked provider. No source edits, issue filing or repair verification was performed. ## Triage assessment - Impact level: P2 - Assessment status: supported - Impact basis: The executed production shared service loses a known virtual-register count and downgrades payload coverage for the demonstrated ordinary PTX prototype; the impact is bounded metadata loss with retained source evidence. - Workaround status: available - Workaround basis: Both executions retained PTX text exactly, including the supported four-register declaration, so the affected fact can be recovered manually without changing or executing the target; a correct structured total still requires a parser fix.

bsakaguchi closed this issue

2026-10-31 00:46:02 +00:00

bsakaguchi referenced this issue

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

[PR #4010] [MERGED] fix(browser): confirm service-worker cancellation cleanup #5989

bsakaguchi referenced this issue

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

[PR #4034] [MERGED] fix(ci): reject false browser environment cancellation passes #6005