This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

4. Benchmark FIPS-compliant alternatives before settling on SHA-256 — on 64-bit systems sha512 is often faster, and security doesn't matter here, only collision resistance and speed.

To clarify: All FIPS-compliant algorithms should be considered and benchmarked. We are not bound to the 256-bit key size. Either shorter or longer is OK, as long as it is fast and has a low chance to collide.

@leofang

Addressed your requested changes.
Changes made:

• removed cuda_core/tests/test_program_cache_fips.py
• added usedforsecurity=False at both hash call sites
• added inline FIPS comments at both hash call sites
• benchmarked FIPS-approved alternatives locally and switched both sites consistently to sha384
• cleaned up the remaining stale test wording so it now says 384-bit digest

• benchmarked FIPS-approved alternatives locally and switched both sites consistently to sha384

Could you share your benchmark script?

Sure. I turned the local benchmark into a standalone stdlib-only script that mirrors both hash sites plus the coupled end-to-end path where make_program_cache_key() feeds FileStreamProgramCache._path_for_key().

On this x86_64 host, the SHA-2 family clearly beat the SHA-3 family for these workloads. For the coupled end-to-end cases, sha384 and sha512 were effectively the two leaders, and sha384 edged out or matched the others closely enough that I kept it consistently at both sites while also keeping the key shorter than sha512.

A representative run from python3 scripts/bench_program_cache_hashes.py --loops 5000 --repeat 2:

• end_to_end_cpp: sha384 10,657.7 ns/op, sha256 11,034.8 ns/op, sha512 11,043.4 ns/op
• end_to_end_ptx: sha384 9,526.4 ns/op, sha512 9,642.1 ns/op, sha256 9,714.3 ns/op

#!/usr/bin/env python3
# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
#
# SPDX-License-Identifier: Apache-2.0

"""Benchmark FIPS-approved hashlib candidates for cuda.core program-cache use.

This mirrors the two relevant call sites:

* ``FileStreamProgramCache._path_for_key()``: hash a cache key to a stable
  filename component via ``hexdigest()``.
* ``make_program_cache_key()``: incrementally build the digest from labeled
  payload chunks and return ``digest()``.

The benchmark is intentionally stdlib-only so reviewers can run it directly.
"""

from __future__ import annotations

import argparse
import hashlib
import inspect
import statistics
import sys
import time
from dataclasses import dataclass
from typing import Callable

_DEFAULT_ALGORITHMS = (
    "sha224",
    "sha256",
    "sha384",
    "sha512",
    "sha3_224",
    "sha3_256",
    "sha3_384",
    "sha3_512",
)


@dataclass(frozen=True)
class HashCase:
    name: str
    runner: Callable[[Callable[..., object]], None]


def _supports_usedforsecurity(constructor: Callable[..., object]) -> bool:
    try:
        signature = inspect.signature(constructor)
    except (TypeError, ValueError):
        return False
    return "usedforsecurity" in signature.parameters


def _make_constructor(name: str) -> Callable[..., object]:
    constructor = getattr(hashlib, name)
    if _supports_usedforsecurity(constructor):
        return lambda data=b"": constructor(data, usedforsecurity=False)
    return constructor


def _file_stream_case(name: str, key: bytes) -> HashCase:
    def _runner(constructor: Callable[..., object]) -> None:
        constructor(key).hexdigest()

    return HashCase(name, _runner)


def _program_cache_case(name: str, payloads: tuple[tuple[str, bytes], ...]) -> HashCase:
    def _runner(constructor: Callable[..., object]) -> None:
        hasher = constructor()
        for label, payload in payloads:
            hasher.update(label.encode("ascii"))
            hasher.update(len(payload).to_bytes(8, "big"))
            hasher.update(payload)
        hasher.digest()

    return HashCase(name, _runner)


def _end_to_end_case(name: str, payloads: tuple[tuple[str, bytes], ...]) -> HashCase:
    def _runner(constructor: Callable[..., object]) -> None:
        hasher = constructor()
        for label, payload in payloads:
            hasher.update(label.encode("ascii"))
            hasher.update(len(payload).to_bytes(8, "big"))
            hasher.update(payload)
        key = hasher.digest()
        constructor(key).hexdigest()

    return HashCase(name, _runner)


