From 39a91d4db96ca13224c25936330ea84fe78970e5 Mon Sep 17 00:00:00 2001 From: morluto <76467478+morluto@users.noreply.github.com> Date: Sat, 24 Oct 2026 01:24:34 +0800 Subject: [PATCH] fix(gpu): preserve declaration grammar and partial resource coverage Share AMDGPU triple grammar between bundle and note owners so modern target spellings participate in processor/feature conflict detection. Propagate CUBIN register conflicts into object, target and report coverage. Normalize adjacent PTX dot keywords in the lexer, removing the downstream pointer splitter. Parse post-type alignment while distinguishing parameter and pointee layout, preserving names and byte provenance when layout remains unknown. Validated production-boundary matrices and public CLI/MCP coverage workflows with inert fixtures. Correct the HIP-v3 control's ELF processor declaration. Real GPU, assembler and compiler-produced modern-note behavior remains unverified. Fixes #6844 Fixes #6823 Fixes #6858 Fixes #6868 Fixes #6861 --- docs/gpu-code-objects.md | 19 +- src/gpu/AmdGpuTarget.ts | 6 +- src/gpu/GpuElfKernels.ts | 24 +- src/gpu/PtxKernels.ts | 53 ++-- tests/acceptance/gpuCodeObjectsCli.test.ts | 109 +++++--- tests/boundary/gpu/codeObjects.test.ts | 288 ++++++++++++++------- 6 files changed, 324 insertions(+), 175 deletions(-) diff --git a/docs/gpu-code-objects.md b/docs/gpu-code-objects.md index e3ca4aaa97..0171b1b0c4 100644 --- a/docs/gpu-code-objects.md +++ b/docs/gpu-code-objects.md @@ -32,10 +32,12 @@ unexamined payloads remain `unexamined`. An exact AMDGPU target is normalized at its note owner for comparison with a Clang bundle declaration; raw values remain intact. Unresolved feature comparisons stay `unknown`. -HIP bundle recognition accepts both the legacy `amdgcn` triple and LLVM's -`amdgpu` subarchitecture triples, including declared generic processors such -as `gfx11-generic` and `gfx9-4-generic`. CUDA's legacy fatbinary compression -flag is retained as an unsupported codec with its encoded and decoded extents; +HIP bundle and metadata note recognition accept both the legacy `amdgcn` triple +and LLVM's `amdgpu` subarchitecture triples, including declared generic processors +such as `gfx11-generic` and `gfx9-4-generic`. Notes with the empty or explicit +`unknown` environment use the same processor and feature comparison while +retaining their raw triple spelling. CUDA's legacy fatbinary compression flag is +retained as an unsupported codec with its encoded and decoded extents; compressed bytes are not treated as PTX text. CCOB is recognized by the exact four signature bytes `CCOB`; high-bit-corrupted bytes are not decoded as ASCII and accepted as a bundle. CCOB v1 has no stored total size, so a following @@ -100,7 +102,9 @@ Kernel representations have different meanings: parameters are excluded from entry virtual-register counts. Parameter names remain available when pointer qualifiers are understood or an unsupported qualifier prevents layout decoding. A pointer's `.align` describes the - pointed-to memory and does not change parameter layout alignment. + pointed-to memory and does not change parameter layout alignment. Spaced, + joined and partly joined pointer qualifiers have the same meaning. Parameter + alignment before or after the scalar type retains its layout and name. Parameter names and following declarations are retained when an array extent uses a parenthesized expression or when dimensions are multidimensional, even though those forms do not currently yield a layout. @@ -108,7 +112,8 @@ Kernel representations have different meanings: directive, including directives after entry declarations; unsupported or conflicting target profiles leave the effective target unknown. Scalar register declarations count each bare or percent-prefixed name and - each parameterized range. Unsupported declaration forms or overflowing totals + each parameterized range, including compact keywords such as `.reg.b32`. + Unsupported declaration forms or overflowing totals retain an unknown count and partial decoding coverage. Unsupported parameter layouts and unknown shared-memory declarations likewise retain their known facts, explicit limitations and partial coverage through retained pages. @@ -116,6 +121,8 @@ Kernel representations have different meanings: including V2 kernel-parameter sizes, and native section declarations. Global register declarations also cover compiler versions that leave the older text-section register field empty. + Conflicting register declarations keep the physical count unknown, preserve the + kernel's other facts, and mark object/report/target decoding coverage partial. - AMDGPU reads v3+ MessagePack notes. V3 processor identity comes from the `EF_AMDGPU_MACH` field in ELF `e_flags`; known processor values are selectable even when the note omits `amdhsa.target`. This inventory reports the processor diff --git a/src/gpu/AmdGpuTarget.ts b/src/gpu/AmdGpuTarget.ts index f29c5e9689..1b289aa3cd 100644 --- a/src/gpu/AmdGpuTarget.ts +++ b/src/gpu/AmdGpuTarget.ts @@ -2,11 +2,11 @@ const processor = String.raw`gfx[0-9a-z]+(?:-[0-9a-z]+)*`; const targetFeaturesSource = String.raw`(?:(?::[a-z0-9_-]+[+-])|(?:\+[a-z0-9_-]+))*`; const targetIdSource = `${processor}${targetFeaturesSource}`; const targetId = new RegExp(`^${targetIdSource}$`, "iu"); -const bundleTriple = - /^(?:hip|hipv4)-(?:amdgcn|amdgpu(?:[0-9]+(?:\.[0-9a-f]+)?)?)-amd-amdhsa-/iu; +const targetTripleSource = String.raw`(?:amdgcn|amdgpu(?:[0-9]+(?:\.[0-9a-f]+)?)?)-amd-amdhsa`; +const bundleTriple = new RegExp(`^(?:hip|hipv4)-${targetTripleSource}-`, "iu"); const bundleTarget = new RegExp(`(?:--|-)(${targetIdSource})$`, "iu"); const noteTarget = new RegExp( - `^amdgcn-amd-amdhsa(?:--|-unknown-)((${processor})${targetFeaturesSource})$`, + `^${targetTripleSource}(?:--|-unknown-)((${processor})${targetFeaturesSource})$`, "iu", ); diff --git a/src/gpu/GpuElfKernels.ts b/src/gpu/GpuElfKernels.ts index 968f4e5b46..c944a08f0d 100644 --- a/src/gpu/GpuElfKernels.ts +++ b/src/gpu/GpuElfKernels.ts @@ -201,7 +201,7 @@ const attachCudaResources = ( text: NativeSection, index: CudaMetadataIndex, symbolKey: string, -): void => { +): "complete" | "partial" => { const { source, directory } = index; const registerRows = index.registerFacts.get(symbolKey); const registerValues = new Set(registerRows?.keys()); @@ -222,12 +222,12 @@ const attachCudaResources = ( section_index: text.index, section_name: text.name, }; - if ( + const registerConflict = registerValues.size > 1 || (kernel.registers.physical !== null && sectionRegisters !== 0 && - kernel.registers.physical !== sectionRegisters) - ) { + kernel.registers.physical !== sectionRegisters); + if (registerConflict) { kernel.registers.physical = null; kernel.limitations.push( "Conflicting CUDA register declarations; physical count remains unknown", @@ -248,6 +248,7 @@ const attachCudaResources = ( section_index: shared[0]!.index, section_name: shared[0]!.name, }; + return registerConflict ? "partial" : "complete"; }; const cudaKernels = async ( @@ -314,12 +315,15 @@ const cudaKernels = async ( readElfString(names, symbol.name_offset), site(source, section, offset, width), ); - attachCudaResources( - kernel, - text!, - index, - `${section.index}:${offset / width}`, - ); + if ( + attachCudaResources( + kernel, + text!, + index, + `${section.index}:${offset / width}`, + ) === "partial" + ) + complete = false; const infos = index.infosByText.get(text!.index) ?? []; if (infos.length === 1) { const info = infos[0]!; diff --git a/src/gpu/PtxKernels.ts b/src/gpu/PtxKernels.ts index b1b1c7ccb6..0d43745096 100644 --- a/src/gpu/PtxKernels.ts +++ b/src/gpu/PtxKernels.ts @@ -55,7 +55,16 @@ function* tokens(text: string): Generator { continue; } if (/[.\w$%@]/u.test(text[at]!)) { - while (at < text.length && /[.\w$%@]/u.test(text[at]!)) advance(); + // Dot-prefixed keywords remain separate when their whitespace is omitted. + // Internal dots in names, instruction mnemonics and version literals stay intact. + const keyword = text[at] === "."; + advance(); + while ( + at < text.length && + /[.\w$%@]/u.test(text[at]!) && + (!keyword || text[at] !== ".") + ) + advance(); } else advance(); if (at - position > 65536) throw new BinaryResourceFailure( @@ -71,6 +80,12 @@ const integer = (token: string | undefined): number | null => { const value = Number(token); return Number.isSafeInteger(value) ? value : null; }; +const alignmentValue = (token: string | undefined): number | null => { + const value = integer(token); + return value !== null && value > 0 && Number.isInteger(Math.log2(value)) + ? value + : null; +}; type PtxVersion = { major: number; minor: number }; const ptxVersion = (token: string | undefined): PtxVersion | null => { const match = /^(\d+)\.(\d+)$/u.exec(token ?? ""); @@ -145,21 +160,10 @@ const declaration = ( size: number | null; align: number | null; } => { - const values = (source[0] === ".param" ? source.slice(1) : source).flatMap( - (value) => - value.startsWith(".ptr.") - ? [ - ".ptr", - ...value - .slice(5) - .split(".") - .map((part) => `.${part}`), - ] - : [value], - ); + const values = source[0] === ".param" ? source.slice(1) : source; let at = 0; - const explicit = values[at] === ".align"; - const alignment = explicit ? integer(values[++at]) : null; + let explicit = values[at] === ".align"; + let alignment = explicit ? alignmentValue(values[++at]) : null; if (explicit) at++; const vector = values[at] === ".v2" ? 2 : values[at] === ".v4" ? 4 : 1; if (vector !== 1) at++; @@ -177,17 +181,17 @@ const declaration = ( [".const", ".global", ".local", ".shared"].includes(qualifier) ) { // Address-space qualifiers describe the pointed-to object. - } else if (pointer && qualifier === ".align") { + } else if (qualifier === ".align") { const alignmentToken = values[at]; if (alignmentToken !== undefined && /^[0-9]+$/u.test(alignmentToken)) at++; - const pointeeAlignment = integer(alignmentToken); - if ( - pointeeAlignment === null || - pointeeAlignment <= 0 || - !Number.isInteger(Math.log2(pointeeAlignment)) - ) - layoutKnown = false; + const declaredAlignment = alignmentValue(alignmentToken); + if (declaredAlignment === null) layoutKnown = false; + if (!pointer) { + if (explicit) layoutKnown = false; + explicit = true; + alignment = declaredAlignment; + } } else layoutKnown = false; } const name = identifier.test(values[at] ?? "") ? values[at]! : null; @@ -221,8 +225,7 @@ const declaration = ( : vector === 1 ? scalar : `.v${vector} ${scalar}`; - const parameterAlignment = explicit ? alignment : size; - const align = parameterAlignment; + const align = explicit ? alignment : size; if (dimensions.length === 1) { const dimension = dimensions[0]!, count = dimension.length === 1 ? integer(dimension[0]) : null; diff --git a/tests/acceptance/gpuCodeObjectsCli.test.ts b/tests/acceptance/gpuCodeObjectsCli.test.ts index 0c219b7dc0..b9f5ea8165 100644 --- a/tests/acceptance/gpuCodeObjectsCli.test.ts +++ b/tests/acceptance/gpuCodeObjectsCli.test.ts @@ -246,10 +246,9 @@ cliTest( async ({ cli }) => { const target = "amdgcn-amd-amdhsa--gfx906+xnack", root = await createTestTempDirectory("rea-hip-v3-target-id-"), - bytes = hipFixture( - amdFixture(false, undefined, target), - "hip-amdgcn-amd-amdhsa--gfx906+xnack", - ), + payload = amdFixture(false, undefined, target); + payload.writeUInt32LE(0x2f, 48); // EF_AMDGPU_MACH_AMDGCN_GFX906 + const bytes = hipFixture(payload, "hip-amdgcn-amd-amdhsa--gfx906+xnack"), path = join(root, "feature.bundle"); await writeFile(path, bytes); const response = await cli.run({ @@ -281,17 +280,18 @@ cliTest( ); cliTest( - "keeps PTX pointer names and excludes call prototypes from resources", + "keeps PTX qualifier layouts and compact registers while excluding call prototypes", async ({ cli }) => { const source = ptxFixture .replace(".u64 output", ".u64 .ptr .global output") .replace( ".u32 value", - ".u64 .ptr .global .align 4 ptr_arg, .param .u64 .ptr .align 1 increment_param_0, .param .u64 .ptr.global.align 4 joined_arg, .param .u32 value", + ".u64 .ptr .global .align 4 ptr_arg, .param .u64 .ptr .align 1 increment_param_0, .param .u64 .ptr.global.align 4 joined_arg, .param .u64 .ptr .global.align 16 partly_joined_arg, .param .u64 .ptr.global .align 16 joined_space_arg, .param .b8 .align 8 payload[8], .param .u32 value", ) + .replace(".reg .b32", ".reg.b32") .replace( ".reg .b64 %rd<5>;", - ".reg .b64 %rd<5>; F: .callprototype _ (.reg .b32 _);", + ".reg .b64 %rd<5>; F: .callprototype _ (.reg.b32 _);", ), root = await createTestTempDirectory("rea-ptx-callprototype-"), path = join(root, "prototype.ptx"); @@ -315,13 +315,17 @@ cliTest( ); const report = gpuCodeObjectsSchema.parse(evidence.normalized_result); expect(report.kernels[0]).toMatchObject({ - parameter_count: 5, + parameter_count: 8, + parameter_bytes: 60, parameters: [ - { name: "output", size: 8 }, - { name: "ptr_arg", size: 8 }, - { name: "increment_param_0", size: 8 }, - { name: "joined_arg", size: 8 }, - { name: "value", size: 4 }, + { name: "output", offset: 0, size: 8 }, + { name: "ptr_arg", offset: 8, size: 8 }, + { name: "increment_param_0", offset: 16, size: 8 }, + { name: "joined_arg", offset: 24, size: 8 }, + { name: "partly_joined_arg", offset: 32, size: 8 }, + { name: "joined_space_arg", offset: 40, size: 8 }, + { name: "payload", offset: 48, size: 8 }, + { name: "value", offset: 56, size: 4 }, ], registers: { virtual: [ @@ -331,6 +335,7 @@ cliTest( }, }); expect(report.coverage.payloads).toBe("complete"); + expect(report.ptx_artifacts[0]!.text).toBe(source); }, ); @@ -696,36 +701,40 @@ cliTest( it("preserves incomplete decoder coverage through MCP target and object pages", async () => { const root = await createTestTempDirectory("rea-gpu-partial-"), { call } = await connectLocalToolsMcp(); - for (const missing of [ - "section-names", - "symbols", - "later-amd-note", - "target-conflict", - "ptx-parameter", - "ptx-shared", - ]) { - const bytes = - missing === "later-amd-note" - ? amdFixture(false, Buffer.alloc(3)) - : missing === "target-conflict" - ? cudaFixture("none", [ - { - kind: 2, - target: 80, - bytes: cubinFixture({ architecture: 90 }), - }, - ]) - : missing === "ptx-parameter" - ? Buffer.from( - ptxFixture.replace(".u64 output", ".u64 .vendor output"), - ) - : missing === "ptx-shared" - ? Buffer.from(ptxFixture.replace("scratch[32]", "scratch[]")) - : cubinFixture(), - path = join(root, missing); - if (missing === "section-names") bytes.writeUInt16LE(0, 62); - else if (missing === "symbols") + const cases = { + "section-names": () => { + const bytes = cubinFixture(); + bytes.writeUInt16LE(0, 62); + return bytes; + }, + symbols: () => { + const bytes = cubinFixture(); bytes.writeUInt32LE(1, Number(bytes.readBigUInt64LE(40)) + 3 * 64 + 4); + return bytes; + }, + "later-amd-note": () => amdFixture(false, Buffer.alloc(3)), + "modern-amd-target-conflict": () => + hipFixture( + amdFixture(false, undefined, "amdgpu9.00-amd-amdhsa--gfx900"), + "hipv4-amdgpu11.00-amd-amdhsa--gfx1100", + ), + "target-conflict": () => + cudaFixture("none", [ + { + kind: 2, + target: 80, + bytes: cubinFixture({ architecture: 90 }), + }, + ]), + "register-conflict": () => cubinFixture({ textRegisters: 11 }), + "ptx-parameter": () => + Buffer.from(ptxFixture.replace(".u64 output", ".u64 .vendor output")), + "ptx-shared": () => + Buffer.from(ptxFixture.replace("scratch[32]", "scratch[]")), + }; + for (const [missing, fixture] of Object.entries(cases)) { + const bytes = fixture(), + path = join(root, missing); await writeFile(path, bytes); const reply = await call("inspect_gpu_code_objects", { path, @@ -740,7 +749,11 @@ it("preserves incomplete decoder coverage through MCP target and object pages", expect(targets.items).toEqual([ expect.objectContaining({ target: - missing === "later-amd-note" ? "amdgcn-amd-amdhsa--gfx900" : "sm_80", + missing === "later-amd-note" + ? "amdgcn-amd-amdhsa--gfx900" + : missing === "modern-amd-target-conflict" + ? "gfx1100" + : "sm_80", kernel_count: missing === "symbols" ? 0 : 1, kernel_coverage: "partial", }), @@ -761,6 +774,9 @@ it("preserves incomplete decoder coverage through MCP target and object pages", parseEvidence(objectReply.structuredContent).normalized_result, ); expect(objects.items).toEqual([ + ...(missing === "modern-amd-target-conflict" + ? [expect.objectContaining({ kind: "host", inspection: "complete" })] + : []), expect.objectContaining({ inspection: missing === "later-amd-note" ? "malformed" : "partial", }), @@ -771,6 +787,13 @@ it("preserves incomplete decoder coverage through MCP target and object pages", payload_target: "sm_90", target_consistency: "conflict", }); + if (missing === "modern-amd-target-conflict") + expect(objects.items[1]).toMatchObject({ + kind: "amdgpu", + target: "gfx1100", + payload_target: "amdgpu9.00-amd-amdhsa--gfx900", + target_consistency: "conflict", + }); } }); diff --git a/tests/boundary/gpu/codeObjects.test.ts b/tests/boundary/gpu/codeObjects.test.ts index c7360c89f4..029cdb14b8 100644 --- a/tests/boundary/gpu/codeObjects.test.ts +++ b/tests/boundary/gpu/codeObjects.test.ts @@ -278,10 +278,17 @@ it("keeps AMDGPU processor identity unknown when ELF flags are unassigned", asyn inspection: "inventory", }); }); -it.each(["gfx900:xnack+", "gfx900+xnack", "gfx900:sramecc-:xnack+"])( - "preserves matching AMDGPU note features with ELF processor flags (%s)", - async (target) => { - const noteTarget = `amdgcn-amd-amdhsa--${target}`; +it.each([ + ["amdgcn", "gfx900:xnack+"], + ["amdgcn", "gfx900+xnack"], + ["amdgcn", "gfx900:sramecc-:xnack+"], + ["amdgpu", "gfx900:xnack+"], + ["amdgpu9", "gfx900:xnack+"], + ["amdgpu9.00", "gfx900:sramecc-:xnack+"], +])( + "preserves matching AMDGPU %s note features with ELF processor flags (%s)", + async (triple, target) => { + const noteTarget = `${triple}-amd-amdhsa--${target}`; const payload = amdFixture(false, undefined, noteTarget); payload.writeUInt32LE(0x12c, 48); const result = await inspect( @@ -301,62 +308,70 @@ it.each(["gfx900:xnack+", "gfx900+xnack", "gfx900:sramecc-:xnack+"])( expect(result.value.coverage.payloads).toBe("complete"); }, ); -it("does not choose between conflicting AMDGPU note and ELF processor identities", async () => { - const bytes = amdFixture( - false, - undefined, - "amdgcn-amd-amdhsa--gfx900:xnack+", - ); - bytes.writeUInt32LE(0x141, 48); // EF_AMDGPU_MACH gfx1100 - const result = await inspect(bytes); - if (!result.ok) throw result.error; - expect(result.value.objects[0]).toMatchObject({ - target: null, - target_consistency: "unknown", - inspection: "partial", - }); - expect(result.value.objects[0]!.limitations.join(" ")).toContain( - "processor flags conflict with target note metadata", - ); - expect(result.value.kernels).toHaveLength(1); -}); -it.each(["match", "conflict", "feature-unknown"] as const)( - "compares AMDGPU targets with an explicit unknown environment (%s)", - async (comparison) => { - const payloadTarget = - comparison === "feature-unknown" - ? "amdgcn-amd-amdhsa-unknown-gfx900:xnack+" - : "amdgcn-amd-amdhsa-unknown-gfx900"; - const inventoryTarget = comparison === "conflict" ? "gfx1100" : "gfx900"; - const bytes = hipFixture( - amdFixture(false, undefined, payloadTarget), - `hipv4-amdgcn-amd-amdhsa--${inventoryTarget}`, +it.each(["amdgcn", "amdgpu9.00"])( + "does not choose between conflicting AMDGPU %s note and ELF processor identities", + async (triple) => { + const bytes = amdFixture( + false, + undefined, + `${triple}-amd-amdhsa--gfx900:xnack+`, ); + bytes.writeUInt32LE(0x141, 48); // EF_AMDGPU_MACH gfx1100 const result = await inspect(bytes); if (!result.ok) throw result.error; - const row = result.value.objects.find((value) => value.kind === "amdgpu")!; - expect(row.payload_target).toBe(payloadTarget); + expect(result.value.objects[0]).toMatchObject({ + target: null, + target_consistency: "unknown", + inspection: "partial", + }); + expect(result.value.objects[0]!.limitations.join(" ")).toContain( + "processor flags conflict with target note metadata", + ); expect(result.value.kernels).toHaveLength(1); - if (comparison === "conflict") { - expect(row).toMatchObject({ - target: "gfx1100", - target_consistency: "conflict", - inspection: "partial", - }); - expect(result.value.coverage.payloads).toBe("partial"); - expect(result.value.kernels[0]!.limitations.join(" ")).toContain( - "Conflicting target declarations", - ); - } else { - expect(row).toMatchObject({ - target: "gfx900", - target_consistency: comparison === "match" ? "consistent" : "unknown", - inspection: "complete", - }); - expect(result.value.coverage.payloads).toBe("complete"); - expect(result.value.kernels[0]!.limitations.join(" ")).not.toContain( - "Conflicting target declarations", + }, +); +it.each([ + "amdgcn-amd-amdhsa-unknown-", + "amdgpu9.00-amd-amdhsa--", + "amdgpu9.00-amd-amdhsa-unknown-", +])( + "compares AMDGPU note processor and feature declarations with %s", + async (prefix) => { + for (const comparison of ["match", "conflict", "feature-unknown"]) { + const payloadTarget = `${prefix}gfx900${comparison === "feature-unknown" ? ":xnack+" : ""}`; + const inventoryTarget = comparison === "conflict" ? "gfx1100" : "gfx900"; + const bytes = hipFixture( + amdFixture(false, undefined, payloadTarget), + `hipv4-amdgcn-amd-amdhsa--${inventoryTarget}`, ); + const result = await inspect(bytes); + if (!result.ok) throw result.error; + const row = result.value.objects.find( + (value) => value.kind === "amdgpu", + )!; + expect(row.payload_target).toBe(payloadTarget); + expect(result.value.kernels).toHaveLength(1); + if (comparison === "conflict") { + expect(row).toMatchObject({ + target: "gfx1100", + target_consistency: "conflict", + inspection: "partial", + }); + expect(result.value.coverage.payloads).toBe("partial"); + expect(result.value.kernels[0]!.limitations.join(" ")).toContain( + "Conflicting target declarations", + ); + } else { + expect(row).toMatchObject({ + target: "gfx900", + target_consistency: comparison === "match" ? "consistent" : "unknown", + inspection: "complete", + }); + expect(result.value.coverage.payloads).toBe("complete"); + expect(result.value.kernels[0]!.limitations.join(" ")).not.toContain( + "Conflicting target declarations", + ); + } } }, ); @@ -482,6 +497,41 @@ it("retains unknown PTX virtual counts and partial payload coverage", async () = count: null, }); }); +it.each(["spaced", "compact", "mixed"])( + "counts %s PTX register keywords while retaining byte provenance", + async (spacing) => { + const directive = spacing === "spaced" ? ".reg " : ".reg"; + const header = ".version 8.8\n.target sm_80\n// GPU 寄存器 🎯\n"; + const entry = `.entry k() { ${directive}.b32 %r<4>; ${spacing === "mixed" ? ".reg " : directive}.b32 r0, r1, %r4; ${directive}.b64 %rd<2>; ret; }`; + const source = header + entry; + const result = await inspect(Buffer.from(source)); + if (!result.ok) throw result.error; + expect(result.value.objects[0]!.inspection).toBe("complete"); + expect(result.value.coverage.payloads).toBe("complete"); + expect(result.value.ptx_artifacts[0]!.text).toBe(source); + const kernel = result.value.kernels[0]!; + expect(kernel.registers.virtual).toEqual([ + { type: ".b32", count: 7 }, + { type: ".b64", count: 2 }, + ]); + expect(kernel.location).toMatchObject({ + offset: Buffer.byteLength(header), + size: Buffer.byteLength(entry), + }); + }, +); +it("keeps unsupported compact PTX register types partial", async () => { + const result = await inspect( + Buffer.from(ptxFixture.replace(".reg .b32", ".reg.v2.b32")), + ); + if (!result.ok) throw result.error; + expect(result.value.objects[0]!.inspection).toBe("partial"); + expect(result.value.coverage.payloads).toBe("partial"); + expect(result.value.kernels[0]!.registers.virtual).toEqual([ + { type: ".v2 .b32", count: null }, + { type: ".b64", count: 5 }, + ]); +}); it.each([64, 65])( "counts all %d scalar names in one PTX register list", async (count) => { @@ -587,24 +637,74 @@ it.each([ } }, ); -it("keeps supported PTX pointer qualifiers complete through the provider", async () => { - const bytes = Buffer.from( - ptxFixture.replace(".u64 output", ".u64 .ptr .global output"), - ); - const result = await inspect(bytes); - if (!result.ok) throw result.error; - expect(result.value.objects[0]!.inspection).toBe("complete"); - expect(result.value.coverage.payloads).toBe("complete"); - expect(result.value.ptx_artifacts[0]!.text).toBe(bytes.toString()); - expect(result.value.kernels[0]).toMatchObject({ - parameter_count: 2, - parameter_bytes: 12, - parameters: [ - { name: "output", offset: 0, size: 8 }, - { name: "value", offset: 8, size: 4 }, - ], - }); -}); +it.each([ + ".ptr .global", + ".ptr .global .align 16", + ".ptr.global.align 16", + ".ptr .global.align 16", + ".ptr.global .align 16", +])( + "keeps PTX pointer qualifiers %s complete through the provider", + async (qualifiers) => { + const bytes = Buffer.from( + ptxFixture.replace(".u64 output", `.u64 ${qualifiers} output`), + ); + const result = await inspect(bytes); + if (!result.ok) throw result.error; + expect(result.value.objects[0]!.inspection).toBe("complete"); + expect(result.value.coverage.payloads).toBe("complete"); + expect(result.value.ptx_artifacts[0]!.text).toBe(bytes.toString()); + expect(result.value.kernels[0]).toMatchObject({ + parameter_count: 2, + parameter_bytes: 12, + parameters: [ + { name: "output", offset: 0, size: 8 }, + { name: "value", offset: 8, size: 4 }, + ], + }); + }, +); +it.each([".align 8 .b8", ".b8 .align 8"])( + "retains PTX array parameter identity and layout with %s", + async (type) => { + const bytes = Buffer.from( + ptxFixture.replace(".u64 output", `${type} payload[8]`), + ); + const result = await inspect(bytes); + if (!result.ok) throw result.error; + expect(result.value.objects[0]!.inspection).toBe("complete"); + expect(result.value.coverage.payloads).toBe("complete"); + expect(result.value.ptx_artifacts[0]!.text).toBe(bytes.toString()); + expect(result.value.kernels[0]).toMatchObject({ + parameter_count: 2, + parameter_bytes: 12, + parameters: [ + { name: "payload", type: ".b8", offset: 0, size: 8 }, + { name: "value", offset: 8, size: 4 }, + ], + }); + }, +); +it.each([".b8 .align 3", ".align 8 .b8 .align 8"])( + "retains PTX names and partial layout for unsupported alignment %s", + async (type) => { + const bytes = Buffer.from( + ptxFixture.replace(".u64 output", `${type} payload[8]`), + ); + const result = await inspect(bytes); + if (!result.ok) throw result.error; + expect(result.value.objects[0]!.inspection).toBe("partial"); + expect(result.value.coverage.payloads).toBe("partial"); + expect(result.value.kernels[0]).toMatchObject({ + parameter_count: 2, + parameter_bytes: null, + parameters: [ + { name: "payload", type: ".b8", offset: null, size: null }, + { name: "value", offset: null, size: 4 }, + ], + }); + }, +); it.each(["whitespace", "comment"] as const)( "recognizes standalone PTX after more than 64 KiB of leading %s", async (kind) => { @@ -700,21 +800,33 @@ it.each(["section-names", "symbols"] as const)( ); }, ); -it("preserves parameters for long mangled names and rejects conflicting register counts", async () => { - const name = "kernel_" + "argument".repeat(64); - const result = await inspect( - cubinFixture({ name, textRegisters: 9, globalRegisters: 10 }), - ); - if (!result.ok) throw result.error; - expect(result.value.kernels[0]).toMatchObject({ - name, - parameter_count: 2, - registers: { physical: null }, - }); - expect(result.value.kernels[0]!.limitations.join(" ")).toContain( - "Conflicting", - ); -}); +it.each([9, 10])( + "preserves parameters and register coverage for text count %d", + async (textRegisters) => { + const name = "kernel_" + "argument".repeat(64); + const result = await inspect( + cubinFixture({ name, textRegisters, globalRegisters: 10 }), + ); + if (!result.ok) throw result.error; + const conflict = textRegisters !== 10; + expect(result.value.kernels[0]).toMatchObject({ + name, + parameter_count: 2, + registers: { physical: conflict ? null : 10 }, + }); + expect(result.value.objects[0]!.inspection).toBe( + conflict ? "partial" : "complete", + ); + expect(result.value.coverage.payloads).toBe( + conflict ? "partial" : "complete", + ); + expect( + result.value.kernels[0]!.limitations.some((value) => + value.includes("Conflicting"), + ), + ).toBe(conflict); + }, +); it("iterates flat concatenated fatbins beyond the genuine container nesting limit", async () => { const result = await inspect( Buffer.concat(Array.from({ length: 20 }, () => cudaFixture())),