.version 6.5
.target sm_30
.address_size 64

.visible .entry activemask(
	.param .u64 input,
	.param .u64 output
)
{
    .reg .u64 	    out_addr;
    .reg .b32 	    temp;

    ld.param.u64 	out_addr, [output];

	activemask.b32  temp;
    st.u32          [out_addr], temp;
	ret;
}