[GH-ISSUE #6664] [Bug] Parenthesized PTX array dimensions discard following kernel parameters #4044

New issue

Closed

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

bsakaguchi commented

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

Owner

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

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 application service reports one parameter for k(.param .b8 data[(4)], .param .u32 count), discarding the valid trailing count declaration. It returns parameter_count: 1 and only data, with partial payload coverage. The equivalent data[4] control returns both parameters with count 2. A data[2+2] control also preserves both declarations, even though its unsupported array layout correctly remains unknown. The failure concerns parameter-list structure and known trailing facts, independently of whether constant-expression sizes are evaluated.

Reproduction

From the source checkout with its locked dependencies, save this executed probe as $WORK/ptx-parameter-probe.mts, where $WORK is an existing disposable directory containing fixtures and evidence subdirectories:

import {writeFile} from 'node:fs/promises';
import {join,resolve} from 'node:path';
import {pathToFileURL} from 'node:url';
const repo=resolve(process.argv[2]),work=resolve(process.argv[3]);
const {GpuCodeObjectsProvider}=await import(pathToFileURL(join(repo,'src/gpu/GpuCodeObjectsProvider.ts')).href);
const {GpuCodeObjectsService}=await import(pathToFileURL(join(repo,'src/application/gpu/GpuCodeObjectsService.ts')).href);
const cases={control:'.param .b8 data[4], .param .u32 count', expression:'.param .b8 data[(4)], .param .u32 count',compound:'.param .b8 data[(2+2)], .param .u32 count',nonparenthesized:'.param .b8 data[2+2], .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){console.log(JSON.stringify({name,ok:false,error:result.error}));continue;}
 const r=result.value.normalized_result;
 console.log(JSON.stringify({name,ok:true,coverage:r.coverage,objects:r.objects.map(x=>({inspection:x.inspection,limitations:x.limitations})),kernels:r.kernels.map(x=>({name:x.name,parameter_count:x.parameter_count,parameters:x.parameters,parameter_bytes:x.parameter_bytes,registers:x.registers.virtual,limitations:x.limitations})),retained_text:r.ptx_artifacts[0]?.text===text}));
}}finally{await service.close();}

Executed command, with the original disposable directory represented by $WORK:

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

Each of the four fixtures was inspected once. Exit status was 0 and stderr was empty. The source service only reads the inert PTX files and constructs Evidence; no target code runs.

Expected behavior

Preserve both parameter declarations, including the known name and four-byte scalar type/size of count. Unsupported expression evaluation may leave the array size, later offsets and aggregate size unknown, but must not erase later parameters or report an incorrect known count. NVIDIA PTX ISA 8.8 permits constant expressions as array dimensions, and its constant-expression grammar includes parentheses. REA's PTX contract requires retaining known facts for unsupported parameter layouts. The nonparenthesized expression control demonstrates that this preservation is already supported.

Impact and scope

A valid small standalone PTX signature loses its trailing parameter in structured Evidence and reports a definite wrong count. Partial coverage and exact retained PTX text survive, so this is bounded metadata loss rather than an undetected complete result. Real compiler prevalence, embedded PTX and earlier versions were not tested. The retained source permits manual recovery of the missing declaration without changing artifact identity.

REA version

Package metadata 6.7.0; source revision 19520d00087bc92b5c1020918ea3b7c064768bb4. The source service was executed; installed published-package behavior was not tested.

Runtime and platform

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

Target and artifact context

Four synthetic UTF-8 PTX 8.8 files targeting sm_80 with 64-bit addresses. Each declares one entry, a four-byte byte array, one trailing .u32 parameter, two virtual .b32 registers, and ret. The cases vary only the array dimension spelling: 4, (4), (2+2) and 2+2. Validity is grounded in the published grammar; no independent assembler was run.

Interface

Other

Earliest failure stage

Native analysis or reconstruction

Relevant output or evidence

Selected fields from the observed service results:

control [4]:              complete; parameter_count=2; names=[data,count]; sizes=[4,4]; parameter_bytes=8
expression [(4)]:         partial;  parameter_count=1; names=[data];       sizes=[null]; parameter_bytes=null
compound [(2+2)]:         partial;  parameter_count=1; names=[data];       sizes=[null]; parameter_bytes=null
nonparenthesized [2+2]:   partial;  parameter_count=2; names=[data,count]; sizes=[null,4]; parameter_bytes=null

All four service results preserved the exact input PTX and the virtual register count 2. The expression cases only add the limitation Unsupported PTX parameter declarations or overflowing layouts leave parameter sizes/offsets and the total byte count unknown. No limitation says that following parameter declarations were omitted or that the returned count is uncertain.

