.version 6.5
.target sm_30
.address_size 64

.global .s32 foobar[4] = {1};

.visible .entry global_array(
	.param .u64 input,
	.param .u64 output
)
{
	.reg .u64 	        in_addr;
    .reg .u64 	        out_addr;
    .reg .u32 	        temp;

	mov.u64 	        in_addr, foobar;
    ld.param.u64 	    out_addr, [output];

    ld.global.u32       temp, [in_addr];
    st.global.u32       [out_addr], temp;
	ret;
}