[GH-ISSUE #6680] [Bug] Multidimensional PTX parameters lose their declared names #4058

Closed
opened 2026-10-31 00:46:26 +00:00 by bsakaguchi · 0 comments
Owner

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

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 production source GPU application service returns name: null for payload in .entry k(.param .b8 payload[2][3], .param .u32 count). It retains two parameter records and the name count, with partial coverage and unknown array layout. The equivalent one-dimensional payload[6] control retains both names. An unsupported constant-expression control, payload[2+4], also retains both names while keeping its layout unknown.

This report concerns the lost identifier, independently of whether multidimensional layout calculation is supported. No installed CLI or published-package runtime reproduction is claimed.

Reproduction

Use the source checkout at the revision below with locked dependencies installed. Set WORK to an existing disposable directory containing fixtures and evidence subdirectories. Save this executed probe as $WORK/ptx-array-names.mts:

import { writeFile } from 'node:fs/promises';
import { resolve, join } from 'node:path';
import { pathToFileURL } from 'node:url';
const repo = resolve(process.argv[2]), work = resolve(process.argv[3]);
const load = (p:string) => import(pathToFileURL(join(repo,p)).href);
const { GpuCodeObjectsProvider } = await load('src/gpu/GpuCodeObjectsProvider.ts');
const { GpuCodeObjectsService } = await load('src/application/gpu/GpuCodeObjectsService.ts');
const cases = {
 literal: '.param .b8 payload[6], .param .u32 count',
 multidimensional: '.param .b8 payload[2][3], .param .u32 count',
 unsupported_expression: '.param .b8 payload[2+4], .param .u32 count',
};
const service = new GpuCodeObjectsService(new GpuCodeObjectsProvider());
console.log(JSON.stringify({ node:process.version,platform:process.platform,arch:process.arch }));
try {
 for (const [name, params] of Object.entries(cases)) {
  const text = `.version 8.8\n.target sm_80\n.address_size 64\n.visible .entry k(${params}) { .reg .b32 %r<2>; ret; }\n`;
  const path = join(work,'fixtures',name+'.ptx');
  await writeFile(path,text);
  const result = await service.inspect({ path,analysis:'kernels',detail:'complete' });
  await writeFile(join(work,'evidence',name+'.json'),JSON.stringify(result,null,2));
  if (!result.ok) throw result.error;
  const report = result.value.normalized_result;
  console.log(JSON.stringify({ name,coverage:report.coverage,objects:report.objects.map(({inspection,limitations}:any)=>({inspection,limitations})),kernels:report.kernels.map(({name,parameter_count,parameter_bytes,parameters,limitations}:any)=>({name,parameter_count,parameter_bytes,parameters,limitations})),retained_text:report.ptx_artifacts[0]?.text===text }));
 }
} finally { await service.close(); }

From the checkout, the executed command was (the disposable directory is represented by $WORK):

node --import ./node_modules/tsx/dist/loader.mjs "$WORK/ptx-array-names.mts" "$PWD" "$WORK" > "$WORK/evidence/ptx-array-names.stdout.jsonl" 2> "$WORK/evidence/ptx-array-names.stderr.txt"

Each of the three inputs was inspected once. The process exited 0 and stderr was empty. The probe writes only its inert local fixtures/results and reads them through the real shared service and provider. No target code runs.

Expected behavior

Retain the declared payload identifier in the first parameter row. Unsupported layout interpretation may leave size, offset and total bytes unknown with partial coverage, as the expression control already demonstrates.

NVIDIA PTX ISA 8.8 permits multidimensional array declarations. REA's PTX contract says unsupported parameter layouts preserve known facts and explicit limitations. The identifier is known without evaluating the array extent.

Impact and scope

A valid multidimensional declaration loses its name from machine-readable kernel parameter facts. Consumers must interpret the retained PTX manually to recover it. The observed loss is bounded: the parameter count, trailing scalar name/type/size and exact original PTX survive, and REA reports partial coverage. Three small standalone fixtures were tested; real-producer frequency, embedded containers and transport execution were not measured. Source input files were not modified by inspection.

REA version

Source revision 8c4f96ef4759fd78da66c4b9726825ca71d7ee0c; package metadata 6.7.0. This is a source-main finding. Latest published tag rea-agents-6.7.0 resolves to 27dc42d10249f0410904a75c0997842eb03b931c and has no src/gpu, src/application/gpu or src/domain/gpu paths. Published-package behavior was not tested.

Runtime and platform

Node.js v26.12.1; Linux 7.2.8-9-omarchy x86_64; repository-locked tsx loader. No CUDA compiler/assembler, GPU, Hopper or Ghidra was used.

Target and artifact context

Three synthetic UTF-8 PTX 8.8 files targeting sm_80, with 64-bit addresses, one two-parameter entry, two virtual .b32 registers and ret. Only the array dimension spelling changes: [6], [2][3], or [2+4]. Syntax validity is grounded in the published PTX array grammar; an independent assembler was not run.

Interface

Other

Earliest failure stage

Native analysis or reconstruction

Relevant output or evidence

Selected fields from the observed results:

literal [6]:             names=[payload,count]; sizes=[6,4];    parameter_count=2; parameter_bytes=12; payloads=complete
multidimensional [2][3]: names=[null,count];    sizes=[null,4]; parameter_count=2; parameter_bytes=null; payloads=partial
expression [2+4]:        names=[payload,count]; sizes=[null,4]; parameter_count=2; parameter_bytes=null; payloads=partial
All three: retained_text=true; process exit=0; stderr empty

The multidimensional and expression cases both add only this layout limitation:

Unsupported PTX parameter declarations or overflowing layouts leave parameter sizes/offsets and the total byte count unknown.
  • Expected source: the published PTX grammar and pinned REA preservation contract linked above.
  • Failure source: the executed production source-service comparison and complete probe above.
  • Evidence provenance: observed service results; grammar applicability and parser diagnosis are separately source-derived.
  • Local verification: reproduced.
  • Reproduction completeness: complete for the static application-service/provider boundary; independent assembly, installed CLI and MCP were not executed.

Typed outcome or recovery expected

Return name: "payload" for the first parameter. Preserve the second parameter and exact retained text. If multidimensional layout stays unsupported, retain honest unknown layout fields, partial coverage and the existing limitation. No new error type is necessary.

Analysis or suspected cause

Source-derived: PtxKernels.ts:193-215 scans backward from the last ] and stops at its matching [. For payload[2][3], this is the second dimension. It then chooses the preceding token, the first dimension's ], as the name candidate, producing null. The parser does not continue past all array suffixes to the actual identifier. This occurs before the separate supported-layout decision.

Additional context

The supplied open-issue corpus and live all-state searches covered PTX multidimensional, PTX "array" "name", and PTX "dimensions". No matching report was found. Related #6664 loses following parameter declarations because a parenthesized dimension prematurely terminates the parameter list. This fixture has no parentheses and retains both declarations. Closed #6606 concerns pointer qualifiers being mistaken for names; this fixture contains no pointer qualifier and reaches the independent array-suffix scan.

NOT_ENOUGH_INFO: first affected source revision, independently assembled output and real-world frequency were not established. No fix or restoration was tested. Prepared with AI assistance; no tracked source was changed.

Triage assessment

  • Impact level: P2
  • Assessment status: supported
  • Impact basis: A valid multidimensional PTX parameter loses its identifier from structured kernel facts through the production source service, while one-dimensional and unknown-layout controls retain that identifier. The affected result is explicitly partial and the original text survives.
  • Workaround status: available
  • Workaround basis: The observed complete Evidence retains the exact source text including payload[2][3], allowing manual identifier recovery without changing the artifact.
Originally created by @coygeek on GitHub (Oct 23, 2026). Original GitHub issue: https://github.com/morluto/rea/issues/6680 ## 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 production source GPU application service returns `name: null` for `payload` in `.entry k(.param .b8 payload[2][3], .param .u32 count)`. It retains two parameter records and the name `count`, with partial coverage and unknown array layout. The equivalent one-dimensional `payload[6]` control retains both names. An unsupported constant-expression control, `payload[2+4]`, also retains both names while keeping its layout unknown. This report concerns the lost identifier, independently of whether multidimensional layout calculation is supported. No installed CLI or published-package runtime reproduction is claimed. ## Reproduction Use the source checkout at the revision below with locked dependencies installed. Set `WORK` to an existing disposable directory containing `fixtures` and `evidence` subdirectories. Save this executed probe as `$WORK/ptx-array-names.mts`: ```ts import { writeFile } from 'node:fs/promises'; import { resolve, join } from 'node:path'; import { pathToFileURL } from 'node:url'; const repo = resolve(process.argv[2]), work = resolve(process.argv[3]); const load = (p:string) => import(pathToFileURL(join(repo,p)).href); const { GpuCodeObjectsProvider } = await load('src/gpu/GpuCodeObjectsProvider.ts'); const { GpuCodeObjectsService } = await load('src/application/gpu/GpuCodeObjectsService.ts'); const cases = { literal: '.param .b8 payload[6], .param .u32 count', multidimensional: '.param .b8 payload[2][3], .param .u32 count', unsupported_expression: '.param .b8 payload[2+4], .param .u32 count', }; const service = new GpuCodeObjectsService(new GpuCodeObjectsProvider()); console.log(JSON.stringify({ node:process.version,platform:process.platform,arch:process.arch })); try { for (const [name, params] of Object.entries(cases)) { const text = `.version 8.8\n.target sm_80\n.address_size 64\n.visible .entry k(${params}) { .reg .b32 %r<2>; ret; }\n`; const path = join(work,'fixtures',name+'.ptx'); await writeFile(path,text); const result = await service.inspect({ path,analysis:'kernels',detail:'complete' }); await writeFile(join(work,'evidence',name+'.json'),JSON.stringify(result,null,2)); if (!result.ok) throw result.error; const report = result.value.normalized_result; console.log(JSON.stringify({ name,coverage:report.coverage,objects:report.objects.map(({inspection,limitations}:any)=>({inspection,limitations})),kernels:report.kernels.map(({name,parameter_count,parameter_bytes,parameters,limitations}:any)=>({name,parameter_count,parameter_bytes,parameters,limitations})),retained_text:report.ptx_artifacts[0]?.text===text })); } } finally { await service.close(); } ``` From the checkout, the executed command was (the disposable directory is represented by `$WORK`): ```sh node --import ./node_modules/tsx/dist/loader.mjs "$WORK/ptx-array-names.mts" "$PWD" "$WORK" > "$WORK/evidence/ptx-array-names.stdout.jsonl" 2> "$WORK/evidence/ptx-array-names.stderr.txt" ``` Each of the three inputs was inspected once. The process exited 0 and stderr was empty. The probe writes only its inert local fixtures/results and reads them through the real shared service and provider. No target code runs. ## Expected behavior Retain the declared `payload` identifier in the first parameter row. Unsupported layout interpretation may leave size, offset and total bytes unknown with partial coverage, as the expression control already demonstrates. NVIDIA PTX ISA 8.8 permits [multidimensional array declarations](https://docs.nvidia.com/cuda/archive/12.9.1/parallel-thread-execution/index.html#array-declarations). REA's [PTX contract](https://github.com/morluto/rea/blob/8c4f96ef4759fd78da66c4b9726825ca71d7ee0c/docs/gpu-code-objects.md#L94-L98) says unsupported parameter layouts preserve known facts and explicit limitations. The identifier is known without evaluating the array extent. ## Impact and scope A valid multidimensional declaration loses its name from machine-readable kernel parameter facts. Consumers must interpret the retained PTX manually to recover it. The observed loss is bounded: the parameter count, trailing scalar name/type/size and exact original PTX survive, and REA reports partial coverage. Three small standalone fixtures were tested; real-producer frequency, embedded containers and transport execution were not measured. Source input files were not modified by inspection. ## REA version Source revision `8c4f96ef4759fd78da66c4b9726825ca71d7ee0c`; package metadata `6.7.0`. This is a source-main finding. Latest published tag `rea-agents-6.7.0` resolves to `27dc42d10249f0410904a75c0997842eb03b931c` and has no `src/gpu`, `src/application/gpu` or `src/domain/gpu` paths. Published-package behavior was not tested. ## Runtime and platform Node.js v26.12.1; Linux 7.2.8-9-omarchy x86_64; repository-locked tsx loader. No CUDA compiler/assembler, GPU, Hopper or Ghidra was used. ## Target and artifact context Three synthetic UTF-8 PTX 8.8 files targeting `sm_80`, with 64-bit addresses, one two-parameter entry, two virtual `.b32` registers and `ret`. Only the array dimension spelling changes: `[6]`, `[2][3]`, or `[2+4]`. Syntax validity is grounded in the published PTX array grammar; an independent assembler was not run. ## Interface Other ## Earliest failure stage Native analysis or reconstruction ## Relevant output or evidence Selected fields from the observed results: ```text literal [6]: names=[payload,count]; sizes=[6,4]; parameter_count=2; parameter_bytes=12; payloads=complete multidimensional [2][3]: names=[null,count]; sizes=[null,4]; parameter_count=2; parameter_bytes=null; payloads=partial expression [2+4]: names=[payload,count]; sizes=[null,4]; parameter_count=2; parameter_bytes=null; payloads=partial All three: retained_text=true; process exit=0; stderr empty ``` The multidimensional and expression cases both add only this layout limitation: ```text Unsupported PTX parameter declarations or overflowing layouts leave parameter sizes/offsets and the total byte count unknown. ``` - Expected source: the published PTX grammar and pinned REA preservation contract linked above. - Failure source: the executed production source-service comparison and complete probe above. - Evidence provenance: observed service results; grammar applicability and parser diagnosis are separately source-derived. - Local verification: reproduced. - Reproduction completeness: complete for the static application-service/provider boundary; independent assembly, installed CLI and MCP were not executed. ## Typed outcome or recovery expected Return `name: "payload"` for the first parameter. Preserve the second parameter and exact retained text. If multidimensional layout stays unsupported, retain honest unknown layout fields, partial coverage and the existing limitation. No new error type is necessary. ## Analysis or suspected cause Source-derived: [PtxKernels.ts:193-215](https://github.com/morluto/rea/blob/8c4f96ef4759fd78da66c4b9726825ca71d7ee0c/src/gpu/PtxKernels.ts#L193-L215) scans backward from the last `]` and stops at its matching `[`. For `payload[2][3]`, this is the second dimension. It then chooses the preceding token, the first dimension's `]`, as the name candidate, producing null. The parser does not continue past all array suffixes to the actual identifier. This occurs before the separate supported-layout decision. ## Additional context The supplied open-issue corpus and live all-state searches covered `PTX multidimensional`, `PTX "array" "name"`, and `PTX "dimensions"`. No matching report was found. Related [#6664](https://github.com/morluto/rea/issues/6664) loses following parameter declarations because a parenthesized dimension prematurely terminates the parameter list. This fixture has no parentheses and retains both declarations. Closed [#6606](https://github.com/morluto/rea/issues/6606) concerns pointer qualifiers being mistaken for names; this fixture contains no pointer qualifier and reaches the independent array-suffix scan. NOT_ENOUGH_INFO: first affected source revision, independently assembled output and real-world frequency were not established. No fix or restoration was tested. Prepared with AI assistance; no tracked source was changed. ## Triage assessment - Impact level: P2 - Assessment status: supported - Impact basis: A valid multidimensional PTX parameter loses its identifier from structured kernel facts through the production source service, while one-dimensional and unknown-layout controls retain that identifier. The affected result is explicitly partial and the original text survives. - Workaround status: available - Workaround basis: The observed complete Evidence retains the exact source text including `payload[2][3]`, allowing manual identifier recovery without changing the artifact.
Sign in to join this conversation.
No milestone
No project
No assignees
1 participant
Notifications
Due date
The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set.

Reference
morluto/rea#4058
No description provided.