.version 6.5
.target sm_30
.address_size 64

.visible .entry pred_not(
	.param .u64 input,
	.param .u64 output
)
{
	.reg .u64 	    in_addr;
    .reg .u64 	    out_addr;
    .reg .u64 	    temp;
    .reg .u64 	    temp2;
    .reg .u64 	    temp3;
    .reg .pred 	    pred;

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

    ld.u64          temp, [in_addr];
    ld.u64          temp2, [in_addr + 8];
	setp.lt.u64		pred, temp, temp2;
    not.pred        pred, pred;
    @pred mov.u64   temp3, 1;
    @!pred mov.u64  temp3, 2;
    st.u64          [out_addr], temp3;
	ret;
}
