.version 6.5
.target sm_30
.address_size 64

.extern .shared .align 4 .b32 shared_mem[];

.func incr_shared_2_global()
{
    .reg .u64 	    temp;
    ld.shared.u64   temp, [shared_mem];
    add.u64         temp, temp, 2;
    st.shared.u64    [shared_mem], temp;
    ret;
}


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

	ld.param.u64 	    in_addr, [input];
    ld.param.u64 	    out_addr, [output];

    ld.global.u64       temp, [in_addr];
    st.shared.u64       [shared_mem], temp;
    call incr_shared_2_global;
    ld.shared.u64       temp, [shared_mem];
    st.global.u64       [out_addr], temp;
	ret;
}