[GH-ISSUE #6719] [Bug] Modern AMDGPU triple spellings make HIP bundles lose readable kernels #4087

New issue

Closed

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

bsakaguchi commented

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

Owner

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

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 service classifies a HIP entry named hipv4-amdgpu9.00-amd-amdhsa--gfx900 as kind: "unknown", despite retaining target gfx900. Kernel analysis marks it unsupported, reports partial payload coverage, and returns no kernels. An otherwise identical bundle named hipv4-amdgcn-amd-amdhsa--gfx900 returns kernel kernel with complete payload coverage. The modern-spelling entry also loses its kernel when enclosed in a valid CCOB v3 Zstd container.

Reproduction

Save the following complete fixture/service probe as $WORK/gpu-hip-triple-repro.mts in a disposable directory. From a prepared source checkout with locked dependencies, run the command below. The existing ELF/bundle builders generate an inert ELF64 AMDGPU payload with a v4 MessagePack kernel declaration; all three cases share exactly the same payload bytes. The script changes the bundle architecture spelling, then adds a CCOB v3 wrapper for the third case.

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,ccobFixture}=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,1],'amdhsa.target':'amdgcn-amd-amdhsa--gfx900','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);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]=2;
const service=new GpuCodeObjectsService(new GpuCodeObjectsProvider());
try{for(const [name,bytes] of [
 ['amdgcn',hipFixture(payload,'hipv4-amdgcn-amd-amdhsa--gfx900')],
 ['amdgpu',hipFixture(payload,'hipv4-amdgpu9.00-amd-amdhsa--gfx900')],
 ['ccob-amdgpu',ccobFixture(3,'zstd',hipFixture(payload,'hipv4-amdgpu9.00-amd-amdhsa--gfx900'))]
] as const){
 const path=join(work,`${name}.hip`);await writeFile(path,bytes);
 const result=await service.inspect({path,analysis:'kernels',detail:'complete'});if(!result.ok)throw result.error;
 const r=result.value.normalized_result;
 console.log(JSON.stringify({name,coverage:r.coverage,objects:r.objects.filter(({bundle_id}:any)=>bundle_id!==null&&!bundle_id.startsWith('host-')).map(({kind,bundle_id,target,inspection,limitations}:any)=>({kind,bundle_id,target,inspection,limitations})),kernels:r.kernels.map(({name}:any)=>name)}));
}}finally{await service.close();}

Command actually run:

node node_modules/tsx/dist/cli.mjs "$WORK/gpu-hip-triple-repro.mts" "$PWD" "$WORK"

The final standalone probe inspected each of its three inputs once and exited 0. A preliminary combined service probe produced the same three kernel-analysis outcomes before the standalone reproduction was written. No target code executes.

Expected behavior

Recognize the supported AMDGPU architecture spelling, retain the AMDGPU kind, and inspect the existing MessagePack kernel declaration. Published LLVM 23.1.0 target-triple documentation identifies amdgpu with a subarchitecture as the current architecture spelling, amdgcn as its legacy alias, and amdgpu9.00 as the subarchitecture for gfx900. Its bundle-entry grammar combines offload kind, target triple, and target ID. REA's docs/gpu-code-objects.md advertises Clang HIP bundles, CCOB v3, and AMDGPU v3+ MessagePack kernel metadata.

Impact and scope

A valid architecture spelling prevents the existing kernel reader from seeing otherwise readable metadata. Target and raw bundle ID survive, and payload coverage explicitly becomes partial. The demonstrated scope is uncompressed HIP binary bundles and a CCOB v3 Zstd wrapper; other subarchitectures, real producer frequency, runtime loadability, and practical workarounds were not tested. This is architecture-spelling recognition, not the separate generic processor-name or feature-suffix parsing defects.

REA version

Source main 8c4f96e, package version 6.7.0. Latest published tag rea-agents-6.7.0 contains no src/gpu tree. This is a main-only source baseline, not a demonstrated released-package regression.