def _sample_cases() -> tuple[HashCase, ...]:
    file_stream_key = bytes.fromhex("ab" * 48)
    long_file_stream_key = (b"cuda-core-cache-key-" * 128)[:4096]

    source = b"""
extern "C" __global__ void saxpy(float a, const float* x, float* y) {
    const int i = blockIdx.x * blockDim.x + threadIdx.x;
    y[i] = a * x[i] + y[i];
}
""".strip()
    ptx = b"""
.version 8.0
.target sm_90
.address_size 64
.visible .entry saxpy() { ret; }
""".strip()
    option_bytes = (
        b"name='saxpy'",
        b"arch='sm_90'",
        b"max_register_count=None",
        b"time=False",
        b"link_time_optimization=False",
        b"debug=False",
        b"lineinfo=False",
        b"ftz=None",
        b"prec_div=None",
        b"prec_sqrt=None",
        b"fma=None",
        b"split_compile=None",
        b"ptxas_options=None",
        b"no_cache=False",
    )
    names = (b"saxpy", b"_Z5saxpyv")
    extra_digest = bytes.fromhex("cd" * 32)

    cpp_payloads = (
        ("schema", b"2"),
        ("nvrtc", b"13.2"),
        ("code_type", b"c++"),
        ("target_type", b"cubin"),
        ("code", source),
        ("option_count", str(len(option_bytes)).encode("ascii")),
        *tuple(("option", item) for item in option_bytes),
        ("names_count", str(len(names)).encode("ascii")),
        *tuple(("name", item) for item in names),
        ("options_name", b"saxpy"),
        ("extra_digest", extra_digest),
    )
    ptx_payloads = (
        ("schema", b"2"),
        ("linker", b"nvJitLink-13.2"),
        ("code_type", b"ptx"),
        ("target_type", b"cubin"),
        ("code", ptx),
        ("option_count", str(len(option_bytes)).encode("ascii")),
        *tuple(("option", item) for item in option_bytes),
        ("names_count", b"0"),
        ("extra_digest", extra_digest),
    )

    return (
        _file_stream_case("file_stream_key_48b", file_stream_key),
        _file_stream_case("file_stream_key_4k", long_file_stream_key),
        _program_cache_case("program_cache_cpp", cpp_payloads),
        _program_cache_case("program_cache_ptx", ptx_payloads),
        _end_to_end_case("end_to_end_cpp", cpp_payloads),
        _end_to_end_case("end_to_end_ptx", ptx_payloads),
    )


def _benchmark_case(
    case: HashCase,
    constructor: Callable[..., object],
    *,
    loops: int,
    repeat: int,
) -> tuple[float, float]:
    samples_ns: list[float] = []
    for _ in range(repeat):
        start = time.perf_counter_ns()
        for _ in range(loops):
            case.runner(constructor)
        elapsed = time.perf_counter_ns() - start
        samples_ns.append(elapsed / loops)
    return statistics.mean(samples_ns), min(samples_ns)


def _format_ns(value: float) -> str:
    return f"{value:,.1f}"


def _write_line(text: str = "") -> None:
    sys.stdout.write(text + "\n")


def main() -> None:
    parser = argparse.ArgumentParser(description=__doc__)
    parser.add_argument(
        "--loops",
        type=int,
        default=200_000,
        help="Iterations per repeat for each algorithm/case pair.",
    )
    parser.add_argument(
        "--repeat",
        type=int,
        default=7,
        help="Independent timing repeats for each algorithm/case pair.",
    )
    parser.add_argument(
        "--algorithms",
        nargs="+",
        default=list(_DEFAULT_ALGORITHMS),
        help="hashlib algorithm names to benchmark.",
    )
    args = parser.parse_args()

    cases = _sample_cases()
    widths = {
        "algorithm": max(len("Algorithm"), max(len(name) for name in args.algorithms)),
        "case": max(len(case.name) for case in cases),
    }

    _write_line(
        f"{'Algorithm':<{widths['algorithm']}}  "
        f"{'Case':<{widths['case']}}  {'mean ns/op':>12}  {'best ns/op':>12}"
    )
    _write_line("-" * (widths["algorithm"] + widths["case"] + 28))

    for algorithm in args.algorithms:
        constructor = _make_constructor(algorithm)
        for case in cases:
            mean_ns, best_ns = _benchmark_case(case, constructor, loops=args.loops, repeat=args.repeat)
            _write_line(
                f"{algorithm:<{widths['algorithm']}}  "
                f"{case.name:<{widths['case']}}  "
                f"{_format_ns(mean_ns):>12}  {_format_ns(best_ns):>12}"
            )


