API tools faq
paste
Guest User

default sm_10

a guest
Jun 12th, 2012
288
0
Never
Add comment
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
ASM (NASM) 19.31 KB | None | 0 0
raw download clone embed print report
  1.     .version 1.4
  2.     .target sm_10, map_f64_to_f32
  3.     // compiled with /usr/local/gpu/cuda-toolkit-4.1.28/cuda/open64/lib//be
  4.     // nvopencc 4.1 built on 2012-01-12
  5.  
  6.     //-----------------------------------------------------------
  7.     // Compiling /tmp/tmpxft_00002172_00000000-9_d2q9.cpp3.i (/tmp/ccBI#.tyRWSL)
  8.     //-----------------------------------------------------------
  9.  
  10.     //-----------------------------------------------------------
  11.     // Options:
  12.     //-----------------------------------------------------------
  13.     //  Target:ptx, ISA:sm_10, Endian:little, Pointer Size:64
  14.     //  -O3 (Optimization level)
  15.     //  -g0 (Debug level)
  16.     //  -m2 (Report advisories)
  17.     //-----------------------------------------------------------
  18.  
  19.     .file   1   "<command-line>"
  20.     .file   2   "/tmp/tmpxft_00002172_00000000-8_d2q9.cudafe2.gpu"
  21.     .file   3   "d2q9.cu"
  22.     .file   4   "/cm/shared/apps/gcc/4.4.6/lib/gcc/x86_64-unknown-linux-gnu/4.4.6/include/stddef.h"
  23.     .file   5   "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/crt/device_runtime.h"
  24.     .file   6   "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/host_defines.h"
  25.     .file   7   "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/builtin_types.h"
  26.     .file   8   "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/device_types.h"
  27.     .file   9   "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/driver_types.h"
  28.     .file   10  "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/surface_types.h"
  29.     .file   11  "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/texture_types.h"
  30.     .file   12  "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/vector_types.h"
  31.     .file   13  "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/device_launch_parameters.h"
  32.     .file   14  "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/crt/storage_class.h"
  33.     .file   15  "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/common_functions.h"
  34.     .file   16  "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/math_functions.h"
  35.     .file   17  "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/math_constants.h"
  36.     .file   18  "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/device_functions.h"
  37.     .file   19  "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/sm_11_atomic_functions.h"
  38.     .file   20  "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/sm_12_atomic_functions.h"
  39.     .file   21  "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/sm_13_double_functions.h"
  40.     .file   22  "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/sm_20_atomic_functions.h"
  41.     .file   23  "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/sm_20_intrinsics.h"
  42.     .file   24  "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/surface_functions.h"
  43.     .file   25  "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/texture_fetch_functions.h"
  44.     .file   26  "/usr/local/gpu/cuda-toolkit-4.1.28/cuda/bin/../include/math_functions_dbl_ptx1.h"
  45.  
  46.  
  47.     .entry _Z9collisionPK7t_speedPKcPS_Pfi (
  48.         .param .u64 __cudaparm__Z9collisionPK7t_speedPKcPS_Pfi_cells,
  49.         .param .u64 __cudaparm__Z9collisionPK7t_speedPKcPS_Pfi_obstacles,
  50.         .param .u64 __cudaparm__Z9collisionPK7t_speedPKcPS_Pfi_tmpy_cells,
  51.         .param .u64 __cudaparm__Z9collisionPK7t_speedPKcPS_Pfi_av_vels,
  52.         .param .s32 __cudaparm__Z9collisionPK7t_speedPKcPS_Pfi_ii)
  53.     {
  54.     .reg .u32 %r<30>;
  55.     .reg .u64 %rd<21>;
  56.     .reg .f32 %f<124>;
  57.     .reg .f64 %fd<127>;
  58.     .reg .pred %p<11>;
  59.     .shared .align 4 .b8 __cuda___cuda_local_var_15345_35_non_const_sdata36[128];
  60.     .loc    3   36  0
  61. $LDWbegin__Z9collisionPK7t_speedPKcPS_Pfi:
  62.     .loc    3   49  0
  63.     cvt.u32.u16     %r1, %ctaid.x;
  64.     cvt.u32.u16     %r2, %ntid.x;
  65.     mul.lo.u32  %r3, %r2, %r1;
  66.     cvt.u32.u16     %r4, %tid.x;
  67.     add.u32     %r5, %r4, %r3;
  68.     cvt.s64.s32     %rd1, %r5;
  69.     mul.wide.s32    %rd2, %r5, 36;
  70.     ld.param.u64    %rd3, [__cudaparm__Z9collisionPK7t_speedPKcPS_Pfi_cells];
  71.     add.u64     %rd4, %rd3, %rd2;
  72.     ld.global.f32   %f1, [%rd4+0];
  73.     mov.f32     %f2, %f1;
  74.     ld.global.f32   %f3, [%rd4+4];
  75.     mov.f32     %f4, %f3;
  76.     ld.global.f32   %f5, [%rd4+8];
  77.     mov.f32     %f6, %f5;
  78.     ld.global.f32   %f7, [%rd4+12];
  79.     mov.f32     %f8, %f7;
  80.     ld.global.f32   %f9, [%rd4+16];
  81.     mov.f32     %f10, %f9;
  82.     ld.global.f32   %f11, [%rd4+20];
  83.     mov.f32     %f12, %f11;
  84.     ld.global.f32   %f13, [%rd4+24];
  85.     mov.f32     %f14, %f13;
  86.     ld.global.f32   %f15, [%rd4+28];
  87.     mov.f32     %f16, %f15;
  88.     ld.global.f32   %f17, [%rd4+32];
  89.     mov.f32     %f18, %f17;
  90.     .loc    3   50  0
  91.     ld.param.u64    %rd5, [__cudaparm__Z9collisionPK7t_speedPKcPS_Pfi_obstacles];
  92.     add.u64     %rd6, %rd1, %rd5;
  93.     ld.global.s8    %r6, [%rd6+0];
  94.     mov.s32     %r7, 0;
  95.     setp.eq.s32     %p1, %r6, %r7;
  96.     mov.s32     %r8, 458129845;
  97.     mov.s32     %r9, 0;
  98.     setp.lt.s32     %p2, %r5, %r9;
  99.     abs.s32     %r10, %r5;
  100.     mul.hi.u32  %r11, %r10, %r8;
  101.     shr.s32     %r12, %r11, 5;
  102.     @%p2 sub.s32    %r12, %r9, %r12;
  103.     mov.s32     %r13, %r12;
  104.     mul.lo.s32  %r14, %r13, 300;
  105.     sub.s32     %r15, %r5, %r14;
  106.     mov.u32     %r16, 0;
  107.     setp.ne.s32     %p3, %r15, %r16;
  108.     @%p3 bra    $Lt_0_11522;
  109.     .loc    3   57  0
  110.     @!%p1 bra   $Lt_0_11522;
  111.     mov.f32     %f19, 0fb8690454;       // -5.55556e-05
  112.     add.f32     %f20, %f7, %f19;
  113.     mov.f32     %f21, 0f00000000;       // 0
  114.     setp.gt.f32     %p4, %f20, %f21;
  115.     @!%p4 bra   $Lt_0_11522;
  116.     mov.f32     %f22, 0fb7690454;       // -1.38889e-05
  117.     add.f32     %f23, %f13, %f22;
  118.     mov.f32     %f24, 0f00000000;       // 0
  119.     setp.gt.f32     %p5, %f23, %f24;
  120.     @!%p5 bra   $Lt_0_11522;
  121.     mov.f32     %f25, 0fb7690454;       // -1.38889e-05
  122.     add.f32     %f26, %f15, %f25;
  123.     mov.f32     %f27, 0f00000000;       // 0
  124.     setp.gt.f32     %p6, %f26, %f27;
  125.     @!%p6 bra   $Lt_0_11522;
  126.     .loc    3   62  0
  127.     mov.f32     %f28, 0f38690454;       // 5.55556e-05
  128.     add.f32     %f3, %f3, %f28;
  129.     mov.f32     %f4, %f3;
  130.     .loc    3   63  0
  131.     mov.f32     %f29, 0f37690454;       // 1.38889e-05
  132.     add.f32     %f11, %f11, %f29;
  133.     mov.f32     %f12, %f11;
  134.     .loc    3   64  0
  135.     mov.f32     %f30, 0f37690454;       // 1.38889e-05
  136.     add.f32     %f17, %f17, %f30;
  137.     mov.f32     %f18, %f17;
  138.     .loc    3   65  0
  139.     mov.f32     %f7, %f20;
  140.     mov.f32     %f8, %f7;
  141.     .loc    3   66  0
  142.     mov.f32     %f13, %f23;
  143.     mov.f32     %f14, %f13;
  144.     .loc    3   67  0
  145.     mov.f32     %f15, %f26;
  146.     mov.f32     %f16, %f15;
  147. $Lt_0_11522:
  148. $L_0_6914:
  149. $Lt_0_7938:
  150.     @!%p1 bra   $Lt_0_8706;
  151.     .loc    3   97  0
  152.     add.f32     %f31, %f3, %f1;
  153.     add.f32     %f32, %f11, %f3;
  154.     add.f32     %f33, %f13, %f7;
  155.     add.f32     %f34, %f11, %f5;
  156.     add.f32     %f35, %f15, %f9;
  157.     add.f32     %f36, %f31, %f5;
  158.     add.f32     %f37, %f32, %f17;
  159.     add.f32     %f38, %f33, %f15;
  160.     add.f32     %f39, %f34, %f13;
  161.     add.f32     %f40, %f35, %f17;
  162.     add.f32     %f41, %f36, %f7;
  163.     sub.f32     %f42, %f37, %f38;
  164.     sub.f32     %f43, %f39, %f40;
  165.     add.f32     %f44, %f41, %f9;
  166.     mul.f32     %f45, %f42, %f42;
  167.     mul.f32     %f46, %f43, %f43;
  168.     add.f32     %f47, %f44, %f11;
  169.     cvt.f64.f32     %fd1, %f45;
  170.     cvt.f64.f32     %fd2, %f46;
  171.     add.f32     %f48, %f47, %f13;
  172.     add.f32     %f49, %f48, %f15;
  173.     add.f32     %f50, %f49, %f17;
  174.     mul.f32     %f51, %f50, %f50;
  175.     mov.f32     %f52, 0fbf59999a;       // -0.85
  176.     mul.f32     %f53, %f1, %f52;
  177.     cvt.f64.f32     %fd3, %f53;
  178.     mov.f32     %f54, 0f3f9dddde;       // 1.23333
  179.     div.full.f32    %f55, %f54, %f50;
  180.     cvt.f64.f32     %fd4, %f55;
  181.     cvt.f64.f32     %fd5, %f51;
  182.     mov.f64     %fd6, 0d3fe5555555555555;   // 0.666667
  183.     mul.f64     %fd7, %fd5, %fd6;
  184.     sub.f64     %fd8, %fd7, %fd1;
  185.     sub.f64     %fd9, %fd8, %fd2;
  186.     mad.rn.f64  %fd10, %fd4, %fd9, %fd3;
  187.     cvt.rn.f32.f64  %f56, %fd10;
  188.     mov.f32     %f57, %f56;
  189.     .loc    3   98  0
  190.     cvt.f64.f32     %fd11, %f42;
  191.     cvt.f64.f32     %fd12, %f50;
  192.     mov.f32     %f58, 0fbf59999a;       // -0.85
  193.     mul.f32     %f59, %f3, %f58;
  194.     cvt.f64.f32     %fd13, %f59;
  195.     add.f32     %f60, %f45, %f45;
  196.     sub.f32     %f61, %f60, %f46;
  197.     cvt.f64.f32     %fd14, %f61;
  198.     add.f64     %fd15, %fd11, %fd11;
  199.     mad.rn.f64  %fd16, %fd12, %fd15, %fd14;
  200.     cvt.rn.f32.f64  %f62, %fd16;
  201.     div.full.f32    %f63, %f62, %f51;
  202.     cvt.f64.f32     %fd17, %f63;
  203.     mov.f64     %fd18, 0d3fe5555555555555;  // 0.666667
  204.     add.f64     %fd19, %fd17, %fd18;
  205.     mov.f64     %fd20, 0d3fd3bbbbc2777778;  // 0.308333
  206.     mul.f64     %fd21, %fd12, %fd20;
  207.     mad.rn.f64  %fd22, %fd19, %fd21, %fd13;
  208.     cvt.rn.f32.f64  %f64, %fd22;
  209.     .loc    3   99  0
  210.     cvt.f64.f32     %fd23, %f5;
  211.     mul.f32     %f65, %f43, %f50;
  212.     mov.f32     %f66, 0f3de38e39;       // 0.111111
  213.     mul.f32     %f67, %f50, %f66;
  214.     cvt.f64.f32     %fd24, %f67;
  215.     add.f32     %f68, %f46, %f65;
  216.     cvt.f64.f32     %fd25, %f68;
  217.     mov.f64     %fd26, 0d3fe0000000000000;  // 0.5
  218.     mul.f64     %fd27, %fd1, %fd26;
  219.     sub.f64     %fd28, %fd25, %fd27;
  220.     mov.f64     %fd29, 0d4008000000000000;  // 3
  221.     mul.f64     %fd30, %fd28, %fd29;
  222.     cvt.rn.f32.f64  %f69, %fd30;
  223.     div.full.f32    %f70, %f69, %f51;
  224.     cvt.f64.f32     %fd31, %f70;
  225.     mov.f64     %fd32, 0d3ff0000000000000;  // 1
  226.     add.f64     %fd33, %fd31, %fd32;
  227.     mul.f64     %fd34, %fd24, %fd33;
  228.     sub.f64     %fd35, %fd34, %fd23;
  229.     mov.f64     %fd36, 0d3ffd9999a0000000;  // 1.85
  230.     mad.rn.f64  %fd37, %fd35, %fd36, %fd23;
  231.     cvt.rn.f32.f64  %f71, %fd37;
  232.     .loc    3   100 0
  233.     cvt.f64.f32     %fd38, %f7;
  234.     mul.f32     %f72, %f42, %f50;
  235.     sub.f32     %f73, %f45, %f72;
  236.     cvt.f64.f32     %fd39, %f73;
  237.     mov.f64     %fd40, 0d3fe0000000000000;  // 0.5
  238.     mul.f64     %fd41, %fd2, %fd40;
  239.     sub.f64     %fd42, %fd39, %fd41;
  240.     mov.f64     %fd43, 0d4008000000000000;  // 3
  241.     mul.f64     %fd44, %fd42, %fd43;
  242.     cvt.rn.f32.f64  %f74, %fd44;
  243.     div.full.f32    %f75, %f74, %f51;
  244.     cvt.f64.f32     %fd45, %f75;
  245.     mov.f64     %fd46, 0d3ff0000000000000;  // 1
  246.     add.f64     %fd47, %fd45, %fd46;
  247.     mul.f64     %fd48, %fd24, %fd47;
  248.     sub.f64     %fd49, %fd48, %fd38;
  249.     mov.f64     %fd50, 0d3ffd9999a0000000;  // 1.85
  250.     mad.rn.f64  %fd51, %fd49, %fd50, %fd38;
  251.     cvt.rn.f32.f64  %f76, %fd51;
  252.     add.f32     %f77, %f56, %f64;
  253.     add.f32     %f78, %f71, %f77;
  254.     add.f32     %f79, %f76, %f78;
  255.     .loc    3   101 0
  256.     cvt.f64.f32     %fd52, %f9;
  257.     sub.f32     %f80, %f46, %f65;
  258.     cvt.f64.f32     %fd53, %f80;
  259.     mov.f64     %fd54, 0d3fe0000000000000;  // 0.5
  260.     mul.f64     %fd55, %fd1, %fd54;
  261.     sub.f64     %fd56, %fd53, %fd55;
  262.     mov.f64     %fd57, 0d4008000000000000;  // 3
  263.     mul.f64     %fd58, %fd56, %fd57;
  264.     cvt.rn.f32.f64  %f81, %fd58;
  265.     div.full.f32    %f82, %f81, %f51;
  266.     cvt.f64.f32     %fd59, %f82;
  267.     mov.f64     %fd60, 0d3ff0000000000000;  // 1
  268.     add.f64     %fd61, %fd59, %fd60;
  269.     mul.f64     %fd62, %fd24, %fd61;
  270.     sub.f64     %fd63, %fd62, %fd52;
  271.     mov.f64     %fd64, 0d3ffd9999a0000000;  // 1.85
  272.     mad.rn.f64  %fd65, %fd63, %fd64, %fd52;
  273.     cvt.rn.f32.f64  %f83, %fd65;
  274.     add.f32     %f84, %f79, %f83;
  275.     .loc    18  1319    0
  276.     add.f32     %f85, %f42, %f43;
  277.     mul.f32     %f86, %f42, %f43;
  278.     cvt.f64.f32     %fd66, %f86;
  279.     mov.f64     %fd67, 0d4008000000000000;  // 3
  280.     mul.f64     %fd68, %fd12, %fd67;
  281.     add.f32     %f87, %f45, %f46;
  282.     cvt.f64.f32     %fd69, %f87;
  283.     mov.f64     %fd70, 0d4008000000000000;  // 3
  284.     mad.rn.f64  %fd71, %fd66, %fd70, %fd69;
  285.     mov.f64     %fd72, 0d4008000000000000;  // 3
  286.     mul.f64     %fd73, %fd71, %fd72;
  287.     cvt.f64.f32     %fd74, %f85;
  288.     mad.rn.f64  %fd75, %fd74, %fd68, %fd73;
  289.     cvt.rn.f32.f64  %f88, %fd75;
  290.     div.full.f32    %f89, %f88, %f51;
  291.     cvt.f64.f32     %fd76, %f89;
  292.     .loc    3   102 0
  293.     cvt.f64.f32     %fd77, %f11;
  294.     mov.f32     %f90, 0f3ce38e39;       // 0.0277778
  295.     mul.f32     %f91, %f50, %f90;
  296.     cvt.f64.f32     %fd78, %f91;
  297.     mov.f64     %fd79, 0d3ff0000000000000;  // 1
  298.     add.f64     %fd80, %fd76, %fd79;
  299.     mul.f64     %fd81, %fd78, %fd80;
  300.     sub.f64     %fd82, %fd81, %fd77;
  301.     mov.f64     %fd83, 0d3ffd9999a0000000;  // 1.85
  302.     mad.rn.f64  %fd84, %fd82, %fd83, %fd77;
  303.     cvt.rn.f32.f64  %f92, %fd84;
  304.     mov.f32     %f93, %f92;
  305.     add.f32     %f94, %f92, %f84;
  306.     .loc    18  1319    0
  307.     mov.f64     %fd85, 0d4022000000000000;  // 9
  308.     mul.f64     %fd86, %fd66, %fd85;
  309.     mov.f64     %fd87, 0d4008000000000000;  // 3
  310.     mul.f64     %fd88, %fd1, %fd87;
  311.     sub.f64     %fd89, %fd88, %fd86;
  312.     mov.f64     %fd90, 0d4008000000000000;  // 3
  313.     mad.rn.f64  %fd91, %fd2, %fd90, %fd89;
  314.     sub.f32     %f95, %f43, %f42;
  315.     cvt.f64.f32     %fd92, %f95;
  316.     mad.rn.f64  %fd93, %fd92, %fd68, %fd91;
  317.     cvt.rn.f32.f64  %f96, %fd93;
  318.     div.full.f32    %f97, %f96, %f51;
  319.     cvt.f64.f32     %fd94, %f97;
  320.     .loc    3   103 0
  321.     cvt.f64.f32     %fd95, %f13;
  322.     mov.f64     %fd96, 0d3ff0000000000000;  // 1
  323.     add.f64     %fd97, %fd94, %fd96;
  324.     mul.f64     %fd98, %fd78, %fd97;
  325.     sub.f64     %fd99, %fd98, %fd95;
  326.     mov.f64     %fd100, 0d3ffd9999a0000000; // 1.85
  327.     mad.rn.f64  %fd101, %fd99, %fd100, %fd95;
  328.     cvt.rn.f32.f64  %f98, %fd101;
  329.     mov.f32     %f99, %f98;
  330.     add.f32     %f100, %f98, %f94;
  331.     .loc    18  1319    0
  332.     mov.f64     %fd102, 0d4008000000000000; // 3
  333.     mad.rn.f64  %fd103, %fd1, %fd102, %fd86;
  334.     mov.f64     %fd104, 0d4008000000000000; // 3
  335.     mad.rn.f64  %fd105, %fd2, %fd104, %fd103;
  336.     neg.f32     %f101, %f85;
  337.     cvt.f64.f32     %fd106, %f101;
  338.     mad.rn.f64  %fd107, %fd106, %fd68, %fd105;
  339.     cvt.rn.f32.f64  %f102, %fd107;
  340.     div.full.f32    %f103, %f102, %f51;
  341.     cvt.f64.f32     %fd108, %f103;
  342.     .loc    3   104 0
  343.     cvt.f64.f32     %fd109, %f15;
  344.     mov.f64     %fd110, 0d3ff0000000000000; // 1
  345.     add.f64     %fd111, %fd108, %fd110;
  346.     mul.f64     %fd112, %fd78, %fd111;
  347.     sub.f64     %fd113, %fd112, %fd109;
  348.     mov.f64     %fd114, 0d3ffd9999a0000000; // 1.85
  349.     mad.rn.f64  %fd115, %fd113, %fd114, %fd109;
  350.     cvt.rn.f32.f64  %f104, %fd115;
  351.     mov.f32     %f105, %f104;
  352.     add.f32     %f106, %f104, %f100;
  353.     .loc    18  1319    0
  354.     sub.f32     %f107, %f42, %f43;
  355.     cvt.f64.f32     %fd116, %f107;
  356.     mad.rn.f64  %fd117, %fd116, %fd68, %fd91;
  357.     cvt.rn.f32.f64  %f108, %fd117;
  358.     div.full.f32    %f109, %f108, %f51;
  359.     cvt.f64.f32     %fd118, %f109;
  360.     .loc    3   105 0
  361.     cvt.f64.f32     %fd119, %f17;
  362.     mov.f64     %fd120, 0d3ff0000000000000; // 1
  363.     add.f64     %fd121, %fd118, %fd120;
  364.     mul.f64     %fd122, %fd78, %fd121;
  365.     sub.f64     %fd123, %fd122, %fd119;
  366.     mov.f64     %fd124, 0d3ffd9999a0000000; // 1.85
  367.     mad.rn.f64  %fd125, %fd123, %fd124, %fd119;
  368.     cvt.rn.f32.f64  %f110, %fd125;
  369.     mov.f32     %f111, %f110;
  370.     add.f32     %f112, %f110, %f106;
  371.     .loc    3   107 0
  372.     add.f32     %f113, %f92, %f64;
  373.     add.f32     %f114, %f110, %f113;
  374.     add.f32     %f115, %f98, %f76;
  375.     add.f32     %f116, %f104, %f115;
  376.     sub.f32     %f117, %f114, %f116;
  377.     div.full.f32    %f118, %f117, %f112;
  378.     bra.uni     $Lt_0_8450;
  379. $Lt_0_8706:
  380.     .loc    3   117 0
  381.     mov.f32     %f57, %f1;
  382.     .loc    3   118 0
  383.     mov.f32     %f76, %f3;
  384.     .loc    3   119 0
  385.     mov.f32     %f83, %f5;
  386.     .loc    3   120 0
  387.     mov.f32     %f64, %f7;
  388.     .loc    3   121 0
  389.     mov.f32     %f71, %f9;
  390.     .loc    3   122 0
  391.     mov.f32     %f105, %f11;
  392.     .loc    3   123 0
  393.     mov.f32     %f111, %f13;
  394.     .loc    3   124 0
  395.     mov.f32     %f93, %f15;
  396.     .loc    3   125 0
  397.     mov.f32     %f99, %f17;
  398.     .loc    3   126 0
  399.     mov.f32     %f118, 0f00000000;      // 0
  400. $Lt_0_8450:
  401.     mov.u64     %rd7, __cuda___cuda_local_var_15345_35_non_const_sdata36;
  402.     .loc    3   133 0
  403.     ld.param.u64    %rd8, [__cudaparm__Z9collisionPK7t_speedPKcPS_Pfi_tmpy_cells];
  404.     add.u64     %rd9, %rd8, %rd2;
  405.     st.global.f32   [%rd9+0], %f57;
  406.     st.global.f32   [%rd9+4], %f64;
  407.     st.global.f32   [%rd9+8], %f71;
  408.     st.global.f32   [%rd9+12], %f76;
  409.     st.global.f32   [%rd9+16], %f83;
  410.     st.global.f32   [%rd9+20], %f93;
  411.     st.global.f32   [%rd9+24], %f99;
  412.     st.global.f32   [%rd9+28], %f105;
  413.     st.global.f32   [%rd9+32], %f111;
  414.     .loc    3   137 0
  415.     cvt.s32.u16     %r17, %tid.x;
  416.     cvt.s64.s32     %rd10, %r17;
  417.     mul.wide.s32    %rd11, %r17, 4;
  418.     add.u64     %rd12, %rd7, %rd11;
  419.     st.shared.f32   [%rd12+0], %f118;
  420.     .loc    3   141 0
  421.     bar.sync    0;
  422.     mov.s32     %r18, 16;
  423. $Lt_0_9474:
  424.  //<loop> Loop body line 141, nesting depth: 1, estimated iterations: unknown
  425.     setp.le.s32     %p7, %r18, %r17;
  426.     @%p7 bra    $Lt_0_9730;
  427.     .loc    3   149 0
  428.     ld.shared.f32   %f119, [%rd12+0];
  429.     add.s32     %r19, %r18, %r17;
  430.     cvt.s64.s32     %rd13, %r19;
  431.     mul.wide.s32    %rd14, %r19, 4;
  432.     add.u64     %rd15, %rd7, %rd14;
  433.     ld.shared.f32   %f120, [%rd15+0];
  434.     add.f32     %f121, %f119, %f120;
  435.     st.shared.f32   [%rd12+0], %f121;
  436. $Lt_0_9730:
  437.     .loc    3   151 0
  438.     bar.sync    0;
  439.     .loc    3   143 0
  440.     shr.s32     %r20, %r18, 31;
  441.     mov.s32     %r21, 1;
  442.     and.b32     %r22, %r20, %r21;
  443.     add.s32     %r23, %r22, %r18;
  444.     shr.s32     %r18, %r23, 1;
  445.     mov.u32     %r24, 0;
  446.     setp.gt.s32     %p8, %r18, %r24;
  447.     @%p8 bra    $Lt_0_9474;
  448.     mov.u32     %r25, 0;
  449.     setp.ne.s32     %p9, %r17, %r25;
  450.     @%p9 bra    $Lt_0_10498;
  451.     .loc    3   154 0
  452.     ld.shared.f32   %f122, [__cuda___cuda_local_var_15345_35_non_const_sdata36+0];
  453.     ld.param.u64    %rd16, [__cudaparm__Z9collisionPK7t_speedPKcPS_Pfi_av_vels];
  454.     ld.param.s32    %r26, [__cudaparm__Z9collisionPK7t_speedPKcPS_Pfi_ii];
  455.     mul.lo.s32  %r27, %r26, 1875;
  456.     add.u32     %r28, %r1, %r27;
  457.     cvt.u64.u32     %rd17, %r28;
  458.     mul.wide.u32    %rd18, %r28, 4;
  459.     add.u64     %rd19, %rd16, %rd18;
  460.     st.global.f32   [%rd19+0], %f122;
  461. $Lt_0_10498:
  462.     .loc    3   156 0
  463.     exit;
  464. $LDWend__Z9collisionPK7t_speedPKcPS_Pfi:
  465.     } // _Z9collisionPK7t_speedPKcPS_Pfi
  466.  
  467.     .entry _Z9propagateP7t_speedPKS_ (
  468.         .param .u64 __cudaparm__Z9propagateP7t_speedPKS__cells,
  469.         .param .u64 __cudaparm__Z9propagateP7t_speedPKS__tmpy_cells)
  470.     {
  471.     .reg .u16 %rh<4>;
  472.     .reg .u32 %r<52>;
  473.     .reg .u64 %rd<32>;
  474.     .reg .f32 %f<15>;
  475.     .reg .pred %p<7>;
  476.     .loc    3   161 0
  477. $LDWbegin__Z9propagateP7t_speedPKS_:
  478.     .loc    3   204 0
  479.     mov.u16     %rh1, %ctaid.x;
  480.     mov.u16     %rh2, %ntid.x;
  481.     mul.wide.u16    %r1, %rh1, %rh2;
  482.     cvt.u32.u16     %r2, %tid.x;
  483.     add.u32     %r3, %r2, %r1;
  484.     cvt.rn.f32.s32  %f1, %r3;
  485.     mov.s32     %r4, 458129845;
  486.     mov.s32     %r5, 0;
  487.     setp.lt.s32     %p1, %r3, %r5;
  488.     abs.s32     %r6, %r3;
  489.     mul.hi.u32  %r7, %r6, %r4;
  490.     shr.s32     %r8, %r7, 5;
  491.     @%p1 sub.s32    %r8, %r5, %r8;
  492.     mov.s32     %r9, %r8;
  493.     mul.lo.s32  %r10, %r9, 300;
  494.     sub.s32     %r11, %r3, %r10;
  495.     mov.f32     %f2, 0f43960000;        // 300
  496.     div.full.f32    %f3, %f1, %f2;
  497.     cvt.rmi.f32.f32     %f4, %f3;
  498.     cvt.rzi.s32.f32     %r12, %f4;
  499.     mul.lo.s32  %r13, %r12, 300;
  500.     add.s32     %r14, %r11, %r13;
  501.     cvt.s64.s32     %rd1, %r14;
  502.     mul.wide.s32    %rd2, %r14, 36;
  503.     ld.param.u64    %rd3, [__cudaparm__Z9propagateP7t_speedPKS__tmpy_cells];
  504.     add.u64     %rd4, %rd3, %rd2;
  505.     ld.global.f32   %f5, [%rd4+0];
  506.     ld.global.f32   %f6, [%rd4+4];
  507.     ld.global.f32   %f7, [%rd4+8];
  508.     ld.global.f32   %f8, [%rd4+12];
  509.     ld.global.f32   %f9, [%rd4+16];
  510.     ld.global.f32   %f10, [%rd4+20];
  511.     ld.global.f32   %f11, [%rd4+24];
  512.     ld.global.f32   %f12, [%rd4+28];
  513.     ld.global.f32   %f13, [%rd4+32];
  514.     .loc    3   210 0
  515.     ld.param.u64    %rd5, [__cudaparm__Z9propagateP7t_speedPKS__cells];
  516.     add.u64     %rd6, %rd2, %rd5;
  517.     st.global.f32   [%rd6+0], %f5;
  518.     .loc    3   211 0
  519.     add.s32     %r15, %r11, 1;
  520.     mov.s32     %r16, 458129845;
  521.     mov.s32     %r17, 0;
  522.     setp.lt.s32     %p2, %r15, %r17;
  523.     abs.s32     %r18, %r15;
  524.     mul.hi.u32  %r19, %r18, %r16;
  525.     shr.s32     %r20, %r19, 5;
  526.     @%p2 sub.s32    %r20, %r17, %r20;
  527.     mov.s32     %r21, %r20;
  528.     mul.lo.s32  %r22, %r21, 300;
  529.     sub.s32     %r23, %r15, %r22;
  530.     add.s32     %r24, %r23, %r13;
  531.     cvt.s64.s32     %rd7, %r24;
  532.     mul.wide.s32    %rd8, %r24, 36;
  533.     add.u64     %rd9, %rd5, %rd8;
  534.     st.global.f32   [%rd9+4], %f6;
  535.     .loc    3   212 0
  536.     add.s32     %r25, %r12, 1;
  537.     mov.s32     %r26, 1374389535;
  538.     mov.s32     %r27, 0;
  539.     setp.lt.s32     %p3, %r25, %r27;
  540.     abs.s32     %r28, %r25;
  541.     mul.hi.u32  %r29, %r28, %r26;
  542.     shr.s32     %r30, %r29, 6;
  543.     @%p3 sub.s32    %r30, %r27, %r30;
  544.     mov.s32     %r31, %r30;
  545.     mul.lo.s32  %r32, %r31, 200;
  546.     sub.s32     %r33, %r25, %r32;
  547.     mul.lo.s32  %r34, %r33, 300;
  548.     add.s32     %r35, %r11, %r34;
  549.     cvt.s64.s32     %rd10, %r35;
  550.     mul.wide.s32    %rd11, %r35, 36;
  551.     add.u64     %rd12, %rd5, %rd11;
  552.     st.global.f32   [%rd12+8], %f7;
  553.     .loc    3   213 0
  554.     mov.s32     %r36, 0;
  555.     setp.ne.s32     %p4, %r11, %r36;
  556.     sub.s32     %r37, %r11, 1;
  557.     mov.s32     %r38, 299;
  558.     selp.s32    %r39, %r37, %r38, %p4;
  559.     add.s32     %r40, %r39, %r13;
  560.     cvt.s64.s32     %rd13, %r40;
  561.     mul.wide.s32    %rd14, %r40, 36;
  562.     add.u64     %rd15, %rd5, %rd14;
  563.     st.global.f32   [%rd15+12], %f8;
  564.     .loc    3   214 0
  565.     mov.s32     %r41, 0;
  566.     setp.ne.s32     %p5, %r12, %r41;
  567.     sub.s32     %r42, %r12, 1;
  568.     mov.s32     %r43, 199;
  569.     selp.s32    %r44, %r42, %r43, %p5;
  570.     mul.lo.s32  %r45, %r44, 300;
  571.     add.s32     %r46, %r11, %r45;
  572.     cvt.s64.s32     %rd16, %r46;
  573.     mul.wide.s32    %rd17, %r46, 36;
  574.     add.u64     %rd18, %rd5, %rd17;
  575.     st.global.f32   [%rd18+16], %f9;
  576.     .loc    3   215 0
  577.     add.s32     %r47, %r23, %r34;
  578.     cvt.s64.s32     %rd19, %r47;
  579.     mul.wide.s32    %rd20, %r47, 36;
  580.     add.u64     %rd21, %rd5, %rd20;
  581.     st.global.f32   [%rd21+20], %f10;
  582.     .loc    3   216 0
  583.     add.s32     %r48, %r39, %r34;
  584.     cvt.s64.s32     %rd22, %r48;
  585.     mul.wide.s32    %rd23, %r48, 36;
  586.     add.u64     %rd24, %rd5, %rd23;
  587.     st.global.f32   [%rd24+24], %f11;
  588.     .loc    3   217 0
  589.     add.s32     %r49, %r39, %r45;
  590.     cvt.s64.s32     %rd25, %r49;
  591.     mul.wide.s32    %rd26, %r49, 36;
  592.     add.u64     %rd27, %rd5, %rd26;
  593.     st.global.f32   [%rd27+28], %f12;
  594.     .loc    3   218 0
  595.     add.s32     %r50, %r23, %r45;
  596.     cvt.s64.s32     %rd28, %r50;
  597.     mul.wide.s32    %rd29, %r50, 36;
  598.     add.u64     %rd30, %rd5, %rd29;
  599.     st.global.f32   [%rd30+32], %f13;
  600.     .loc    3   221 0
  601.     exit;
  602. $LDWend__Z9propagateP7t_speedPKS_:
  603.     } // _Z9propagateP7t_speedPKS_
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!