Runtime and platform

Node.js v26.12.1, Linux x64, repository-locked tsx and MessagePack dependency. No GPU, HIP compiler, ROCm runtime, Hopper, Ghidra, built CLI, or MCP transport was used.

Target and artifact context

Self-made Clang HIP bundles containing the same 592-byte ELF64 little-endian EM_AMDGPU payload for gfx900, HSA ABI version 2/code object v4, and one zero-argument MessagePack kernel declaration. The newer bundle ID uses the documented amdgpu9.00 triple architecture. These are static format-boundary fixtures, not executable compiler-produced kernels or GPU-loadability evidence.

Interface

Other

Earliest failure stage

Native analysis or reconstruction

Relevant output or evidence

Selected fields from the actual source-service output:

{"name":"amdgcn","kind":"amdgpu","target":"gfx900","inspection":"complete","payloads":"complete","kernels":["kernel"]}
{"name":"amdgpu","kind":"unknown","target":"gfx900","inspection":"unsupported","payloads":"partial","kernels":[]}
{"name":"ccob-amdgpu","kind":"unknown","target":"gfx900","inspection":"unsupported","payloads":"partial","kernels":[]}

Both failing entries contain limitation Payload kind is outside kernel metadata coverage. The host placeholder is omitted from these projections; it remains complete. The CCOB wrapper itself decompresses successfully and is complete, isolating the failure to the enclosed entry's classification.

Evidence provenance: observed production-service results; architecture validity is separately source-derived from released LLVM documentation.
Local verification: reproduced.
Reproduction completeness: complete for the synthetic static format/service boundary. Compiler production, GPU loadability, and CLI/MCP transport were not run.

The complete request reaches GpuCodeObjectsService, its input schema, ordinary-file provider admission, hashing, inventory, and output validation. Public CLI/MCP dispatch uses that same shared service.

Typed outcome or recovery expected

All three cases should retain the same kernel facts, with the modern entry classified as AMDGPU. A proposed restoration check should compare the unchanged modern-spelling and legacy-alias fixtures, including the CCOB wrapper, rather than rewriting the input ID. No fix was made.

Analysis or suspected cause

Source-derived: ClangOffloadBundle.ts:65 recognizes HIP payloads only when the bundle ID begins with hip-amdgcn- or hipv4-amdgcn-. The valid modern architecture fails that regular expression. Target extraction still succeeds independently. GpuInventory.ts only dispatches cubin/amdgpu kinds to the existing ELF kernel reader and therefore classifies this payload as outside coverage. The source permalink is given in REA version.

Additional context

Related all-state issue bodies and the initial open snapshot were reviewed. Exact all-state search for amdgpu9.00 found no report. #6704 concerns hyphens inside generic processor names; #6613 concerns v3 plus-feature target suffixes; neither changes the bundle triple architecture used here. The submission checkbox stays unchecked for the filing agent's immediate duplicate check. Independent semantic review accepted this finding before this draft was written. No runtime recheck or report review followed draft creation.

Triage assessment

Impact level: P2
Assessment status: supported
Impact basis: Supported HIP and compressed-bundle paths lose a readable kernel solely because of a documented AMDGPU architecture spelling; target identity survives and incomplete payload coverage remains visible.
Workaround status: unverified
Workaround basis: No safe practical workflow for inspecting an existing affected artifact through REA was demonstrated. Rewriting its bundle ID would change the analyzed artifact, and the legacy-alias control establishes causality rather than a verified workaround.