Expected source: the NVIDIA grammar and pinned REA contract linked above.
Failure source: the executed source-service comparison and complete probe above.
Evidence provenance: observed.
Local verification: reproduced.
Reproduction completeness: complete at the source application-service/provider boundary; CLI/MCP transport and independent assembly were not run.

Typed outcome or recovery expected

Return both parameter records and count 2. If expression layout stays unsupported, retain unknown layout fields and partial coverage while preserving the trailing scalar declaration. A proposed restoration check uses the same four cases: parenthesized forms must retain the same names, count and scalar size as the nonparenthesized expression control; the ordinary literal control must remain complete.

Analysis or suspected cause

Source-derived: PtxKernels.ts:267-289 terminates the parameter list on the first ) without tracking nested parentheses inside array dimensions. The subsequent header loop consumes the remaining ], comma and trailing parameter until the body brace. The surviving parameter array then supplies the incorrect known count.

Additional context

Open/closed searches covered PTX, PTX parentheses, PTX "constant expression", and PTX "parameter list", plus the supplied open corpus. Related closed #6493 concerns entries with no parameter list; #6594 concerns PTX 1.3 body parameters. #6606 concerns pointer-qualified parameter names. The present fixtures contain an explicit modern list, no pointer attributes, no pragma and no call prototype. No matching report was found in the open and closed searches refreshed before submission. First affected release is NOT_ENOUGH_INFO. Prepared with AI assistance.

Triage assessment

Impact level: P2
Assessment status: supported
Impact basis: The production source service demonstrably omits a declared parameter and reports an incorrect known count for a valid small PTX signature; partial coverage and full source Evidence bound the consequence.
Workaround status: available
Workaround basis: The observed complete Evidence retains the exact input PTX including the missing scalar declaration, allowing manual signature recovery without modifying the input.

