Large valued arguments cause local memory overhead

```
julia> using CUDA
^[[A^[[A
julia> kernel(a,data) = (@inbounds a[threadIdx().x] = data[threadIdx().x]; return)
kernel (generic function with 1 method)

julia> CUDA.code_llvm(kernel, Tuple{CuDeviceArray{Int,1,AS.Global}, NTuple{256,Int}}; debuginfo=:none, kernel=true)

define dso_local ptx_kernel void @_Z17julia_kernel_216913CuDeviceArrayI5Int64Li1E6GlobalE5TupleI5Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int64E({ [1 x i64], i64 }* nocapture nonnull readonly byval dereferenceable(16), [256 x i64]* nocapture nonnull readonly byval dereferenceable(2048)) local_unnamed_addr {
top:
%2 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
%3 = zext i32 %2 to i64
%4 = getelementptr inbounds [256 x i64], [256 x i64]* %1, i64 0, i64 %3
%5 = getelementptr inbounds { [1 x i64], i64 }, { [1 x i64], i64 }* %0, i64 0, i32 1
%6 = bitcast i64* %5 to i64 addrspace(1)**
%7 = load i64 addrspace(1)*, i64 addrspace(1)** %6, align 8
%8 = load i64, i64* %4, align 8
%9 = getelementptr i64, i64 addrspace(1)* %7, i64 %3
store i64 %8, i64 addrspace(1)* %9, align 8
ret void
}

julia> CUDA.code_ptx(kernel, Tuple{CuDeviceArray{Int,1,AS.Global}, NTuple{32,Int}}; kernel=true)
//
// Generated by LLVM NVPTX Back-End
//

.version 6.3
.target sm_75
.address_size 64

// .globl _Z17julia_kernel_217213CuDeviceArrayI5Int64Li1E6GlobalE5TupleI5Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int64E // -- Begin function _Z17julia_kernel_217213CuDeviceArrayI5Int64Li1E6GlobalE5TupleI5Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int64E
.weak .global .align 8 .u64 exception_flag;
// @_Z17julia_kernel_217213CuDeviceArrayI5Int64Li1E6GlobalE5TupleI5Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int64E
.visible .entry _Z17julia_kernel_217213CuDeviceArrayI5Int64Li1E6GlobalE5TupleI5Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int64E(
.param .align 8 .b8 _Z17julia_kernel_217213CuDeviceArrayI5Int64Li1E6GlobalE5TupleI5Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int64E_param_0[16],
.param .align 8 .b8 _Z17julia_kernel_217213CuDeviceArrayI5Int64Li1E6GlobalE5TupleI5Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int64E_param_1[256]
)
{
.local .align 8 .b8 __local_depot0[256];
.reg .b64 %SP;
.reg .b64 %SPL;
.reg .b32 %r<2>;
.reg .b64 %rd<42>;

// %bb.0: // %top
mov.u64 %SPL, __local_depot0;
mov.b64 %rd1, _Z17julia_kernel_217213CuDeviceArrayI5Int64Li1E6GlobalE5TupleI5Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int64E_param_0;
mov.b64 %rd2, _Z17julia_kernel_217213CuDeviceArrayI5Int64Li1E6GlobalE5TupleI5Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int645Int64E_param_1;
add.u64 %rd4, %SPL, 0;
ld.param.u64 %rd5, [%rd2+248];
ld.param.u64 %rd6, [%rd2+240];
ld.param.u64 %rd7, [%rd2+232];
ld.param.u64 %rd8, [%rd2+224];
ld.param.u64 %rd9, [%rd2+216];
ld.param.u64 %rd10, [%rd2+208];
ld.param.u64 %rd11, [%rd2+200];
ld.param.u64 %rd12, [%rd2+192];
ld.param.u64 %rd13, [%rd2+184];
ld.param.u64 %rd14, [%rd2+176];
ld.param.u64 %rd15, [%rd2+168];
ld.param.u64 %rd16, [%rd2+160];
ld.param.u64 %rd17, [%rd2+152];
ld.param.u64 %rd18, [%rd2+144];
ld.param.u64 %rd19, [%rd2+136];
ld.param.u64 %rd20, [%rd2+128];
ld.param.u64 %rd21, [%rd2+120];
ld.param.u64 %rd22, [%rd2+112];
ld.param.u64 %rd23, [%rd2+104];
ld.param.u64 %rd24, [%rd2+96];
ld.param.u64 %rd25, [%rd2+88];
ld.param.u64 %rd26, [%rd2+80];
ld.param.u64 %rd27, [%rd2+72];
ld.param.u64 %rd28, [%rd2+64];
ld.param.u64 %rd29, [%rd2+56];
ld.param.u64 %rd30, [%rd2+48];
ld.param.u64 %rd31, [%rd2+40];
ld.param.u64 %rd32, [%rd2+32];
ld.param.u64 %rd33, [%rd2+24];
ld.param.u64 %rd34, [%rd2+16];
ld.param.u64 %rd35, [%rd2+8];
ld.param.u64 %rd36, [%rd2];
st.local.u64 [%rd4], %rd36;
st.local.u64 [%rd4+8], %rd35;
st.local.u64 [%rd4+16], %rd34;
st.local.u64 [%rd4+24], %rd33;
st.local.u64 [%rd4+32], %rd32;
st.local.u64 [%rd4+40], %rd31;
st.local.u64 [%rd4+48], %rd30;
st.local.u64 [%rd4+56], %rd29;
st.local.u64 [%rd4+64], %rd28;
st.local.u64 [%rd4+72], %rd27;
st.local.u64 [%rd4+80], %rd26;
st.local.u64 [%rd4+88], %rd25;
st.local.u64 [%rd4+96], %rd24;
st.local.u64 [%rd4+104], %rd23;
st.local.u64 [%rd4+112], %rd22;
st.local.u64 [%rd4+120], %rd21;
st.local.u64 [%rd4+128], %rd20;
st.local.u64 [%rd4+136], %rd19;
st.local.u64 [%rd4+144], %rd18;
st.local.u64 [%rd4+152], %rd17;
st.local.u64 [%rd4+160], %rd16;
st.local.u64 [%rd4+168], %rd15;
st.local.u64 [%rd4+176], %rd14;
st.local.u64 [%rd4+184], %rd13;
st.local.u64 [%rd4+192], %rd12;
st.local.u64 [%rd4+200], %rd11;
st.local.u64 [%rd4+208], %rd10;
st.local.u64 [%rd4+216], %rd9;
st.local.u64 [%rd4+224], %rd8;
st.local.u64 [%rd4+232], %rd7;
st.local.u64 [%rd4+240], %rd6;
st.local.u64 [%rd4+248], %rd5;
ld.param.u64 %rd37, [%rd1+8];
mov.u32 %r1, %tid.x;
mul.wide.u32 %rd38, %r1, 8;
add.s64 %rd39, %rd4, %rd38;
ld.local.u64 %rd40, [%rd39];
add.s64 %rd41, %rd37, %rd38;
st.global.u64 [%rd41], %rd40;
ret;
// -- End function
}
```