PTX: Conditional reference-to-value rewrite

[maleadt]

We should not unconditionally rewrite all parameters to be passed by-value, as there's actually a limit:

__global__ function parameters are passed to the device via constant memory and are limited to 32,764 bytes starting with Volta, and 4 KB on older architectures.

Not doing the rewrite requires that CUDA.jl allocates a box on the GPU that contains the argument. This also needs to be a stream-ordered allocation, in order to respect the kernel's lifetime.