Pad allocations and shared memory to whole 32-bit words

[maleadt]

PTX has few 8- and 16-bit atomics, so LLVM implements them as a compare-and-swap on the containing 32-bit word. When the value sits in the last partial word of an array, that word extends past the end of the allocation. This never faults (allocations are 256-byte aligned), but compute-sanitizer reports it as an out-of-bounds access and aborts the kernel, leaving a sticky error:

========= Invalid __global__ read of size 4 bytes
=========     Access to 0xb00000000 is out of bounds
=========     and is inside the nearest allocation at 0xb00000000 of size 2 bytes

This hits any sub-word atomic that goes through LLVM: Int8/Int16 read-modify-writes and CAS from UnsafeAtomics, Atomix or KernelAbstractions on every architecture (current LLVM never emits atom.cas.b16), and CUDA.jl's own 16-bit atomics before sm_70 or BFloat16 add before sm_90. It applies to global memory and to static and dynamic shared memory alike. NVCC doesn't run into this because it emits atom.cas.b16, which ptxas widens after the sanitizer's view of the access; libcu++'s cuda::atomic_ref<short> has the same problem we do.

Like Metal.jl did in JuliaGPU/Metal.jl#977, this pads pool allocations, static shared memory arrays and the dynamic shared memory size of a launch to a whole number of 32-bit words. That costs no memory: the stream-ordered pool hands out 512-byte blocks even for 2-byte requests, and shared memory is allocated in units of at least 128 bytes. The new tests make the sanitizer job abort without this change.

Caveats:

• Memory that doesn't come from the pool isn't padded (unsafe_wrap, the low-level CUDA.alloc, device-side malloc). The kernel programming docs now say so.
• Memory statistics now count the padded size, e.g. CUDA.@allocated CuArray{Int8}(undef, 1) reports 4 bytes.
• A launch whose dynamic shared memory size equals an odd-sized FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES limit will now exceed it by up to 3 bytes.

With this in place, #3350 shouldn't need its native atom.cas.b16 paths to keep the sanitizer happy, and 8-bit atomics (which #3350 enables through UnsafeAtomics and which have no native fallback) are covered.