[GH-ISSUE #6613] [Bug] HIP v3 feature-suffixed bundle targets become unknown and cannot be selected #3998

New issue

Closed

opened 2026-10-31 00:45:55 +00:00 by bsakaguchi · 0 comments

bsakaguchi commented

2026-10-31 00:45:55 +00:00

Owner

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

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 production source-service probe recognizes a HIP code-object-v3 bundle entry named hip-amdgcn-amd-amdhsa--gfx906+xnack, but returns target: null. Its kernel is present without a target filter, then disappears with targets: ["gfx906+xnack"], where the object becomes not-selected. The same fixture shape with the no-feature gfx906 declaration retains its target and kernel both with and without the exact filter.

The shared service and provider ran with real files, input admission, hashing, static parsing, schema validation and Evidence construction. CLI/MCP transport, a legacy HIP compiler and GPU execution were not run.

Reproduction

With the source checkout and locked dependencies, save the following as $WORK/probe.mts in a disposable directory. From the checkout run node node_modules/tsx/dist/cli.mjs "$WORK/probe.mts" "$PWD" "$WORK". The probe imports the repository's existing ELF/bundle fixture builders, constructs a v3 MessagePack note, and inspects the two inert fixtures through the production shared service. It writes only beneath the supplied work directory. The final fixture form was inspected once per target/filter combination; earlier preliminary fixtures showed the same signal before their headers were tightened.

import {pathToFileURL} from 'node:url';
import {join,resolve} from 'node:path';
import {writeFile} from 'node:fs/promises';
import {createRequire} from 'node:module';
const repo=resolve(process.argv[2]),work=resolve(process.argv[3]);
const load=(p:string)=>import(pathToFileURL(join(repo,p)).href);
const {encode}=createRequire(join(repo,'package.json'))('@msgpack/msgpack');
const {elfFixture,hipFixture}=await load('tests/fixtures/gpu/codeObjects.ts');
const {GpuCodeObjectsService}=await load('src/application/gpu/GpuCodeObjectsService.ts');
const {GpuCodeObjectsProvider}=await load('src/gpu/GpuCodeObjectsProvider.ts');
const metadata=Buffer.from(encode({'amdhsa.version':[1,0],'amdhsa.kernels':[{
 '.name':'kernel','.symbol':'kernel.kd','.args':[],'.kernarg_segment_size':0,
 '.group_segment_fixed_size':0,'.private_segment_fixed_size':0,'.kernarg_segment_align':4,
 '.wavefront_size':64,'.sgpr_count':6,'.vgpr_count':2,'.max_flat_workgroup_size':256
}]}));
const note=Buffer.alloc(20+Math.ceil(metadata.length/4)*4);
note.writeUInt32LE(7,0);note.writeUInt32LE(metadata.length,4);note.writeUInt32LE(32,8);note.write('AMDGPU\0',12);metadata.copy(note,20);
const payload=elfFixture(224,[{name:'.note',type:7,bytes:note}]);
payload[8]=1;payload.writeUInt16LE(1,16);payload.writeUInt32LE(0x12f,48);
const service=new GpuCodeObjectsService(new GpuCodeObjectsProvider());
try {
 for(const target of ['gfx906','gfx906+xnack']){
  payload.writeUInt32LE(target.includes('+')?0x12f:0x2f,48);
  const bytes=hipFixture(payload,`hip-amdgcn-amd-amdhsa--${target}`),path=join(work,`${target}.hip`);
  await writeFile(path,bytes);
  for(const filter of [false,true]){
   const result=await service.inspect({path,analysis:'kernels',detail:'complete',...(filter?{targets:[target]}:{})});
   if(!result.ok)throw result.error;
   const r=result.value.normalized_result;
   console.log(JSON.stringify({target,filter,coverage:r.coverage,objects:r.objects.map(({kind,target,bundle_id,inspection,limitations}:any)=>({kind,target,bundle_id,inspection,limitations})),kernels:r.kernels.map(({name,parameter_count}:any)=>({name,parameter_count}))}));
  }
 }
}finally{await service.close();}

Expected behavior

Preserve the declared v3 target, including its feature suffix, and allow exact selection of that inventoried declaration. LLVM's published code-object-v2/v3 bundle-ID grammar explicitly uses a processor followed by +<target-feature> components. Its v3 metadata definition has no amdhsa.target requirement; that field was introduced by the v4 changes. REA documents Clang HIP bundle inspection, v3+ MessagePack metadata and target-based selection in docs/gpu-code-objects.md.

Impact and scope

The observed target declaration is lost and an exact selected-kernel request returns no kernels despite a readable metadata declaration. Both coverage fields correctly signal partial coverage for the affected input; this report does not claim silent complete coverage. The effect is bounded to the demonstrated supported v3 feature spelling. Current hipv4 colon-suffixed IDs and real legacy compiler prevalence were not tested. Omitting the target filter returns the kernel, and the raw bundle_id remains available for manual attribution; both were directly observed. No target bytes are executed or modified.

REA version

Package 6.7.0; source revision 21b836a32710a575357715a798beadd0f3976b8a. The source package and revision are the inspection baseline; no installed launcher was used.

Runtime and platform

Node.js v26.12.1, Linux x64; repository-locked tsx and MessagePack dependency. No Hopper, Ghidra, CUDA or ROCm runtime was used.

Target and artifact context

Synthetic Clang HIP binary bundles containing an ELF64 little-endian EM_AMDGPU relocatable object, HSA ABI version 1 (code object v3), gfx906 architecture flags, and one MessagePack kernel declaration with zero arguments. The feature case sets the documented v3 XNACK flag 0x100 and uses gfx906+xnack; the control omits that flag and feature. The probe is a static format-boundary fixture, not a loadable compiled kernel or a claim of GPU runtime validity.

Interface

Other

Earliest failure stage

Native analysis or reconstruction

Relevant output or evidence

Selected fields from the observed production-service output:

[
  {"target":"gfx906","filter":false,"object_target":"gfx906","inspection":"complete","kernel_names":["kernel"],"discovery":"complete","payloads":"complete"},
  {"target":"gfx906","filter":true,"object_target":"gfx906","inspection":"complete","kernel_names":["kernel"],"discovery":"complete","payloads":"complete"},
  {"target":"gfx906+xnack","filter":false,"object_target":null,"inspection":"complete","kernel_names":["kernel"],"discovery":"partial","payloads":"partial"},
  {"target":"gfx906+xnack","filter":true,"object_target":null,"inspection":"not-selected","kernel_names":[],"discovery":"partial","payloads":"partial"}
]

These are a field projection of the actual four JSON result lines, with the host entry omitted. Every result retains the original bundle_id; no object limitation explains the discarded feature spelling.

Evidence provenance: observed service results; grammar applicability and parser diagnosis are separately source-derived.
Local verification: reproduced.
Reproduction completeness: complete for the static format/service boundary; no legacy producer or public transport execution is claimed.

Typed outcome or recovery expected

The feature-suffixed object's target should remain selectable and retain kernel kernel. A restoration check should compare that exact declaration and its filtered result with the no-feature control, without renaming the target or replacing its artifact.

Analysis or suspected cause

Source-derived: src/gpu/ClangOffloadBundle.ts:65 recognizes the hip-amdgcn- entry kind, but its target expression at line 72 accepts an alphanumeric processor and optional colon suffix only. The plus suffix therefore leaves the target null. src/gpu/GpuInventory.ts:453 attempts note-based recovery, but src/gpu/GpuElfKernels.ts:402 only reads amdhsa.target, which is absent in the valid v3 map. The exact selection predicate in src/gpu/GpuInventory.ts:119 excludes the now-null target. This is reached through the service input schema and provider without a guard rejecting the fixture or request.

Additional context

The open snapshot's GPU/PTX/CUBIN issue titles and relevant bodies were inspected; no report of v3 plus-suffixed HIP target loss was found. Open and closed issues were searched before submission. NOT_ENOUGH_INFO: first affected release and real-world frequency were not investigated.

Triage assessment

Impact level: P2
Assessment status: supported
Impact basis: A documented HIP v3 target spelling loses its inventoried target and cannot select its readable kernel, while the otherwise comparable no-feature control succeeds.
Workaround status: available
Workaround basis: Omitting the exact target filter returned the kernel in the observed feature fixture, and the retained raw bundle ID permits manual target attribution without changing input bytes.

Originally created by @coygeek on GitHub (Oct 23, 2026). Original GitHub issue: https://github.com/morluto/rea/issues/6613 ## 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 production source-service probe recognizes a HIP code-object-v3 bundle entry named `hip-amdgcn-amd-amdhsa--gfx906+xnack`, but returns `target: null`. Its kernel is present without a target filter, then disappears with `targets: ["gfx906+xnack"]`, where the object becomes `not-selected`. The same fixture shape with the no-feature `gfx906` declaration retains its target and kernel both with and without the exact filter. The shared service and provider ran with real files, input admission, hashing, static parsing, schema validation and Evidence construction. CLI/MCP transport, a legacy HIP compiler and GPU execution were not run. ## Reproduction With the source checkout and locked dependencies, save the following as `$WORK/probe.mts` in a disposable directory. From the checkout run `node node_modules/tsx/dist/cli.mjs "$WORK/probe.mts" "$PWD" "$WORK"`. The probe imports the repository's existing ELF/bundle fixture builders, constructs a v3 MessagePack note, and inspects the two inert fixtures through the production shared service. It writes only beneath the supplied work directory. The final fixture form was inspected once per target/filter combination; earlier preliminary fixtures showed the same signal before their headers were tightened. ```ts import {pathToFileURL} from 'node:url'; import {join,resolve} from 'node:path'; import {writeFile} from 'node:fs/promises'; import {createRequire} from 'node:module'; const repo=resolve(process.argv[2]),work=resolve(process.argv[3]); const load=(p:string)=>import(pathToFileURL(join(repo,p)).href); const {encode}=createRequire(join(repo,'package.json'))('@msgpack/msgpack'); const {elfFixture,hipFixture}=await load('tests/fixtures/gpu/codeObjects.ts'); const {GpuCodeObjectsService}=await load('src/application/gpu/GpuCodeObjectsService.ts'); const {GpuCodeObjectsProvider}=await load('src/gpu/GpuCodeObjectsProvider.ts'); const metadata=Buffer.from(encode({'amdhsa.version':[1,0],'amdhsa.kernels':[{ '.name':'kernel','.symbol':'kernel.kd','.args':[],'.kernarg_segment_size':0, '.group_segment_fixed_size':0,'.private_segment_fixed_size':0,'.kernarg_segment_align':4, '.wavefront_size':64,'.sgpr_count':6,'.vgpr_count':2,'.max_flat_workgroup_size':256 }]})); const note=Buffer.alloc(20+Math.ceil(metadata.length/4)*4); note.writeUInt32LE(7,0);note.writeUInt32LE(metadata.length,4);note.writeUInt32LE(32,8);note.write('AMDGPU\0',12);metadata.copy(note,20); const payload=elfFixture(224,[{name:'.note',type:7,bytes:note}]); payload[8]=1;payload.writeUInt16LE(1,16);payload.writeUInt32LE(0x12f,48); const service=new GpuCodeObjectsService(new GpuCodeObjectsProvider()); try { for(const target of ['gfx906','gfx906+xnack']){ payload.writeUInt32LE(target.includes('+')?0x12f:0x2f,48); const bytes=hipFixture(payload,`hip-amdgcn-amd-amdhsa--${target}`),path=join(work,`${target}.hip`); await writeFile(path,bytes); for(const filter of [false,true]){ const result=await service.inspect({path,analysis:'kernels',detail:'complete',...(filter?{targets:[target]}:{})}); if(!result.ok)throw result.error; const r=result.value.normalized_result; console.log(JSON.stringify({target,filter,coverage:r.coverage,objects:r.objects.map(({kind,target,bundle_id,inspection,limitations}:any)=>({kind,target,bundle_id,inspection,limitations})),kernels:r.kernels.map(({name,parameter_count}:any)=>({name,parameter_count}))})); } } }finally{await service.close();} ``` ## Expected behavior Preserve the declared v3 target, including its feature suffix, and allow exact selection of that inventoried declaration. LLVM's [published code-object-v2/v3 bundle-ID grammar](https://releases.llvm.org/14.0.0/docs/AMDGPUUsage.html#code-object-v2-to-v3-target-id) explicitly uses a processor followed by `+<target-feature>` components. Its [v3 metadata definition](https://releases.llvm.org/14.0.0/docs/AMDGPUUsage.html#code-object-v3-metadata) has no `amdhsa.target` requirement; that field was introduced by the v4 changes. REA documents Clang HIP bundle inspection, v3+ MessagePack metadata and target-based selection in `docs/gpu-code-objects.md`. ## Impact and scope The observed target declaration is lost and an exact selected-kernel request returns no kernels despite a readable metadata declaration. Both coverage fields correctly signal partial coverage for the affected input; this report does not claim silent complete coverage. The effect is bounded to the demonstrated supported v3 feature spelling. Current `hipv4` colon-suffixed IDs and real legacy compiler prevalence were not tested. Omitting the target filter returns the kernel, and the raw `bundle_id` remains available for manual attribution; both were directly observed. No target bytes are executed or modified. ## REA version Package 6.7.0; source revision `21b836a32710a575357715a798beadd0f3976b8a`. The source package and revision are the inspection baseline; no installed launcher was used. ## Runtime and platform Node.js v26.12.1, Linux x64; repository-locked tsx and MessagePack dependency. No Hopper, Ghidra, CUDA or ROCm runtime was used. ## Target and artifact context Synthetic Clang HIP binary bundles containing an ELF64 little-endian EM_AMDGPU relocatable object, HSA ABI version 1 (code object v3), `gfx906` architecture flags, and one MessagePack kernel declaration with zero arguments. The feature case sets the documented v3 XNACK flag `0x100` and uses `gfx906+xnack`; the control omits that flag and feature. The probe is a static format-boundary fixture, not a loadable compiled kernel or a claim of GPU runtime validity. ## Interface Other ## Earliest failure stage Native analysis or reconstruction ## Relevant output or evidence Selected fields from the observed production-service output: ```json [ {"target":"gfx906","filter":false,"object_target":"gfx906","inspection":"complete","kernel_names":["kernel"],"discovery":"complete","payloads":"complete"}, {"target":"gfx906","filter":true,"object_target":"gfx906","inspection":"complete","kernel_names":["kernel"],"discovery":"complete","payloads":"complete"}, {"target":"gfx906+xnack","filter":false,"object_target":null,"inspection":"complete","kernel_names":["kernel"],"discovery":"partial","payloads":"partial"}, {"target":"gfx906+xnack","filter":true,"object_target":null,"inspection":"not-selected","kernel_names":[],"discovery":"partial","payloads":"partial"} ] ``` These are a field projection of the actual four JSON result lines, with the host entry omitted. Every result retains the original `bundle_id`; no object limitation explains the discarded feature spelling. - Evidence provenance: observed service results; grammar applicability and parser diagnosis are separately source-derived. - Local verification: reproduced. - Reproduction completeness: complete for the static format/service boundary; no legacy producer or public transport execution is claimed. ## Typed outcome or recovery expected The feature-suffixed object's target should remain selectable and retain kernel `kernel`. A restoration check should compare that exact declaration and its filtered result with the no-feature control, without renaming the target or replacing its artifact. ## Analysis or suspected cause Source-derived: `src/gpu/ClangOffloadBundle.ts:65` recognizes the `hip-amdgcn-` entry kind, but its target expression at line 72 accepts an alphanumeric processor and optional colon suffix only. The plus suffix therefore leaves the target null. `src/gpu/GpuInventory.ts:453` attempts note-based recovery, but `src/gpu/GpuElfKernels.ts:402` only reads `amdhsa.target`, which is absent in the valid v3 map. The exact selection predicate in `src/gpu/GpuInventory.ts:119` excludes the now-null target. This is reached through the service input schema and provider without a guard rejecting the fixture or request. ## Additional context The open snapshot's GPU/PTX/CUBIN issue titles and relevant bodies were inspected; no report of v3 plus-suffixed HIP target loss was found. Open and closed issues were searched before submission. NOT_ENOUGH_INFO: first affected release and real-world frequency were not investigated. ## Triage assessment - Impact level: P2 - Assessment status: supported - Impact basis: A documented HIP v3 target spelling loses its inventoried target and cannot select its readable kernel, while the otherwise comparable no-feature control succeeds. - Workaround status: available - Workaround basis: Omitting the exact target filter returned the kernel in the observed feature fixture, and the retained raw bundle ID permits manual target attribution without changing input bytes.

bsakaguchi closed this issue

2026-10-31 00:45:55 +00:00

bsakaguchi referenced this issue

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

[PR #3998] [MERGED] fix(javascript): retain static paths through TypeScript wrappers #5983