| { |
| "asm": [ |
| { |
| "labels": [], |
| "source": null, |
| "text": "//" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": "// Generated by NVIDIA NVVM Compiler" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": "//" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": "// Compiler Build ID: CL-23083092" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": "// Cuda compilation tools, release 9.1, V9.1.85" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": "// Based on LLVM 3.4svn" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": "//" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": "" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": ".version 6.1" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": ".target sm_30" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": ".address_size 64" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": "" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": " // .globl _Z6vecAddPfS_S_i" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": "" |
| }, |
| { |
| "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": " .reg .pred %p<2>;" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": " .reg .f32 %f<4>;" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": " .reg .b32 %r<6>;" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": " .reg .b64 %rd<11>;" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": "" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": " ld.param.u64 %rd1, [_Z6vecAddPfS_S_i_param_0];" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": " ld.param.u64 %rd2, [_Z6vecAddPfS_S_i_param_1];" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": " ld.param.u64 %rd3, [_Z6vecAddPfS_S_i_param_2];" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": " ld.param.u32 %r2, [_Z6vecAddPfS_S_i_param_3];" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": " .loc 1 9 12" |
| }, |
| { |
| "labels": [], |
| "source": { |
| "column": 12, |
| "file": "/tmp/moo.cu", |
| "line": 9 |
| }, |
| "text": " mov.u32 %r3, %ctaid.x;" |
| }, |
| { |
| "labels": [], |
| "source": { |
| "column": 12, |
| "file": "/tmp/moo.cu", |
| "line": 9 |
| }, |
| "text": " mov.u32 %r4, %ntid.x;" |
| }, |
| { |
| "labels": [], |
| "source": { |
| "column": 12, |
| "file": "/tmp/moo.cu", |
| "line": 9 |
| }, |
| "text": " mov.u32 %r5, %tid.x;" |
| }, |
| { |
| "labels": [], |
| "source": { |
| "column": 12, |
| "file": "/tmp/moo.cu", |
| "line": 9 |
| }, |
| "text": " mad.lo.s32 %r1, %r4, %r3, %r5;" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": " .loc 1 12 5" |
| }, |
| { |
| "labels": [], |
| "source": { |
| "column": 5, |
| "file": "/tmp/moo.cu", |
| "line": 12 |
| }, |
| "text": " setp.ge.s32 %p1, %r1, %r2;" |
| }, |
| { |
| "labels": [ |
| { |
| "name": "BB0_2", |
| "range": { |
| "endCol": 30, |
| "startCol": 25 |
| } |
| } |
| ], |
| "source": { |
| "column": 5, |
| "file": "/tmp/moo.cu", |
| "line": 12 |
| }, |
| "text": " @%p1 bra BB0_2;" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": "" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": " .loc 1 9 12" |
| }, |
| { |
| "labels": [], |
| "source": { |
| "column": 12, |
| "file": "/tmp/moo.cu", |
| "line": 9 |
| }, |
| "text": " cvta.to.global.u64 %rd4, %rd1;" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": " .loc 1 13 9" |
| }, |
| { |
| "labels": [], |
| "source": { |
| "column": 9, |
| "file": "/tmp/moo.cu", |
| "line": 13 |
| }, |
| "text": " mul.wide.s32 %rd5, %r1, 4;" |
| }, |
| { |
| "labels": [], |
| "source": { |
| "column": 9, |
| "file": "/tmp/moo.cu", |
| "line": 13 |
| }, |
| "text": " add.s64 %rd6, %rd4, %rd5;" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": " .loc 1 9 12" |
| }, |
| { |
| "labels": [], |
| "source": { |
| "column": 12, |
| "file": "/tmp/moo.cu", |
| "line": 9 |
| }, |
| "text": " cvta.to.global.u64 %rd7, %rd2;" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": " .loc 1 13 9" |
| }, |
| { |
| "labels": [], |
| "source": { |
| "column": 9, |
| "file": "/tmp/moo.cu", |
| "line": 13 |
| }, |
| "text": " add.s64 %rd8, %rd7, %rd5;" |
| }, |
| { |
| "labels": [], |
| "source": { |
| "column": 9, |
| "file": "/tmp/moo.cu", |
| "line": 13 |
| }, |
| "text": " ld.global.f32 %f1, [%rd8];" |
| }, |
| { |
| "labels": [], |
| "source": { |
| "column": 9, |
| "file": "/tmp/moo.cu", |
| "line": 13 |
| }, |
| "text": " ld.global.f32 %f2, [%rd6];" |
| }, |
| { |
| "labels": [], |
| "source": { |
| "column": 9, |
| "file": "/tmp/moo.cu", |
| "line": 13 |
| }, |
| "text": " add.f32 %f3, %f2, %f1;" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": " .loc 1 9 12" |
| }, |
| { |
| "labels": [], |
| "source": { |
| "column": 12, |
| "file": "/tmp/moo.cu", |
| "line": 9 |
| }, |
| "text": " cvta.to.global.u64 %rd9, %rd3;" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": " .loc 1 13 9" |
| }, |
| { |
| "labels": [], |
| "source": { |
| "column": 9, |
| "file": "/tmp/moo.cu", |
| "line": 13 |
| }, |
| "text": " add.s64 %rd10, %rd9, %rd5;" |
| }, |
| { |
| "labels": [], |
| "source": { |
| "column": 9, |
| "file": "/tmp/moo.cu", |
| "line": 13 |
| }, |
| "text": " st.global.f32 [%rd10], %f3;" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": "" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": "BB0_2:" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": " .loc 1 14 1" |
| }, |
| { |
| "labels": [], |
| "source": { |
| "column": 1, |
| "file": "/tmp/moo.cu", |
| "line": 14 |
| }, |
| "text": " ret;" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": "}" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": "" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": " .file 1 \"/tmp/moo.cu\", 1525722349, 2016" |
| }, |
| { |
| "labels": [], |
| "source": null, |
| "text": " .file 2 \"/opt/compiler-explorer/gcc-6.4.0/include/c++/6.4.0/cmath\", 1517261630, 47676" |
| } |
| ], |
| "labelDefinitions": { |
| "BB0_2": 58, |
| "entry": 15 |
| } |
| } |