Originally created by @coygeek on GitHub (Oct 23, 2026). Original GitHub issue: https://github.com/morluto/rea/issues/6719 ## 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 service classifies a HIP entry named `hipv4-amdgpu9.00-amd-amdhsa--gfx900` as `kind: "unknown"`, despite retaining target `gfx900`. Kernel analysis marks it `unsupported`, reports partial payload coverage, and returns no kernels. An otherwise identical bundle named `hipv4-amdgcn-amd-amdhsa--gfx900` returns kernel `kernel` with complete payload coverage. The modern-spelling entry also loses its kernel when enclosed in a valid CCOB v3 Zstd container. ## Reproduction Save the following complete fixture/service probe as `$WORK/gpu-hip-triple-repro.mts` in a disposable directory. From a prepared source checkout with locked dependencies, run the command below. The existing ELF/bundle builders generate an inert ELF64 AMDGPU payload with a v4 MessagePack kernel declaration; all three cases share exactly the same payload bytes. The script changes the bundle architecture spelling, then adds a CCOB v3 wrapper for the third case. ```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,ccobFixture}=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,1],'amdhsa.target':'amdgcn-amd-amdhsa--gfx900','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);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]=2; const service=new GpuCodeObjectsService(new GpuCodeObjectsProvider()); try{for(const [name,bytes] of [ ['amdgcn',hipFixture(payload,'hipv4-amdgcn-amd-amdhsa--gfx900')], ['amdgpu',hipFixture(payload,'hipv4-amdgpu9.00-amd-amdhsa--gfx900')], ['ccob-amdgpu',ccobFixture(3,'zstd',hipFixture(payload,'hipv4-amdgpu9.00-amd-amdhsa--gfx900'))] ] as const){ const path=join(work,`${name}.hip`);await writeFile(path,bytes); const result=await service.inspect({path,analysis:'kernels',detail:'complete'});if(!result.ok)throw result.error; const r=result.value.normalized_result; console.log(JSON.stringify({name,coverage:r.coverage,objects:r.objects.filter(({bundle_id}:any)=>bundle_id!==null&&!bundle_id.startsWith('host-')).map(({kind,bundle_id,target,inspection,limitations}:any)=>({kind,bundle_id,target,inspection,limitations})),kernels:r.kernels.map(({name}:any)=>name)})); }}finally{await service.close();} ``` Command actually run: ```sh node node_modules/tsx/dist/cli.mjs "$WORK/gpu-hip-triple-repro.mts" "$PWD" "$WORK" ``` The final standalone probe inspected each of its three inputs once and exited 0. A preliminary combined service probe produced the same three kernel-analysis outcomes before the standalone reproduction was written. No target code executes. ## Expected behavior Recognize the supported AMDGPU architecture spelling, retain the AMDGPU kind, and inspect the existing MessagePack kernel declaration. Published [LLVM 23.1.0 target-triple documentation](https://releases.llvm.org/23.1.0/docs/AMDGPUUsage.html#target-triples) identifies `amdgpu` with a subarchitecture as the current architecture spelling, `amdgcn` as its legacy alias, and `amdgpu9.00` as the subarchitecture for `gfx900`. Its [bundle-entry grammar](https://github.com/llvm/llvm-project/blob/llvmorg-23.1.0/clang/docs/ClangOffloadBundler.rst#bundle-entry-id) combines offload kind, target triple, and target ID. REA's `docs/gpu-code-objects.md` advertises Clang HIP bundles, CCOB v3, and AMDGPU v3+ MessagePack kernel metadata. ## Impact and scope A valid architecture spelling prevents the existing kernel reader from seeing otherwise readable metadata. Target and raw bundle ID survive, and payload coverage explicitly becomes partial. The demonstrated scope is uncompressed HIP binary bundles and a CCOB v3 Zstd wrapper; other subarchitectures, real producer frequency, runtime loadability, and practical workarounds were not tested. This is architecture-spelling recognition, not the separate generic processor-name or feature-suffix parsing defects. ## REA version Source main [8c4f96e](https://github.com/morluto/rea/blob/8c4f96ef4759fd78da66c4b9726825ca71d7ee0c/src/gpu/ClangOffloadBundle.ts#L60-L74), package version 6.7.0. Latest published tag `rea-agents-6.7.0` contains no `src/gpu` tree. This is a main-only source baseline, not a demonstrated released-package regression. ## Runtime and platform Node.js v26.12.1, Linux x64, repository-locked tsx and MessagePack dependency. No GPU, HIP compiler, ROCm runtime, Hopper, Ghidra, built CLI, or MCP transport was used. ## Target and artifact context Self-made Clang HIP bundles containing the same 592-byte ELF64 little-endian `EM_AMDGPU` payload for `gfx900`, HSA ABI version 2/code object v4, and one zero-argument MessagePack kernel declaration. The newer bundle ID uses the documented `amdgpu9.00` triple architecture. These are static format-boundary fixtures, not executable compiler-produced kernels or GPU-loadability evidence. ## Interface Other ## Earliest failure stage Native analysis or reconstruction ## Relevant output or evidence Selected fields from the actual source-service output: ```json {"name":"amdgcn","kind":"amdgpu","target":"gfx900","inspection":"complete","payloads":"complete","kernels":["kernel"]} {"name":"amdgpu","kind":"unknown","target":"gfx900","inspection":"unsupported","payloads":"partial","kernels":[]} {"name":"ccob-amdgpu","kind":"unknown","target":"gfx900","inspection":"unsupported","payloads":"partial","kernels":[]} ``` Both failing entries contain limitation `Payload kind is outside kernel metadata coverage`. The host placeholder is omitted from these projections; it remains complete. The CCOB wrapper itself decompresses successfully and is complete, isolating the failure to the enclosed entry's classification. - Evidence provenance: observed production-service results; architecture validity is separately source-derived from released LLVM documentation. - Local verification: reproduced. - Reproduction completeness: complete for the synthetic static format/service boundary. Compiler production, GPU loadability, and CLI/MCP transport were not run. The complete request reaches `GpuCodeObjectsService`, its input schema, ordinary-file provider admission, hashing, inventory, and output validation. Public CLI/MCP dispatch uses that same shared service. ## Typed outcome or recovery expected All three cases should retain the same kernel facts, with the modern entry classified as AMDGPU. A proposed restoration check should compare the unchanged modern-spelling and legacy-alias fixtures, including the CCOB wrapper, rather than rewriting the input ID. No fix was made. ## Analysis or suspected cause Source-derived: `ClangOffloadBundle.ts:65` recognizes HIP payloads only when the bundle ID begins with `hip-amdgcn-` or `hipv4-amdgcn-`. The valid modern architecture fails that regular expression. Target extraction still succeeds independently. `GpuInventory.ts` only dispatches `cubin`/`amdgpu` kinds to the existing ELF kernel reader and therefore classifies this payload as outside coverage. The source permalink is given in REA version. ## Additional context Related all-state issue bodies and the initial open snapshot were reviewed. Exact all-state search for `amdgpu9.00` found no report. #6704 concerns hyphens inside generic processor names; #6613 concerns v3 plus-feature target suffixes; neither changes the bundle triple architecture used here. The submission checkbox stays unchecked for the filing agent's immediate duplicate check. Independent semantic review accepted this finding before this draft was written. No runtime recheck or report review followed draft creation. ## Triage assessment - Impact level: P2 - Assessment status: supported - Impact basis: Supported HIP and compressed-bundle paths lose a readable kernel solely because of a documented AMDGPU architecture spelling; target identity survives and incomplete payload coverage remains visible. - Workaround status: unverified - Workaround basis: No safe practical workflow for inspecting an existing affected artifact through REA was demonstrated. Rewriting its bundle ID would change the analyzed artifact, and the legacy-alias control establishes causality rather than a verified workaround.

bsakaguchi closed this issue

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

bsakaguchi referenced this issue

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

[PR #4087] [CLOSED] fix(mcp): normalize advertised boolean literals for Gemini compatibility (#4025) #6047