cuda-intrinsics-gen: unsafe families without a safety renderer inherit the packed-atomic contract #1152

Summary

validate_renderable says every unsafe record has a dedicated family safety renderer, but it only
checks a hand-maintained family-name allowlist (crates/cuda-intrinsics-gen/src/render.rs:445-484).
It does not verify that render_raw_abi has a matching safety-dispatch arm. An allowed family
without an arm reaches the terminal else at render.rs:5173-5178 and inherits the packed-atomic
contract:

} else {
    output.push_str(
        "/// `addr` must designate four writable bytes in global memory and be naturally aligned to four bytes.\n\
         /// Do not overlap this operation with a whole-word atomic or any non-atomic access to either 16-bit lane.\n\
         /// Racing atomics must use scopes that include each other; this relaxed GPU-scope operation is not atomic with host/system access.\n",
    );
}

That text is correct for the two packed_atomic records. It is also emitted for six CLC records in
crates/cuda-intrinsics/src/generated/abi_v1.rs:641-731:

The four query functions have no pointer or addr. The two request functions write a naturally
aligned 16-byte response in .shared::cta and signal an mbarrier, not a four-byte global atomic.

Reproduction

Commit 18cc1888 on safety-guard-red adds two invariants over rendered raw-ABI safety blocks.
Against the unfixed renderer, 298 tests pass and exactly these two fail:

unsafe safety docs reference `addr` without a pointer argument: clc_query_get_first_ctaid_x, clc_query_get_first_ctaid_y, clc_query_get_first_ctaid_z, clc_query_is_canceled
unsafe safety blocks shared across families: [{"clc", "packed_atomic"}]

The generated reference independently contradicts the inherited prose:
intrinsics/generated-reference.md:62-65 records all four clc_query_* operations as
pure; memory none.

Why this is a guard defect

Editing the allowlist is sufficient to pass validation; no corresponding renderer arm is required.
render_raw_abi, the guard, the fallback, the CLC records, and the CLC allowlist entry arrived
together in 6a9a6373 (PR #406). The same failure mode surfaced during review of PR #1138 when
barrier_reduction was added to the allowlist without a safety arm.

Impact and next step

The impact is bounded: cuda_device::clc has separate family-specific prose and is the recommended
path. However, intrinsics/generated-reference.md:1227-1250 calls cuda_intrinsics::clc::* the
public source path, and crates/cuda-intrinsics/src/generated/mod.rs:93-100 publicly re-exports it.

A two-commit RED/GREEN fix is ready on safety-guard-red: 18cc1888 adds the failing invariants;
101b1a01 makes rendering fail closed and corrects the CLC contracts.

The remaining maintainer decision is whether those contracts should also spell out object lifetime
through asynchronous completion and proxy-fence ordering when a response is reused across loop
iterations.

No release tag contains 6a9a6373; I did not find a published cuda-intrinsics crate containing
this code.