if __name__ == "__main__":
    main()

Benchmark: FIPS-approved hash algorithms for program cache

I re-ran the benchmark with proper parameters (--loops 200000 --repeat 7) on an AMD Threadripper PRO 3975WX (x86_64 with SHA-NI). The original run used --loops 5000 --repeat 2, which is far too few samples to draw conclusions — the sha384 "win" in the original results was noise.

Results (end-to-end case, best ns/op)

All FIPS-approved algorithms available in hashlib were tested:

Analysis

• SHA-NI hardware acceleration makes sha256/sha224/sha1 the fastest group at ~6,500 ns. SHA-NI is present on most modern AMD (Zen+) and Intel (Ice Lake+) CPUs. It only accelerates SHA-256 and SHA-1 — not SHA-512 or SHA-3.
• sha384 and sha512 are ~15% slower because they use 64-bit software operations with no hardware acceleration on x86_64.
• SHA-3 family is 25-50% slower — expected, as SHA-3 (Keccak) is designed for hardware, not software on x86_64.
• sha1 is equally fast but only provides 160-bit / 2^80 collision resistance. No reason to pick it when sha256 is the same speed with 2^128.

Recommendation

Please switch back from sha384 to sha256. It is:

• Fastest (tied with sha224/sha1, thanks to SHA-NI)
• 256-bit collision resistance (2^128) — more than sufficient
• Universally available in every Python build
• FIPS-approved

When running the benchmark, please use the default parameters (--loops 200000 --repeat 7) — --loops 5000 --repeat 2 does not produce statistically meaningful results.

Follow-up on the hash choice: the earlier sha384 switch came from an undersampled local run. After the maintainer reran the benchmark with --loops 200000 --repeat 7, the result on SHA-NI x86_64 was that sha256 is the fastest FIPS-approved practical choice for this program-cache workload.

I have switched both runtime call sites back to sha256(..., usedforsecurity=False) consistently, kept _KEY_SCHEMA_VERSION = 2, and updated the cache-key/file-stream wording and digest-width tests back to 32 bytes.

For reproducibility, the PR now carries scripts/bench_program_cache_hashes.py as a stdlib-only review/support artifact with the broader algorithm set and the maintainer-style default benchmark parameters.

Please kindly push your local changes, thanks!

/ok to test 404db9a

Following up here because while I was rereading your feedback on the other PR and then looking more closely at this branch, I noticed these two Windows MCDM failures were not coming from the FIPS hash change itself. They are both hitting the same pinned-memory setup path with a CUDA_ERROR_OUT_OF_MEMORY during PinnedMemoryResource(...) construction, so I updated the two NUMA tests to use the existing create_pinned_memory_resource_or_xfail(...) helper that the test suite already uses for known Windows MCDM mempool setup failures. I also pushed the formatter-only cleanup needed for pre-commit.ci.

Following up here because while I was rereading your feedback on the other PR and then looking more closely at this branch, I noticed these two Windows MCDM failures were not coming from the FIPS hash change itself. They are both hitting the same pinned-memory setup path with a CUDA_ERROR_OUT_OF_MEMORY during PinnedMemoryResource(...) construction, so I updated the two NUMA tests to use the existing create_pinned_memory_resource_or_xfail(...) helper that the test suite already uses for known Windows MCDM mempool setup failures. I also pushed the formatter-only cleanup needed for pre-commit.ci.

