.version 8.4
.target sm_75
.address_size 64

.visible .func add_u16(
    .param .b64 out_c,
    .param .align 16 .b8 in_a[16],
    .param .align 16 .b8 in_b[16]
) {
    .reg .u64 %ra<2>, %rb<2>;
    .reg .b64 %rdc;
    
    ld.param.b64 %rdc, [out_c];
    
    ld.param.v2.u64 {%ra1, %ra0}, [in_a];
    ld.param.v2.u64 {%rb1, %rb0}, [in_b];

    add.cc.u64 %ra0, %ra0, %rb0;
    addc.u64 %ra1, %ra1, %rb1;

    st.v2.u64 [%rdc], {%ra1, %ra0};

    ret;
}

.visible .func sub_u16(
    .param .b64 out_c,
    .param .align 16 .b8 in_a[16],
    .param .align 16 .b8 in_b[16]
) {
    .reg .u64 %ra<2>, %rb<2>;
    .reg .b64 %rdc;
    
    ld.param.b64 %rdc, [out_c];
    
    ld.param.v2.u64 {%ra1, %ra0}, [in_a];
    ld.param.v2.u64 {%rb1, %rb0}, [in_b];

    add.cc.u64 %ra0, %ra0, %rb0;
    addc.u64 %ra1, %ra1, %rb1;

    st.v2.u64 [%rdc], {%ra1, %ra0};

    ret;
}


.visible .func add_u32(
    .param .b64 out_c,
    .param .align 16 .b8 in_a[32],
    .param .align 16 .b8 in_b[32]
) {
    .reg .u64 %ra<4>, %rb<4>;
    .reg .b64 %rdc;
    
    ld.param.b64 %rdc, [out_c];
    
    ld.param.v2.u64 {%ra3, %ra2}, [in_a];
    ld.param.v2.u64 {%ra1, %ra0}, [in_a + 16];
    ld.param.v2.u64 {%rb3, %rb2}, [in_b];
    ld.param.v2.u64 {%rb1, %rb0}, [in_b + 16];

    add.cc.u64 %ra0, %ra0, %rb0;
    addc.cc.u64 %ra1, %ra1, %rb1;
    addc.cc.u64 %ra2, %ra2, %rb2;
    addc.u64 %ra3, %ra3, %rb3;

    st.v2.u64 [%rdc], {%ra3, %ra2};
    st.v2.u64 [%rdc + 16], {%ra1, %ra0};

    ret;
}

.visible .func sub_u32(
    .param .b64 out_c,
    .param .align 16 .b8 in_a[32],
    .param .align 16 .b8 in_b[32]
) {
    .reg .u64 %ra<4>, %rb<4>;
    .reg .b64 %rdc;
    
    ld.param.b64 %rdc, [out_c];
    
    ld.param.v2.u64 {%ra3, %ra2}, [in_a];
    ld.param.v2.u64 {%ra1, %ra0}, [in_a + 16];
    ld.param.v2.u64 {%rb3, %rb2}, [in_b];
    ld.param.v2.u64 {%rb1, %rb0}, [in_b + 16];

    sub.cc.u64 %ra0, %ra0, %rb0;
    subc.cc.u64 %ra1, %ra1, %rb1;
    subc.cc.u64 %ra2, %ra2, %rb2;
    subc.u64 %ra3, %ra3, %rb3;

    st.v2.u64 [%rdc], {%ra3, %ra2};
    st.v2.u64 [%rdc + 16], {%ra1, %ra0};

    ret;
}