[GH-ISSUE #6469] [Bug] A 65-name PTX register list marks the object malformed and drops kernel declarations #3879

New issue

Closed

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

bsakaguchi commented

2026-10-31 00:44:52 +00:00

Owner

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

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 public inspect-gpu-code-objects CLI reports a standalone PTX object as malformed and returns no kernels when its single entry declares 65 scalar .b32 registers in one comma-separated list. The command exits 0 with Malformed PTX resource declaration and partial payload coverage. A 64-name list works. A three-entry fixture retains the entry before the long list, but loses the affected entry and the entry after it. Original PTX text remains retained.

Reproduction

Prerequisites: run from the source checkout described under the version/environment field, with its pinned Node toolchain, locked dependencies, and compiled CLI already available. No install, build, CUDA compiler, driver, or target execution is part of this reproduction. The following executed block generates all inputs, gives each CLI child only an isolated HOME/USERPROFILE/XDG/TMPDIR environment plus PATH, and removes its temporary directory in finally. All client/provider override variables are absent from the child environment.

node --input-type=module <<'NODE'
import { mkdtempSync, mkdirSync, writeFileSync, rmSync } from 'node:fs';
import { tmpdir } from 'node:os';
import { join } from 'node:path';
import { spawnSync } from 'node:child_process';
const work = mkdtempSync(join(tmpdir(), 'gpu-register-repro-'));
const env = { PATH: process.env.PATH };
for (const [key, dir] of Object.entries({ HOME: 'home', USERPROFILE: 'home', XDG_CONFIG_HOME: 'config', XDG_CACHE_HOME: 'cache', XDG_DATA_HOME: 'data', XDG_STATE_HOME: 'state', XDG_RUNTIME_DIR: 'run', TMPDIR: 'tmp' })) {
  env[key] = join(work, dir);
  mkdirSync(env[key], { recursive: true, mode: 0o700 });
}
function cli(args) {
  const p = spawnSync(process.execPath, ['scripts/rea.mjs', ...args, '--json'], { env, encoding: 'utf8', timeout: 45000 });
  if (p.error || p.status !== 0) throw new Error(JSON.stringify({ status: p.status, error: p.error?.message, stderr: p.stderr }));
  return JSON.parse(p.stdout);
}
try {
  const names = n => Array.from({ length: n }, (_, i) => `%r${i}`).join(', ');
  const entry = (name, body) => `.visible .entry ${name}() { ${body} ret; }\n`;
  const fixtures = {
    regs64: entry('k', `.reg .b32 ${names(64)};`),
    regs65: entry('k', `.reg .b32 ${names(65)};`),
    regs66: entry('k', `.reg .b32 ${names(66)};`),
    split65: entry('k', `.reg .b32 ${names(64)}; .reg .b32 %r64;`),
    range65: entry('k', '.reg .b32 %r<65>;'),
    multi65: entry('before', '') + entry('affected', `.reg .b32 ${names(65)};`) + entry('after', ''),
  };
  for (const [name, body] of Object.entries(fixtures)) {
    const text = '.version 8.8\n.target sm_80\n.address_size 64\n' + body;
    const path = join(work, `${name}.ptx`);
    writeFileSync(path, text);
    const evidence = cli(['inspect-gpu-code-objects', path, '--analysis', 'kernels', '--detail', 'complete']);
    const r = evidence.normalized_result;
    console.log(JSON.stringify({ name, exit: 0, inspection: r.objects[0].inspection, limitations: r.objects[0].limitations, kernels: r.kernels.map(k => ({ name: k.name, virtual: k.registers.virtual })), payloads: r.coverage.payloads, retainedTextMatches: r.ptx_artifacts[0].text === text }));
    if (name === 'regs65') {
      for (const collection of ['ptx', 'kernels']) {
        const request = { source: { kind: 'inline', evidence }, view: { kind: 'gpu-code-objects', collection, ...(collection === 'ptx' ? { object_id: r.objects[0].id, offset: 0, limit: 4096 } : {}) } };
        const requestPath = join(work, `${collection}-view.json`);
        writeFileSync(requestPath, JSON.stringify(request));
        const page = cli(['inspect-analysis-view', requestPath]).normalized_result;
        console.log(JSON.stringify({ view: collection, exit: 0, ...(collection === 'ptx' ? { textMatches: page.ptx?.text === text } : {}), items: page.items }));
      }
    }
  }
} finally {
  rmSync(work, { recursive: true, force: true });
}
NODE

Expected behavior

Scalar register lists must retain the declared kernel and count each name. The 65-name input should produce one k kernel with virtual .b32 count 65 and complete payload coverage, just as the split-declaration and range controls do. A necessary resource refusal must preserve known facts and explicitly identify the resource limitation; declaration length alone must not turn valid PTX into malformed syntax or discard later entries.

Impact and scope

P2: bounded loss of structured kernel facts and incorrect format classification for the executed 65- and 66-name standalone lists. All observed failing runs returned partial coverage, so they did not claim complete analysis. The earlier entry survives in the three-entry case; not every kernel in every affected artifact is necessarily lost. The original PTX text survives in complete Evidence and the executed PTX analysis view returned that exact text, so an analyst can inspect/count the small declaration manually without changing artifact identity. The kernels analysis view remained empty: paging cannot recover structured rows that were never retained. Split declarations and the range form worked as controls, but rewriting a source artifact changes its bytes, digest, and coordinates and does not restore structured facts for the original artifact. Embedded containers and other declaration forms were not tested. No security boundary crossing was identified.

REA version

Source main at e32a0494d4; package version and node scripts/rea.mjs --version both 6.7.0. This used the checkout's existing compiled runtime; a separate source-service execution confirmed the same failure. The latest published tag is rea-agents-6.7.0, but a source-tree check found no GPU reader at that tag. Published-package runtime reproduction was not run, and no released-package failure is claimed.

Runtime and platform

Node.js v24.18.0; macOS 27.0.1 arm64; existing compiled source-checkout runtime. No CUDA/ROCm tool, GPU driver, or provider engine was invoked.

Target and artifact context

Generated standalone ASCII PTX with .version 8.8, .target sm_80, .address_size 64, and empty-parameter .visible .entry bodies ending in ret;. The fixtures and all required inputs are generated inline above.

Interface

CLI

Earliest failure stage

Native analysis or reconstruction

Relevant output or evidence

The executed block printed these fields for the boundary and controls; every command exited 0:

regs64: complete; k virtual .b32=64; payloads complete
regs65: malformed; "Malformed PTX resource declaration"; kernels []; payloads partial
regs66: malformed; "Malformed PTX resource declaration"; kernels []; payloads partial
split65: complete; k virtual .b32=65; payloads complete
range65: complete; k virtual .b32=65; payloads complete
multi65: malformed; kernels [before]; payloads partial
PTX view: original text matches; kernels view: items []

Contract evidence: docs/gpu-code-objects.md:81-88 promises per-name/range virtual-register counts and retained unknown facts with partial coverage for unsupported/overflowing declarations; its resource-limit discussion at lines 140-151 promises retained facts and explicit uncertainty, with no 128-token syntax allowance. NVIDIA's PTX ISA 8.8 specification, sections 3.1, 5.4.1, and 5.4.6, permits percent-prefixed identifiers, demonstrates comma-separated register declarations, and describes 100 virtual .b32 registers. Its instruction examples also mix percent-prefixed scalar names and ranges in a .reg .b32 list. The count 65 introduces no different syntax or platform register-allocation requirement; no compiler or target execution was performed.

The same input was also executed against the source service and returned the same malformed diagnostic and zero kernels. All 11 existing PTX parser tests passed with one worker; that result does not cover the long-list case.

Typed outcome or recovery expected

Retain the entry and its scalar register total. If decoding is genuinely limited, preserve the known entry declarations and mark the affected resource facts unknown with an explicit limitation; do not mislabel this list as malformed PTX.

Analysis or suspected cause

Source-derived: src/gpu/PtxKernels.ts:259-268 tests the collected resource-declaration token count against 128 before parsing individual register names. A 64-name declaration contributes 128 tokens including the type and separators; the 65-name case has 130 and fails at the final name. src/gpu/GpuInventory.ts:265-278 maps that format failure to malformed. The generator stops at the failing entry, explaining the observed loss of later entries. No implementation change is proposed.

Additional context

Restoration check: Rerun the same generated inputs without rewriting the failing artifact. The 65-name single-entry input should retain k with virtual .b32 count 65 and complete payload coverage, and the three-entry fixture should retain before, affected, and after in order with the affected count 65. The 64-name, split65, and range65 controls must keep their observed counts and complete coverage; the original PTX bytes must stay retained. If a real admission constraint prevents complete decoding, report that constraint and unknown resource facts explicitly while preserving the entries, rather than classifying this scalar list as malformed. Do not weaken existing malformed/truncated-input rejection to satisfy this check.

Read-only searches of open and closed issues found the related GPU feature requests #3674 and #5345, but no duplicate declaration-length failure. The documented scalar-counting behavior is already implemented on main; this report does not request SASS, GPU execution, or other future CUDA features. The last known good revision is NOT_ENOUGH_INFO; only the specified main baseline was runtime-tested. Prepared with AI assistance.

Originally created by @coygeek on GitHub (Oct 23, 2026). Original GitHub issue: https://github.com/morluto/rea/issues/6469 ## 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 public `inspect-gpu-code-objects` CLI reports a standalone PTX object as `malformed` and returns no kernels when its single entry declares 65 scalar `.b32` registers in one comma-separated list. The command exits 0 with `Malformed PTX resource declaration` and partial payload coverage. A 64-name list works. A three-entry fixture retains the entry before the long list, but loses the affected entry and the entry after it. Original PTX text remains retained. ## Reproduction Prerequisites: run from the source checkout described under the version/environment field, with its pinned Node toolchain, locked dependencies, and compiled CLI already available. No install, build, CUDA compiler, driver, or target execution is part of this reproduction. The following executed block generates all inputs, gives each CLI child only an isolated HOME/USERPROFILE/XDG/TMPDIR environment plus PATH, and removes its temporary directory in `finally`. All client/provider override variables are absent from the child environment. ```sh node --input-type=module <<'NODE' import { mkdtempSync, mkdirSync, writeFileSync, rmSync } from 'node:fs'; import { tmpdir } from 'node:os'; import { join } from 'node:path'; import { spawnSync } from 'node:child_process'; const work = mkdtempSync(join(tmpdir(), 'gpu-register-repro-')); const env = { PATH: process.env.PATH }; for (const [key, dir] of Object.entries({ HOME: 'home', USERPROFILE: 'home', XDG_CONFIG_HOME: 'config', XDG_CACHE_HOME: 'cache', XDG_DATA_HOME: 'data', XDG_STATE_HOME: 'state', XDG_RUNTIME_DIR: 'run', TMPDIR: 'tmp' })) { env[key] = join(work, dir); mkdirSync(env[key], { recursive: true, mode: 0o700 }); } function cli(args) { const p = spawnSync(process.execPath, ['scripts/rea.mjs', ...args, '--json'], { env, encoding: 'utf8', timeout: 45000 }); if (p.error || p.status !== 0) throw new Error(JSON.stringify({ status: p.status, error: p.error?.message, stderr: p.stderr })); return JSON.parse(p.stdout); } try { const names = n => Array.from({ length: n }, (_, i) => `%r${i}`).join(', '); const entry = (name, body) => `.visible .entry ${name}() { ${body} ret; }\n`; const fixtures = { regs64: entry('k', `.reg .b32 ${names(64)};`), regs65: entry('k', `.reg .b32 ${names(65)};`), regs66: entry('k', `.reg .b32 ${names(66)};`), split65: entry('k', `.reg .b32 ${names(64)}; .reg .b32 %r64;`), range65: entry('k', '.reg .b32 %r<65>;'), multi65: entry('before', '') + entry('affected', `.reg .b32 ${names(65)};`) + entry('after', ''), }; for (const [name, body] of Object.entries(fixtures)) { const text = '.version 8.8\n.target sm_80\n.address_size 64\n' + body; const path = join(work, `${name}.ptx`); writeFileSync(path, text); const evidence = cli(['inspect-gpu-code-objects', path, '--analysis', 'kernels', '--detail', 'complete']); const r = evidence.normalized_result; console.log(JSON.stringify({ name, exit: 0, inspection: r.objects[0].inspection, limitations: r.objects[0].limitations, kernels: r.kernels.map(k => ({ name: k.name, virtual: k.registers.virtual })), payloads: r.coverage.payloads, retainedTextMatches: r.ptx_artifacts[0].text === text })); if (name === 'regs65') { for (const collection of ['ptx', 'kernels']) { const request = { source: { kind: 'inline', evidence }, view: { kind: 'gpu-code-objects', collection, ...(collection === 'ptx' ? { object_id: r.objects[0].id, offset: 0, limit: 4096 } : {}) } }; const requestPath = join(work, `${collection}-view.json`); writeFileSync(requestPath, JSON.stringify(request)); const page = cli(['inspect-analysis-view', requestPath]).normalized_result; console.log(JSON.stringify({ view: collection, exit: 0, ...(collection === 'ptx' ? { textMatches: page.ptx?.text === text } : {}), items: page.items })); } } } } finally { rmSync(work, { recursive: true, force: true }); } NODE ``` ## Expected behavior Scalar register lists must retain the declared kernel and count each name. The 65-name input should produce one `k` kernel with virtual `.b32` count 65 and complete payload coverage, just as the split-declaration and range controls do. A necessary resource refusal must preserve known facts and explicitly identify the resource limitation; declaration length alone must not turn valid PTX into malformed syntax or discard later entries. ## Impact and scope P2: bounded loss of structured kernel facts and incorrect format classification for the executed 65- and 66-name standalone lists. All observed failing runs returned partial coverage, so they did not claim complete analysis. The earlier entry survives in the three-entry case; not every kernel in every affected artifact is necessarily lost. The original PTX text survives in complete Evidence and the executed PTX analysis view returned that exact text, so an analyst can inspect/count the small declaration manually without changing artifact identity. The kernels analysis view remained empty: paging cannot recover structured rows that were never retained. Split declarations and the range form worked as controls, but rewriting a source artifact changes its bytes, digest, and coordinates and does not restore structured facts for the original artifact. Embedded containers and other declaration forms were not tested. No security boundary crossing was identified. ## REA version Source main at e32a0494d42dedd5e79ca341d20c53f07479d9bb; package version and `node scripts/rea.mjs --version` both 6.7.0. This used the checkout's existing compiled runtime; a separate source-service execution confirmed the same failure. The latest published tag is `rea-agents-6.7.0`, but a source-tree check found no GPU reader at that tag. Published-package runtime reproduction was not run, and no released-package failure is claimed. ## Runtime and platform Node.js v24.18.0; macOS 27.0.1 arm64; existing compiled source-checkout runtime. No CUDA/ROCm tool, GPU driver, or provider engine was invoked. ## Target and artifact context Generated standalone ASCII PTX with `.version 8.8`, `.target sm_80`, `.address_size 64`, and empty-parameter `.visible .entry` bodies ending in `ret;`. The fixtures and all required inputs are generated inline above. ## Interface CLI ## Earliest failure stage Native analysis or reconstruction ## Relevant output or evidence The executed block printed these fields for the boundary and controls; every command exited 0: ```text regs64: complete; k virtual .b32=64; payloads complete regs65: malformed; "Malformed PTX resource declaration"; kernels []; payloads partial regs66: malformed; "Malformed PTX resource declaration"; kernels []; payloads partial split65: complete; k virtual .b32=65; payloads complete range65: complete; k virtual .b32=65; payloads complete multi65: malformed; kernels [before]; payloads partial PTX view: original text matches; kernels view: items [] ``` Contract evidence: `docs/gpu-code-objects.md:81-88` promises per-name/range virtual-register counts and retained unknown facts with partial coverage for unsupported/overflowing declarations; its resource-limit discussion at lines 140-151 promises retained facts and explicit uncertainty, with no 128-token syntax allowance. NVIDIA's [PTX ISA 8.8 specification](https://docs.nvidia.com/cuda/archive/12.9.1/parallel-thread-execution/index.html#variable-declarations), sections 3.1, 5.4.1, and 5.4.6, permits percent-prefixed identifiers, demonstrates comma-separated register declarations, and describes 100 virtual `.b32` registers. Its instruction examples also mix percent-prefixed scalar names and ranges in a `.reg .b32` list. The count 65 introduces no different syntax or platform register-allocation requirement; no compiler or target execution was performed. The same input was also executed against the source service and returned the same malformed diagnostic and zero kernels. All 11 existing PTX parser tests passed with one worker; that result does not cover the long-list case. ## Typed outcome or recovery expected Retain the entry and its scalar register total. If decoding is genuinely limited, preserve the known entry declarations and mark the affected resource facts unknown with an explicit limitation; do not mislabel this list as malformed PTX. ## Analysis or suspected cause Source-derived: `src/gpu/PtxKernels.ts:259-268` tests the collected resource-declaration token count against 128 before parsing individual register names. A 64-name declaration contributes 128 tokens including the type and separators; the 65-name case has 130 and fails at the final name. `src/gpu/GpuInventory.ts:265-278` maps that format failure to `malformed`. The generator stops at the failing entry, explaining the observed loss of later entries. No implementation change is proposed. ## Additional context Restoration check: Rerun the same generated inputs without rewriting the failing artifact. The 65-name single-entry input should retain `k` with virtual `.b32` count 65 and complete payload coverage, and the three-entry fixture should retain `before`, `affected`, and `after` in order with the affected count 65. The 64-name, split65, and range65 controls must keep their observed counts and complete coverage; the original PTX bytes must stay retained. If a real admission constraint prevents complete decoding, report that constraint and unknown resource facts explicitly while preserving the entries, rather than classifying this scalar list as malformed. Do not weaken existing malformed/truncated-input rejection to satisfy this check. Read-only searches of open and closed issues found the related GPU feature requests #3674 and #5345, but no duplicate declaration-length failure. The documented scalar-counting behavior is already implemented on main; this report does not request SASS, GPU execution, or other future CUDA features. The last known good revision is NOT_ENOUGH_INFO; only the specified main baseline was runtime-tested. Prepared with AI assistance.

bsakaguchi closed this issue

2026-10-31 00:44:52 +00:00

bsakaguchi referenced this issue

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

[PR #4014] [CLOSED] fix(browser): compare retained script source identities #5990