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
2026-10-25 22:08:11 -04:00 · 2026-10-24 01:24:34 +08:00
parent d137a4cdde
commit 39a91d4db9
6 changed files with 324 additions and 175 deletions
@@ -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
@@ -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",
 );

@@ -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]!;
@@ -55,7 +55,16 @@ function* tokens(text: string): Generator<Token> {
      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;
@@ -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",
+      });
  }
 });

@@ -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())),