API tools faq
paste
Guest User

sm_20

a guest
Jun 12th, 2012
224
0
Never
Add comment
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
ASM (NASM) 15.55 KB | None | 0 0
raw download clone embed print report
  1. //
  2. // Generated by NVIDIA NVVM Compiler
  3. // Compiler built on Thu Jan 12 22:46:01 2012 (1326408361)
  4. // Cuda compilation tools, release 4.1, V0.2.1221
  5. //
  6.  
  7. .version 3.0
  8. .target sm_20
  9. .address_size 64
  10.  
  11.     .file   1 "/tmp/tmpxft_000021b9_00000000-9_d2q9.cpp3.i"
  12.     .file   2 "d2q9.cu"
  13.     .file   3 "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/device_functions.h"
  14.     .file   4 "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/nvvm/ci_include.h"
  15. // __cuda_local_var_17168_35_non_const_sdata has been demoted
  16.  
  17. .entry _Z9collisionPK7t_speedPKcPS_Pfi(
  18.     .param .u64 _Z9collisionPK7t_speedPKcPS_Pfi_param_0,
  19.     .param .u64 _Z9collisionPK7t_speedPKcPS_Pfi_param_1,
  20.     .param .u64 _Z9collisionPK7t_speedPKcPS_Pfi_param_2,
  21.     .param .u64 _Z9collisionPK7t_speedPKcPS_Pfi_param_3,
  22.     .param .u32 _Z9collisionPK7t_speedPKcPS_Pfi_param_4
  23. )
  24. {
  25.     .reg .f32   %f<139>;
  26.     .reg .f64   %fd<97>;
  27.     .reg .pred  %p<13>;
  28.     .reg .s32   %r<68>;
  29.     .reg .s64   %rl<21>;
  30.     .reg .s16   %rc<2>;
  31.     // demoted variable
  32.     .shared .align 4 .b8 __cuda_local_var_17168_35_non_const_sdata[128];
  33.  
  34.     ld.param.u64    %rl4, [_Z9collisionPK7t_speedPKcPS_Pfi_param_0];
  35.     ld.param.u64    %rl5, [_Z9collisionPK7t_speedPKcPS_Pfi_param_1];
  36.     cvta.to.global.u64  %rl6, %rl5;
  37.     cvta.to.global.u64  %rl7, %rl4;
  38.     .loc 2 44 1
  39.     mov.u32     %r3, %ctaid.x;
  40.     mov.u32     %r4, %ntid.x;
  41.     mov.u32     %r5, %tid.x;
  42.     mad.lo.s32  %r6, %r3, %r4, %r5;
  43.     .loc 2 49 1
  44.     cvt.s64.s32     %rl8, %r6;
  45.     mad.lo.s64  %rl9, %rl8, 36, %rl7;
  46.     ld.global.f32   %f127, [%rl9+32];
  47.     ld.global.f32   %f8, [%rl9+28];
  48.     ld.global.f32   %f7, [%rl9+24];
  49.     ld.global.f32   %f124, [%rl9+20];
  50.     ld.global.f32   %f5, [%rl9+16];
  51.     ld.global.f32   %f4, [%rl9+12];
  52.     ld.global.f32   %f131, [%rl9+8];
  53.     ld.global.f32   %f132, [%rl9+4];
  54.     ld.global.f32   %f137, [%rl9];
  55.     .loc 2 50 1
  56.     add.s64     %rl10, %rl6, %rl8;
  57.     .loc 2 54 1
  58.     mul.hi.s32  %r16, %r6, 458129845;
  59.     shr.u32     %r17, %r16, 31;
  60.     shr.s32     %r18, %r16, 5;
  61.     add.s32     %r19, %r18, %r17;
  62.     mul.lo.s32  %r20, %r19, 300;
  63.     sub.s32     %r21, %r6, %r20;
  64.     setp.eq.s32     %p2, %r21, 0;
  65.     .loc 2 50 1
  66.     ld.global.u8    %rc1, [%rl10];
  67.     {
  68.     .reg .s16   %temp1;
  69.     .reg .s16   %temp2;
  70.     cvt.s16.s8  %temp1, %rc1;
  71.     mov.b16     %temp2, 0;
  72.     cvt.s16.s8  %temp2, %temp2;
  73.     setp.eq.s16     %p1, %temp1, %temp2;
  74.     }
  75.     and.pred    %p3, %p2, %p1;
  76.     .loc 2 54 1
  77.     @!%p3 bra   BB0_4;
  78.  
  79.     add.f32     %f10, %f4, 0fB8690454;
  80.     .loc 2 57 1
  81.     setp.leu.f32    %p4, %f10, 0f00000000;
  82.     @%p4 bra    BB0_4;
  83.  
  84.     add.f32     %f11, %f7, 0fB7690454;
  85.     .loc 2 57 1
  86.     setp.leu.f32    %p5, %f11, 0f00000000;
  87.     @%p5 bra    BB0_4;
  88.  
  89.     add.f32     %f12, %f8, 0fB7690454;
  90.     .loc 2 57 1
  91.     setp.gt.f32     %p6, %f12, 0f00000000;
  92.     @%p6 bra    BB0_5;
  93.  
  94. BB0_4:
  95.     mov.f32     %f123, %f7;
  96.     mov.f32     %f130, %f8;
  97.     mov.f32     %f136, %f4;
  98.     bra.uni     BB0_6;
  99.  
  100. BB0_5:
  101.     .loc 2 62 1
  102.     add.f32     %f132, %f132, 0f38690454;
  103.     .loc 2 63 1
  104.     add.f32     %f124, %f124, 0f37690454;
  105.     .loc 2 64 1
  106.     add.f32     %f127, %f127, 0f37690454;
  107.     mov.f32     %f123, %f11;
  108.     mov.f32     %f130, %f12;
  109.     mov.f32     %f136, %f10;
  110.  
  111. BB0_6:
  112.     mov.f32     %f19, %f136;
  113.     mov.f32     %f18, %f130;
  114.     mov.f32     %f16, %f127;
  115.     mov.f32     %f122, %f123;
  116.     .loc 2 73 1
  117.     add.f32     %f49, %f137, 0f00000000;
  118.     add.f32     %f50, %f49, %f132;
  119.     add.f32     %f51, %f50, %f131;
  120.     add.f32     %f52, %f51, %f19;
  121.     add.f32     %f53, %f52, %f5;
  122.     add.f32     %f54, %f53, %f124;
  123.     add.f32     %f55, %f54, %f122;
  124.     add.f32     %f56, %f55, %f18;
  125.     add.f32     %f22, %f56, %f16;
  126.     .loc 2 75 1
  127.     add.f32     %f57, %f132, %f124;
  128.     add.f32     %f58, %f57, %f16;
  129.     add.f32     %f59, %f19, %f122;
  130.     add.f32     %f60, %f59, %f18;
  131.     sub.f32     %f61, %f58, %f60;
  132.     .loc 2 82 1
  133.     add.f32     %f62, %f131, %f124;
  134.     add.f32     %f63, %f62, %f122;
  135.     add.f32     %f64, %f5, %f18;
  136.     add.f32     %f65, %f64, %f16;
  137.     sub.f32     %f23, %f63, %f65;
  138.     .loc 2 88 1
  139.     mul.f32     %f24, %f23, %f23;
  140.     .loc 2 89 1
  141.     mul.f32     %f25, %f61, %f61;
  142.     .loc 2 90 1
  143.     mul.f32     %f26, %f61, %f23;
  144.     .loc 2 91 1
  145.     mul.f32     %f27, %f23, %f22;
  146.     .loc 2 92 1
  147.     mul.f32     %f28, %f22, %f22;
  148.     .loc 2 94 1
  149.     @%p1 bra    BB0_8;
  150.  
  151.     mov.f32     %f138, 0f00000000;
  152.     mov.f32     %f126, %f16;
  153.     mov.f32     %f129, %f18;
  154.     mov.f32     %f133, %f5;
  155.     mov.f32     %f135, %f19;
  156.     bra.uni     BB0_9;
  157.  
  158. BB0_8:
  159.     .loc 2 97 1
  160.     mul.f32     %f67, %f137, 0fBF59999A;
  161.     cvt.f64.f32     %fd1, %f67;
  162.     mov.f32     %f68, 0f3F9DDDDE;
  163.     .loc 3 1311 3
  164.     div.rn.f32  %f69, %f68, %f22;
  165.     .loc 2 97 1
  166.     cvt.f64.f32     %fd2, %f69;
  167.     cvt.f64.f32     %fd3, %f28;
  168.     mov.f64     %fd4, 0d3FE5555555555555;
  169.     .loc 2 97 1
  170.     cvt.f64.f32     %fd5, %f25;
  171.     .loc 2 99 1
  172.     neg.f64     %fd6, %fd5;
  173.     .loc 2 97 1
  174.     fma.rn.f64  %fd7, %fd3, %fd4, %fd6;
  175.     cvt.f64.f32     %fd8, %f24;
  176.     sub.f64     %fd9, %fd7, %fd8;
  177.     fma.rn.f64  %fd10, %fd2, %fd9, %fd1;
  178.     cvt.rn.f32.f64  %f137, %fd10;
  179.     add.f32     %f70, %f137, 0f00000000;
  180.     .loc 2 98 1
  181.     mul.f32     %f71, %f132, 0fBF59999A;
  182.     cvt.f64.f32     %fd11, %f71;
  183.     cvt.f64.f32     %fd12, %f22;
  184.     mul.f64     %fd13, %fd12, 0d3FD3BBBBC2777778;
  185.     add.f32     %f72, %f25, %f25;
  186.     sub.f32     %f73, %f72, %f24;
  187.     cvt.f64.f32     %fd14, %f73;
  188.     cvt.f64.f32     %fd15, %f61;
  189.     add.f64     %fd16, %fd15, %fd15;
  190.     fma.rn.f64  %fd17, %fd16, %fd12, %fd14;
  191.     .loc 3 1328 3
  192.     div.rn.f64  %fd18, %fd17, %fd3;
  193.     .loc 2 98 1
  194.     add.f64     %fd19, %fd18, 0d3FE5555555555555;
  195.     fma.rn.f64  %fd20, %fd13, %fd19, %fd11;
  196.     cvt.rn.f32.f64  %f30, %fd20;
  197.     add.f32     %f79, %f70, %f30;
  198.     .loc 2 99 1
  199.     mul.f32     %f80, %f22, 0f3DE38E39;
  200.     cvt.f64.f32     %fd21, %f80;
  201.     add.f32     %f81, %f27, %f24;
  202.     cvt.f64.f32     %fd22, %f81;
  203.     mov.f64     %fd23, 0d3FE0000000000000;
  204.     .loc 2 99 1
  205.     fma.rn.f64  %fd24, %fd6, %fd23, %fd22;
  206.     mul.f64     %fd25, %fd24, 0d4008000000000000;
  207.     .loc 3 1328 3
  208.     div.rn.f64  %fd26, %fd25, %fd3;
  209.     .loc 2 99 1
  210.     add.f64     %fd27, %fd26, 0d3FF0000000000000;
  211.     cvt.f64.f32     %fd28, %f131;
  212.     neg.f64     %fd29, %fd28;
  213.     fma.rn.f64  %fd30, %fd21, %fd27, %fd29;
  214.     fma.rn.f64  %fd31, %fd30, 0d3FFD9999A0000000, %fd28;
  215.     cvt.rn.f32.f64  %f31, %fd31;
  216.     add.f32     %f82, %f79, %f31;
  217.     .loc 2 100 1
  218.     neg.f32     %f83, %f61;
  219.     neg.f32     %f84, %f22;
  220.     fma.rn.f32  %f85, %f84, %f61, %f25;
  221.     cvt.f64.f32     %fd32, %f85;
  222.     neg.f64     %fd33, %fd8;
  223.     fma.rn.f64  %fd34, %fd33, %fd23, %fd32;
  224.     mul.f64     %fd35, %fd34, 0d4008000000000000;
  225.     .loc 3 1328 3
  226.     div.rn.f64  %fd36, %fd35, %fd3;
  227.     .loc 2 100 1
  228.     add.f64     %fd37, %fd36, 0d3FF0000000000000;
  229.     cvt.f64.f32     %fd38, %f19;
  230.     neg.f64     %fd39, %fd38;
  231.     fma.rn.f64  %fd40, %fd21, %fd37, %fd39;
  232.     fma.rn.f64  %fd41, %fd40, 0d3FFD9999A0000000, %fd38;
  233.     cvt.rn.f32.f64  %f132, %fd41;
  234.     add.f32     %f86, %f82, %f132;
  235.     sub.f32     %f87, %f24, %f27;
  236.     .loc 2 101 1
  237.     cvt.f64.f32     %fd42, %f87;
  238.     fma.rn.f64  %fd43, %fd6, %fd23, %fd42;
  239.     mul.f64     %fd44, %fd43, 0d4008000000000000;
  240.     .loc 3 1328 3
  241.     div.rn.f64  %fd45, %fd44, %fd3;
  242.     .loc 2 101 1
  243.     add.f64     %fd46, %fd45, 0d3FF0000000000000;
  244.     cvt.f64.f32     %fd47, %f5;
  245.     neg.f64     %fd48, %fd47;
  246.     fma.rn.f64  %fd49, %fd21, %fd46, %fd48;
  247.     fma.rn.f64  %fd50, %fd49, 0d3FFD9999A0000000, %fd47;
  248.     cvt.rn.f32.f64  %f131, %fd50;
  249.     add.f32     %f88, %f86, %f131;
  250.     .loc 2 102 1
  251.     mul.f32     %f89, %f22, 0f3CE38E39;
  252.     cvt.f64.f32     %fd51, %f89;
  253.     mul.f64     %fd52, %fd12, 0d4008000000000000;
  254.     add.f32     %f90, %f61, %f23;
  255.     cvt.f64.f32     %fd53, %f90;
  256.     mul.f64     %fd54, %fd52, %fd53;
  257.     add.f32     %f91, %f25, %f24;
  258.     cvt.f64.f32     %fd55, %f91;
  259.     cvt.f64.f32     %fd56, %f26;
  260.     fma.rn.f64  %fd57, %fd56, 0d4008000000000000, %fd55;
  261.     fma.rn.f64  %fd58, %fd57, 0d4008000000000000, %fd54;
  262.     .loc 3 1328 3
  263.     div.rn.f64  %fd59, %fd58, %fd3;
  264.     .loc 2 102 1
  265.     add.f64     %fd60, %fd59, 0d3FF0000000000000;
  266.     cvt.f64.f32     %fd61, %f124;
  267.     neg.f64     %fd62, %fd61;
  268.     fma.rn.f64  %fd63, %fd51, %fd60, %fd62;
  269.     fma.rn.f64  %fd64, %fd63, 0d3FFD9999A0000000, %fd61;
  270.     cvt.rn.f32.f64  %f34, %fd64;
  271.     add.f32     %f92, %f88, %f34;
  272.     sub.f32     %f93, %f23, %f61;
  273.     .loc 2 103 1
  274.     cvt.f64.f32     %fd65, %f93;
  275.     mul.f64     %fd66, %fd5, 0d4008000000000000;
  276.     mul.f64     %fd67, %fd56, 0d4022000000000000;
  277.     mov.f64     %fd68, 0d4022000000000000;
  278.     .loc 2 103 1
  279.     neg.f64     %fd69, %fd56;
  280.     fma.rn.f64  %fd70, %fd69, %fd68, %fd66;
  281.     fma.rn.f64  %fd71, %fd8, 0d4008000000000000, %fd70;
  282.     fma.rn.f64  %fd72, %fd52, %fd65, %fd71;
  283.     .loc 3 1328 3
  284.     div.rn.f64  %fd73, %fd72, %fd3;
  285.     .loc 2 103 1
  286.     add.f64     %fd74, %fd73, 0d3FF0000000000000;
  287.     cvt.f64.f32     %fd75, %f122;
  288.     neg.f64     %fd76, %fd75;
  289.     fma.rn.f64  %fd77, %fd51, %fd74, %fd76;
  290.     fma.rn.f64  %fd78, %fd77, 0d3FFD9999A0000000, %fd75;
  291.     cvt.rn.f32.f64  %f35, %fd78;
  292.     add.f32     %f94, %f92, %f35;
  293.     .loc 2 104 1
  294.     sub.f32     %f95, %f83, %f23;
  295.     cvt.f64.f32     %fd79, %f95;
  296.     fma.rn.f64  %fd80, %fd5, 0d4008000000000000, %fd67;
  297.     fma.rn.f64  %fd81, %fd8, 0d4008000000000000, %fd80;
  298.     fma.rn.f64  %fd82, %fd52, %fd79, %fd81;
  299.     .loc 3 1328 3
  300.     div.rn.f64  %fd83, %fd82, %fd3;
  301.     .loc 2 104 1
  302.     add.f64     %fd84, %fd83, 0d3FF0000000000000;
  303.     cvt.f64.f32     %fd85, %f18;
  304.     neg.f64     %fd86, %fd85;
  305.     fma.rn.f64  %fd87, %fd51, %fd84, %fd86;
  306.     fma.rn.f64  %fd88, %fd87, 0d3FFD9999A0000000, %fd85;
  307.     cvt.rn.f32.f64  %f124, %fd88;
  308.     add.f32     %f96, %f94, %f124;
  309.     .loc 2 105 1
  310.     sub.f32     %f97, %f61, %f23;
  311.     cvt.f64.f32     %fd89, %f97;
  312.     fma.rn.f64  %fd90, %fd52, %fd89, %fd71;
  313.     .loc 3 1328 3
  314.     div.rn.f64  %fd91, %fd90, %fd3;
  315.     .loc 2 105 1
  316.     add.f64     %fd92, %fd91, 0d3FF0000000000000;
  317.     cvt.f64.f32     %fd93, %f16;
  318.     neg.f64     %fd94, %fd93;
  319.     fma.rn.f64  %fd95, %fd51, %fd92, %fd94;
  320.     fma.rn.f64  %fd96, %fd95, 0d3FFD9999A0000000, %fd93;
  321.     cvt.rn.f32.f64  %f122, %fd96;
  322.     add.f32     %f98, %f96, %f122;
  323.     .loc 2 107 1
  324.     add.f32     %f99, %f30, %f34;
  325.     add.f32     %f100, %f99, %f122;
  326.     add.f32     %f101, %f132, %f35;
  327.     add.f32     %f102, %f101, %f124;
  328.     sub.f32     %f103, %f100, %f102;
  329.     .loc 3 1311 3
  330.     div.rn.f32  %f104, %f103, %f98;
  331.     .loc 2 107 1
  332.     add.f32     %f138, %f104, 0f00000000;
  333.     mov.f32     %f126, %f35;
  334.     mov.f32     %f129, %f34;
  335.     mov.f32     %f133, %f31;
  336.     mov.f32     %f135, %f30;
  337.  
  338. BB0_9:
  339.     ld.param.u64    %rl19, [_Z9collisionPK7t_speedPKcPS_Pfi_param_2];
  340.     cvta.to.global.u64  %rl11, %rl19;
  341.     .loc 2 133 1
  342.     mad.lo.s64  %rl13, %rl8, 36, %rl11;
  343.     st.global.f32   [%rl13+32], %f122;
  344.     st.global.f32   [%rl13+28], %f124;
  345.     st.global.f32   [%rl13+24], %f126;
  346.     st.global.f32   [%rl13+20], %f129;
  347.     st.global.f32   [%rl13+16], %f131;
  348.     st.global.f32   [%rl13+12], %f132;
  349.     st.global.f32   [%rl13+8], %f133;
  350.     st.global.f32   [%rl13+4], %f135;
  351.     st.global.f32   [%rl13], %f137;
  352.     .loc 2 44 1
  353.     mov.u32     %r66, %tid.x;
  354.     mul.wide.s32    %rl14, %r66, 4;
  355.     mov.u64     %rl15, __cuda_local_var_17168_35_non_const_sdata;
  356.     add.s64     %rl3, %rl15, %rl14;
  357.     .loc 2 137 1
  358.     st.shared.f32   [%rl3], %f138;
  359.     .loc 2 141 1
  360.     bar.sync    0;
  361.     .loc 2 44 1
  362.     mov.u32     %r67, %tid.x;
  363.     .loc 2 145 1
  364.     setp.gt.s32     %p7, %r67, 15;
  365.     @%p7 bra    BB0_11;
  366.  
  367.     .loc 2 148 1
  368.     ld.shared.f32   %f105, [%rl3];
  369.     .loc 2 147 1
  370.     ld.shared.f32   %f106, [%rl3+64];
  371.     .loc 2 149 1
  372.     add.f32     %f107, %f105, %f106;
  373.     st.shared.f32   [%rl3], %f107;
  374.  
  375. BB0_11:
  376.     .loc 2 151 1
  377.     bar.sync    0;
  378.     .loc 2 44 1
  379.     mov.u32     %r65, %tid.x;
  380.     .loc 2 145 1
  381.     setp.gt.s32     %p8, %r65, 7;
  382.     @%p8 bra    BB0_13;
  383.  
  384.     .loc 2 148 1
  385.     ld.shared.f32   %f108, [%rl3];
  386.     .loc 2 147 1
  387.     ld.shared.f32   %f109, [%rl3+32];
  388.     .loc 2 149 1
  389.     add.f32     %f110, %f108, %f109;
  390.     st.shared.f32   [%rl3], %f110;
  391.  
  392. BB0_13:
  393.     .loc 2 151 1
  394.     bar.sync    0;
  395.     .loc 2 44 1
  396.     mov.u32     %r64, %tid.x;
  397.     .loc 2 145 1
  398.     setp.gt.s32     %p9, %r64, 3;
  399.     @%p9 bra    BB0_15;
  400.  
  401.     .loc 2 148 1
  402.     ld.shared.f32   %f111, [%rl3];
  403.     .loc 2 147 1
  404.     ld.shared.f32   %f112, [%rl3+16];
  405.     .loc 2 149 1
  406.     add.f32     %f113, %f111, %f112;
  407.     st.shared.f32   [%rl3], %f113;
  408.  
  409. BB0_15:
  410.     .loc 2 151 1
  411.     bar.sync    0;
  412.     .loc 2 44 1
  413.     mov.u32     %r63, %tid.x;
  414.     .loc 2 145 1
  415.     setp.gt.s32     %p10, %r63, 1;
  416.     @%p10 bra   BB0_17;
  417.  
  418.     .loc 2 148 1
  419.     ld.shared.f32   %f114, [%rl3];
  420.     .loc 2 147 1
  421.     ld.shared.f32   %f115, [%rl3+8];
  422.     .loc 2 149 1
  423.     add.f32     %f116, %f114, %f115;
  424.     st.shared.f32   [%rl3], %f116;
  425.  
  426. BB0_17:
  427.     .loc 2 151 1
  428.     bar.sync    0;
  429.     .loc 2 44 1
  430.     mov.u32     %r62, %tid.x;
  431.     .loc 2 145 1
  432.     setp.gt.s32     %p11, %r62, 0;
  433.     @%p11 bra   BB0_19;
  434.  
  435.     .loc 2 148 1
  436.     ld.shared.f32   %f117, [%rl3];
  437.     .loc 2 147 1
  438.     ld.shared.f32   %f118, [%rl3+4];
  439.     .loc 2 149 1
  440.     add.f32     %f119, %f117, %f118;
  441.     st.shared.f32   [%rl3], %f119;
  442.  
  443. BB0_19:
  444.     .loc 2 151 1
  445.     bar.sync    0;
  446.     .loc 2 44 1
  447.     mov.u32     %r61, %tid.x;
  448.     .loc 2 153 1
  449.     setp.eq.s32     %p12, %r61, 0;
  450.     @%p12 bra   BB0_21;
  451.  
  452.     .loc 2 156 1
  453.     ret;
  454.  
  455. BB0_21:
  456.     .loc 2 154 1
  457.     ld.shared.f32   %f120, [__cuda_local_var_17168_35_non_const_sdata];
  458.     ld.param.u32    %r59, [_Z9collisionPK7t_speedPKcPS_Pfi_param_4];
  459.     .loc 2 44 1
  460.     mov.u32     %r60, %ctaid.x;
  461.     .loc 2 154 1
  462.     mad.lo.s32  %r57, %r59, 1875, %r60;
  463.     ld.param.u64    %rl20, [_Z9collisionPK7t_speedPKcPS_Pfi_param_3];
  464.     cvta.to.global.u64  %rl16, %rl20;
  465.     .loc 2 154 1
  466.     mul.wide.u32    %rl17, %r57, 4;
  467.     add.s64     %rl18, %rl16, %rl17;
  468.     st.global.f32   [%rl18], %f120;
  469.     .loc 2 156 1
  470.     ret;
  471. }
  472.  
  473. .entry _Z9propagateP7t_speedPKS_(
  474.     .param .u64 _Z9propagateP7t_speedPKS__param_0,
  475.     .param .u64 _Z9propagateP7t_speedPKS__param_1
  476. )
  477. {
  478.     .reg .f32   %f<13>;
  479.     .reg .pred  %p<3>;
  480.     .reg .s32   %r<57>;
  481.     .reg .s64   %rl<24>;
  482.  
  483.  
  484.     ld.param.u64    %rl1, [_Z9propagateP7t_speedPKS__param_0];
  485.     ld.param.u64    %rl2, [_Z9propagateP7t_speedPKS__param_1];
  486.     cvta.to.global.u64  %rl3, %rl1;
  487.     cvta.to.global.u64  %rl4, %rl2;
  488.     .loc 2 164 1
  489.     mov.u32     %r1, %ctaid.x;
  490.     mov.u32     %r2, %ntid.x;
  491.     mov.u32     %r3, %tid.x;
  492.     mad.lo.s32  %r4, %r1, %r2, %r3;
  493.     .loc 2 167 1
  494.     mul.hi.s32  %r5, %r4, 458129845;
  495.     shr.u32     %r6, %r5, 31;
  496.     shr.s32     %r7, %r5, 5;
  497.     add.s32     %r8, %r7, %r6;
  498.     mul.lo.s32  %r9, %r8, 300;
  499.     sub.s32     %r10, %r4, %r9;
  500.     .loc 2 169 1
  501.     cvt.rn.f32.s32  %f1, %r4;
  502.     .loc 3 1311 3
  503.     div.rn.f32  %f2, %f1, 0f43960000;
  504.     .loc 4 373 5
  505.     cvt.rmi.f32.f32     %f3, %f2;
  506.     .loc 4 1343 5
  507.     cvt.rzi.s32.f32     %r11, %f3;
  508.     .loc 2 204 1
  509.     mad.lo.s32  %r12, %r11, 300, %r10;
  510.     cvt.s64.s32     %rl5, %r12;
  511.     mad.lo.s64  %rl6, %rl5, 36, %rl4;
  512.     ld.global.f32   %f4, [%rl6+32];
  513.     ld.global.f32   %f5, [%rl6+28];
  514.     ld.global.f32   %f6, [%rl6+24];
  515.     ld.global.f32   %f7, [%rl6+20];
  516.     ld.global.f32   %f8, [%rl6+16];
  517.     ld.global.f32   %f9, [%rl6+12];
  518.     ld.global.f32   %f10, [%rl6+8];
  519.     ld.global.f32   %f11, [%rl6+4];
  520.     ld.global.f32   %f12, [%rl6];
  521.     .loc 2 205 1
  522.     add.s32     %r22, %r11, 1;
  523.     mul.hi.s32  %r23, %r22, 1374389535;
  524.     shr.u32     %r24, %r23, 31;
  525.     shr.s32     %r25, %r23, 6;
  526.     add.s32     %r26, %r25, %r24;
  527.     mul.lo.s32  %r27, %r26, 200;
  528.     sub.s32     %r28, %r22, %r27;
  529.     .loc 2 206 1
  530.     add.s32     %r29, %r10, 1;
  531.     mul.hi.s32  %r30, %r29, 458129845;
  532.     shr.u32     %r31, %r30, 31;
  533.     shr.s32     %r32, %r30, 5;
  534.     add.s32     %r33, %r32, %r31;
  535.     mul.lo.s32  %r34, %r33, 300;
  536.     sub.s32     %r35, %r29, %r34;
  537.     .loc 2 207 1
  538.     setp.eq.s32     %p1, %r11, 0;
  539.     add.s32     %r36, %r10, -1;
  540.     .loc 2 208 1
  541.     setp.eq.s32     %p2, %r10, 0;
  542.     selp.b32    %r37, 299, %r36, %p2;
  543.     mad.lo.s64  %rl7, %rl5, 36, %rl3;
  544.     .loc 2 210 1
  545.     st.global.f32   [%rl7], %f12;
  546.     .loc 2 211 1
  547.     mad.lo.s32  %r39, %r11, 300, %r35;
  548.     cvt.s64.s32     %rl8, %r39;
  549.     mad.lo.s64  %rl9, %rl8, 36, %rl3;
  550.     .loc 2 211 1
  551.     st.global.f32   [%rl9+4], %f11;
  552.     .loc 2 212 1
  553.     mad.lo.s32  %r41, %r28, 300, %r10;
  554.     cvt.s64.s32     %rl10, %r41;
  555.     mad.lo.s64  %rl11, %rl10, 36, %rl3;
  556.     .loc 2 212 1
  557.     st.global.f32   [%rl11+8], %f10;
  558.     .loc 2 213 1
  559.     mad.lo.s32  %r43, %r11, 300, %r37;
  560.     cvt.s64.s32     %rl12, %r43;
  561.     mad.lo.s64  %rl13, %rl12, 36, %rl3;
  562.     .loc 2 213 1
  563.     st.global.f32   [%rl13+12], %f9;
  564.     mad.lo.s32  %r45, %r11, 300, -300;
  565.     selp.b32    %r46, 59700, %r45, %p1;
  566.     .loc 2 214 1
  567.     add.s32     %r47, %r46, %r10;
  568.     cvt.s64.s32     %rl14, %r47;
  569.     mad.lo.s64  %rl15, %rl14, 36, %rl3;
  570.     .loc 2 214 1
  571.     st.global.f32   [%rl15+16], %f8;
  572.     .loc 2 215 1
  573.     mad.lo.s32  %r49, %r28, 300, %r35;
  574.     cvt.s64.s32     %rl16, %r49;
  575.     mad.lo.s64  %rl17, %rl16, 36, %rl3;
  576.     .loc 2 215 1
  577.     st.global.f32   [%rl17+20], %f7;
  578.     .loc 2 216 1
  579.     mad.lo.s32  %r51, %r28, 300, %r37;
  580.     cvt.s64.s32     %rl18, %r51;
  581.     mad.lo.s64  %rl19, %rl18, 36, %rl3;
  582.     .loc 2 216 1
  583.     st.global.f32   [%rl19+24], %f6;
  584.     .loc 2 217 1
  585.     add.s32     %r53, %r46, %r37;
  586.     cvt.s64.s32     %rl20, %r53;
  587.     mad.lo.s64  %rl21, %rl20, 36, %rl3;
  588.     .loc 2 217 1
  589.     st.global.f32   [%rl21+28], %f5;
  590.     .loc 2 218 1
  591.     add.s32     %r55, %r46, %r35;
  592.     cvt.s64.s32     %rl22, %r55;
  593.     mad.lo.s64  %rl23, %rl22, 36, %rl3;
  594.     .loc 2 218 1
  595.     st.global.f32   [%rl23+32], %f4;
  596.     .loc 2 221 2
  597.     ret;
  598. }
Advertisement
Add Comment
Please, Sign In to add comment
Public Pastes
We use cookies for various purposes including analytics. By continuing to use Pastebin, you agree to our use of cookies as described in the Cookies Policy.  OK, I Understand
Not a member of Pastebin yet?
Sign Up, it unlocks many cool features!