Kernel arguments with `#[repr(align(16))]` cause misaligned access / `IllegalAddress`

### Summary I encountered runtime `IllegalAddress` errors when passing 16-byte aligned structs (e.g., structs containing `u128` or `#[repr(align(16))]`) by value to CUDA kernels. The root cause is in `rustc_codegen_nvvm/src/abi.rs`. The current logic forces `PassMode::Direct` for all ADTs. However, for the NVPTX target, LLVM often defaults to 8-byte alignment for generic structs passed directly, even if the Rust type requires 16-byte alignment. This results in the Host packing arguments at 16-byte boundaries, but the Device (PTX) reading them at 8-byte boundaries, leading to misaligned reads and crashes. ### Reproduction Case Define a struct with 16-byte alignment and pass it to a kernel: ``` #[repr(C, align(16))] pub struct AlignedStruct { a: u64, b: u64, } #[cuda_std::kernel] pub unsafe fn my_kernel(input: AlignedStruct) { // Accessing input.a or input.b causes IllegalAddress // because the kernel expects 8-byte alignment but the pointer // generated by the driver is 16-byte aligned (or vice versa depending on padding). } ``` ### Issue In `rustc_codegen_nvvm/src/abi.r`s, the function `readjust_fn_abi` currently does this: ``` // Current logic if arg.layout.ty.is_adt() { // ... arg.mode = PassMode::Direct(ArgAttributes::new()); } ``` This generates PTX kernel parameters like `.param .align 8 .b8 input[16]`. However, because the struct contains aligned data, the kernel code may generate vector loads (e.g.,` ld.global.v2.u64` or 128-bit loads) which trap if the address is not 16-byte aligned.