diff --git a/docs/gpu-code-objects.md b/docs/gpu-code-objects.md index 7cd3a8b5c5..0c2e14fac1 100644 --- a/docs/gpu-code-objects.md +++ b/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 diff --git a/src/gpu/ClangOffloadBundle.ts b/src/gpu/ClangOffloadBundle.ts index fdac9390da..8e9da156e1 100644 --- a/src/gpu/ClangOffloadBundle.ts +++ b/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; } diff --git a/src/gpu/CudaFatbinary.ts b/src/gpu/CudaFatbinary.ts index f64d6f6569..1747e973b1 100644 --- a/src/gpu/CudaFatbinary.ts +++ b/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 diff --git a/src/gpu/CudaInfo.test.ts b/src/gpu/CudaInfo.test.ts index 55c402110d..43e7a2db35 100644 --- a/src/gpu/CudaInfo.test.ts +++ b/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; diff --git a/src/gpu/CudaInfo.ts b/src/gpu/CudaInfo.ts index 74a3128b8e..12e5e08b9f 100644 --- a/src/gpu/CudaInfo.ts +++ b/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, }; diff --git a/src/gpu/GpuElfKernels.ts b/src/gpu/GpuElfKernels.ts index 8e6d16393b..5592a83713 100644 --- a/src/gpu/GpuElfKernels.ts +++ b/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; diff --git a/src/gpu/GpuInventory.ts b/src/gpu/GpuInventory.ts index cebc7a06c9..85d3beaf6b 100644 --- a/src/gpu/GpuInventory.ts +++ b/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 ?? diff --git a/src/gpu/PtxKernels.test.ts b/src/gpu/PtxKernels.test.ts index 548d23275a..25256a06f2 100644 --- a/src/gpu/PtxKernels.test.ts +++ b/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( diff --git a/src/gpu/PtxKernels.ts b/src/gpu/PtxKernels.ts index 438351ccb3..a145af2e8f 100644 --- a/src/gpu/PtxKernels.ts +++ b/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; } diff --git a/tests/acceptance/gpuCodeObjectsCli.test.ts b/tests/acceptance/gpuCodeObjectsCli.test.ts index b35d88b890..4c15e242a2 100644 --- a/tests/acceptance/gpuCodeObjectsCli.test.ts +++ b/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[]")) diff --git a/tests/fixtures/gpu/codeObjects.ts b/tests/fixtures/gpu/codeObjects.ts index ab6bfeb94a..8f6a3f6134 100644 --- a/tests/fixtures/gpu/codeObjects.ts +++ b/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;