Originally created by @coygeek on GitHub (Oct 23, 2026). Original GitHub issue: https://github.com/morluto/rea/issues/6664 ## 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 application service reports one parameter for `k(.param .b8 data[(4)], .param .u32 count)`, discarding the valid trailing `count` declaration. It returns `parameter_count: 1` and only `data`, with partial payload coverage. The equivalent `data[4]` control returns both parameters with count 2. A `data[2+2]` control also preserves both declarations, even though its unsupported array layout correctly remains unknown. The failure concerns parameter-list structure and known trailing facts, independently of whether constant-expression sizes are evaluated. ## Reproduction From the source checkout with its locked dependencies, save this executed probe as `$WORK/ptx-parameter-probe.mts`, where `$WORK` is an existing disposable directory containing `fixtures` and `evidence` subdirectories: ```ts import {writeFile} from 'node:fs/promises'; import {join,resolve} from 'node:path'; import {pathToFileURL} from 'node:url'; const repo=resolve(process.argv[2]),work=resolve(process.argv[3]); const {GpuCodeObjectsProvider}=await import(pathToFileURL(join(repo,'src/gpu/GpuCodeObjectsProvider.ts')).href); const {GpuCodeObjectsService}=await import(pathToFileURL(join(repo,'src/application/gpu/GpuCodeObjectsService.ts')).href); const cases={control:'.param .b8 data[4], .param .u32 count', expression:'.param .b8 data[(4)], .param .u32 count',compound:'.param .b8 data[(2+2)], .param .u32 count',nonparenthesized:'.param .b8 data[2+2], .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){console.log(JSON.stringify({name,ok:false,error:result.error}));continue;} const r=result.value.normalized_result; console.log(JSON.stringify({name,ok:true,coverage:r.coverage,objects:r.objects.map(x=>({inspection:x.inspection,limitations:x.limitations})),kernels:r.kernels.map(x=>({name:x.name,parameter_count:x.parameter_count,parameters:x.parameters,parameter_bytes:x.parameter_bytes,registers:x.registers.virtual,limitations:x.limitations})),retained_text:r.ptx_artifacts[0]?.text===text})); }}finally{await service.close();} ``` Executed command, with the original disposable directory represented by `$WORK`: ```sh node --import ./node_modules/tsx/dist/loader.mjs "$WORK/ptx-parameter-probe.mts" "$PWD" "$WORK" > "$WORK/evidence/ptx-parameter-stdout.jsonl" 2> "$WORK/evidence/ptx-parameter-stderr.txt" ``` Each of the four fixtures was inspected once. Exit status was 0 and stderr was empty. The source service only reads the inert PTX files and constructs Evidence; no target code runs. ## Expected behavior Preserve both parameter declarations, including the known name and four-byte scalar type/size of `count`. Unsupported expression evaluation may leave the array size, later offsets and aggregate size unknown, but must not erase later parameters or report an incorrect known count. NVIDIA PTX ISA 8.8 permits [constant expressions as array dimensions](https://docs.nvidia.com/cuda/archive/12.9.1/parallel-thread-execution/index.html#array-declarations), and its [constant-expression grammar](https://docs.nvidia.com/cuda/archive/12.9.1/parallel-thread-execution/index.html#constant-expressions) includes parentheses. REA's [PTX contract](https://github.com/morluto/rea/blob/19520d00087bc92b5c1020918ea3b7c064768bb4/docs/gpu-code-objects.md#L79-L91) requires retaining known facts for unsupported parameter layouts. The nonparenthesized expression control demonstrates that this preservation is already supported. ## Impact and scope A valid small standalone PTX signature loses its trailing parameter in structured Evidence and reports a definite wrong count. Partial coverage and exact retained PTX text survive, so this is bounded metadata loss rather than an undetected complete result. Real compiler prevalence, embedded PTX and earlier versions were not tested. The retained source permits manual recovery of the missing declaration without changing artifact identity. ## REA version Package metadata `6.7.0`; source revision `19520d00087bc92b5c1020918ea3b7c064768bb4`. The source service was executed; installed published-package behavior was not tested. ## Runtime and platform Node.js v24.22.1; Linux 7.2.8-9-omarchy, x86_64; repository-locked tsx loader. No CUDA/ROCm compiler, driver, GPU, Hopper or Ghidra was used. ## Target and artifact context Four synthetic UTF-8 PTX 8.8 files targeting `sm_80` with 64-bit addresses. Each declares one entry, a four-byte byte array, one trailing `.u32` parameter, two virtual `.b32` registers, and `ret`. The cases vary only the array dimension spelling: `4`, `(4)`, `(2+2)` and `2+2`. Validity is grounded in the published grammar; no independent assembler was run. ## Interface Other ## Earliest failure stage Native analysis or reconstruction ## Relevant output or evidence Selected fields from the observed service results: ```text control [4]: complete; parameter_count=2; names=[data,count]; sizes=[4,4]; parameter_bytes=8 expression [(4)]: partial; parameter_count=1; names=[data]; sizes=[null]; parameter_bytes=null compound [(2+2)]: partial; parameter_count=1; names=[data]; sizes=[null]; parameter_bytes=null nonparenthesized [2+2]: partial; parameter_count=2; names=[data,count]; sizes=[null,4]; parameter_bytes=null ``` All four service results preserved the exact input PTX and the virtual register count 2. The expression cases only add the limitation `Unsupported PTX parameter declarations or overflowing layouts leave parameter sizes/offsets and the total byte count unknown.` No limitation says that following parameter declarations were omitted or that the returned count is uncertain. - Expected source: the NVIDIA grammar and pinned REA contract linked above. - Failure source: the executed source-service comparison and complete probe above. - Evidence provenance: observed. - Local verification: reproduced. - Reproduction completeness: complete at the source application-service/provider boundary; CLI/MCP transport and independent assembly were not run. ## Typed outcome or recovery expected Return both parameter records and count 2. If expression layout stays unsupported, retain unknown layout fields and partial coverage while preserving the trailing scalar declaration. A proposed restoration check uses the same four cases: parenthesized forms must retain the same names, count and scalar size as the nonparenthesized expression control; the ordinary literal control must remain complete. ## Analysis or suspected cause Source-derived: [PtxKernels.ts:267-289](https://github.com/morluto/rea/blob/19520d00087bc92b5c1020918ea3b7c064768bb4/src/gpu/PtxKernels.ts#L267-L289) terminates the parameter list on the first `)` without tracking nested parentheses inside array dimensions. The subsequent header loop consumes the remaining `]`, comma and trailing parameter until the body brace. The surviving parameter array then supplies the incorrect known count. ## Additional context Open/closed searches covered `PTX`, `PTX parentheses`, `PTX "constant expression"`, and `PTX "parameter list"`, plus the supplied open corpus. Related closed #6493 concerns entries with no parameter list; #6594 concerns PTX 1.3 body parameters. #6606 concerns pointer-qualified parameter names. The present fixtures contain an explicit modern list, no pointer attributes, no pragma and no call prototype. No matching report was found in the open and closed searches refreshed before submission. First affected release is NOT_ENOUGH_INFO. Prepared with AI assistance. ## Triage assessment - Impact level: P2 - Assessment status: supported - Impact basis: The production source service demonstrably omits a declared parameter and reports an incorrect known count for a valid small PTX signature; partial coverage and full source Evidence bound the consequence. - Workaround status: available - Workaround basis: The observed complete Evidence retains the exact input PTX including the missing scalar declaration, allowing manual signature recovery without modifying the input.

bsakaguchi closed this issue

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

bsakaguchi referenced this issue

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

[PR #4044] [CLOSED] fix(rizin): preserve trailing whitespace before the completion marker #6014