Top-K repeats driver lookup, metadata allocations, and launch attributes

Top-K repeats host setup work before every kernel launch: `make_tensor_map` resolves `cuTensorMapEncodeTiled`, rank-three descriptor metadata is built with short-lived vectors, and the Top-K launchers set attributes that `ku::launch_kernel` immediately sets again.

At revision `8e70df7`, FP32 and non-cluster BF16 set `cudaFuncAttributeMaxDynamicSharedMemorySize` twice for the same kernel/value. The large-cluster BF16 path also repeats `cudaFuncAttributeNonPortableClusterSizeAllowed=1`. The shared launch helper already checks both calls, so the wrapper copies are redundant.

A host harness using the actual baseline headers, with a recording mock resolver/encoder, counted eight metadata allocations and one lookup per rank-three call. A fixed-array/cached-symbol candidate counted zero of each after initialization. The mock verifies the arguments passed to the encoder; it does not model CUDA encoding cost. Separately, an Nsight Systems trace of the real runtime lookup recorded 110,001 resolver calls for repeated lookup versus one for a cached-pointer probe (one initial lookup, 10,000 warmup calls, 100,000 measured calls).

The proposed scope is to cache only the stable driver entry point, use fixed-size metadata storage while rebuilding every descriptor, and leave checked attribute setup with the common launcher. Preserve the vector helper interface and failed-initialization retry. Tensor-descriptor caching and kernel changes are outside this proposal.

Relevant precedents are [PyTorch's cached driver API table](https://github.com/pytorch/pytorch/blob/main/c10/cuda/driver_api.cpp) and [CUTLASS's fixed-size TMA metadata arrays](https://github.com/NVIDIA/cutlass/blob/main/include/cute/atom/copy_traits_sm90_tma.hpp). These support the approach; they do not establish a DeepSelect speedup.

Start at make_tensor_map, the rank-three descriptor construction, and the Top-K launchers that call ku::launch_kernel. Check the actual baseline headers and use the recording mock resolver/encoder to verify one cached lookup, zero metadata allocations after initialization, preserved encoder arguments, and retry after failed initialization; confirm redundant attribute setup is removed without changing tensor-descriptor or kernel behavior.

Written by the indexing model from the issue text.