[GH-ISSUE #6483] [Bug] omit AMDGPU target conflicts for LLVM's explicit-environment note spelling #3888

New issue

Closed

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

bsakaguchi commented

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

Owner

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

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

inspect-gpu-code-objects --analysis kernels does not identify the base-architecture conflict between a Clang bundle declaring gfx1100 and an embedded AMDGPU note declaring amdgcn-amd-amdhsa-unknown-gfx900. The CLI reports target_consistency: "unknown", object inspection: "complete", report coverage.payloads: "complete", complete target kernel_coverage, and no conflict limitation or kernel conflict annotation. A control changing only the synthetic note spelling to amdgcn-amd-amdhsa--gfx900 correctly reports conflict and partial coverage.

A freshly compiled LLVM gfx900 object emits the explicit unknown spelling and reproduces the same failure when placed in the conflicting bundle fixture. Both declarations remain visible in the result. unknown explicitly leaves agreement unresolved; it does not mean that the targets matched. The defect is the missing documented base-conflict handling and coverage downgrade.

Reproduction

Run from the source checkout with its pinned Node toolchain, locked dependencies, and compiled runtime. The following complete script creates synthetic controls, compiles a fresh OpenCL gfx900 object, wraps each payload in a one-entry Clang bundle, and invokes the public CLI. It also checks the target-summary projection. The producer replay used LLVM clang and llvm-readelf 23.1.3; put those executables on PATH, or set CLANG_BIN and LLVM_READELF_BIN to their executable paths. No GPU, ROCm libraries, target execution, external provider, or network is required. Child processes receive a clean environment with home/config/cache/data/state/temp directories inside the disposable work directory; the script removes that directory afterward.

node --input-type=module <<'NODE'
import fs from 'node:fs';
import path from 'node:path';
import os from 'node:os';
import { createRequire } from 'node:module';
import { spawnSync } from 'node:child_process';
const repo=process.cwd();
const work=fs.mkdtempSync(path.join(os.tmpdir(),'rea-amdgpu-'));
const clang=process.env.CLANG_BIN??'clang';
const readelf=process.env.LLVM_READELF_BIN??'llvm-readelf';
const {encode}=createRequire(path.join(repo,'package.json'))('@msgpack/msgpack');
try {
const env={PATH:process.env.PATH,HOME:path.join(work,'home'),TMPDIR:path.join(work,'tmp'),XDG_CONFIG_HOME:path.join(work,'config'),XDG_CACHE_HOME:path.join(work,'cache'),XDG_DATA_HOME:path.join(work,'data'),XDG_STATE_HOME:path.join(work,'state'),DSH_HOME:path.join(work,'dsh'),DSH_AGENTS_HOME:path.join(work,'dsh-agents')};
for(const [k,v] of Object.entries(env))if(k!=='PATH')fs.mkdirSync(v,{recursive:true});
function run(exe,args){const r=spawnSync(exe,args,{cwd:repo,env,encoding:'utf8'});if(r.status!==0)throw Error(JSON.stringify({exe,args,status:r.status,stderr:r.stderr,stdout:r.stdout}));return r.stdout;}
function object(target){
 const meta=Buffer.from(encode({'amdhsa.version':[1,2],'amdhsa.target':target,'amdhsa.kernels':[{'.name':'add_one','.symbol':'add_one.kd','.args':[],'.kernarg_segment_size':0,'.group_segment_fixed_size':0,'.private_segment_fixed_size':0,'.sgpr_count':10,'.vgpr_count':2}]}));
 const note=Buffer.alloc(20+Math.ceil(meta.length/4)*4);note.writeUInt32LE(7,0);note.writeUInt32LE(meta.length,4);note.writeUInt32LE(32,8);note.write('AMDGPU\0',12);meta.copy(note,20);
 const names=Buffer.from('\0.shstrtab\0.note\0');const noteOffset=Math.ceil((64+names.length)/8)*8;const table=Math.ceil((noteOffset+note.length)/8)*8;
 const b=Buffer.alloc(table+3*64);b.set([0x7f,69,76,70,2,1,1,0x40]);b.writeUInt16LE(1,16);b.writeUInt16LE(224,18);b.writeUInt32LE(1,20);b.writeBigUInt64LE(BigInt(table),40);b.writeUInt32LE(0x12c,48);b.writeUInt16LE(64,52);b.writeUInt16LE(64,58);b.writeUInt16LE(3,60);b.writeUInt16LE(1,62);names.copy(b,64);note.copy(b,noteOffset);
 for(const [i,name,type,off,size,align] of [[1,1,3,64,names.length,1],[2,11,7,noteOffset,note.length,4]]){const h=table+i*64;b.writeUInt32LE(name,h);b.writeUInt32LE(type,h+4);b.writeBigUInt64LE(BigInt(off),h+24);b.writeBigUInt64LE(BigInt(size),h+32);b.writeBigUInt64LE(BigInt(align),h+48);}return b;
}
function bundle(payload,target){const id=Buffer.from('hipv4-amdgcn-amd-amdhsa--'+target);const b=Buffer.alloc(56+id.length);b.write('__CLANG_OFFLOAD_BUNDLE__');b.writeBigUInt64LE(1n,24);b.writeBigUInt64LE(BigInt(b.length),32);b.writeBigUInt64LE(BigInt(payload.length),40);b.writeBigUInt64LE(BigInt(id.length),48);id.copy(b,56);return Buffer.concat([b,payload]);}
const results=[];
function inspect(label,bytes,target){const file=path.join(work,label+'.bundle');fs.writeFileSync(file,bundle(bytes,target));const stdout=run(process.execPath,['scripts/rea.mjs','inspect-gpu-code-objects',file,'--analysis','kernels','--format','json']);fs.writeFileSync(path.join(work,label+'.json'),stdout);const r=JSON.parse(stdout).normalized_result;const row=r.objects.find(x=>x.kind==='amdgpu');results.push({label,target:row.target,payload_target:row.payload_target,target_consistency:row.target_consistency,inspection:row.inspection,coverage:r.coverage,limitations:row.limitations,kernels:r.kernels.map(k=>({name:k.name,limitations:k.limitations})),targets:r.targets});}
for(const [label,note] of [['empty','amdgcn-amd-amdhsa--gfx900'],['unknown','amdgcn-amd-amdhsa-unknown-gfx900']])for(const target of ['gfx900','gfx1100'])inspect(label+'-'+target,object(note),target);
inspect('feature-unknown-control',object('amdgcn-amd-amdhsa--gfx900:xnack+'),'gfx900');
fs.writeFileSync(path.join(work,'kernel.cl'),'__kernel void add_one(__global int *out, int v) { out[0] = v + 1; }\n');
fs.writeFileSync(path.join(work,'clang-version.txt'),run(clang,['--version']));
run(clang,['-target','amdgcn-amd-amdhsa','-mcpu=gfx900','-x','cl','-cl-std=CL2.0','-nogpulib','-O2','-c',path.join(work,'kernel.cl'),'-o',path.join(work,'gfx900.o')]);
fs.writeFileSync(path.join(work,'llvm-notes.txt'),run(readelf,['-h','-n',path.join(work,'gfx900.o')]));
for(const target of ['gfx900','gfx1100'])inspect('real-'+target,fs.readFileSync(path.join(work,'gfx900.o')),target);
const standalone=path.join(work,'gfx900.o');const st=run(process.execPath,['scripts/rea.mjs','inspect-gpu-code-objects',standalone,'--analysis','kernels','--format','json']);fs.writeFileSync(path.join(work,'standalone.json'),st);
fs.writeFileSync(path.join(work,'results.json'),JSON.stringify(results,null,2));
for(const r of results) console.log(JSON.stringify({case:r.label,target:r.target,payload_target:r.payload_target,target_consistency:r.target_consistency,inspection:r.inspection,payloads:r.coverage.payloads,limitations:r.limitations}));
console.log(fs.readFileSync(path.join(work,'llvm-notes.txt'),'utf8').split('\n').find(line=>line.startsWith('amdhsa.target:')));
for(const label of ['empty-gfx1100','unknown-gfx1100','real-gfx1100']) {
 const text=run(process.execPath,['scripts/rea.mjs','inspect-gpu-code-objects',path.join(work,label+'.bundle'),'--analysis','kernels','--detail','summary','--view','{"collection":"targets"}','--evidence_output',path.join(work,label+'-evidence.json'),'--format','json']);
 const view=JSON.parse(text).summary.normalized_result;
 console.log(JSON.stringify({case:label,targets:view.items.map(row=>({target:row.target,kernel_coverage:row.kernel_coverage}))}));
}
} finally { fs.rmSync(work,{recursive:true,force:true}); }
NODE

Every inspection exited 0. The seven-case matrix reproduced once in the retained replay and once again with the exact inline script above. Relevant compact output is summarized below; unknown-gfx1100 and real-gfx1100 are the two mismatch failures, and empty-gfx1100 is the discriminating control.

Expected behavior

For the explicit-environment gfx900 note inside the gfx1100 bundle, retain both raw declarations and report target_consistency: "conflict", partial object/report/target kernel coverage, and the conflict limitation on the object and affected kernel. Matching exact gfx900 declarations should report consistent. Preserve unknown for genuinely unresolved feature comparisons, such as a plain gfx900 bundle versus a gfx900:xnack+ note; those are different from the demonstrated base-architecture conflict.

Impact and scope

This is a bounded integrity defect in the static consistency and coverage results for the tested explicit-environment target spelling. Kernel facts still parse, and both raw target declarations remain available. A caller can manually compare the retained base architectures; that workaround was checked for both mismatch cases and yields gfx1100 != gfx900. It does not establish general feature compatibility. No data mutation or runtime execution failure was observed, and no claim is made about every LLVM/HIP release or environment spelling.

REA version

6.7.0, source revision 6dff37594c

Runtime and platform

Node v24.18.0; macOS 27.0.1 arm64; Homebrew clang 23.1.3 and llvm-readelf 23.1.3 for fixture production and independent note inspection. This identifies the tested producer release and does not claim it is the newest LLVM release.

Target and artifact context

Synthetic ELF64 little-endian EM_AMDGPU objects containing an AMDGPU type-32 MessagePack note, plus a freshly LLVM-produced ELF64 relocatable gfx900 object carrying metadata version [1,2]. The standalone real object parses successfully with target_consistency: "single-source". The Clang bundle wrappers are constructed fixtures, including the intentionally conflicting outer declaration; they are not claimed to be hipcc-produced bundles.

Interface

CLI

Earliest failure stage

Native analysis or reconstruction

Relevant output or evidence

Expected source: docs/gpu-code-objects.md:23-33 requires independent raw declarations, normalized AMDGPU comparison, partial object/report/target coverage and kernel annotations for conflicting base architectures, while unresolved feature comparisons remain unknown. The released LLVM AMDGPU usage guide defines the triple's Environment component and the amdhsa.target triple-plus-target-ID syntax in its Target Triples and Code Object V4 Metadata sections; subsequent metadata versions inherit that target field. The exact explicit-unknown spelling is also directly established by the producer output below.
Failure source: Executed public CLI output from the complete inline reproduction. Source breadcrumbs are src/gpu/GpuElfKernels.ts:354-360 for note-owner normalization and src/gpu/GpuInventory.ts:391-435 / :471-498 for target comparison and coverage propagation; these explain the observed path without substituting for runtime evidence.
Evidence provenance: observed; the primary violation is independently replayed CLI output, with source-derived explanation identified below.
Local verification: reproduced.
Reproduction completeness: complete.

Case	Inventory target	Payload target	Consistency	Object inspection	Report payload coverage	Target kernel coverage
Synthetic empty-environment match	gfx900	amdgcn-amd-amdhsa--gfx900	consistent	complete	complete	not requested
Synthetic empty-environment mismatch	gfx1100	amdgcn-amd-amdhsa--gfx900	conflict	partial	partial	partial
Synthetic explicit-unknown match	gfx900	amdgcn-amd-amdhsa-unknown-gfx900	unknown	complete	complete	not requested
Synthetic explicit-unknown mismatch	gfx1100	amdgcn-amd-amdhsa-unknown-gfx900	unknown	complete	complete	complete
Real LLVM object match	gfx900	amdgcn-amd-amdhsa-unknown-gfx900	unknown	complete	complete	not requested
Real LLVM object mismatch	gfx1100	amdgcn-amd-amdhsa-unknown-gfx900	unknown	complete	complete	complete
Same-base feature control	gfx900	amdgcn-amd-amdhsa--gfx900:xnack+	unknown	complete	complete	not requested

The empty-environment mismatch has the following object limitation, also appended to its kernel. Both explicit-unknown mismatch cases have an empty object limitation list and no equivalent kernel conflict annotation; their ordinary static-metadata warning remains present.

Conflicting target declarations: inventory gfx1100; payload amdgcn-amd-amdhsa--gfx900. Target agreement is not established.

Independent llvm-readelf -h -n observation on the freshly compiled object:

Machine: EM_AMDGPU
Flags: 0x12c, gfx900, xnack
amdhsa.target:   amdgcn-amd-amdhsa-unknown-gfx900
amdhsa.version:
  - 1
  - 2

Restoration check: The synthetic and real explicit-unknown mismatches must produce the same base-conflict state, partial object/report/target kernel coverage, and conflict annotations as the empty-environment mismatch. Both exact gfx900 match cases must be consistent, while the same-base feature control must retain its unresolved agreement. Preserve the raw note and inventory values and the parsed kernel facts.

Analysis or suspected cause

Source-derived explanation: the note owner derives a comparison key only from the empty-environment spelling. The explicit-unknown spelling retains its raw value but yields a null key, so the inventory comparison remains unknown and returns no conflict to the coverage/kernel propagation path. The real object and its successful standalone inspection rule out an exclusively synthetic or unsupported-note explanation. This report does not require accepting arbitrary trailing gfx strings or changing intentional unresolved-feature semantics.

Additional context

Open and closed duplicate searches covered AMDGPU, amdhsa, target_consistency, GPU conflict, HIP target, and gfx900, plus the supplied full issue snapshot. Related issue #5345 is the completed initial static CUDA/HIP feature, and issue #3674 is its broader open CUDA analysis umbrella. Their bodies and comments do not report this target-normalization/conflict failure or its restoration.

Triage assessment

Impact level: P2
Assessment status: supported
Impact basis: The demonstrated supported base mismatch loses documented conflict, partial-coverage and kernel-annotation signals; this affects static analysis results while retaining the underlying declarations and kernel facts. No severe harm, data loss or execution failure was observed.
Workaround status: available
Workaround basis: The CLI retains inventory gfx1100 and payload amdgcn-amd-amdhsa-unknown-gfx900. Manual base-architecture comparison was executed on both mismatch results and detects gfx1100 versus gfx900 without modifying inputs or code; feature compatibility remains outside that workaround.

Originally created by @coygeek on GitHub (Oct 23, 2026). Original GitHub issue: https://github.com/morluto/rea/issues/6483 ## 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 `inspect-gpu-code-objects --analysis kernels` does not identify the base-architecture conflict between a Clang bundle declaring `gfx1100` and an embedded AMDGPU note declaring `amdgcn-amd-amdhsa-unknown-gfx900`. The CLI reports `target_consistency: "unknown"`, object `inspection: "complete"`, report `coverage.payloads: "complete"`, complete target `kernel_coverage`, and no conflict limitation or kernel conflict annotation. A control changing only the synthetic note spelling to `amdgcn-amd-amdhsa--gfx900` correctly reports `conflict` and partial coverage. A freshly compiled LLVM gfx900 object emits the explicit `unknown` spelling and reproduces the same failure when placed in the conflicting bundle fixture. Both declarations remain visible in the result. `unknown` explicitly leaves agreement unresolved; it does not mean that the targets matched. The defect is the missing documented base-conflict handling and coverage downgrade. ## Reproduction Run from the source checkout with its pinned Node toolchain, locked dependencies, and compiled runtime. The following complete script creates synthetic controls, compiles a fresh OpenCL gfx900 object, wraps each payload in a one-entry Clang bundle, and invokes the public CLI. It also checks the target-summary projection. The producer replay used LLVM clang and llvm-readelf 23.1.3; put those executables on PATH, or set `CLANG_BIN` and `LLVM_READELF_BIN` to their executable paths. No GPU, ROCm libraries, target execution, external provider, or network is required. Child processes receive a clean environment with home/config/cache/data/state/temp directories inside the disposable work directory; the script removes that directory afterward. ```sh node --input-type=module <<'NODE' import fs from 'node:fs'; import path from 'node:path'; import os from 'node:os'; import { createRequire } from 'node:module'; import { spawnSync } from 'node:child_process'; const repo=process.cwd(); const work=fs.mkdtempSync(path.join(os.tmpdir(),'rea-amdgpu-')); const clang=process.env.CLANG_BIN??'clang'; const readelf=process.env.LLVM_READELF_BIN??'llvm-readelf'; const {encode}=createRequire(path.join(repo,'package.json'))('@msgpack/msgpack'); try { const env={PATH:process.env.PATH,HOME:path.join(work,'home'),TMPDIR:path.join(work,'tmp'),XDG_CONFIG_HOME:path.join(work,'config'),XDG_CACHE_HOME:path.join(work,'cache'),XDG_DATA_HOME:path.join(work,'data'),XDG_STATE_HOME:path.join(work,'state'),DSH_HOME:path.join(work,'dsh'),DSH_AGENTS_HOME:path.join(work,'dsh-agents')}; for(const [k,v] of Object.entries(env))if(k!=='PATH')fs.mkdirSync(v,{recursive:true}); function run(exe,args){const r=spawnSync(exe,args,{cwd:repo,env,encoding:'utf8'});if(r.status!==0)throw Error(JSON.stringify({exe,args,status:r.status,stderr:r.stderr,stdout:r.stdout}));return r.stdout;} function object(target){ const meta=Buffer.from(encode({'amdhsa.version':[1,2],'amdhsa.target':target,'amdhsa.kernels':[{'.name':'add_one','.symbol':'add_one.kd','.args':[],'.kernarg_segment_size':0,'.group_segment_fixed_size':0,'.private_segment_fixed_size':0,'.sgpr_count':10,'.vgpr_count':2}]})); const note=Buffer.alloc(20+Math.ceil(meta.length/4)*4);note.writeUInt32LE(7,0);note.writeUInt32LE(meta.length,4);note.writeUInt32LE(32,8);note.write('AMDGPU\0',12);meta.copy(note,20); const names=Buffer.from('\0.shstrtab\0.note\0');const noteOffset=Math.ceil((64+names.length)/8)*8;const table=Math.ceil((noteOffset+note.length)/8)*8; const b=Buffer.alloc(table+3*64);b.set([0x7f,69,76,70,2,1,1,0x40]);b.writeUInt16LE(1,16);b.writeUInt16LE(224,18);b.writeUInt32LE(1,20);b.writeBigUInt64LE(BigInt(table),40);b.writeUInt32LE(0x12c,48);b.writeUInt16LE(64,52);b.writeUInt16LE(64,58);b.writeUInt16LE(3,60);b.writeUInt16LE(1,62);names.copy(b,64);note.copy(b,noteOffset); for(const [i,name,type,off,size,align] of [[1,1,3,64,names.length,1],[2,11,7,noteOffset,note.length,4]]){const h=table+i*64;b.writeUInt32LE(name,h);b.writeUInt32LE(type,h+4);b.writeBigUInt64LE(BigInt(off),h+24);b.writeBigUInt64LE(BigInt(size),h+32);b.writeBigUInt64LE(BigInt(align),h+48);}return b; } function bundle(payload,target){const id=Buffer.from('hipv4-amdgcn-amd-amdhsa--'+target);const b=Buffer.alloc(56+id.length);b.write('__CLANG_OFFLOAD_BUNDLE__');b.writeBigUInt64LE(1n,24);b.writeBigUInt64LE(BigInt(b.length),32);b.writeBigUInt64LE(BigInt(payload.length),40);b.writeBigUInt64LE(BigInt(id.length),48);id.copy(b,56);return Buffer.concat([b,payload]);} const results=[]; function inspect(label,bytes,target){const file=path.join(work,label+'.bundle');fs.writeFileSync(file,bundle(bytes,target));const stdout=run(process.execPath,['scripts/rea.mjs','inspect-gpu-code-objects',file,'--analysis','kernels','--format','json']);fs.writeFileSync(path.join(work,label+'.json'),stdout);const r=JSON.parse(stdout).normalized_result;const row=r.objects.find(x=>x.kind==='amdgpu');results.push({label,target:row.target,payload_target:row.payload_target,target_consistency:row.target_consistency,inspection:row.inspection,coverage:r.coverage,limitations:row.limitations,kernels:r.kernels.map(k=>({name:k.name,limitations:k.limitations})),targets:r.targets});} for(const [label,note] of [['empty','amdgcn-amd-amdhsa--gfx900'],['unknown','amdgcn-amd-amdhsa-unknown-gfx900']])for(const target of ['gfx900','gfx1100'])inspect(label+'-'+target,object(note),target); inspect('feature-unknown-control',object('amdgcn-amd-amdhsa--gfx900:xnack+'),'gfx900'); fs.writeFileSync(path.join(work,'kernel.cl'),'__kernel void add_one(__global int *out, int v) { out[0] = v + 1; }\n'); fs.writeFileSync(path.join(work,'clang-version.txt'),run(clang,['--version'])); run(clang,['-target','amdgcn-amd-amdhsa','-mcpu=gfx900','-x','cl','-cl-std=CL2.0','-nogpulib','-O2','-c',path.join(work,'kernel.cl'),'-o',path.join(work,'gfx900.o')]); fs.writeFileSync(path.join(work,'llvm-notes.txt'),run(readelf,['-h','-n',path.join(work,'gfx900.o')])); for(const target of ['gfx900','gfx1100'])inspect('real-'+target,fs.readFileSync(path.join(work,'gfx900.o')),target); const standalone=path.join(work,'gfx900.o');const st=run(process.execPath,['scripts/rea.mjs','inspect-gpu-code-objects',standalone,'--analysis','kernels','--format','json']);fs.writeFileSync(path.join(work,'standalone.json'),st); fs.writeFileSync(path.join(work,'results.json'),JSON.stringify(results,null,2)); for(const r of results) console.log(JSON.stringify({case:r.label,target:r.target,payload_target:r.payload_target,target_consistency:r.target_consistency,inspection:r.inspection,payloads:r.coverage.payloads,limitations:r.limitations})); console.log(fs.readFileSync(path.join(work,'llvm-notes.txt'),'utf8').split('\n').find(line=>line.startsWith('amdhsa.target:'))); for(const label of ['empty-gfx1100','unknown-gfx1100','real-gfx1100']) { const text=run(process.execPath,['scripts/rea.mjs','inspect-gpu-code-objects',path.join(work,label+'.bundle'),'--analysis','kernels','--detail','summary','--view','{"collection":"targets"}','--evidence_output',path.join(work,label+'-evidence.json'),'--format','json']); const view=JSON.parse(text).summary.normalized_result; console.log(JSON.stringify({case:label,targets:view.items.map(row=>({target:row.target,kernel_coverage:row.kernel_coverage}))})); } } finally { fs.rmSync(work,{recursive:true,force:true}); } NODE ``` Every inspection exited 0. The seven-case matrix reproduced once in the retained replay and once again with the exact inline script above. Relevant compact output is summarized below; `unknown-gfx1100` and `real-gfx1100` are the two mismatch failures, and `empty-gfx1100` is the discriminating control. ## Expected behavior For the explicit-environment `gfx900` note inside the `gfx1100` bundle, retain both raw declarations and report `target_consistency: "conflict"`, partial object/report/target kernel coverage, and the conflict limitation on the object and affected kernel. Matching exact `gfx900` declarations should report `consistent`. Preserve `unknown` for genuinely unresolved feature comparisons, such as a plain `gfx900` bundle versus a `gfx900:xnack+` note; those are different from the demonstrated base-architecture conflict. ## Impact and scope This is a bounded integrity defect in the static consistency and coverage results for the tested explicit-environment target spelling. Kernel facts still parse, and both raw target declarations remain available. A caller can manually compare the retained base architectures; that workaround was checked for both mismatch cases and yields `gfx1100 != gfx900`. It does not establish general feature compatibility. No data mutation or runtime execution failure was observed, and no claim is made about every LLVM/HIP release or environment spelling. ## REA version 6.7.0, source revision 6dff37594c0bc942c9748b34a7a028d223e88a6d ## Runtime and platform Node v24.18.0; macOS 27.0.1 arm64; Homebrew clang 23.1.3 and llvm-readelf 23.1.3 for fixture production and independent note inspection. This identifies the tested producer release and does not claim it is the newest LLVM release. ## Target and artifact context Synthetic ELF64 little-endian EM_AMDGPU objects containing an AMDGPU type-32 MessagePack note, plus a freshly LLVM-produced ELF64 relocatable gfx900 object carrying metadata version `[1,2]`. The standalone real object parses successfully with `target_consistency: "single-source"`. The Clang bundle wrappers are constructed fixtures, including the intentionally conflicting outer declaration; they are not claimed to be hipcc-produced bundles. ## Interface CLI ## Earliest failure stage Native analysis or reconstruction ## Relevant output or evidence - Expected source: `docs/gpu-code-objects.md:23-33` requires independent raw declarations, normalized AMDGPU comparison, partial object/report/target coverage and kernel annotations for conflicting base architectures, while unresolved feature comparisons remain unknown. The released [LLVM AMDGPU usage guide](https://github.com/llvm/llvm-project/blob/llvmorg-23.1.3/llvm/docs/AMDGPUUsage.rst) defines the triple's Environment component and the `amdhsa.target` triple-plus-target-ID syntax in its Target Triples and Code Object V4 Metadata sections; subsequent metadata versions inherit that target field. The exact explicit-unknown spelling is also directly established by the producer output below. - Failure source: Executed public CLI output from the complete inline reproduction. Source breadcrumbs are `src/gpu/GpuElfKernels.ts:354-360` for note-owner normalization and `src/gpu/GpuInventory.ts:391-435` / `:471-498` for target comparison and coverage propagation; these explain the observed path without substituting for runtime evidence. - Evidence provenance: observed; the primary violation is independently replayed CLI output, with source-derived explanation identified below. - Local verification: reproduced. - Reproduction completeness: complete. | Case | Inventory target | Payload target | Consistency | Object inspection | Report payload coverage | Target kernel coverage | | --- | --- | --- | --- | --- | --- | --- | | Synthetic empty-environment match | gfx900 | amdgcn-amd-amdhsa--gfx900 | consistent | complete | complete | not requested | | Synthetic empty-environment mismatch | gfx1100 | amdgcn-amd-amdhsa--gfx900 | conflict | partial | partial | partial | | Synthetic explicit-unknown match | gfx900 | amdgcn-amd-amdhsa-unknown-gfx900 | unknown | complete | complete | not requested | | Synthetic explicit-unknown mismatch | gfx1100 | amdgcn-amd-amdhsa-unknown-gfx900 | unknown | complete | complete | complete | | Real LLVM object match | gfx900 | amdgcn-amd-amdhsa-unknown-gfx900 | unknown | complete | complete | not requested | | Real LLVM object mismatch | gfx1100 | amdgcn-amd-amdhsa-unknown-gfx900 | unknown | complete | complete | complete | | Same-base feature control | gfx900 | amdgcn-amd-amdhsa--gfx900:xnack+ | unknown | complete | complete | not requested | The empty-environment mismatch has the following object limitation, also appended to its kernel. Both explicit-unknown mismatch cases have an empty object limitation list and no equivalent kernel conflict annotation; their ordinary static-metadata warning remains present. ```text Conflicting target declarations: inventory gfx1100; payload amdgcn-amd-amdhsa--gfx900. Target agreement is not established. ``` Independent `llvm-readelf -h -n` observation on the freshly compiled object: ```text Machine: EM_AMDGPU Flags: 0x12c, gfx900, xnack amdhsa.target: amdgcn-amd-amdhsa-unknown-gfx900 amdhsa.version: - 1 - 2 ``` - Restoration check: The synthetic and real explicit-unknown mismatches must produce the same base-conflict state, partial object/report/target kernel coverage, and conflict annotations as the empty-environment mismatch. Both exact gfx900 match cases must be consistent, while the same-base feature control must retain its unresolved agreement. Preserve the raw note and inventory values and the parsed kernel facts. ## Analysis or suspected cause Source-derived explanation: the note owner derives a comparison key only from the empty-environment spelling. The explicit-unknown spelling retains its raw value but yields a null key, so the inventory comparison remains `unknown` and returns no conflict to the coverage/kernel propagation path. The real object and its successful standalone inspection rule out an exclusively synthetic or unsupported-note explanation. This report does not require accepting arbitrary trailing `gfx` strings or changing intentional unresolved-feature semantics. ## Additional context Open and closed duplicate searches covered `AMDGPU`, `amdhsa`, `target_consistency`, `GPU conflict`, `HIP target`, and `gfx900`, plus the supplied full issue snapshot. Related [issue #5345](https://github.com/morluto/rea/issues/5345) is the completed initial static CUDA/HIP feature, and [issue #3674](https://github.com/morluto/rea/issues/3674) is its broader open CUDA analysis umbrella. Their bodies and comments do not report this target-normalization/conflict failure or its restoration. ## Triage assessment - Impact level: P2 - Assessment status: supported - Impact basis: The demonstrated supported base mismatch loses documented conflict, partial-coverage and kernel-annotation signals; this affects static analysis results while retaining the underlying declarations and kernel facts. No severe harm, data loss or execution failure was observed. - Workaround status: available - Workaround basis: The CLI retains inventory gfx1100 and payload amdgcn-amd-amdhsa-unknown-gfx900. Manual base-architecture comparison was executed on both mismatch results and detects gfx1100 versus gfx900 without modifying inputs or code; feature compatibility remains outside that workaround.

bsakaguchi closed this issue

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

bsakaguchi referenced this issue

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

[PR #4430] [MERGED] fix(firmware): reject staged files cut off at the per-file limit #6238