[BUG] Validate the complete cp.async source transfer range

### Prerequisite

- [x] I have searched for related cp.async issues.

### Problem

LegalizeSafeMemoryAccess validates the base index of a direct low-level T.ptx_cp_async source access, but not the full transfer range. A source base can be valid while base + num_elems - 1 is out of bounds.

This can issue an illegal global-memory read or copy adjacent memory into shared memory. The scope is direct/low-level ptx_cp_async callers; the higher-level async-copy lowering already predicates tested tail copies.

### Minimal counterexample

For a logical fp16 source of length 9:

```text
source: A[0 ... 8]
copy: start=A[8], count=8

previous check: A[8] is valid -> unguarded cp.async
required check: A[8 ... 15] is valid -> predicate/zero-fill the invalid tail
```

The base branch lowers this call with three arguments, meaning no predicate is attached.

### Expected behavior

A cp.async source transfer must be legal only when its final transferred element remains within the flattened source buffer. Otherwise legalization should attach the false predicate that causes the existing cp.async lowering to zero-fill rather than read out of bounds.

### Proposed fix

Compute the flattened physical source index and validate its final transferred element against the flattened buffer extent. Keep the existing base-index checks, and deduplicate the equivalent upper-bound predicate for one-element transfers. A regression test covers the valid-base/invalid-range case.

Start at LegalizeSafeMemoryAccess and trace the direct, low-level cp.async source-access path, then review the existing cp.async lowering and regression test mentioned in the issue. Verify that validation covers the final transferred element, preserves base-index checks, and attaches the false predicate for an invalid tail so the transfer zero-fills instead of reading out of bounds.

Written by the indexing model from the issue text.