morluto/rea

Fork 0

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

[GH-ISSUE #6651] [Bug] Entry-scoped PTX pragmas discard body resource declarations #4031

New issue

Closed

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

bsakaguchi commented

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

Owner

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

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

A valid entry-scoped .pragma "nounroll"; makes static PTX inspection stop the kernel declaration before its body. The returned k kernel retains its one input parameter but loses its three virtual .b32 registers, reports shared_bytes: null instead of 16, and has a source extent ending at the pragma semicolon. Object and payload coverage become partial. Removing only the header pragma, or moving it inside the body, returns the resources completely.

Reproduction

Run from the source checkout below with its installed locked dependencies and pinned Node on PATH. This is the executed shared-application-service probe, not a claimed CLI or MCP execution. The optional GPU_EVIDENCE_DIRECTORY selects where the disposable fixtures/results are retained; omit it to use /tmp.

node --import tsx --input-type=module <<'NODE'
import { mkdtemp, writeFile } from 'node:fs/promises';
import { join, resolve } from 'node:path';
import { pathToFileURL } from 'node:url';
import { tmpdir } from 'node:os';
const root = process.cwd();
const { GpuCodeObjectsProvider } = await import(pathToFileURL(resolve(root, 'src/gpu/GpuCodeObjectsProvider.ts')).href);
const { GpuCodeObjectsService } = await import(pathToFileURL(resolve(root, 'src/application/gpu/GpuCodeObjectsService.ts')).href);
const work = await mkdtemp(join(process.env.GPU_EVIDENCE_DIRECTORY ?? tmpdir(), 'ptx-pragma-'));
const preamble = '.version 8.8\n.target sm_80\n.address_size 64\n';
const body = '{ .reg .b32 %r<3>; .shared .align 4 .b8 scratch[16]; ret; }';
const cases = {
  plain: '.visible .entry k(.param .u32 input) ' + body,
  headerPragma: '.visible .entry k(.param .u32 input) .pragma "nounroll"; ' + body,
  bodyPragma: '.visible .entry k(.param .u32 input) { .pragma "nounroll"; .reg .b32 %r<3>; .shared .align 4 .b8 scratch[16]; ret; }',
  declarationOnly: '.extern .entry k(.param .u32 input);',
};
for (const [name, declaration] of Object.entries(cases)) {
  const path = join(work, name + '.ptx');
  const text = preamble + declaration + '\n';
  await writeFile(path, text);
  const service = new GpuCodeObjectsService(new GpuCodeObjectsProvider());
  try {
    const result = await service.inspect({ path, analysis: 'kernels', detail: 'complete' });
    if (!result.ok) throw result.error;
    const report = result.value.normalized_result;
    await writeFile(join(work, name + '.json'), JSON.stringify(result.value, null, 2));
    console.log(JSON.stringify({ name, inspection: report.objects[0].inspection, payloads: report.coverage.payloads, retainedTextMatches: report.ptx_artifacts[0].text === text, kernels: report.kernels.map(k => ({ name: k.name, parameter_count: k.parameter_count, parameter_bytes: k.parameter_bytes, virtual: k.registers.virtual, shared_bytes: k.shared_bytes, source_size: k.location.size, limitations: k.limitations })) }));
  } finally { await service.close(); }
}
NODE

Each of the four cases was inspected once. The process exited 0. No fixture code was executed.

Expected behavior

Retain the same register declarations and entry-local static shared bytes with an entry-scoped pragma as without it, and include the actual body in the kernel source extent. NVIDIA explicitly allows entry-scoped pragmas, including the .entry foo .pragma "nounroll"; { ... } spelling; pragma strings have no impact on PTX semantics. REA's documented PTX representation reports virtual registers and supported entry-local static shared declarations. The same body is already handled completely in the two controls.

Impact and scope

One valid standalone PTX definition lost structured resource declarations in the executed sample. The returned kernel remains present, its parameter count/bytes remain correct, and coverage is partial rather than falsely complete. The complete Evidence retained the exact original PTX text in all four cases, so manual reading of these small resource declarations is an available workaround. No embedded PTX, real compiler-produced artifact, physical resource allocation, CUDA compilation, or GPU execution was tested.

REA version

Source revision 23def04096869b85c4d3c14a73ed09168483c9f6; package metadata 6.7.0. The source service at that revision was executed. The installed published-package runtime was not tested.

Runtime and platform

Node.js v24.22.1; Linux 7.2.8-9-omarchy x86_64. No CUDA/ROCm compiler, driver, GPU, Hopper, or Ghidra was used.

Target and artifact context

Synthetic inert standalone UTF-8 PTX 8.8, .target sm_80, .address_size 64, one .u32 entry parameter, a .b32 register range of three, and a 16-byte static shared array. All fixtures are generated by the reproduction.

Interface

Other

Earliest failure stage

Native analysis or reconstruction

Relevant output or evidence

Selected fields from the observed service results:

plain:           complete / complete; virtual=[{type:".b32",count:3}]; shared_bytes=16; source_size=87
headerPragma:    partial  / partial;  virtual=[]; shared_bytes=null; source_size=47
bodyPragma:      complete / complete; virtual=[{type:".b32",count:3}]; shared_bytes=16; source_size=107
declarationOnly: partial  / partial;  virtual=[]; shared_bytes=null; source_size=28

The two state columns are object inspection and report coverage.payloads. Every result retained parameter_count: 1, parameter_bytes: 4, and the exact original PTX text. The header-pragma case reports Unsupported, dynamic or overflowing PTX shared declarations, or an absent entry body, leave entry-local static shared bytes unknown. Although the body is present, its resource facts are the same as the separate declaration-only control.

Expected source: NVIDIA's pragma grammar and REA's PTX representation linked above.
Failure source: executed GpuCodeObjectsService.inspect results and the inline fixtures; source-derived explanation below.
Evidence provenance: observed.
Local verification: reproduced.
Reproduction completeness: complete; validity is grounded in NVIDIA's published grammar and controls, not an assembler run.

Typed outcome or recovery expected

Return k with virtual .b32 count 3, shared_bytes: 16, and an extent through its closing brace, with complete inspection for this supported body. Keep genuine body-less external declarations distinguishable.

Analysis or suspected cause

Source-derived: the entry-header loop stops at the first semicolon even when it terminates a valid entry-scoped pragma. Body parsing only runs when that loop stops at {. The outer loop subsequently skips tokens within braces, so the real body does not recover the lost facts.

Additional context

Restoration check: the same four inputs should retain their current parameter facts and original text; headerPragma should additionally match plain for virtual-register count and shared bytes and extend through its actual body. declarationOnly must continue to represent its genuinely absent body as unknown rather than fabricate zero resource declarations.

Read-only all-state GitHub issue searches for PTX pragma returned no match; the broader PTX header results and the supplied open-issue bodies were also checked. Related open #6626 concerns .callprototype syntax inside the body, not the header; this fixture has no call prototype. NOT_ENOUGH_INFO: first affected release and real-producer prevalence were not investigated. Prepared with AI assistance.

Triage assessment

Impact level: P2
Assessment status: supported
Impact basis: A valid bounded PTX definition loses structured virtual-register and shared-memory declarations, while its ordinary parameter facts and exact source text survive and coverage becomes partial.
Workaround status: available
Workaround basis: Every observed result retains the exact original PTX text in complete Evidence, so an analyst can read the small .reg and .shared declarations manually without modifying artifact identity. Rewriting the pragma was only a control, not the workaround.

Originally created by @coygeek on GitHub (Oct 23, 2026). Original GitHub issue: https://github.com/morluto/rea/issues/6651 ## 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 A valid entry-scoped `.pragma "nounroll";` makes static PTX inspection stop the kernel declaration before its body. The returned `k` kernel retains its one input parameter but loses its three virtual `.b32` registers, reports `shared_bytes: null` instead of 16, and has a source extent ending at the pragma semicolon. Object and payload coverage become `partial`. Removing only the header pragma, or moving it inside the body, returns the resources completely. ## Reproduction Run from the source checkout below with its installed locked dependencies and pinned Node on `PATH`. This is the executed shared-application-service probe, not a claimed CLI or MCP execution. The optional `GPU_EVIDENCE_DIRECTORY` selects where the disposable fixtures/results are retained; omit it to use `/tmp`. ```sh node --import tsx --input-type=module <<'NODE' import { mkdtemp, writeFile } from 'node:fs/promises'; import { join, resolve } from 'node:path'; import { pathToFileURL } from 'node:url'; import { tmpdir } from 'node:os'; const root = process.cwd(); const { GpuCodeObjectsProvider } = await import(pathToFileURL(resolve(root, 'src/gpu/GpuCodeObjectsProvider.ts')).href); const { GpuCodeObjectsService } = await import(pathToFileURL(resolve(root, 'src/application/gpu/GpuCodeObjectsService.ts')).href); const work = await mkdtemp(join(process.env.GPU_EVIDENCE_DIRECTORY ?? tmpdir(), 'ptx-pragma-')); const preamble = '.version 8.8\n.target sm_80\n.address_size 64\n'; const body = '{ .reg .b32 %r<3>; .shared .align 4 .b8 scratch[16]; ret; }'; const cases = { plain: '.visible .entry k(.param .u32 input) ' + body, headerPragma: '.visible .entry k(.param .u32 input) .pragma "nounroll"; ' + body, bodyPragma: '.visible .entry k(.param .u32 input) { .pragma "nounroll"; .reg .b32 %r<3>; .shared .align 4 .b8 scratch[16]; ret; }', declarationOnly: '.extern .entry k(.param .u32 input);', }; for (const [name, declaration] of Object.entries(cases)) { const path = join(work, name + '.ptx'); const text = preamble + declaration + '\n'; await writeFile(path, text); const service = new GpuCodeObjectsService(new GpuCodeObjectsProvider()); try { const result = await service.inspect({ path, analysis: 'kernels', detail: 'complete' }); if (!result.ok) throw result.error; const report = result.value.normalized_result; await writeFile(join(work, name + '.json'), JSON.stringify(result.value, null, 2)); console.log(JSON.stringify({ name, inspection: report.objects[0].inspection, payloads: report.coverage.payloads, retainedTextMatches: report.ptx_artifacts[0].text === text, kernels: report.kernels.map(k => ({ name: k.name, parameter_count: k.parameter_count, parameter_bytes: k.parameter_bytes, virtual: k.registers.virtual, shared_bytes: k.shared_bytes, source_size: k.location.size, limitations: k.limitations })) })); } finally { await service.close(); } } NODE ``` Each of the four cases was inspected once. The process exited 0. No fixture code was executed. ## Expected behavior Retain the same register declarations and entry-local static shared bytes with an entry-scoped pragma as without it, and include the actual body in the kernel source extent. NVIDIA explicitly allows [entry-scoped pragmas](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#performance-tuning-directives-pragma), including the `.entry foo .pragma "nounroll"; { ... }` spelling; pragma strings have no impact on PTX semantics. REA's [documented PTX representation](https://github.com/morluto/rea/blob/23def04096869b85c4d3c14a73ed09168483c9f6/docs/gpu-code-objects.md#L79-L91) reports virtual registers and supported entry-local static shared declarations. The same body is already handled completely in the two controls. ## Impact and scope One valid standalone PTX definition lost structured resource declarations in the executed sample. The returned kernel remains present, its parameter count/bytes remain correct, and coverage is partial rather than falsely complete. The complete Evidence retained the exact original PTX text in all four cases, so manual reading of these small resource declarations is an available workaround. No embedded PTX, real compiler-produced artifact, physical resource allocation, CUDA compilation, or GPU execution was tested. ## REA version Source revision `23def04096869b85c4d3c14a73ed09168483c9f6`; package metadata 6.7.0. The source service at that revision was executed. The installed published-package runtime was not tested. ## Runtime and platform Node.js v24.22.1; Linux 7.2.8-9-omarchy x86_64. No CUDA/ROCm compiler, driver, GPU, Hopper, or Ghidra was used. ## Target and artifact context Synthetic inert standalone UTF-8 PTX 8.8, `.target sm_80`, `.address_size 64`, one `.u32` entry parameter, a `.b32` register range of three, and a 16-byte static shared array. All fixtures are generated by the reproduction. ## Interface Other ## Earliest failure stage Native analysis or reconstruction ## Relevant output or evidence Selected fields from the observed service results: ```text plain: complete / complete; virtual=[{type:".b32",count:3}]; shared_bytes=16; source_size=87 headerPragma: partial / partial; virtual=[]; shared_bytes=null; source_size=47 bodyPragma: complete / complete; virtual=[{type:".b32",count:3}]; shared_bytes=16; source_size=107 declarationOnly: partial / partial; virtual=[]; shared_bytes=null; source_size=28 ``` The two state columns are object `inspection` and report `coverage.payloads`. Every result retained `parameter_count: 1`, `parameter_bytes: 4`, and the exact original PTX text. The header-pragma case reports `Unsupported, dynamic or overflowing PTX shared declarations, or an absent entry body, leave entry-local static shared bytes unknown.` Although the body is present, its resource facts are the same as the separate declaration-only control. - Expected source: NVIDIA's pragma grammar and REA's PTX representation linked above. - Failure source: executed `GpuCodeObjectsService.inspect` results and the inline fixtures; source-derived explanation below. - Evidence provenance: observed. - Local verification: reproduced. - Reproduction completeness: complete; validity is grounded in NVIDIA's published grammar and controls, not an assembler run. ## Typed outcome or recovery expected Return `k` with virtual `.b32` count 3, `shared_bytes: 16`, and an extent through its closing brace, with complete inspection for this supported body. Keep genuine body-less external declarations distinguishable. ## Analysis or suspected cause Source-derived: [the entry-header loop](https://github.com/morluto/rea/blob/23def04096869b85c4d3c14a73ed09168483c9f6/src/gpu/PtxKernels.ts#L267-L294) stops at the first semicolon even when it terminates a valid entry-scoped pragma. Body parsing only runs when that loop stops at `{`. The outer loop subsequently [skips tokens within braces](https://github.com/morluto/rea/blob/23def04096869b85c4d3c14a73ed09168483c9f6/src/gpu/PtxKernels.ts#L207-L221), so the real body does not recover the lost facts. ## Additional context Restoration check: the same four inputs should retain their current parameter facts and original text; `headerPragma` should additionally match `plain` for virtual-register count and shared bytes and extend through its actual body. `declarationOnly` must continue to represent its genuinely absent body as unknown rather than fabricate zero resource declarations. Read-only all-state GitHub issue searches for `PTX pragma` returned no match; the broader `PTX header` results and the supplied open-issue bodies were also checked. Related open [#6626](https://github.com/morluto/rea/issues/6626) concerns `.callprototype` syntax inside the body, not the header; this fixture has no call prototype. NOT_ENOUGH_INFO: first affected release and real-producer prevalence were not investigated. Prepared with AI assistance. ## Triage assessment - Impact level: P2 - Assessment status: supported - Impact basis: A valid bounded PTX definition loses structured virtual-register and shared-memory declarations, while its ordinary parameter facts and exact source text survive and coverage becomes partial. - Workaround status: available - Workaround basis: Every observed result retains the exact original PTX text in complete Evidence, so an analyst can read the small .reg and .shared declarations manually without modifying artifact identity. Rewriting the pragma was only a control, not the workaround.

bsakaguchi closed this issue

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

bsakaguchi referenced this issue

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

[PR #4031] feat(artifacts): select extraction entries and output destination #6002

bsakaguchi referenced this issue

2026-11-01 21:10:57 +00:00

[GH-ISSUE #3796] [Tracker] Open feature requests by area (October 2026) #1471

bsakaguchi referenced this issue

2026-11-01 21:11:07 +00:00

[GH-ISSUE #5654] [Tracker] Open PRs that can't merge yet: conflicts and red CI (October 2026) #3425