I didn't compare byte-for-byte, but it looks like my #2096. Where did you see the error? Here in the CI? (My PR addresses test failures outside our CI; found by our QA team).

This is a scope-creep and please revert the recent changes. Irrelevant test failures will be triaged by the team and if any action is needed we'll let you know. In fact, the OOM issue is being addressed (#2096).

I think it's referring to https://github.com/NVIDIA/cuda-python/actions/runs/25950797754/job/76288908185.

I think it's referring to https://github.com/NVIDIA/cuda-python/actions/runs/25950797754/job/76288908185.

Oh awesome, so we're seeing the same flakiness that QA observed also in our CI now.

Yes, that was the CI job I was referring to. I agree this should stay out of #2087, so I reverted the unrelated Windows MCDM test changes and pushed the branch back to the FIPS-only changeset in e84d143a9a. I will leave the OOM issue to #2096 and the team triage path.

I think it's referring to https://github.com/NVIDIA/cuda-python/actions/runs/25950797754/job/76288908185.

Oh awesome, so we're seeing the same flakiness that QA observed also in our CI now.

Yes. I added dual-GPU runners yesterday (#2090; see also team channel).

@rwgk @leofang Thanks, that helps clarify it. My read is that the CI failure looks more like the same Windows MCDM pinned-memory OOM behavior QA saw than something specific to the NUMA test logic itself.

A possible fault boundary to check in #2096 is whether the instability is already present at bare PinnedMemoryResource(...) construction on that Windows H100 MCDM lane, versus something introduced later by the NUMA-specific test path. If that is reproducible in isolation, it would point more toward platform/driver/runtime allocation limits or resource initialization behavior on Windows MCDM than the test logic itself.

A simple validation path might be:

1. isolate just PinnedMemoryResource(...) creation in the same CI lane
2. compare behavior with and without the NUMA-related parameters
3. record whether the failure tracks the environment/driver state rather than the individual test body

I reverted my unrelated test changes from #2087, so I am not proposing to address it in this PR. I would be glad to help on the CI side if useful, although I may be limited by contributor permissions on some of the workflow/debugging path.

@aryanputta could you please force-push to keep 404db9a as the last commit? We were about to merge but all the changes would require re-running the full CI, which is a waste of resource (even for NVIDIA, GPUs are scarce resources). If you revert, we just need to re-run the failing pipeline.

@rwgk @leofang I did a follow-up CLI pass just to understand the failure surface more concretely. I am not proposing a change on #2087, only sharing what I found in case it is useful for #2096 or future runner triage.

What I checked:

• pulled the failing log for Test win-64 / Python 3.14, CUDA 13.2.1 (wheels), GPU h100 (x2) (MCDM)
• compared it against a passing Windows MCDM lane on l4 and a passing dual-GPU Windows lane on t4 (x2)
• checked the current cuda.core test call sites around the failing tests

What seems to be happening:

• the only real test failures in that H100 MCDM lane are test_pinned_mr_numa_id_default_no_ipc and test_pinned_mr_numa_id_explicit
• both fail immediately at bare PinnedMemoryResource(...) construction with CUDA_ERROR_OUT_OF_MEMORY, before reaching the NUMA assertions themselves
• neighboring pinned-memory tests in the same file are otherwise green or Windows-skipped, so the failure boundary looks more like constructor/setup on that platform path than the later NUMA assertion logic
• the repo already has the shared Windows MCDM mempool OOM xfail path in create_pinned_memory_resource_or_xfail(...), and some other pinned-memory tests already go through it, but these two NUMA tests still call PinnedMemoryResource(...) directly

So my current read is that this is less “the NUMA logic is broken” and more “these two tests are still exposing the raw constructor OOM on the new H100 x2 MCDM lane.”

I could not do a full local repro on my side because this environment is missing cuda.pathfinder, so my testing here was log-driven plus code-path comparison. If there is a preferred way you want contributors to investigate this kind of CI-only lane issue without stepping on scope, I would appreciate the advice. Happy to help on the analysis side where contributor permissions allow.

/ok to test 6a9c413