1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18
.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; }