//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-35059454
// Cuda compilation tools, release 12.6, V12.6.85
// Based on NVVM 7.0.1
//
.version 8.5
.target sm_89, debug
.address_size 64
// .globl vector_add_scalar
.visible .entry vector_add_scalar(
.param .u64 vector_add_scalar_param_0,
.param .f64 vector_add_scalar_param_1,
.param .u32 vector_add_scalar_param_2
)
{
.reg .pred %p<3>;
.reg .b32 %r<7>;
.reg .f64 %fd<4>;
.reg .b64 %rd<5>;
.loc 1 4 0
$L__func_begin0:
.loc 1 4 0
ld.param.u64 %rd1, [vector_add_scalar_param_0];
ld.param.f64 %fd1, [vector_add_scalar_param_1];
ld.param.u32 %r2, [vector_add_scalar_param_2];
$L__tmp0:
.loc 1 5 5
mov.u32 %r3, %ctaid.x;
mov.u32 %r4, %ntid.x;
mul.lo.s32 %r5, %r3, %r4;
mov.u32 %r6, %tid.x;
add.s32 %r1, %r5, %r6;
$L__tmp1:
.loc 1 6 5
setp.lt.s32 %p1, %r1, %r2;
not.pred %p2, %p1;
@%p2 bra $L__BB0_2;
bra.uni $L__BB0_1;
$L__BB0_1:
$L__tmp2:
.loc 1 7 9
cvt.s64.s32 %rd2, %r1;
shl.b64 %rd3, %rd2, 3;
add.s64 %rd4, %rd1, %rd3;
ld.f64 %fd2, [%rd4];
add.f64 %fd3, %fd2, %fd1;
st.f64 [%rd4], %fd3;
bra.uni $L__BB0_2;
$L__tmp3:
$L__BB0_2:
.loc 1 9 1
ret;
$L__tmp4:
$L__func_end0:
}
.file 1 "/home/reatank/MOCHA/hce-superopt/ptx/samples/small_kernels/vector_add_scalar/vector_add_scalar.cu"
.section .debug_abbrev
{
.b8 1
.b8 17
.b8 1
.b8 37
.b8 8
.b8 19
.b8 5
.b8 3
.b8 8
.b8 16
.b8 6
.b8 27
.b8 8
.b8 17
.b8 1
.b8 0
.b8 0
.b8 2
.b8 46
.b8 1
.b8 17
.b8 1
.b8 18
.b8 1
.b8 64
.b8 10
.b8 135,64
.b8 8
.b8 3
.b8 8
.b8 58
.b8 11
.b8 59
.b8 11
.b8 73
.b8 19
.b8 63
.b8 12
.b8 0
.b8 0
.b8 3
.b8 5
.b8 0
.b8 2
.b8 10
.b8 51
.b8 11
.b8 3
.b8 8
.b8 58
.b8 11
.b8 59
.b8 11
.b8 73
.b8 19
.b8 0
.b8 0
.b8 4
.b8 11
.b8 1
.b8 17
.b8 1
.b8 18
.b8 1
.b8 0
.b8 0
.b8 5
.b8 52
.b8 0
.b8 2
.b8 10
.b8 51
.b8 11
.b8 3
.b8 8
.b8 58
.b8 11
.b8 59
.b8 11
.b8 73
.b8 19
.b8 0
.b8 0
.b8 6
.b8 59
.b8 0
.b8 3
.b8 8
.b8 0
.b8 0
.b8 7
.b8 15
.b8 0
.b8 73
.b8 19
.b8 51
.b8 6
.b8 0
.b8 0
.b8 8
.b8 36
.b8 0
.b8 3
.b8 8
.b8 62
.b8 11
.b8 11
.b8 11
.b8 0
.b8 0
.b8 0
}
.section .debug_info
{
.b32 332
.b8 2
.b8 0
.b32 .debug_abbrev
.b8 8
.b8 1
.b8 108,103,101,110,102,101,58,32,69,68,71,32,54,46,54
.b8 0
.b8 4
.b8 0
.b8 118,101,99,116,111,114,95,97,100,100,95,115,99,97,108,97,114,46,99,117
.b8 0
.b32 .debug_line
.b8 47,104,111,109,101,47,114,101,97,116,97,110,107,47,77,79,67,72,65,47,104,99,101,45,115,117,112,101,114,111,112,116,47,112,116,120,47,115,97,109
.b8 112,108,101,115,47,115,109,97,108,108,95,107,101,114,110,101,108,115,47,118,101,99,116,111,114,95,97,100,100,95,115,99,97,108,97,114
.b8 0
.b64 0
.b8 2
.b64 $L__func_begin0
.b64 $L__func_end0
.b8 1
.b8 156
.b8 118,101,99,116,111,114,95,97,100,100,95,115,99,97,108,97,114
.b8 0
.b8 118,101,99,116,111,114,95,97,100,100,95,115,99,97,108,97,114
.b8 0
.b8 1
.b8 4
.b32 303
.b8 1
.b8 3
.b8 9
.b8 3
.b64 vector_add_scalar_param_0
.b8 7
.b8 100,97,116,97
.b8 0
.b8 1
.b8 4
.b32 309
.b8 3
.b8 6
.b8 144
.b8 177
.b8 200
.b8 153
.b8 171
.b8 2
.b8 2
.b8 97,108,112,104,97
.b8 0
.b8 1
.b8 4
.b32 318
.b8 3
.b8 9
.b8 3
.b64 vector_add_scalar_param_2
.b8 7
.b8 110
.b8 0
.b8 1
.b8 4
.b32 328
.b8 4
.b64 $L__tmp0
.b64 $L__tmp4
.b8 5
.b8 5
.b8 144
.b8 177
.b8 228
.b8 149
.b8 1
.b8 2
.b8 105,100,120
.b8 0
.b8 1
.b8 5
.b32 328
.b8 0
.b8 0
.b8 6
.b8 118,111,105,100
.b8 0
.b8 7
.b32 318
.b32 12
.b8 8
.b8 100,111,117,98,108,101
.b8 0
.b8 4
.b8 8
.b8 8
.b8 105,110,116
.b8 0
.b8 5
.b8 4
.b8 0
}
.section .debug_macinfo
{
.b8 0
}