fix(gpu): preserve parameter and target declarations

Decode observed V2 CUDA parameters, retain qualified PTX parameter names and entry resource counts, and preserve family and feature target declarations through selection and consistency checks.

Fixes #6624
Fixes #6625
Fixes #6626
Fixes #6618
Fixes #6613
Fixes #6606

docs/gpu-code-objects.md

@@ -27,7 +27,7 @@ selected fields.
 Conflicting base architectures retain both values, mark object/report/target
 coverage partial and annotate the affected kernels. CUBIN flags establish a base
 architecture: `base-only` explicitly leaves feature-profile agreement unknown,
-including `a` variants. A standalone declaration is `single-source`, and
+including `a` and `f` variants. A standalone declaration is `single-source`, and
 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`.
@@ -43,8 +43,8 @@ target remains unexamined; an empty filtered result does not establish that
 no enclosed code declares that target. Discovering every inner declaration
 requires decoding those payloads.
 
-HIP IDs with one or two separators before the architecture retain the declared
-target. If a recognized AMDGPU bundle ID does not expose a target, the same
+HIP bundle IDs retain the declared processor and v2-v3 `+feature` suffixes for
+exact selection. If a recognized AMDGPU bundle ID does not expose a target, the same
 bounded note-discovery owner used for standalone AMDGPU files establishes it
 before filtering. Those directory/note reads share the inspection work budget;
 successive queries can select the target returned by inventory.
@@ -83,14 +83,20 @@ Kernel representations have different meanings:
   shared memory, dynamic allocation or physical device resources.
   PTX 1.0–1.3 declares entry parameters in the body; PTX 1.4 and later uses the
   entry parameter list. A `.target compute_xx` spelling is retained as written;
-  NVIDIA PTX defines it as a synonym for `.target sm_xx`.
+  NVIDIA PTX defines it as a synonym for `.target sm_xx` when comparing its
+  target identity with the fatbin declaration. `.callprototype` formal
+  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.
   Scalar register declarations count each bare or percent-prefixed name and
   each parameterized range. 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.
-- CUBIN identifies CUDA entry symbols and reads numeric `nv.info` records and
-  native section declarations. Global register declarations also cover compiler
+- CUBIN identifies CUDA entry symbols and reads numeric `nv.info` records,
+  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.
 - AMDGPU reads v3+ MessagePack notes. Argument counts include hidden ABI arguments
   when present; SGPR/VGPR counts and fixed group/private segment sizes retain

src/gpu/ClangOffloadBundle.ts

@@ -69,7 +69,9 @@ export const readClangBundleDirectory = async (
     });
     row.bundle_id = name;
     row.target =
-      /-(?:-)?(gfx[a-zA-Z0-9]+(?::[^\s]+)?)$/u.exec(name)?.[1] ?? null;
+      /-(?:-)?(gfx[a-zA-Z0-9]+(?:(?::[^\s]+)|(?:\+[a-zA-Z0-9_-]+))*)$/u.exec(
+        name,
+      )?.[1] ?? null;
     objects.push(row);
     position += 24 + nameSize;
   }

src/gpu/CudaFatbinary.ts

@@ -47,7 +47,13 @@ export async function* readCudaFatbinary(
         kind: kind === 1 ? "ptx" : kind === 2 ? "cubin" : "unknown",
         metadata: { offset: position, size: width },
       });
-      row.target = `sm_${h.readUInt32LE(28)}${(flags & 0x100000n) !== 0n ? "a" : ""}`;
+      const architectureSuffix =
+        (flags & 0x100000n) !== 0n
+          ? "a"
+          : (flags & 0x200000n) !== 0n
+            ? "f"
+            : "";
+      row.target = `sm_${h.readUInt32LE(28)}${architectureSuffix}`;
       row.raw_flags = `0x${flags.toString(16)}`;
       row.compression =
         (flags & 0x8000n) !== 0n && (flags & 0x2000n) === 0n

src/gpu/CudaInfo.test.ts

@@ -18,6 +18,27 @@ it("retains observed argument ranges while reporting inconsistent area declarati
   });
   expect(read.limitations.join(" ")).toContain("disagree");
 });
+it("decodes V2 kernel parameter sizes without the legacy packed-size shift", () => {
+  const read = readCudaParameters(
+    Buffer.from("0319204e04450c000000000000000000204e0000", "hex"),
+    true,
+  );
+  expect(read).toMatchObject({
+    parameter_count: 1,
+    parameter_bytes: 20000,
+    parameters: [{ ordinal: 0, offset: 0, size: 20000 }],
+    limitations: [],
+  });
+});
+it("marks malformed V2 parameter payloads incomplete", () => {
+  const read = readCudaParameters(
+    Buffer.from("0319204e044508000000000000000000", "hex"),
+    true,
+  );
+  expect(read.parameter_count).toBeNull();
+  expect(read.parameter_bytes).toBe(20000);
+  expect(read.limitations.join(" ")).toContain("unsupported indices/sizes");
+});
 it("bounds dense numeric record streams without retaining a record array", () => {
   const bytes = Buffer.alloc(4 * 262145);
   for (let offset = 0; offset < bytes.length; offset += 4) bytes[offset] = 1;

src/gpu/CudaInfo.ts

@@ -89,7 +89,11 @@ export const readCudaParameters = (
   for (const record of readCudaInfo(bytes, little, budget)) {
     if (record.format === 3 && record.attribute === 0x19)
       byteDeclarations.add(record.value);
-    if (record.format !== 4 || record.attribute !== 0x17) continue;
+    if (
+      record.format !== 4 ||
+      (record.attribute !== 0x17 && record.attribute !== 0x45)
+    )
+      continue;
     if (record.data.length !== 12) {
       complete = false;
       continue;
@@ -100,7 +104,7 @@ export const readCudaParameters = (
       ordinal: r.u16(4),
       name: null,
       offset: r.u16(6),
-      size: r.u32(8) >>> 18,
+      size: record.attribute === 0x45 ? r.u32(8) : r.u32(8) >>> 18,
       type: null,
       hidden: null,
     };

src/gpu/GpuElfKernels.ts

@@ -381,7 +381,7 @@ export const readAmdElfMetadata = async (
     const targetMatch =
       target === null
         ? null
-        : /^amdgcn-amd-amdhsa-(?:-(gfx[a-zA-Z0-9]+(?::[^\s]+)?)|unknown-(gfx[a-zA-Z0-9]+(?::[^\s]+)?))$/u.exec(
+        : /^amdgcn-amd-amdhsa-(?:-(gfx[a-zA-Z0-9]+(?:(?::[^\s]+)|(?:\+[a-zA-Z0-9_-]+))*)|unknown-(gfx[a-zA-Z0-9]+(?:(?::[^\s]+)|(?:\+[a-zA-Z0-9_-]+))*))$/u.exec(
             target,
           );
     const target_key = targetMatch?.[1] ?? targetMatch?.[2] ?? null;

src/gpu/GpuInventory.ts

@@ -47,6 +47,10 @@ import {
 } from "./GpuElfKernels.js";
 import { readPtxKernels, readPtxTarget } from "./PtxKernels.js";
 
+/** PTX calls compute_XX a synonym for the corresponding sm_XX target. */
+const targetIdentity = (target: string): string =>
+  target.replace(/^compute_(\d+[a-z]?)$/u, "sm_$1");
+
 /** File declarations only. Every child range stays in its named parent coordinate space. */
 export class GpuInventory {
   readonly objects: GpuObject[] = [];
@@ -400,7 +404,8 @@ export class GpuInventory {
         const inner = /^sm_([0-9]+)$/u.exec(payload.key)?.[1];
         if (base !== undefined && inner !== undefined)
           state = base === inner ? "base-only" : "conflict";
-      } else if (outer === payload.key) state = "consistent";
+      } else if (targetIdentity(outer) === targetIdentity(payload.key))
+        state = "consistent";
       else if (
         row.kind === "amdgpu" &&
         outer.split(":")[0] === payload.key.split(":")[0]
@@ -441,6 +446,7 @@ export class GpuInventory {
         return source.read(row.location.offset + offset, size);
       },
     };
+
     try {
       const native =
         directory ??

src/gpu/PtxKernels.test.ts

@@ -16,7 +16,7 @@ it("recognizes directives outside comments and strings and preserves UTF-8 byte
     { type: ".b32", count: 3 },
   ]);
 });
-it("derives vector/array parameter declarations and leaves unknown qualifiers unresolved", () => {
+it("derives vector/array and pointer parameters while preserving their names", () => {
   const rows = [
     ...readPtxKernels(
       ".entry k(.param .v2 .u32 x, .param .align 16 .b8 a[32], .param .u64 .ptr .global y) {ret;}",
@@ -28,15 +28,63 @@ it("derives vector/array parameter declarations and leaves unknown qualifiers un
   ).toEqual([
     [0, 8],
     [16, 32],
-    [null, null],
+    [48, 8],
   ]);
-  expect(rows[0]!.kernel.parameter_bytes).toBeNull();
-  expect(rows[0]!.coverage).toBe("partial");
-  expect(rows[0]!.kernel.limitations.join(" ")).toContain(
-    "parameter declarations",
-  );
+  expect(rows[0]!.kernel.parameters[2]!.name).toBe("y");
+  expect(rows[0]!.kernel.parameter_bytes).toBe(56);
+  expect(rows[0]!.coverage).toBe("complete");
   expect(rows[0]!.kernel.parameter_count).toBe(3);
 });
+it("retains a parameter name when a qualifier is unknown", () => {
+  const [row] = [
+    ...readPtxKernels(
+      ".entry k(.param .u64 .vendor_specific value) { ret; }",
+      "object",
+    ),
+  ];
+  expect(row!.kernel.parameters).toEqual([
+    expect.objectContaining({ name: "value", size: null, offset: null }),
+  ]);
+  expect(row!.coverage).toBe("partial");
+});
+it.each([".ptr .global .align 4", ".ptr .align 1", ".ptr.global.align 4"])(
+  "parses PTX pointer qualifiers %s without treating pointee alignment as parameter alignment",
+  (qualifiers) => {
+    const [row] = [
+      ...readPtxKernels(
+        `.entry k(.param .u64 ${qualifiers} ptr_arg) { ret; }`,
+        "object",
+      ),
+    ];
+    expect(row!.kernel.parameters).toEqual([
+      expect.objectContaining({ name: "ptr_arg", offset: 0, size: 8 }),
+    ]);
+    expect(row!.kernel.parameter_bytes).toBe(8);
+    expect(row!.coverage).toBe("complete");
+  },
+);
+it("preserves pointer names but withholds layout for unknown pointer qualifiers", () => {
+  const [row] = [
+    ...readPtxKernels(
+      ".entry k(.param .u64 .ptr .mystery .align 4 ptr_arg) { ret; }",
+      "object",
+    ),
+  ];
+  expect(row!.kernel.parameters).toEqual([
+    expect.objectContaining({ name: "ptr_arg", offset: null, size: null }),
+  ]);
+  expect(row!.coverage).toBe("partial");
+});
+it("excludes call-prototype formals from entry virtual register counts", () => {
+  const [row] = [
+    ...readPtxKernels(
+      ".entry k() { .reg .b32 %r<4>; F: .callprototype _ (.reg .b32 _); ret; }",
+      "object",
+    ),
+  ];
+  expect(row!.kernel.registers.virtual).toEqual([{ type: ".b32", count: 4 }]);
+  expect(row!.coverage).toBe("complete");
+});
 it("keeps dynamic and compound shared declarations unknown", () => {
   const [row] = [
     ...readPtxKernels(

src/gpu/PtxKernels.ts

@@ -153,10 +153,66 @@ const declaration = (
   const vector = values[at] === ".v2" ? 2 : values[at] === ".v4" ? 4 : 1;
   if (vector !== 1) at++;
   const scalar = values[at++],
-    unit = sizeOf(scalar),
-    candidate = values[at++];
-  const name =
+    unit = sizeOf(scalar);
+  const qualifiers = values.slice(at).flatMap((value) =>
+    value.startsWith(".ptr.")
+      ? [
+          ".ptr",
+          ...value
+            .slice(5)
+            .split(".")
+            .map((part) => `.${part}`),
+        ]
+      : [value],
+  );
+  let qualifierAt = 0,
+    pointer = false,
+    pointerLayoutKnown = true;
+  while (qualifiers[qualifierAt]?.startsWith(".")) {
+    const qualifier = qualifiers[qualifierAt++]!;
+    if (qualifier === ".ptr") {
+      if (pointer) pointerLayoutKnown = false;
+      pointer = true;
+    } else if (
+      pointer &&
+      [".const", ".global", ".local", ".shared"].includes(qualifier)
+    ) {
+      // Address-space qualifiers describe the pointed-to object.
+    } else if (pointer && qualifier === ".align") {
+      const alignmentToken = qualifiers[qualifierAt];
+      if (alignmentToken !== undefined && /^[0-9]+$/u.test(alignmentToken))
+        qualifierAt++;
+      const pointeeAlignment = integer(alignmentToken);
+      if (
+        pointeeAlignment === null ||
+        pointeeAlignment <= 0 ||
+        !Number.isInteger(Math.log2(pointeeAlignment))
+      )
+        pointerLayoutKnown = false;
+    } else pointerLayoutKnown = false;
+  }
+  const candidate = qualifiers[qualifierAt];
+  const candidateIndex =
+    candidate === undefined ? -1 : values.indexOf(candidate, at);
+  let name =
     candidate !== undefined && identifier.test(candidate) ? candidate : null;
+  if (!pointerLayoutKnown) name = null;
+  // Unknown qualifiers can carry operands; preserve a unique trailing name but
+  // do not derive a byte layout from an unrecognized declaration.
+  if (!pointerLayoutKnown) {
+    const nameToken = qualifiers
+      .slice(qualifierAt)
+      .reverse()
+      .find((value) => identifier.test(value));
+    name = nameToken ?? null;
+  }
+  const namePosition =
+    candidateIndex >= 0
+      ? candidateIndex
+      : name === null
+        ? values.length
+        : values.lastIndexOf(name);
+  at = namePosition + 1;
   let size = unit === null ? null : unit * vector;
   const type =
     scalar === undefined
@@ -178,6 +234,7 @@ const declaration = (
   }
   if (
     at !== values.length ||
+    !pointerLayoutKnown ||
     name === null ||
     align === null ||
     align <= 0 ||
@@ -326,6 +383,14 @@ export function* readPtxKernels(
           end = t.end;
           continue;
         }
+        // Call prototypes contain formal .reg/.param declarations, not entry resources.
+        if (t.value === ".callprototype") {
+          for (t = next(); t.value !== ";"; t = next()) {
+            if (t.value === "}") rejectGpu("Malformed PTX call prototype");
+          }
+          end = t.end;
+          continue;
+        }
         if (t.value !== ".reg" && t.value !== ".shared") continue;
         if (t.value === ".reg") {
           const { type, count } = registerDeclaration(next);
@@ -453,7 +518,7 @@ export const readPtxTarget = (text: string): string | null => {
         const value = stream.next();
         target =
           !value.done &&
-          /^(?:sm|compute)_[0-9]+[a-z]?$/u.test(value.value.value)
+          /^(?:sm_[0-9]+(?:a|f)?|compute_[0-9]+)$/u.test(value.value.value)
             ? value.value.value
             : null;
       }

tests/acceptance/gpuCodeObjectsCli.test.ts

@@ -19,6 +19,7 @@ import {
   ptxFixture,
   amdFixture,
   cubinFixture,
+  hipFixture,
 } from "../fixtures/gpu/codeObjects.js";
 
 cliTest(
@@ -63,6 +64,197 @@ cliTest(
   },
 );
 
+cliTest(
+  "compares equivalent PTX targets and retains family-specific CUDA targets",
+  async ({ cli }) => {
+    const root = await createTestTempDirectory("rea-gpu-target-spelling-");
+    for (const [name, target, payloadTarget, ptx, flags] of [
+      [
+        "compute",
+        "sm_80",
+        "compute_80",
+        ptxFixture.replace(".target sm_80", ".target compute_80"),
+        0,
+      ],
+      [
+        "family",
+        "sm_100f",
+        "sm_100f",
+        ptxFixture.replace(".target sm_80", ".target sm_100f"),
+        0x200000,
+      ],
+    ] as const) {
+      const bytes = cudaFixture("none", [
+          {
+            kind: 1,
+            target: name === "compute" ? 80 : 100,
+            bytes: Buffer.from(ptx),
+            flags,
+          },
+        ]),
+        path = join(root, name);
+      await writeFile(path, bytes);
+      const response = await cli.run({
+        arguments: [
+          "inspect-gpu-code-objects",
+          path,
+          "--analysis",
+          "kernels",
+          "--detail",
+          "complete",
+          "--json",
+        ],
+      });
+      expect(response.exitCode, response.stderr).toBe(0);
+      const evidence = parseEvidence(
+        toolContract("inspect_gpu_code_objects").outputSchema.parse(
+          response.json,
+        ),
+      );
+      const report = gpuCodeObjectsSchema.parse(evidence.normalized_result);
+      expect(report.objects[0]).toMatchObject({
+        target,
+        payload_target: payloadTarget,
+        target_consistency: "consistent",
+      });
+      expect(report.coverage.payloads).toBe("complete");
+    }
+  },
+);
+
+cliTest(
+  "reads CUDA V2 parameter records through the public CLI",
+  async ({ cli }) => {
+    const root = await createTestTempDirectory("rea-cubin-v2-parameter-"),
+      bytes = cudaFixture("none", [
+        {
+          kind: 2,
+          target: 80,
+          bytes: cubinFixture({ architecture: 80, v2ParameterSize: 20000 }),
+        },
+      ]),
+      path = join(root, "large-parameter.cubin");
+    await writeFile(path, bytes);
+    const response = await cli.run({
+      arguments: [
+        "inspect-gpu-code-objects",
+        path,
+        "--analysis",
+        "kernels",
+        "--detail",
+        "complete",
+        "--json",
+      ],
+    });
+    expect(response.exitCode, response.stderr).toBe(0);
+    const evidence = parseEvidence(
+      toolContract("inspect_gpu_code_objects").outputSchema.parse(
+        response.json,
+      ),
+    );
+    const report = gpuCodeObjectsSchema.parse(evidence.normalized_result);
+    expect(report.kernels[0]).toMatchObject({
+      parameter_count: 1,
+      parameter_bytes: 20000,
+      parameters: [{ ordinal: 0, offset: 0, size: 20000 }],
+    });
+    expect(report.coverage.payloads).toBe("complete");
+  },
+);
+
+cliTest(
+  "preserves HIP v3 feature target IDs for exact selection",
+  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",
+      ),
+      path = join(root, "feature.bundle");
+    await writeFile(path, bytes);
+    const response = await cli.run({
+      arguments: [
+        "inspect-gpu-code-objects",
+        path,
+        "--analysis",
+        "kernels",
+        "--targets",
+        JSON.stringify(["gfx906+xnack"]),
+        "--detail",
+        "complete",
+        "--json",
+      ],
+    });
+    expect(response.exitCode, response.stderr).toBe(0);
+    const evidence = parseEvidence(
+      toolContract("inspect_gpu_code_objects").outputSchema.parse(
+        response.json,
+      ),
+    );
+    const report = gpuCodeObjectsSchema.parse(evidence.normalized_result);
+    expect(report.objects.find((row) => row.kind === "amdgpu")!.target).toBe(
+      "gfx906+xnack",
+    );
+    expect(report.kernels).toHaveLength(1);
+    expect(report.coverage.payloads).toBe("complete");
+  },
+);
+
+cliTest(
+  "keeps PTX pointer names and excludes call prototypes from resources",
+  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",
+        )
+        .replace(
+          ".reg .b64 %rd<5>;",
+          ".reg .b64 %rd<5>; F: .callprototype _ (.reg .b32 _);",
+        ),
+      root = await createTestTempDirectory("rea-ptx-callprototype-"),
+      path = join(root, "prototype.ptx");
+    await writeFile(path, source);
+    const response = await cli.run({
+      arguments: [
+        "inspect-gpu-code-objects",
+        path,
+        "--analysis",
+        "kernels",
+        "--detail",
+        "complete",
+        "--json",
+      ],
+    });
+    expect(response.exitCode, response.stderr).toBe(0);
+    const evidence = parseEvidence(
+      toolContract("inspect_gpu_code_objects").outputSchema.parse(
+        response.json,
+      ),
+    );
+    const report = gpuCodeObjectsSchema.parse(evidence.normalized_result);
+    expect(report.kernels[0]).toMatchObject({
+      parameter_count: 5,
+      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 },
+      ],
+      registers: {
+        virtual: [
+          { type: ".b32", count: 9 },
+          { type: ".b64", count: 5 },
+        ],
+      },
+    });
+    expect(report.coverage.payloads).toBe("complete");
+  },
+);
+
 cliTest(
   "resolves CUBIN SHN_XINDEX symbols through their ELF extension table",
   async ({ cli }) => {
@@ -343,7 +535,7 @@ it("preserves incomplete decoder coverage through MCP target and object pages",
               ])
             : missing === "ptx-parameter"
               ? Buffer.from(
-                  ptxFixture.replace(".u64 output", ".u64 .ptr .global output"),
+                  ptxFixture.replace(".u64 output", ".u64 .vendor output"),
                 )
               : missing === "ptx-shared"
                 ? Buffer.from(ptxFixture.replace("scratch[32]", "scratch[]"))

tests/fixtures/gpu/codeObjects.ts

@@ -94,6 +94,7 @@ export const cubinFixture = (
     osAbi?: number;
     flags?: number;
     extendedSectionIndex?: boolean;
+    v2ParameterSize?: number;
   } = {},
 ): Buffer => {
   const name = options.name ?? "add_one";
@@ -108,10 +109,24 @@ export const cubinFixture = (
   const global = Buffer.alloc(record.length * (options.globalRecords ?? 1));
   for (let offset = 0; offset < global.length; offset += record.length)
     record.copy(global, offset);
-  const info = Buffer.from(
-    "03190c0004170c00000000000100080000f0110004170c00000000000000000000f02100",
-    "hex",
-  );
+  const info =
+    options.v2ParameterSize === undefined
+      ? Buffer.from(
+          "03190c0004170c00000000000100080000f0110004170c00000000000000000000f02100",
+          "hex",
+        )
+      : (() => {
+          const byteCount = Buffer.alloc(4);
+          byteCount[0] = 3;
+          byteCount[1] = 0x19;
+          byteCount.writeUInt16LE(options.v2ParameterSize!, 2);
+          const v2 = Buffer.alloc(16);
+          v2[0] = 4;
+          v2[1] = 0x45;
+          v2.writeUInt16LE(12, 2);
+          v2.writeUInt32LE(options.v2ParameterSize!, 12);
+          return Buffer.concat([byteCount, v2]);
+        })();
   const paddingSections = options.extendedSectionIndex
     ? Array.from({ length: 65531 }, (_, index) => ({
         name: `.padding.${index}`,
@@ -243,12 +258,12 @@ export const amdFixture = (
 };
 export const cudaFixture = (
   compression: "none" | "zstd" = "none",
-  entries = [
+  entries: { kind: number; target: number; bytes: Buffer; flags?: number }[] = [
     { kind: 1, target: 80, bytes: Buffer.from(ptxFixture) },
     { kind: 2, target: 90, bytes: cubinFixture({ architecture: 90 }) },
   ],
 ): Buffer => {
-  const records = entries.map(({ kind, target, bytes }) => {
+  const records = entries.map(({ kind, target, bytes, flags: entryFlags }) => {
     const compressed = compression === "zstd",
       data = compressed ? zstdCompressSync(bytes) : bytes;
     const result = Buffer.alloc(64 + Math.ceil(data.length / 8) * 8);
@@ -258,7 +273,10 @@ export const cudaFixture = (
     result.writeBigUInt64LE(BigInt(result.length - 64), 8);
     result.writeUInt32LE(compressed ? data.length : 0, 16);
     result.writeUInt32LE(target, 28);
-    result.writeBigUInt64LE(compressed ? 0x8011n : 0x11n, 40);
+    result.writeBigUInt64LE(
+      (compressed ? 0x8011n : 0x11n) | BigInt(entryFlags ?? 0),
+      40,
+    );
     result.writeBigUInt64LE(compressed ? BigInt(bytes.length) : 0n, 56);
     data.copy(result, 64);
     return result;