blob: 44e474ba6ad7d8c895d0831ee81c46f7f7f33de8 [file] [log] [blame] [raw]
{
"asm": [
{
"labels": [],
"source": null,
"text": "//"
},
{
"labels": [],
"source": null,
"text": "// Generated by LLVM NVPTX Back-End"
},
{
"labels": [],
"source": null,
"text": "//"
},
{
"labels": [],
"source": null,
"text": ""
},
{
"labels": [],
"source": null,
"text": ".version 6.1"
},
{
"labels": [],
"source": null,
"text": ".target sm_35"
},
{
"labels": [],
"source": null,
"text": ".address_size 64"
},
{
"labels": [],
"source": null,
"text": ""
},
{
"labels": [],
"source": null,
"text": " // .globl cudaMalloc"
},
{
"labels": [],
"source": null,
"text": ".global .align 1 .b8 blockIdx[1];"
},
{
"labels": [],
"source": null,
"text": ".global .align 1 .b8 blockDim[1];"
},
{
"labels": [],
"source": null,
"text": ".global .align 1 .b8 threadIdx[1];"
},
{
"labels": [],
"source": null,
"text": ""
},
{
"labels": [],
"source": null,
"text": ".visible .func (.param .b32 func_retval0) cudaMalloc("
},
{
"labels": [],
"source": null,
"text": " .param .b64 cudaMalloc_param_0,"
},
{
"labels": [],
"source": null,
"text": " .param .b64 cudaMalloc_param_1"
},
{
"labels": [],
"source": null,
"text": ")"
},
{
"labels": [],
"source": null,
"text": "{"
},
{
"labels": [],
"source": null,
"text": " .local .align 8 .b8 __local_depot0[16];"
},
{
"labels": [],
"source": null,
"text": " .reg .b64 %SP;"
},
{
"labels": [],
"source": null,
"text": " .reg .b64 %SPL;"
},
{
"labels": [],
"source": null,
"text": " .reg .b32 %r<2>;"
},
{
"labels": [],
"source": null,
"text": " .reg .b64 %rd<3>;"
},
{
"labels": [],
"source": null,
"text": ""
},
{
"labels": [],
"source": null,
"text": " mov.u64 %SPL, __local_depot0;"
},
{
"labels": [],
"source": null,
"text": " cvta.local.u64 %SP, %SPL;"
},
{
"labels": [],
"source": null,
"text": " ld.param.u64 %rd2, [cudaMalloc_param_1];"
},
{
"labels": [],
"source": null,
"text": " ld.param.u64 %rd1, [cudaMalloc_param_0];"
},
{
"labels": [],
"source": null,
"text": " st.u64 [%SP+0], %rd1;"
},
{
"labels": [],
"source": null,
"text": " st.u64 [%SP+8], %rd2;"
},
{
"labels": [],
"source": null,
"text": " mov.u32 %r1, 30;"
},
{
"labels": [],
"source": null,
"text": " st.param.b32 [func_retval0+0], %r1;"
},
{
"labels": [],
"source": null,
"text": " ret;"
},
{
"labels": [],
"source": null,
"text": "}"
},
{
"labels": [],
"source": null,
"text": ""
},
{
"labels": [],
"source": null,
"text": " // .globl cudaFuncGetAttributes"
},
{
"labels": [],
"source": null,
"text": ".visible .func (.param .b32 func_retval0) cudaFuncGetAttributes("
},
{
"labels": [],
"source": null,
"text": " .param .b64 cudaFuncGetAttributes_param_0,"
},
{
"labels": [],
"source": null,
"text": " .param .b64 cudaFuncGetAttributes_param_1"
},
{
"labels": [],
"source": null,
"text": ")"
},
{
"labels": [],
"source": null,
"text": "{"
},
{
"labels": [],
"source": null,
"text": " .local .align 8 .b8 __local_depot1[16];"
},
{
"labels": [],
"source": null,
"text": " .reg .b64 %SP;"
},
{
"labels": [],
"source": null,
"text": " .reg .b64 %SPL;"
},
{
"labels": [],
"source": null,
"text": " .reg .b32 %r<2>;"
},
{
"labels": [],
"source": null,
"text": " .reg .b64 %rd<3>;"
},
{
"labels": [],
"source": null,
"text": ""
},
{
"labels": [],
"source": null,
"text": " mov.u64 %SPL, __local_depot1;"
},
{
"labels": [],
"source": null,
"text": " cvta.local.u64 %SP, %SPL;"
},
{
"labels": [],
"source": null,
"text": " ld.param.u64 %rd2, [cudaFuncGetAttributes_param_1];"
},
{
"labels": [],
"source": null,
"text": " ld.param.u64 %rd1, [cudaFuncGetAttributes_param_0];"
},
{
"labels": [],
"source": null,
"text": " st.u64 [%SP+0], %rd1;"
},
{
"labels": [],
"source": null,
"text": " st.u64 [%SP+8], %rd2;"
},
{
"labels": [],
"source": null,
"text": " mov.u32 %r1, 30;"
},
{
"labels": [],
"source": null,
"text": " st.param.b32 [func_retval0+0], %r1;"
},
{
"labels": [],
"source": null,
"text": " ret;"
},
{
"labels": [],
"source": null,
"text": "}"
},
{
"labels": [],
"source": null,
"text": ""
},
{
"labels": [],
"source": null,
"text": " // .globl cudaDeviceGetAttribute"
},
{
"labels": [],
"source": null,
"text": ".visible .func (.param .b32 func_retval0) cudaDeviceGetAttribute("
},
{
"labels": [],
"source": null,
"text": " .param .b64 cudaDeviceGetAttribute_param_0,"
},
{
"labels": [],
"source": null,
"text": " .param .b32 cudaDeviceGetAttribute_param_1,"
},
{
"labels": [],
"source": null,
"text": " .param .b32 cudaDeviceGetAttribute_param_2"
},
{
"labels": [],
"source": null,
"text": ")"
},
{
"labels": [],
"source": null,
"text": "{"
},
{
"labels": [],
"source": null,
"text": " .local .align 8 .b8 __local_depot2[16];"
},
{
"labels": [],
"source": null,
"text": " .reg .b64 %SP;"
},
{
"labels": [],
"source": null,
"text": " .reg .b64 %SPL;"
},
{
"labels": [],
"source": null,
"text": " .reg .b32 %r<4>;"
},
{
"labels": [],
"source": null,
"text": " .reg .b64 %rd<2>;"
},
{
"labels": [],
"source": null,
"text": ""
},
{
"labels": [],
"source": null,
"text": " mov.u64 %SPL, __local_depot2;"
},
{
"labels": [],
"source": null,
"text": " cvta.local.u64 %SP, %SPL;"
},
{
"labels": [],
"source": null,
"text": " ld.param.u32 %r2, [cudaDeviceGetAttribute_param_2];"
},
{
"labels": [],
"source": null,
"text": " ld.param.u32 %r1, [cudaDeviceGetAttribute_param_1];"
},
{
"labels": [],
"source": null,
"text": " ld.param.u64 %rd1, [cudaDeviceGetAttribute_param_0];"
},
{
"labels": [],
"source": null,
"text": " st.u64 [%SP+0], %rd1;"
},
{
"labels": [],
"source": null,
"text": " st.u32 [%SP+8], %r1;"
},
{
"labels": [],
"source": null,
"text": " st.u32 [%SP+12], %r2;"
},
{
"labels": [],
"source": null,
"text": " mov.u32 %r3, 30;"
},
{
"labels": [],
"source": null,
"text": " st.param.b32 [func_retval0+0], %r3;"
},
{
"labels": [],
"source": null,
"text": " ret;"
},
{
"labels": [],
"source": null,
"text": "}"
},
{
"labels": [],
"source": null,
"text": ""
},
{
"labels": [],
"source": null,
"text": " // .globl cudaGetDevice"
},
{
"labels": [],
"source": null,
"text": ".visible .func (.param .b32 func_retval0) cudaGetDevice("
},
{
"labels": [],
"source": null,
"text": " .param .b64 cudaGetDevice_param_0"
},
{
"labels": [],
"source": null,
"text": ")"
},
{
"labels": [],
"source": null,
"text": "{"
},
{
"labels": [],
"source": null,
"text": " .local .align 8 .b8 __local_depot3[8];"
},
{
"labels": [],
"source": null,
"text": " .reg .b64 %SP;"
},
{
"labels": [],
"source": null,
"text": " .reg .b64 %SPL;"
},
{
"labels": [],
"source": null,
"text": " .reg .b32 %r<2>;"
},
{
"labels": [],
"source": null,
"text": " .reg .b64 %rd<2>;"
},
{
"labels": [],
"source": null,
"text": ""
},
{
"labels": [],
"source": null,
"text": " mov.u64 %SPL, __local_depot3;"
},
{
"labels": [],
"source": null,
"text": " cvta.local.u64 %SP, %SPL;"
},
{
"labels": [],
"source": null,
"text": " ld.param.u64 %rd1, [cudaGetDevice_param_0];"
},
{
"labels": [],
"source": null,
"text": " st.u64 [%SP+0], %rd1;"
},
{
"labels": [],
"source": null,
"text": " mov.u32 %r1, 30;"
},
{
"labels": [],
"source": null,
"text": " st.param.b32 [func_retval0+0], %r1;"
},
{
"labels": [],
"source": null,
"text": " ret;"
},
{
"labels": [],
"source": null,
"text": "}"
},
{
"labels": [],
"source": null,
"text": ""
},
{
"labels": [],
"source": null,
"text": " // .globl cudaOccupancyMaxActiveBlocksPerMultiprocessor"
},
{
"labels": [],
"source": null,
"text": ".visible .func (.param .b32 func_retval0) cudaOccupancyMaxActiveBlocksPerMultiprocessor("
},
{
"labels": [],
"source": null,
"text": " .param .b64 cudaOccupancyMaxActiveBlocksPerMultiprocessor_param_0,"
},
{
"labels": [],
"source": null,
"text": " .param .b64 cudaOccupancyMaxActiveBlocksPerMultiprocessor_param_1,"
},
{
"labels": [],
"source": null,
"text": " .param .b32 cudaOccupancyMaxActiveBlocksPerMultiprocessor_param_2,"
},
{
"labels": [],
"source": null,
"text": " .param .b64 cudaOccupancyMaxActiveBlocksPerMultiprocessor_param_3"
},
{
"labels": [],
"source": null,
"text": ")"
},
{
"labels": [],
"source": null,
"text": "{"
},
{
"labels": [],
"source": null,
"text": " .local .align 8 .b8 __local_depot4[32];"
},
{
"labels": [],
"source": null,
"text": " .reg .b64 %SP;"
},
{
"labels": [],
"source": null,
"text": " .reg .b64 %SPL;"
},
{
"labels": [],
"source": null,
"text": " .reg .b32 %r<3>;"
},
{
"labels": [],
"source": null,
"text": " .reg .b64 %rd<4>;"
},
{
"labels": [],
"source": null,
"text": ""
},
{
"labels": [],
"source": null,
"text": " mov.u64 %SPL, __local_depot4;"
},
{
"labels": [],
"source": null,
"text": " cvta.local.u64 %SP, %SPL;"
},
{
"labels": [],
"source": null,
"text": " ld.param.u64 %rd3, [cudaOccupancyMaxActiveBlocksPerMultiprocessor_param_3];"
},
{
"labels": [],
"source": null,
"text": " ld.param.u32 %r1, [cudaOccupancyMaxActiveBlocksPerMultiprocessor_param_2];"
},
{
"labels": [],
"source": null,
"text": " ld.param.u64 %rd2, [cudaOccupancyMaxActiveBlocksPerMultiprocessor_param_1];"
},
{
"labels": [],
"source": null,
"text": " ld.param.u64 %rd1, [cudaOccupancyMaxActiveBlocksPerMultiprocessor_param_0];"
},
{
"labels": [],
"source": null,
"text": " st.u64 [%SP+0], %rd1;"
},
{
"labels": [],
"source": null,
"text": " st.u64 [%SP+8], %rd2;"
},
{
"labels": [],
"source": null,
"text": " st.u32 [%SP+16], %r1;"
},
{
"labels": [],
"source": null,
"text": " st.u64 [%SP+24], %rd3;"
},
{
"labels": [],
"source": null,
"text": " mov.u32 %r2, 30;"
},
{
"labels": [],
"source": null,
"text": " st.param.b32 [func_retval0+0], %r2;"
},
{
"labels": [],
"source": null,
"text": " ret;"
},
{
"labels": [],
"source": null,
"text": "}"
},
{
"labels": [],
"source": null,
"text": ""
},
{
"labels": [],
"source": null,
"text": " // .globl cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags"
},
{
"labels": [],
"source": null,
"text": ".visible .func (.param .b32 func_retval0) cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags("
},
{
"labels": [],
"source": null,
"text": " .param .b64 cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_param_0,"
},
{
"labels": [],
"source": null,
"text": " .param .b64 cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_param_1,"
},
{
"labels": [],
"source": null,
"text": " .param .b32 cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_param_2,"
},
{
"labels": [],
"source": null,
"text": " .param .b64 cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_param_3,"
},
{
"labels": [],
"source": null,
"text": " .param .b32 cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_param_4"
},
{
"labels": [],
"source": null,
"text": ")"
},
{
"labels": [],
"source": null,
"text": "{"
},
{
"labels": [],
"source": null,
"text": " .local .align 8 .b8 __local_depot5[40];"
},
{
"labels": [],
"source": null,
"text": " .reg .b64 %SP;"
},
{
"labels": [],
"source": null,
"text": " .reg .b64 %SPL;"
},
{
"labels": [],
"source": null,
"text": " .reg .b32 %r<4>;"
},
{
"labels": [],
"source": null,
"text": " .reg .b64 %rd<4>;"
},
{
"labels": [],
"source": null,
"text": ""
},
{
"labels": [],
"source": null,
"text": " mov.u64 %SPL, __local_depot5;"
},
{
"labels": [],
"source": null,
"text": " cvta.local.u64 %SP, %SPL;"
},
{
"labels": [],
"source": null,
"text": " ld.param.u32 %r2, [cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_param_4];"
},
{
"labels": [],
"source": null,
"text": " ld.param.u64 %rd3, [cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_param_3];"
},
{
"labels": [],
"source": null,
"text": " ld.param.u32 %r1, [cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_param_2];"
},
{
"labels": [],
"source": null,
"text": " ld.param.u64 %rd2, [cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_param_1];"
},
{
"labels": [],
"source": null,
"text": " ld.param.u64 %rd1, [cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_param_0];"
},
{
"labels": [],
"source": null,
"text": " st.u64 [%SP+0], %rd1;"
},
{
"labels": [],
"source": null,
"text": " st.u64 [%SP+8], %rd2;"
},
{
"labels": [],
"source": null,
"text": " st.u32 [%SP+16], %r1;"
},
{
"labels": [],
"source": null,
"text": " st.u64 [%SP+24], %rd3;"
},
{
"labels": [],
"source": null,
"text": " st.u32 [%SP+32], %r2;"
},
{
"labels": [],
"source": null,
"text": " mov.u32 %r3, 30;"
},
{
"labels": [],
"source": null,
"text": " st.param.b32 [func_retval0+0], %r3;"
},
{
"labels": [],
"source": null,
"text": " ret;"
},
{
"labels": [],
"source": null,
"text": "}"
},
{
"labels": [],
"source": null,
"text": ""
},
{
"labels": [],
"source": null,
"text": " // .globl _Z6vecAddPfS_S_i"
},
{
"labels": [],
"source": null,
"text": ".visible .entry _Z6vecAddPfS_S_i("
},
{
"labels": [],
"source": null,
"text": " .param .u64 _Z6vecAddPfS_S_i_param_0,"
},
{
"labels": [],
"source": null,
"text": " .param .u64 _Z6vecAddPfS_S_i_param_1,"
},
{
"labels": [],
"source": null,
"text": " .param .u64 _Z6vecAddPfS_S_i_param_2,"
},
{
"labels": [],
"source": null,
"text": " .param .u32 _Z6vecAddPfS_S_i_param_3"
},
{
"labels": [],
"source": null,
"text": ")"
},
{
"labels": [],
"source": null,
"text": "{"
},
{
"labels": [],
"source": null,
"text": " .local .align 8 .b8 __local_depot6[32];"
},
{
"labels": [],
"source": null,
"text": " .reg .b64 %SP;"
},
{
"labels": [],
"source": null,
"text": " .reg .b64 %SPL;"
},
{
"labels": [],
"source": null,
"text": " .reg .pred %p<2>;"
},
{
"labels": [],
"source": null,
"text": " .reg .f32 %f<4>;"
},
{
"labels": [],
"source": null,
"text": " .reg .b32 %r<9>;"
},
{
"labels": [],
"source": null,
"text": " .reg .b64 %rd<18>;"
},
{
"labels": [],
"source": null,
"text": ""
},
{
"labels": [],
"source": null,
"text": " mov.u64 %SPL, __local_depot6;"
},
{
"labels": [],
"source": null,
"text": " cvta.local.u64 %SP, %SPL;"
},
{
"labels": [],
"source": null,
"text": " ld.param.u32 %r1, [_Z6vecAddPfS_S_i_param_3];"
},
{
"labels": [],
"source": null,
"text": " ld.param.u64 %rd3, [_Z6vecAddPfS_S_i_param_2];"
},
{
"labels": [],
"source": null,
"text": " ld.param.u64 %rd2, [_Z6vecAddPfS_S_i_param_1];"
},
{
"labels": [],
"source": null,
"text": " ld.param.u64 %rd1, [_Z6vecAddPfS_S_i_param_0];"
},
{
"labels": [],
"source": null,
"text": " cvta.to.global.u64 %rd4, %rd3;"
},
{
"labels": [],
"source": null,
"text": " cvta.global.u64 %rd5, %rd4;"
},
{
"labels": [],
"source": null,
"text": " cvta.to.global.u64 %rd6, %rd2;"
},
{
"labels": [],
"source": null,
"text": " cvta.global.u64 %rd7, %rd6;"
},
{
"labels": [],
"source": null,
"text": " cvta.to.global.u64 %rd8, %rd1;"
},
{
"labels": [],
"source": null,
"text": " cvta.global.u64 %rd9, %rd8;"
},
{
"labels": [],
"source": null,
"text": " st.u64 [%SP+0], %rd9;"
},
{
"labels": [],
"source": null,
"text": " st.u64 [%SP+8], %rd7;"
},
{
"labels": [],
"source": null,
"text": " st.u64 [%SP+16], %rd5;"
},
{
"labels": [],
"source": null,
"text": " st.u32 [%SP+24], %r1;"
},
{
"labels": [],
"source": null,
"text": " mov.u32 %r2, %ctaid.x;"
},
{
"labels": [],
"source": null,
"text": " mov.u32 %r3, %ntid.x;"
},
{
"labels": [],
"source": null,
"text": " mul.lo.s32 %r4, %r2, %r3;"
},
{
"labels": [],
"source": null,
"text": " mov.u32 %r5, %tid.x;"
},
{
"labels": [],
"source": null,
"text": " add.s32 %r6, %r4, %r5;"
},
{
"labels": [],
"source": null,
"text": " st.u32 [%SP+28], %r6;"
},
{
"labels": [],
"source": null,
"text": " ld.u32 %r7, [%SP+28];"
},
{
"labels": [],
"source": null,
"text": " ld.u32 %r8, [%SP+24];"
},
{
"labels": [],
"source": null,
"text": " setp.ge.s32 %p1, %r7, %r8;"
},
{
"labels": [
{
"name": "LBB6_2",
"range": {
"endCol": 31,
"startCol": 25
}
}
],
"source": null,
"text": " @%p1 bra LBB6_2;"
},
{
"labels": [
{
"name": "LBB6_1",
"range": {
"endCol": 31,
"startCol": 25
}
}
],
"source": null,
"text": " bra.uni LBB6_1;"
},
{
"labels": [],
"source": null,
"text": "LBB6_1:"
},
{
"labels": [],
"source": null,
"text": " ld.u64 %rd10, [%SP+0];"
},
{
"labels": [],
"source": null,
"text": " ld.s32 %rd11, [%SP+28];"
},
{
"labels": [],
"source": null,
"text": " shl.b64 %rd12, %rd11, 2;"
},
{
"labels": [],
"source": null,
"text": " add.s64 %rd13, %rd10, %rd12;"
},
{
"labels": [],
"source": null,
"text": " ld.f32 %f1, [%rd13];"
},
{
"labels": [],
"source": null,
"text": " ld.u64 %rd14, [%SP+8];"
},
{
"labels": [],
"source": null,
"text": " add.s64 %rd15, %rd14, %rd12;"
},
{
"labels": [],
"source": null,
"text": " ld.f32 %f2, [%rd15];"
},
{
"labels": [],
"source": null,
"text": " add.rn.f32 %f3, %f1, %f2;"
},
{
"labels": [],
"source": null,
"text": " ld.u64 %rd16, [%SP+16];"
},
{
"labels": [],
"source": null,
"text": " add.s64 %rd17, %rd16, %rd12;"
},
{
"labels": [],
"source": null,
"text": " st.f32 [%rd17], %f3;"
},
{
"labels": [
{
"name": "LBB6_2",
"range": {
"endCol": 31,
"startCol": 25
}
}
],
"source": null,
"text": " bra.uni LBB6_2;"
},
{
"labels": [],
"source": null,
"text": "LBB6_2:"
},
{
"labels": [],
"source": null,
"text": " ret;"
},
{
"labels": [],
"source": null,
"text": "}"
}
],
"labelDefinitions": {
"LBB6_1": 209,
"LBB6_2": 223,
"entry": 167,
"func": 135
}
}