[GH-ISSUE #6858] [Bug] Compact PTX register declarations disappear with complete decoding coverage #4196

New issue

Closed

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

bsakaguchi commented

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

Owner

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

Before submitting

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

Area

Evidence, artifacts, or reconstruction

Observed behavior

Executed through the production GPU service and provider: .reg.b32 %r<4>; .reg.b64 %rd<2>; produces an empty registers.virtual array. The same declarations with whitespace between .reg and the type return .b32: 4 and .b64: 2. Both results mark the PTX object and payload decoding complete. The compact case has no limitation identifying omitted registers.

This report concerns unreleased main, not the published npm package. The fetched published release tag lacks the GPU subsystem.

Reproduction

In the source checkout with locked dependencies installed, save the following as /tmp/compact-registers.mts, then run:

REA_ROOT="$PWD" node --import ./node_modules/tsx/dist/loader.mjs /tmp/compact-registers.mts

import { writeFile } from 'node:fs/promises';
import { dirname, join } from 'node:path';
import { fileURLToPath, pathToFileURL } from 'node:url';
const root=process.env.REA_ROOT!;
const dir=dirname(fileURLToPath(import.meta.url));
const { createGpuCodeObjectsService }=await import(pathToFileURL(`${root}/src/composition/gpu.ts`).href);
const rows=[];
for(const compact of [false,true]){
 const path=join(dir,compact?'compact.ptx':'spaced.ptx');
 const reg=compact?'.reg.b32 %r<4>; .reg.b64 %rd<2>;':'.reg .b32 %r<4>; .reg .b64 %rd<2>;';
 await writeFile(path,`.version 8.8\n.target sm_80\n.address_size 64\n.visible .entry k() { ${reg} ret; }\n`);
 const service=createGpuCodeObjectsService();
 try{
  const result=await service.inspect({path,analysis:'kernels',detail:'complete'});
  if(!result.ok)throw result.error;
  const report=result.value.normalized_result as any;
  rows.push({compact,coverage:report.coverage,objects:report.objects.map((o:any)=>({inspection:o.inspection,target:o.target})),registers:report.kernels[0].registers.virtual,limitations:report.kernels[0].limitations});
 }finally{await service.close();}
}
console.log(JSON.stringify({node:process.version,platform:process.platform,arch:process.arch,rows},null,2));

This two-case source-service probe ran once. An earlier independent five-case source-service matrix also observed the compact-register omission, including a mixed form where four spaced .b32 registers survived but three compact .b32 registers and two compact .b64 registers disappeared. All execution finished before this draft was written.

Expected behavior

Retain the four .b32 and two .b64 virtual register declarations regardless of this spacing choice. If a declaration cannot be decoded, retain explicit uncertainty and partial coverage rather than silently reporting a complete empty inventory. The documented PTX contract promises counts for named and parameterized registers and unknown counts with partial decoding coverage for unsupported forms.

The compact spelling is grounded in real producer source: Dr.Jit's CUDA generator emits .reg.b32, .reg.b64, and other compact register declarations. That producer was source-inspected, not built or executed.

Impact and scope

The reduced case loses all six declared virtual registers while claiming complete decoding. The mixed case undercounts a retained type and omits another type. This affects static metadata and completeness claims; it does not establish anything about physical register allocation or GPU behavior. Observed scope is synthetic standalone PTX with the compact declaration form. Real Dr.Jit-generated artifacts, embedded fatbin delivery, other platforms, and CLI/MCP transport were not executed.

REA version

Source revision 96ea32b7e4e28d3c9cd79a1de4b8e1ba3b4244e9 on unreleased main; package manifest 6.7.0. Published tag rea-agents-6.7.0 was inspected and contains no src/gpu or this operation; no released-package reproduction is claimed.

Runtime and platform

Node.js v24.22.1; Linux x64; locked source tsx loader. No CUDA assembler, GPU, Hopper, or Ghidra was used.

Target and artifact context

Two self-authored PTX 8.8 text fixtures declaring target sm_80 and entry k. They differ only in whitespace between the register directive and type. The service reads files without executing them.

Interface

Other

Earliest failure stage

Native analysis or reconstruction

Relevant output or evidence

Reduced output from the executed two-case probe:

{
  "spaced": {
    "registers": [{"type":".b32","count":4},{"type":".b64","count":2}],
    "object_inspection": "complete",
    "coverage_payloads": "complete"
  },
  "compact": {
    "registers": [],
    "object_inspection": "complete",
    "coverage_payloads": "complete"
  }
}

Evidence provenance: observed source-service results and file fixtures; source-derived producer syntax and cause. The compact kernel's only limitation is the generic warning that virtual declarations are not physical allocations; no unsupported-declaration warning appears. The earlier mixed case returned only [{"type":".b32","count":4}] with complete payload coverage.

Typed outcome or recovery expected

Complete facts should include the declared virtual counts. A reader that cannot interpret the syntax should return partial decoding coverage with an explicit limitation. The spaced control must continue to report the same counts.

Analysis or suspected cause

Source-derived: the tokenizer keeps adjacent dots and word characters in one token, so .reg.b32 does not become a recognized .reg directive. The entry resource scanner skips it. The partial-coverage check examines only retained register rows, so this omission does not lower coverage. No fix was implemented.

Additional context

Known CUBIN register-conflict report #6823 concerns a different representation and failure. Closed #6626 concerns call-prototype formals, and closed #6469 concerns long register lists. They do not describe the compact entry-local declarations here. Related variants with compact shared/entry directives were not promoted to separate reports because their syntax validity is less established and they share the tokenization boundary.

NOT_ENOUGH_INFO: first affected source revision, real generated-artifact prevalence, independent assembler acceptance of these exact fixtures, and behavior in a future published release. Prepared with AI assistance. No verification or execution of this bug occurred after drafting.

Triage assessment

Impact level: P2
Assessment status: supported
Impact basis: Production source-service observations show omitted or undercounted virtual registers with complete decoding coverage on a bounded, producer-grounded PTX spelling. The failure affects analysis facts without executing or modifying target programs.
Workaround status: available
Workaround basis: For the reproduced declarations, inserting whitespace between .reg and the type in an inspection copy produces the expected counts, as demonstrated by the spaced control. This changes the inspected copy's identity; a general automatic rewriting workflow was not tested.

Originally created by @coygeek on GitHub (Oct 23, 2026). Original GitHub issue: https://github.com/morluto/rea/issues/6858 ### 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 Evidence, artifacts, or reconstruction ### Observed behavior Executed through the production GPU service and provider: `.reg.b32 %r<4>; .reg.b64 %rd<2>;` produces an empty `registers.virtual` array. The same declarations with whitespace between `.reg` and the type return `.b32: 4` and `.b64: 2`. Both results mark the PTX object and payload decoding complete. The compact case has no limitation identifying omitted registers. This report concerns unreleased `main`, not the published npm package. The fetched published release tag lacks the GPU subsystem. ### Reproduction In the source checkout with locked dependencies installed, save the following as `/tmp/compact-registers.mts`, then run: ```sh REA_ROOT="$PWD" node --import ./node_modules/tsx/dist/loader.mjs /tmp/compact-registers.mts ``` ```ts import { writeFile } from 'node:fs/promises'; import { dirname, join } from 'node:path'; import { fileURLToPath, pathToFileURL } from 'node:url'; const root=process.env.REA_ROOT!; const dir=dirname(fileURLToPath(import.meta.url)); const { createGpuCodeObjectsService }=await import(pathToFileURL(`${root}/src/composition/gpu.ts`).href); const rows=[]; for(const compact of [false,true]){ const path=join(dir,compact?'compact.ptx':'spaced.ptx'); const reg=compact?'.reg.b32 %r<4>; .reg.b64 %rd<2>;':'.reg .b32 %r<4>; .reg .b64 %rd<2>;'; await writeFile(path,`.version 8.8\n.target sm_80\n.address_size 64\n.visible .entry k() { ${reg} ret; }\n`); const service=createGpuCodeObjectsService(); try{ const result=await service.inspect({path,analysis:'kernels',detail:'complete'}); if(!result.ok)throw result.error; const report=result.value.normalized_result as any; rows.push({compact,coverage:report.coverage,objects:report.objects.map((o:any)=>({inspection:o.inspection,target:o.target})),registers:report.kernels[0].registers.virtual,limitations:report.kernels[0].limitations}); }finally{await service.close();} } console.log(JSON.stringify({node:process.version,platform:process.platform,arch:process.arch,rows},null,2)); ``` This two-case source-service probe ran once. An earlier independent five-case source-service matrix also observed the compact-register omission, including a mixed form where four spaced `.b32` registers survived but three compact `.b32` registers and two compact `.b64` registers disappeared. All execution finished before this draft was written. ### Expected behavior Retain the four `.b32` and two `.b64` virtual register declarations regardless of this spacing choice. If a declaration cannot be decoded, retain explicit uncertainty and partial coverage rather than silently reporting a complete empty inventory. The [documented PTX contract](https://github.com/morluto/rea/blob/96ea32b7e4e28d3c9cd79a1de4b8e1ba3b4244e9/docs/gpu-code-objects.md#L110-L114) promises counts for named and parameterized registers and unknown counts with partial decoding coverage for unsupported forms. The compact spelling is grounded in real producer source: [Dr.Jit's CUDA generator](https://github.com/mitsuba-renderer/drjit-core/blob/1feeb021d24a17a4ffc2b9cab79652cea920021a/src/cuda_eval.cpp#L114-L126) emits `.reg.b32`, `.reg.b64`, and other compact register declarations. That producer was source-inspected, not built or executed. ### Impact and scope The reduced case loses all six declared virtual registers while claiming complete decoding. The mixed case undercounts a retained type and omits another type. This affects static metadata and completeness claims; it does not establish anything about physical register allocation or GPU behavior. Observed scope is synthetic standalone PTX with the compact declaration form. Real Dr.Jit-generated artifacts, embedded fatbin delivery, other platforms, and CLI/MCP transport were not executed. ### REA version Source revision `96ea32b7e4e28d3c9cd79a1de4b8e1ba3b4244e9` on unreleased `main`; package manifest 6.7.0. Published tag `rea-agents-6.7.0` was inspected and contains no `src/gpu` or this operation; no released-package reproduction is claimed. ### Runtime and platform Node.js v24.22.1; Linux x64; locked source tsx loader. No CUDA assembler, GPU, Hopper, or Ghidra was used. ### Target and artifact context Two self-authored PTX 8.8 text fixtures declaring target `sm_80` and entry `k`. They differ only in whitespace between the register directive and type. The service reads files without executing them. ### Interface Other ### Earliest failure stage Native analysis or reconstruction ### Relevant output or evidence Reduced output from the executed two-case probe: ```json { "spaced": { "registers": [{"type":".b32","count":4},{"type":".b64","count":2}], "object_inspection": "complete", "coverage_payloads": "complete" }, "compact": { "registers": [], "object_inspection": "complete", "coverage_payloads": "complete" } } ``` Evidence provenance: observed source-service results and file fixtures; source-derived producer syntax and cause. The compact kernel's only limitation is the generic warning that virtual declarations are not physical allocations; no unsupported-declaration warning appears. The earlier mixed case returned only `[{"type":".b32","count":4}]` with complete payload coverage. ### Typed outcome or recovery expected Complete facts should include the declared virtual counts. A reader that cannot interpret the syntax should return partial decoding coverage with an explicit limitation. The spaced control must continue to report the same counts. ### Analysis or suspected cause Source-derived: the [tokenizer](https://github.com/morluto/rea/blob/96ea32b7e4e28d3c9cd79a1de4b8e1ba3b4244e9/src/gpu/PtxKernels.ts#L57-L66) keeps adjacent dots and word characters in one token, so `.reg.b32` does not become a recognized `.reg` directive. The [entry resource scanner](https://github.com/morluto/rea/blob/96ea32b7e4e28d3c9cd79a1de4b8e1ba3b4244e9/src/gpu/PtxKernels.ts#L419-L421) skips it. The [partial-coverage check](https://github.com/morluto/rea/blob/96ea32b7e4e28d3c9cd79a1de4b8e1ba3b4244e9/src/gpu/PtxKernels.ts#L508-L511) examines only retained register rows, so this omission does not lower coverage. No fix was implemented. ### Additional context Known CUBIN register-conflict report #6823 concerns a different representation and failure. Closed #6626 concerns call-prototype formals, and closed #6469 concerns long register lists. They do not describe the compact entry-local declarations here. Related variants with compact shared/entry directives were not promoted to separate reports because their syntax validity is less established and they share the tokenization boundary. NOT_ENOUGH_INFO: first affected source revision, real generated-artifact prevalence, independent assembler acceptance of these exact fixtures, and behavior in a future published release. Prepared with AI assistance. No verification or execution of this bug occurred after drafting. ## Triage assessment - Impact level: P2 - Assessment status: supported - Impact basis: Production source-service observations show omitted or undercounted virtual registers with complete decoding coverage on a bounded, producer-grounded PTX spelling. The failure affects analysis facts without executing or modifying target programs. - Workaround status: available - Workaround basis: For the reproduced declarations, inserting whitespace between `.reg` and the type in an inspection copy produces the expected counts, as demonstrated by the spaced control. This changes the inspected copy's identity; a general automatic rewriting workflow was not tested.

bsakaguchi closed this issue

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

bsakaguchi referenced this issue

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

[PR #4196] [MERGED] fix(javascript): preserve values captured before later assignments #6119