.version 8.4
.target sm_75
.address_size 64

.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 {%ra0, %ra1}, [in_a];
    ld.param.v2.u64 {%ra2, %ra3}, [in_a + 16];
    ld.param.v2.u64 {%rb0, %rb1}, [in_b];
    ld.param.v2.u64 {%rb2, %rb3}, [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], {%ra0, %ra1};
    st.v2.u64 [%rdc + 16], {%ra2, %ra3};

    ret;
}