Advertisement
dineiar

raytracer_kernel_ptx_nvrtc

Mar 30th, 2020
773
0
Never
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
text 86.20 KB | None | 0 0
  1. //
  2. // Generated by NVIDIA NVVM Compiler
  3. //
  4. // Compiler Build ID: CL-26218862
  5. // Cuda compilation tools, release 10.1, V10.1.168
  6. // Based on LLVM 3.4svn
  7. //
  8.  
  9. .version 6.4
  10. .target sm_30
  11. .address_size 64
  12.  
  13. .extern .func (.param .b32 func_retval0) _Z5powiffi
  14. (
  15. .param .b32 _Z5powiffi_param_0,
  16. .param .b32 _Z5powiffi_param_1
  17. )
  18. ;
  19.  
  20. .func (.param .align 4 .b8 func_retval0[12]) _Z5traceRK4Vec3IfES2_PK6SpherejRKi(
  21. .param .b32 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_0,
  22. .param .b32 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_1,
  23. .param .b32 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_2,
  24. .param .b64 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_3,
  25. .param .b64 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_4,
  26. .param .b32 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_5,
  27. .param .b64 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_6
  28. )
  29. {
  30. .local .align 4 .b8 __local_depot0[32];
  31. .reg .b64 %SP;
  32. .reg .b64 %SPL;
  33. .reg .pred %p<88>;
  34. .reg .b16 %rs<5>;
  35. .reg .f32 %f<480>;
  36. .reg .b32 %r<40>;
  37. .reg .b64 %rd<93>;
  38.  
  39.  
  40. mov.u64 %SPL, __local_depot0;
  41. cvta.local.u64 %SP, %SPL;
  42. ld.param.f32 %f179, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_0];
  43. ld.param.f32 %f180, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_1];
  44. ld.param.f32 %f181, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_2];
  45. ld.param.u64 %rd34, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_3];
  46. ld.param.u64 %rd35, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_4];
  47. ld.param.u32 %r13, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_5];
  48. ld.param.u64 %rd36, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_6];
  49. setp.eq.s32 %p8, %r13, 0;
  50. mov.f32 %f433, 0f4CBEBC20;
  51. mov.u64 %rd83, 0;
  52. @%p8 bra BB0_61;
  53.  
  54. cvta.to.local.u64 %rd39, %rd34;
  55. ld.local.f32 %f1, [%rd39];
  56. ld.local.f32 %f2, [%rd39+4];
  57. ld.local.f32 %f3, [%rd39+8];
  58. and.b32 %r1, %r13, 3;
  59. setp.eq.s32 %p9, %r1, 0;
  60. mov.f32 %f439, 0f00000000;
  61. mov.f32 %f433, 0f4CBEBC20;
  62. mov.u64 %rd83, 0;
  63. mov.u32 %r36, 0;
  64. @%p9 bra BB0_30;
  65.  
  66. setp.eq.s32 %p10, %r1, 1;
  67. mov.f32 %f435, 0f4CBEBC20;
  68. mov.u64 %rd83, 0;
  69. mov.u32 %r35, 0;
  70. mov.f32 %f433, %f435;
  71. @%p10 bra BB0_22;
  72.  
  73. setp.eq.s32 %p11, %r1, 2;
  74. mov.f32 %f431, 0f4CBEBC20;
  75. mov.u64 %rd83, 0;
  76. mov.u32 %r34, 0;
  77. @%p11 bra BB0_4;
  78. bra.uni BB0_5;
  79.  
  80. BB0_4:
  81. mov.f32 %f433, %f431;
  82. bra.uni BB0_14;
  83.  
  84. BB0_5:
  85. ld.f32 %f189, [%rd35+8];
  86. sub.f32 %f4, %f189, %f179;
  87. ld.f32 %f190, [%rd35+12];
  88. sub.f32 %f5, %f190, %f180;
  89. ld.f32 %f191, [%rd35+16];
  90. sub.f32 %f6, %f191, %f181;
  91. mul.f32 %f192, %f5, %f2;
  92. fma.rn.f32 %f193, %f4, %f1, %f192;
  93. fma.rn.f32 %f7, %f6, %f3, %f193;
  94. setp.lt.f32 %p13, %f7, 0f00000000;
  95. mov.pred %p81, -1;
  96. mov.f32 %f188, 0f4CBEBC20;
  97. @%p13 bra BB0_6;
  98. bra.uni BB0_7;
  99.  
  100. BB0_6:
  101. mov.f32 %f428, %f188;
  102. mov.f32 %f429, %f188;
  103. bra.uni BB0_9;
  104.  
  105. BB0_7:
  106. mul.f32 %f196, %f5, %f5;
  107. fma.rn.f32 %f197, %f4, %f4, %f196;
  108. fma.rn.f32 %f198, %f6, %f6, %f197;
  109. mul.f32 %f199, %f7, %f7;
  110. sub.f32 %f8, %f198, %f199;
  111. ld.f32 %f9, [%rd35+24];
  112. setp.gt.f32 %p15, %f8, %f9;
  113. mov.f32 %f428, %f188;
  114. mov.f32 %f429, %f188;
  115. @%p15 bra BB0_9;
  116.  
  117. sub.f32 %f200, %f9, %f8;
  118. sqrt.rn.f32 %f201, %f200;
  119. sub.f32 %f429, %f7, %f201;
  120. add.f32 %f428, %f7, %f201;
  121. mov.pred %p81, 0;
  122.  
  123. BB0_9:
  124. mov.u32 %r34, 1;
  125. @%p81 bra BB0_10;
  126. bra.uni BB0_11;
  127.  
  128. BB0_10:
  129. mov.f32 %f433, %f188;
  130. bra.uni BB0_14;
  131.  
  132. BB0_11:
  133. setp.lt.f32 %p17, %f429, 0f00000000;
  134. selp.f32 %f433, %f428, %f429, %p17;
  135. setp.geu.f32 %p18, %f433, 0f4CBEBC20;
  136. @%p18 bra BB0_12;
  137.  
  138. mov.u64 %rd83, %rd35;
  139. bra.uni BB0_14;
  140.  
  141. BB0_12:
  142. mov.f32 %f433, %f188;
  143.  
  144. BB0_14:
  145. mul.wide.u32 %rd44, %r34, 80;
  146. add.s64 %rd45, %rd35, %rd44;
  147. ld.f32 %f206, [%rd45+8];
  148. sub.f32 %f16, %f206, %f179;
  149. ld.f32 %f207, [%rd45+12];
  150. sub.f32 %f17, %f207, %f180;
  151. ld.f32 %f208, [%rd45+16];
  152. sub.f32 %f18, %f208, %f181;
  153. mul.f32 %f209, %f17, %f2;
  154. fma.rn.f32 %f210, %f16, %f1, %f209;
  155. fma.rn.f32 %f19, %f18, %f3, %f210;
  156. setp.lt.f32 %p20, %f19, 0f00000000;
  157. mov.pred %p82, -1;
  158. @%p20 bra BB0_15;
  159. bra.uni BB0_16;
  160.  
  161. BB0_15:
  162. mov.f32 %f432, %f431;
  163. bra.uni BB0_18;
  164.  
  165. BB0_16:
  166. mul.f32 %f213, %f17, %f17;
  167. fma.rn.f32 %f214, %f16, %f16, %f213;
  168. fma.rn.f32 %f215, %f18, %f18, %f214;
  169. mul.f32 %f216, %f19, %f19;
  170. sub.f32 %f20, %f215, %f216;
  171. ld.f32 %f21, [%rd45+24];
  172. setp.gt.f32 %p22, %f20, %f21;
  173. mov.f32 %f432, %f431;
  174. @%p22 bra BB0_18;
  175.  
  176. sub.f32 %f217, %f21, %f20;
  177. sqrt.rn.f32 %f218, %f217;
  178. sub.f32 %f432, %f19, %f218;
  179. add.f32 %f431, %f19, %f218;
  180. mov.pred %p82, 0;
  181.  
  182. BB0_18:
  183. @%p82 bra BB0_21;
  184.  
  185. setp.lt.f32 %p24, %f432, 0f00000000;
  186. selp.f32 %f26, %f431, %f432, %p24;
  187. setp.geu.f32 %p25, %f26, %f433;
  188. @%p25 bra BB0_21;
  189.  
  190. mov.f32 %f433, %f26;
  191. mov.u64 %rd83, %rd45;
  192.  
  193. BB0_21:
  194. add.s32 %r35, %r34, 1;
  195.  
  196. BB0_22:
  197. cvt.u64.u32 %rd5, %r35;
  198. mul.wide.u32 %rd49, %r35, 80;
  199. add.s64 %rd50, %rd35, %rd49;
  200. add.s64 %rd6, %rd50, 8;
  201. ld.f32 %f221, [%rd50+8];
  202. sub.f32 %f29, %f221, %f179;
  203. ld.f32 %f222, [%rd50+12];
  204. sub.f32 %f30, %f222, %f180;
  205. ld.f32 %f223, [%rd50+16];
  206. sub.f32 %f31, %f223, %f181;
  207. mul.f32 %f224, %f30, %f2;
  208. fma.rn.f32 %f225, %f29, %f1, %f224;
  209. fma.rn.f32 %f32, %f31, %f3, %f225;
  210. setp.lt.f32 %p27, %f32, 0f00000000;
  211. mov.pred %p83, -1;
  212. @%p27 bra BB0_23;
  213. bra.uni BB0_24;
  214.  
  215. BB0_23:
  216. mov.f32 %f436, %f435;
  217. bra.uni BB0_26;
  218.  
  219. BB0_24:
  220. mul.f32 %f228, %f30, %f30;
  221. fma.rn.f32 %f229, %f29, %f29, %f228;
  222. fma.rn.f32 %f230, %f31, %f31, %f229;
  223. mul.f32 %f231, %f32, %f32;
  224. sub.f32 %f33, %f230, %f231;
  225. ld.f32 %f34, [%rd6+16];
  226. setp.gt.f32 %p29, %f33, %f34;
  227. mov.f32 %f436, %f435;
  228. @%p29 bra BB0_26;
  229.  
  230. sub.f32 %f232, %f34, %f33;
  231. sqrt.rn.f32 %f233, %f232;
  232. sub.f32 %f436, %f32, %f233;
  233. add.f32 %f435, %f32, %f233;
  234. mov.pred %p83, 0;
  235.  
  236. BB0_26:
  237. @%p83 bra BB0_29;
  238.  
  239. setp.lt.f32 %p31, %f436, 0f00000000;
  240. selp.f32 %f39, %f435, %f436, %p31;
  241. setp.geu.f32 %p32, %f39, %f433;
  242. @%p32 bra BB0_29;
  243.  
  244. mul.lo.s64 %rd51, %rd5, 80;
  245. add.s64 %rd83, %rd35, %rd51;
  246. mov.f32 %f433, %f39;
  247.  
  248. BB0_29:
  249. cvt.u32.u64 %r20, %rd5;
  250. add.s32 %r36, %r20, 1;
  251. mov.f32 %f439, %f433;
  252.  
  253. BB0_30:
  254. setp.lt.u32 %p33, %r13, 4;
  255. @%p33 bra BB0_31;
  256. bra.uni BB0_32;
  257.  
  258. BB0_31:
  259. mov.f32 %f433, %f439;
  260. bra.uni BB0_61;
  261.  
  262. BB0_32:
  263. cvt.u64.u32 %rd11, %r36;
  264. mul.wide.u32 %rd52, %r36, 80;
  265. add.s64 %rd53, %rd35, %rd52;
  266. add.s64 %rd12, %rd53, 8;
  267. ld.f32 %f236, [%rd53+8];
  268. sub.f32 %f44, %f236, %f179;
  269. ld.f32 %f237, [%rd53+12];
  270. sub.f32 %f45, %f237, %f180;
  271. ld.f32 %f238, [%rd53+16];
  272. sub.f32 %f46, %f238, %f181;
  273. mul.f32 %f239, %f45, %f2;
  274. fma.rn.f32 %f240, %f44, %f1, %f239;
  275. fma.rn.f32 %f47, %f46, %f3, %f240;
  276. setp.lt.f32 %p35, %f47, 0f00000000;
  277. mov.pred %p84, -1;
  278. mov.f32 %f441, 0f4CBEBC20;
  279. @%p35 bra BB0_33;
  280. bra.uni BB0_34;
  281.  
  282. BB0_33:
  283. mov.f32 %f442, %f441;
  284. bra.uni BB0_36;
  285.  
  286. BB0_34:
  287. mul.f32 %f243, %f45, %f45;
  288. fma.rn.f32 %f244, %f44, %f44, %f243;
  289. fma.rn.f32 %f245, %f46, %f46, %f244;
  290. mul.f32 %f246, %f47, %f47;
  291. sub.f32 %f48, %f245, %f246;
  292. ld.f32 %f49, [%rd12+16];
  293. setp.gt.f32 %p37, %f48, %f49;
  294. mov.f32 %f442, %f441;
  295. @%p37 bra BB0_36;
  296.  
  297. sub.f32 %f247, %f49, %f48;
  298. sqrt.rn.f32 %f248, %f247;
  299. sub.f32 %f442, %f47, %f248;
  300. add.f32 %f441, %f47, %f248;
  301. mov.pred %p84, 0;
  302.  
  303. BB0_36:
  304. @%p84 bra BB0_39;
  305.  
  306. setp.lt.f32 %p39, %f442, 0f00000000;
  307. selp.f32 %f54, %f441, %f442, %p39;
  308. setp.geu.f32 %p40, %f54, %f433;
  309. @%p40 bra BB0_39;
  310.  
  311. mul.lo.s64 %rd54, %rd11, 80;
  312. add.s64 %rd83, %rd35, %rd54;
  313. mov.f32 %f433, %f54;
  314.  
  315. BB0_39:
  316. add.s32 %r21, %r36, 1;
  317. cvt.u64.u32 %rd15, %r21;
  318. mul.wide.u32 %rd55, %r21, 80;
  319. add.s64 %rd56, %rd35, %rd55;
  320. add.s64 %rd16, %rd56, 8;
  321. ld.f32 %f251, [%rd56+8];
  322. sub.f32 %f56, %f251, %f179;
  323. ld.f32 %f252, [%rd56+12];
  324. sub.f32 %f57, %f252, %f180;
  325. ld.f32 %f253, [%rd56+16];
  326. sub.f32 %f58, %f253, %f181;
  327. mul.f32 %f254, %f57, %f2;
  328. fma.rn.f32 %f255, %f56, %f1, %f254;
  329. fma.rn.f32 %f59, %f58, %f3, %f255;
  330. setp.lt.f32 %p42, %f59, 0f00000000;
  331. mov.pred %p85, -1;
  332. mov.f32 %f444, 0f4CBEBC20;
  333. @%p42 bra BB0_40;
  334. bra.uni BB0_41;
  335.  
  336. BB0_40:
  337. mov.f32 %f445, %f444;
  338. bra.uni BB0_43;
  339.  
  340. BB0_41:
  341. mul.f32 %f258, %f57, %f57;
  342. fma.rn.f32 %f259, %f56, %f56, %f258;
  343. fma.rn.f32 %f260, %f58, %f58, %f259;
  344. mul.f32 %f261, %f59, %f59;
  345. sub.f32 %f60, %f260, %f261;
  346. ld.f32 %f61, [%rd16+16];
  347. setp.gt.f32 %p44, %f60, %f61;
  348. mov.f32 %f445, %f444;
  349. @%p44 bra BB0_43;
  350.  
  351. sub.f32 %f262, %f61, %f60;
  352. sqrt.rn.f32 %f263, %f262;
  353. sub.f32 %f445, %f59, %f263;
  354. add.f32 %f444, %f59, %f263;
  355. mov.pred %p85, 0;
  356.  
  357. BB0_43:
  358. @%p85 bra BB0_46;
  359.  
  360. setp.lt.f32 %p46, %f445, 0f00000000;
  361. selp.f32 %f66, %f444, %f445, %p46;
  362. setp.geu.f32 %p47, %f66, %f433;
  363. @%p47 bra BB0_46;
  364.  
  365. mul.lo.s64 %rd57, %rd15, 80;
  366. add.s64 %rd83, %rd35, %rd57;
  367. mov.f32 %f433, %f66;
  368.  
  369. BB0_46:
  370. add.s32 %r22, %r36, 2;
  371. cvt.u64.u32 %rd19, %r22;
  372. mul.wide.u32 %rd58, %r22, 80;
  373. add.s64 %rd59, %rd35, %rd58;
  374. add.s64 %rd20, %rd59, 8;
  375. ld.f32 %f266, [%rd59+8];
  376. sub.f32 %f68, %f266, %f179;
  377. ld.f32 %f267, [%rd59+12];
  378. sub.f32 %f69, %f267, %f180;
  379. ld.f32 %f268, [%rd59+16];
  380. sub.f32 %f70, %f268, %f181;
  381. mul.f32 %f269, %f69, %f2;
  382. fma.rn.f32 %f270, %f68, %f1, %f269;
  383. fma.rn.f32 %f71, %f70, %f3, %f270;
  384. setp.lt.f32 %p49, %f71, 0f00000000;
  385. mov.pred %p86, -1;
  386. mov.f32 %f447, 0f4CBEBC20;
  387. @%p49 bra BB0_47;
  388. bra.uni BB0_48;
  389.  
  390. BB0_47:
  391. mov.f32 %f448, %f447;
  392. bra.uni BB0_50;
  393.  
  394. BB0_48:
  395. mul.f32 %f273, %f69, %f69;
  396. fma.rn.f32 %f274, %f68, %f68, %f273;
  397. fma.rn.f32 %f275, %f70, %f70, %f274;
  398. mul.f32 %f276, %f71, %f71;
  399. sub.f32 %f72, %f275, %f276;
  400. ld.f32 %f73, [%rd20+16];
  401. setp.gt.f32 %p51, %f72, %f73;
  402. mov.f32 %f448, %f447;
  403. @%p51 bra BB0_50;
  404.  
  405. sub.f32 %f277, %f73, %f72;
  406. sqrt.rn.f32 %f278, %f277;
  407. sub.f32 %f448, %f71, %f278;
  408. add.f32 %f447, %f71, %f278;
  409. mov.pred %p86, 0;
  410.  
  411. BB0_50:
  412. @%p86 bra BB0_53;
  413.  
  414. setp.lt.f32 %p53, %f448, 0f00000000;
  415. selp.f32 %f78, %f447, %f448, %p53;
  416. setp.geu.f32 %p54, %f78, %f433;
  417. @%p54 bra BB0_53;
  418.  
  419. mul.lo.s64 %rd60, %rd19, 80;
  420. add.s64 %rd83, %rd35, %rd60;
  421. mov.f32 %f433, %f78;
  422.  
  423. BB0_53:
  424. add.s32 %r23, %r36, 3;
  425. cvt.u64.u32 %rd23, %r23;
  426. mul.wide.u32 %rd61, %r23, 80;
  427. add.s64 %rd62, %rd35, %rd61;
  428. add.s64 %rd24, %rd62, 8;
  429. ld.f32 %f281, [%rd62+8];
  430. sub.f32 %f80, %f281, %f179;
  431. ld.f32 %f282, [%rd62+12];
  432. sub.f32 %f81, %f282, %f180;
  433. ld.f32 %f283, [%rd62+16];
  434. sub.f32 %f82, %f283, %f181;
  435. mul.f32 %f284, %f81, %f2;
  436. fma.rn.f32 %f285, %f80, %f1, %f284;
  437. fma.rn.f32 %f83, %f82, %f3, %f285;
  438. setp.lt.f32 %p56, %f83, 0f00000000;
  439. mov.pred %p87, -1;
  440. mov.f32 %f450, 0f4CBEBC20;
  441. @%p56 bra BB0_54;
  442. bra.uni BB0_55;
  443.  
  444. BB0_54:
  445. mov.f32 %f451, %f450;
  446. bra.uni BB0_57;
  447.  
  448. BB0_55:
  449. mul.f32 %f288, %f81, %f81;
  450. fma.rn.f32 %f289, %f80, %f80, %f288;
  451. fma.rn.f32 %f290, %f82, %f82, %f289;
  452. mul.f32 %f291, %f83, %f83;
  453. sub.f32 %f84, %f290, %f291;
  454. ld.f32 %f85, [%rd24+16];
  455. setp.gt.f32 %p58, %f84, %f85;
  456. mov.f32 %f451, %f450;
  457. @%p58 bra BB0_57;
  458.  
  459. sub.f32 %f292, %f85, %f84;
  460. sqrt.rn.f32 %f293, %f292;
  461. sub.f32 %f451, %f83, %f293;
  462. add.f32 %f450, %f83, %f293;
  463. mov.pred %p87, 0;
  464.  
  465. BB0_57:
  466. @%p87 bra BB0_60;
  467.  
  468. setp.lt.f32 %p60, %f451, 0f00000000;
  469. selp.f32 %f90, %f450, %f451, %p60;
  470. setp.geu.f32 %p61, %f90, %f433;
  471. @%p61 bra BB0_60;
  472.  
  473. mul.lo.s64 %rd63, %rd23, 80;
  474. add.s64 %rd83, %rd35, %rd63;
  475. mov.f32 %f433, %f90;
  476.  
  477. BB0_60:
  478. cvt.u32.u64 %r24, %rd11;
  479. add.s32 %r36, %r24, 4;
  480. setp.lt.u32 %p62, %r36, %r13;
  481. @%p62 bra BB0_32;
  482.  
  483. BB0_61:
  484. setp.eq.s64 %p63, %rd83, 0;
  485. mov.f32 %f477, 0f40000000;
  486. mov.f32 %f478, %f477;
  487. mov.f32 %f479, %f477;
  488. @%p63 bra BB0_89;
  489.  
  490. cvta.to.local.u64 %rd64, %rd34;
  491. ld.local.f32 %f93, [%rd64];
  492. ld.local.f32 %f94, [%rd64+4];
  493. ld.local.f32 %f95, [%rd64+8];
  494. fma.rn.f32 %f96, %f433, %f93, %f179;
  495. fma.rn.f32 %f97, %f433, %f94, %f180;
  496. fma.rn.f32 %f98, %f433, %f95, %f181;
  497. add.s64 %rd28, %rd83, 8;
  498. ld.f32 %f297, [%rd83+8];
  499. sub.f32 %f454, %f96, %f297;
  500. ld.f32 %f298, [%rd83+12];
  501. sub.f32 %f455, %f97, %f298;
  502. ld.f32 %f299, [%rd83+16];
  503. sub.f32 %f456, %f98, %f299;
  504. mul.f32 %f300, %f455, %f455;
  505. fma.rn.f32 %f301, %f454, %f454, %f300;
  506. fma.rn.f32 %f102, %f456, %f456, %f301;
  507. setp.leu.f32 %p64, %f102, 0f00000000;
  508. @%p64 bra BB0_64;
  509.  
  510. sqrt.rn.f32 %f302, %f102;
  511. rcp.rn.f32 %f303, %f302;
  512. mul.f32 %f454, %f454, %f303;
  513. mul.f32 %f455, %f455, %f303;
  514. mul.f32 %f456, %f456, %f303;
  515.  
  516. BB0_64:
  517. mul.f32 %f304, %f455, %f94;
  518. fma.rn.f32 %f305, %f454, %f93, %f304;
  519. fma.rn.f32 %f306, %f456, %f95, %f305;
  520. mov.u16 %rs4, 0;
  521. setp.leu.f32 %p65, %f306, 0f00000000;
  522. @%p65 bra BB0_66;
  523.  
  524. neg.f32 %f454, %f454;
  525. neg.f32 %f455, %f455;
  526. neg.f32 %f456, %f456;
  527. mov.u16 %rs4, 1;
  528.  
  529. BB0_66:
  530. ld.f32 %f307, [%rd28+44];
  531. setp.gt.f32 %p66, %f307, 0f00000000;
  532. @%p66 bra BB0_68;
  533.  
  534. ld.f32 %f308, [%rd28+48];
  535. setp.leu.f32 %p67, %f308, 0f00000000;
  536. @%p67 bra BB0_69;
  537.  
  538. BB0_68:
  539. cvta.to.local.u64 %rd65, %rd36;
  540. ld.local.u32 %r25, [%rd65];
  541. setp.lt.s32 %p68, %r25, 5;
  542. @%p68 bra BB0_81;
  543. bra.uni BB0_69;
  544.  
  545. BB0_81:
  546. mul.f32 %f351, %f455, %f94;
  547. fma.rn.f32 %f352, %f454, %f93, %f351;
  548. fma.rn.f32 %f353, %f456, %f95, %f352;
  549. add.f32 %f354, %f353, 0f3F800000;
  550. mov.u32 %r29, 3;
  551. // Callseq Start 0
  552. {
  553. .reg .b32 temp_param_reg;
  554. // <end>}
  555. .param .b32 param0;
  556. st.param.f32 [param0+0], %f354;
  557. .param .b32 param1;
  558. st.param.b32 [param1+0], %r29;
  559. .param .b32 retval0;
  560. call.uni (retval0),
  561. _Z5powiffi,
  562. (
  563. param0,
  564. param1
  565. );
  566. ld.param.f32 %f355, [retval0+0];
  567.  
  568. //{
  569. }// Callseq End 0
  570. fma.rn.f32 %f143, %f355, 0f3F666666, 0f3DCCCCCD;
  571. ld.local.f32 %f356, [%rd64];
  572. ld.local.f32 %f357, [%rd64+4];
  573. mul.f32 %f358, %f455, %f357;
  574. fma.rn.f32 %f359, %f454, %f356, %f358;
  575. ld.local.f32 %f360, [%rd64+8];
  576. fma.rn.f32 %f361, %f456, %f360, %f359;
  577. add.f32 %f362, %f454, %f454;
  578. mul.f32 %f363, %f362, %f361;
  579. add.f32 %f364, %f455, %f455;
  580. mul.f32 %f365, %f364, %f361;
  581. add.f32 %f366, %f456, %f456;
  582. mul.f32 %f367, %f366, %f361;
  583. sub.f32 %f144, %f356, %f363;
  584. sub.f32 %f145, %f357, %f365;
  585. sub.f32 %f146, %f360, %f367;
  586. add.u64 %rd71, %SP, 4;
  587. add.u64 %rd32, %SPL, 4;
  588. st.local.f32 [%rd32], %f144;
  589. st.local.f32 [%rd32+4], %f145;
  590. st.local.f32 [%rd32+8], %f146;
  591. mul.f32 %f368, %f145, %f145;
  592. fma.rn.f32 %f369, %f144, %f144, %f368;
  593. fma.rn.f32 %f147, %f146, %f146, %f369;
  594. setp.leu.f32 %p77, %f147, 0f00000000;
  595. @%p77 bra BB0_83;
  596.  
  597. sqrt.rn.f32 %f370, %f147;
  598. rcp.rn.f32 %f371, %f370;
  599. mul.f32 %f372, %f371, %f144;
  600. st.local.f32 [%rd32], %f372;
  601. mul.f32 %f373, %f371, %f145;
  602. st.local.f32 [%rd32+4], %f373;
  603. mul.f32 %f374, %f371, %f146;
  604. st.local.f32 [%rd32+8], %f374;
  605.  
  606. BB0_83:
  607. mul.f32 %f148, %f454, 0f38D1B717;
  608. add.f32 %f378, %f96, %f148;
  609. mul.f32 %f149, %f455, 0f38D1B717;
  610. add.f32 %f379, %f97, %f149;
  611. mul.f32 %f150, %f456, 0f38D1B717;
  612. add.f32 %f380, %f98, %f150;
  613. ld.local.u32 %r30, [%rd65];
  614. add.s32 %r31, %r30, 1;
  615. add.u64 %rd73, %SP, 0;
  616. add.u64 %rd74, %SPL, 0;
  617. st.local.u32 [%rd74], %r31;
  618. // Callseq Start 1
  619. {
  620. .reg .b32 temp_param_reg;
  621. // <end>}
  622. .param .b32 param0;
  623. st.param.f32 [param0+0], %f378;
  624. .param .b32 param1;
  625. st.param.f32 [param1+0], %f379;
  626. .param .b32 param2;
  627. st.param.f32 [param2+0], %f380;
  628. .param .b64 param3;
  629. st.param.b64 [param3+0], %rd71;
  630. .param .b64 param4;
  631. st.param.b64 [param4+0], %rd35;
  632. .param .b32 param5;
  633. st.param.b32 [param5+0], %r13;
  634. .param .b64 param6;
  635. st.param.b64 [param6+0], %rd73;
  636. .param .align 4 .b8 retval0[12];
  637. call.uni (retval0),
  638. _Z5traceRK4Vec3IfES2_PK6SpherejRKi,
  639. (
  640. param0,
  641. param1,
  642. param2,
  643. param3,
  644. param4,
  645. param5,
  646. param6
  647. );
  648. ld.param.f32 %f151, [retval0+0];
  649. ld.param.f32 %f152, [retval0+4];
  650. ld.param.f32 %f153, [retval0+8];
  651.  
  652. //{
  653. }// Callseq End 1
  654. ld.f32 %f470, [%rd28+44];
  655. mov.f32 %f471, 0f00000000;
  656. setp.eq.f32 %p78, %f470, 0f00000000;
  657. mov.f32 %f472, %f471;
  658. mov.f32 %f473, %f471;
  659. @%p78 bra BB0_87;
  660.  
  661. setp.eq.s16 %p79, %rs4, 0;
  662. selp.f32 %f381, 0f3F68BA2E, 0f3F8CCCCD, %p79;
  663. ld.local.f32 %f382, [%rd64];
  664. ld.local.f32 %f383, [%rd64+4];
  665. mul.f32 %f384, %f455, %f383;
  666. fma.rn.f32 %f385, %f454, %f382, %f384;
  667. ld.local.f32 %f386, [%rd64+8];
  668. fma.rn.f32 %f387, %f456, %f386, %f385;
  669. mul.f32 %f388, %f381, %f381;
  670. mul.f32 %f389, %f387, %f387;
  671. mov.f32 %f390, 0f3F800000;
  672. sub.f32 %f391, %f390, %f389;
  673. mul.f32 %f392, %f388, %f391;
  674. sub.f32 %f393, %f390, %f392;
  675. mul.f32 %f394, %f381, %f387;
  676. neg.f32 %f395, %f394;
  677. sqrt.rn.f32 %f396, %f393;
  678. sub.f32 %f397, %f395, %f396;
  679. mul.f32 %f398, %f454, %f397;
  680. mul.f32 %f399, %f455, %f397;
  681. mul.f32 %f400, %f456, %f397;
  682. fma.rn.f32 %f155, %f381, %f382, %f398;
  683. fma.rn.f32 %f156, %f381, %f383, %f399;
  684. fma.rn.f32 %f157, %f381, %f386, %f400;
  685. add.u64 %rd77, %SP, 20;
  686. add.u64 %rd33, %SPL, 20;
  687. st.local.f32 [%rd33], %f155;
  688. st.local.f32 [%rd33+4], %f156;
  689. st.local.f32 [%rd33+8], %f157;
  690. mul.f32 %f401, %f156, %f156;
  691. fma.rn.f32 %f402, %f155, %f155, %f401;
  692. fma.rn.f32 %f158, %f157, %f157, %f402;
  693. setp.leu.f32 %p80, %f158, 0f00000000;
  694. @%p80 bra BB0_86;
  695.  
  696. sqrt.rn.f32 %f403, %f158;
  697. rcp.rn.f32 %f404, %f403;
  698. mul.f32 %f405, %f404, %f155;
  699. st.local.f32 [%rd33], %f405;
  700. mul.f32 %f406, %f404, %f156;
  701. st.local.f32 [%rd33+4], %f406;
  702. mul.f32 %f407, %f404, %f157;
  703. st.local.f32 [%rd33+8], %f407;
  704.  
  705. BB0_86:
  706. sub.f32 %f408, %f96, %f148;
  707. ld.local.u32 %r32, [%rd65];
  708. add.s32 %r33, %r32, 1;
  709. add.u64 %rd79, %SP, 16;
  710. add.u64 %rd80, %SPL, 16;
  711. st.local.u32 [%rd80], %r33;
  712. sub.f32 %f409, %f98, %f150;
  713. sub.f32 %f410, %f97, %f149;
  714. // Callseq Start 2
  715. {
  716. .reg .b32 temp_param_reg;
  717. // <end>}
  718. .param .b32 param0;
  719. st.param.f32 [param0+0], %f408;
  720. .param .b32 param1;
  721. st.param.f32 [param1+0], %f410;
  722. .param .b32 param2;
  723. st.param.f32 [param2+0], %f409;
  724. .param .b64 param3;
  725. st.param.b64 [param3+0], %rd77;
  726. .param .b64 param4;
  727. st.param.b64 [param4+0], %rd35;
  728. .param .b32 param5;
  729. st.param.b32 [param5+0], %r13;
  730. .param .b64 param6;
  731. st.param.b64 [param6+0], %rd79;
  732. .param .align 4 .b8 retval0[12];
  733. call.uni (retval0),
  734. _Z5traceRK4Vec3IfES2_PK6SpherejRKi,
  735. (
  736. param0,
  737. param1,
  738. param2,
  739. param3,
  740. param4,
  741. param5,
  742. param6
  743. );
  744. ld.param.f32 %f471, [retval0+0];
  745. ld.param.f32 %f472, [retval0+4];
  746. ld.param.f32 %f473, [retval0+8];
  747.  
  748. //{
  749. }// Callseq End 2
  750. ld.f32 %f470, [%rd28+44];
  751.  
  752. BB0_87:
  753. mov.f32 %f411, 0f3F800000;
  754. sub.f32 %f412, %f411, %f143;
  755. mul.f32 %f413, %f412, %f471;
  756. mul.f32 %f414, %f412, %f472;
  757. mul.f32 %f415, %f412, %f473;
  758. mul.f32 %f416, %f413, %f470;
  759. mul.f32 %f417, %f414, %f470;
  760. mul.f32 %f418, %f415, %f470;
  761. fma.rn.f32 %f419, %f143, %f151, %f416;
  762. fma.rn.f32 %f420, %f143, %f152, %f417;
  763. fma.rn.f32 %f421, %f143, %f153, %f418;
  764. ld.f32 %f422, [%rd28+20];
  765. mul.f32 %f474, %f422, %f419;
  766. ld.f32 %f423, [%rd28+24];
  767. mul.f32 %f475, %f420, %f423;
  768. ld.f32 %f424, [%rd28+28];
  769. mul.f32 %f476, %f421, %f424;
  770. bra.uni BB0_88;
  771.  
  772. BB0_69:
  773. mov.f32 %f474, 0f00000000;
  774. mov.f32 %f475, %f474;
  775. mov.f32 %f476, %f474;
  776. @%p8 bra BB0_88;
  777.  
  778. fma.rn.f32 %f115, %f454, 0f38D1B717, %f96;
  779. fma.rn.f32 %f116, %f455, 0f38D1B717, %f97;
  780. fma.rn.f32 %f117, %f456, 0f38D1B717, %f98;
  781. mov.f32 %f476, 0f00000000;
  782. mov.u32 %r38, 0;
  783. mov.f32 %f475, %f476;
  784. mov.f32 %f474, %f476;
  785.  
  786. BB0_71:
  787. cvt.u64.u32 %rd29, %r38;
  788. mul.wide.u32 %rd66, %r38, 80;
  789. add.s64 %rd67, %rd35, %rd66;
  790. add.s64 %rd30, %rd67, 40;
  791. ld.f32 %f121, [%rd67+40];
  792. setp.leu.f32 %p70, %f121, 0f00000000;
  793. @%p70 bra BB0_80;
  794.  
  795. ld.f32 %f315, [%rd30+-32];
  796. sub.f32 %f465, %f315, %f96;
  797. ld.f32 %f316, [%rd30+-28];
  798. sub.f32 %f464, %f316, %f97;
  799. ld.f32 %f317, [%rd30+-24];
  800. sub.f32 %f463, %f317, %f98;
  801. mul.f32 %f318, %f464, %f464;
  802. fma.rn.f32 %f319, %f465, %f465, %f318;
  803. fma.rn.f32 %f125, %f463, %f463, %f319;
  804. setp.leu.f32 %p71, %f125, 0f00000000;
  805. @%p71 bra BB0_74;
  806.  
  807. sqrt.rn.f32 %f320, %f125;
  808. rcp.rn.f32 %f321, %f320;
  809. mul.f32 %f465, %f465, %f321;
  810. mul.f32 %f464, %f464, %f321;
  811. mul.f32 %f463, %f463, %f321;
  812.  
  813. BB0_74:
  814. mov.u32 %r39, 0;
  815.  
  816. BB0_75:
  817. cvt.u32.u64 %r28, %rd29;
  818. setp.eq.s32 %p72, %r28, %r39;
  819. @%p72 bra BB0_78;
  820.  
  821. mul.wide.u32 %rd68, %r39, 80;
  822. add.s64 %rd69, %rd35, %rd68;
  823. add.s64 %rd31, %rd69, 8;
  824. ld.f32 %f322, [%rd69+8];
  825. sub.f32 %f132, %f322, %f115;
  826. ld.f32 %f323, [%rd69+12];
  827. sub.f32 %f133, %f323, %f116;
  828. ld.f32 %f324, [%rd69+16];
  829. sub.f32 %f134, %f324, %f117;
  830. mul.f32 %f325, %f464, %f133;
  831. fma.rn.f32 %f326, %f465, %f132, %f325;
  832. fma.rn.f32 %f135, %f463, %f134, %f326;
  833. setp.lt.f32 %p73, %f135, 0f00000000;
  834. @%p73 bra BB0_78;
  835.  
  836. mul.f32 %f328, %f133, %f133;
  837. fma.rn.f32 %f329, %f132, %f132, %f328;
  838. fma.rn.f32 %f330, %f134, %f134, %f329;
  839. mul.f32 %f331, %f135, %f135;
  840. sub.f32 %f332, %f330, %f331;
  841. ld.f32 %f333, [%rd31+16];
  842. mov.f32 %f466, 0f00000000;
  843. setp.leu.f32 %p74, %f332, %f333;
  844. @%p74 bra BB0_79;
  845.  
  846. BB0_78:
  847. add.s32 %r39, %r39, 1;
  848. setp.lt.u32 %p75, %r39, %r13;
  849. mov.f32 %f466, 0f3F800000;
  850. @%p75 bra BB0_75;
  851.  
  852. BB0_79:
  853. ld.f32 %f335, [%rd28+20];
  854. mul.f32 %f336, %f466, %f335;
  855. ld.f32 %f337, [%rd28+24];
  856. mul.f32 %f338, %f466, %f337;
  857. ld.f32 %f339, [%rd28+28];
  858. mul.f32 %f340, %f466, %f339;
  859. mul.f32 %f341, %f454, %f465;
  860. fma.rn.f32 %f342, %f455, %f464, %f341;
  861. fma.rn.f32 %f343, %f456, %f463, %f342;
  862. mov.f32 %f344, 0f00000000;
  863. max.f32 %f345, %f344, %f343;
  864. mul.f32 %f346, %f336, %f345;
  865. mul.f32 %f347, %f338, %f345;
  866. mul.f32 %f348, %f340, %f345;
  867. ld.f32 %f349, [%rd30+4];
  868. ld.f32 %f350, [%rd30+8];
  869. fma.rn.f32 %f474, %f346, %f121, %f474;
  870. fma.rn.f32 %f475, %f347, %f349, %f475;
  871. fma.rn.f32 %f476, %f348, %f350, %f476;
  872.  
  873. BB0_80:
  874. add.s32 %r38, %r38, 1;
  875. setp.lt.u32 %p76, %r38, %r13;
  876. @%p76 bra BB0_71;
  877.  
  878. BB0_88:
  879. ld.f32 %f425, [%rd28+32];
  880. add.f32 %f477, %f474, %f425;
  881. ld.f32 %f426, [%rd28+36];
  882. add.f32 %f478, %f475, %f426;
  883. ld.f32 %f427, [%rd28+40];
  884. add.f32 %f479, %f476, %f427;
  885.  
  886. BB0_89:
  887. st.param.f32 [func_retval0+0], %f477;
  888. st.param.f32 [func_retval0+4], %f478;
  889. st.param.f32 [func_retval0+8], %f479;
  890. ret;
  891. }
  892.  
  893. // .globl raytrace_kernel
  894. .visible .entry raytrace_kernel(
  895. .param .u32 raytrace_kernel_param_0,
  896. .param .u32 raytrace_kernel_param_1,
  897. .param .u64 raytrace_kernel_param_2,
  898. .param .u64 raytrace_kernel_param_3,
  899. .param .u32 raytrace_kernel_param_4,
  900. .param .f32 raytrace_kernel_param_5,
  901. .param .f32 raytrace_kernel_param_6,
  902. .param .f32 raytrace_kernel_param_7,
  903. .param .f32 raytrace_kernel_param_8
  904. )
  905. {
  906. .local .align 4 .b8 __local_depot1[56];
  907. .reg .b64 %SP;
  908. .reg .b64 %SPL;
  909. .reg .pred %p<263>;
  910. .reg .b16 %rs<13>;
  911. .reg .f32 %f<1399>;
  912. .reg .b32 %r<148>;
  913. .reg .f64 %fd<18>;
  914. .reg .b64 %rd<391>;
  915.  
  916.  
  917. mov.u64 %SPL, __local_depot1;
  918. cvta.local.u64 %SP, %SPL;
  919. ld.param.u32 %r36, [raytrace_kernel_param_0];
  920. ld.param.u32 %r38, [raytrace_kernel_param_1];
  921. ld.param.u64 %rd158, [raytrace_kernel_param_3];
  922. ld.param.u32 %r37, [raytrace_kernel_param_4];
  923. ld.param.f32 %f539, [raytrace_kernel_param_5];
  924. ld.param.f32 %f540, [raytrace_kernel_param_6];
  925. ld.param.f32 %f541, [raytrace_kernel_param_7];
  926. ld.param.f32 %f542, [raytrace_kernel_param_8];
  927. cvta.to.global.u64 %rd1, %rd158;
  928. mov.u32 %r39, %ntid.x;
  929. mov.u32 %r40, %ctaid.x;
  930. mov.u32 %r41, %tid.x;
  931. mad.lo.s32 %r1, %r39, %r40, %r41;
  932. mov.u32 %r42, %ntid.y;
  933. mov.u32 %r43, %ctaid.y;
  934. mov.u32 %r44, %tid.y;
  935. mad.lo.s32 %r2, %r42, %r43, %r44;
  936. setp.ge.u32 %p22, %r2, %r38;
  937. setp.ge.u32 %p23, %r1, %r36;
  938. or.pred %p24, %p22, %p23;
  939. @%p24 bra BB1_271;
  940.  
  941. cvt.rn.f64.s32 %fd1, %r1;
  942. add.f64 %fd2, %fd1, 0d3FE0000000000000;
  943. cvt.f64.f32 %fd3, %f539;
  944. mul.f64 %fd4, %fd3, %fd2;
  945. fma.rn.f64 %fd5, %fd4, 0d4000000000000000, 0dBFF0000000000000;
  946. cvt.f64.f32 %fd6, %f542;
  947. mul.f64 %fd7, %fd6, %fd5;
  948. cvt.f64.f32 %fd8, %f541;
  949. mul.f64 %fd9, %fd8, %fd7;
  950. cvt.rn.f32.f64 %f9, %fd9;
  951. cvt.rn.f64.s32 %fd10, %r2;
  952. add.f64 %fd11, %fd10, 0d3FE0000000000000;
  953. cvt.f64.f32 %fd12, %f540;
  954. mul.f64 %fd13, %fd12, %fd11;
  955. fma.rn.f64 %fd14, %fd12, %fd11, %fd13;
  956. mov.f64 %fd15, 0d3FF0000000000000;
  957. sub.f64 %fd16, %fd15, %fd14;
  958. mul.f64 %fd17, %fd6, %fd16;
  959. cvt.rn.f32.f64 %f8, %fd17;
  960. mul.f32 %f544, %f8, %f8;
  961. fma.rn.f32 %f545, %f9, %f9, %f544;
  962. add.f32 %f3, %f545, 0f3F800000;
  963. mov.f32 %f7, 0fBF800000;
  964. setp.leu.f32 %p25, %f3, 0f00000000;
  965. @%p25 bra BB1_3;
  966.  
  967. sqrt.rn.f32 %f546, %f3;
  968. rcp.rn.f32 %f547, %f546;
  969. mul.f32 %f9, %f9, %f547;
  970. mul.f32 %f8, %f8, %f547;
  971. neg.f32 %f7, %f547;
  972.  
  973. BB1_3:
  974. setp.eq.s32 %p26, %r37, 0;
  975. mov.f32 %f1263, 0f4CBEBC20;
  976. mov.u64 %rd328, 0;
  977. mov.u64 %rd327, %rd328;
  978. @%p26 bra BB1_63;
  979.  
  980. and.b32 %r48, %r37, 3;
  981. mov.f32 %f1263, 0f4CBEBC20;
  982. mov.u64 %rd328, 0;
  983. mov.u32 %r132, 0;
  984. mov.f32 %f1250, 0f00000000;
  985. setp.eq.s32 %p27, %r48, 0;
  986. mov.u64 %rd327, %rd328;
  987. @%p27 bra BB1_32;
  988.  
  989. mov.f32 %f1263, 0f4CBEBC20;
  990. mov.u64 %rd328, 0;
  991. mov.u32 %r131, 0;
  992. and.b32 %r126, %r37, 3;
  993. setp.eq.s32 %p28, %r126, 1;
  994. mov.u64 %rd327, %rd328;
  995. @%p28 bra BB1_24;
  996.  
  997. mov.f32 %f1263, 0f4CBEBC20;
  998. mov.u64 %rd328, 0;
  999. mov.u32 %r130, 0;
  1000. and.b32 %r128, %r37, 3;
  1001. setp.eq.s32 %p29, %r128, 2;
  1002. mov.u64 %rd327, %rd328;
  1003. @%p29 bra BB1_16;
  1004.  
  1005. ld.global.f32 %f10, [%rd1+8];
  1006. ld.global.f32 %f11, [%rd1+12];
  1007. mul.f32 %f555, %f8, %f11;
  1008. fma.rn.f32 %f556, %f9, %f10, %f555;
  1009. ld.global.f32 %f12, [%rd1+16];
  1010. fma.rn.f32 %f13, %f7, %f12, %f556;
  1011. setp.lt.f32 %p31, %f13, 0f00000000;
  1012. mov.pred %p242, -1;
  1013. mov.f32 %f1239, 0f4CBEBC20;
  1014. @%p31 bra BB1_8;
  1015. bra.uni BB1_9;
  1016.  
  1017. BB1_8:
  1018. mov.f32 %f1240, %f1239;
  1019. bra.uni BB1_11;
  1020.  
  1021. BB1_9:
  1022. mov.f32 %f1239, 0f4CBEBC20;
  1023. add.s64 %rd286, %rd1, 8;
  1024. mul.f32 %f559, %f11, %f11;
  1025. fma.rn.f32 %f560, %f10, %f10, %f559;
  1026. fma.rn.f32 %f561, %f12, %f12, %f560;
  1027. mul.f32 %f562, %f13, %f13;
  1028. sub.f32 %f14, %f561, %f562;
  1029. ld.global.f32 %f15, [%rd286+16];
  1030. setp.gt.f32 %p33, %f14, %f15;
  1031. mov.f32 %f1240, %f1239;
  1032. @%p33 bra BB1_11;
  1033.  
  1034. sub.f32 %f563, %f15, %f14;
  1035. sqrt.rn.f32 %f564, %f563;
  1036. sub.f32 %f1240, %f13, %f564;
  1037. add.f32 %f1239, %f13, %f564;
  1038. mov.pred %p242, 0;
  1039.  
  1040. BB1_11:
  1041. mov.u64 %rd328, 0;
  1042. mov.f32 %f1263, 0f4CBEBC20;
  1043. mov.u32 %r130, 1;
  1044. @%p242 bra BB1_12;
  1045. bra.uni BB1_13;
  1046.  
  1047. BB1_12:
  1048. mov.u64 %rd327, %rd328;
  1049. bra.uni BB1_16;
  1050.  
  1051. BB1_13:
  1052. mov.u64 %rd328, 0;
  1053. mov.f32 %f1219, 0f4CBEBC20;
  1054. setp.lt.f32 %p35, %f1240, 0f00000000;
  1055. selp.f32 %f1263, %f1239, %f1240, %p35;
  1056. setp.geu.f32 %p36, %f1263, 0f4CBEBC20;
  1057. @%p36 bra BB1_14;
  1058.  
  1059. mov.u64 %rd328, %rd158;
  1060. mov.u64 %rd327, %rd1;
  1061. bra.uni BB1_16;
  1062.  
  1063. BB1_14:
  1064. mov.u64 %rd327, %rd328;
  1065. mov.f32 %f1263, %f1219;
  1066.  
  1067. BB1_16:
  1068. mov.f32 %f1242, 0f4CBEBC20;
  1069. mul.wide.u32 %rd171, %r130, 80;
  1070. add.s64 %rd172, %rd1, %rd171;
  1071. ld.global.f32 %f22, [%rd172+8];
  1072. ld.global.f32 %f23, [%rd172+12];
  1073. mul.f32 %f569, %f8, %f23;
  1074. fma.rn.f32 %f570, %f9, %f22, %f569;
  1075. ld.global.f32 %f24, [%rd172+16];
  1076. fma.rn.f32 %f25, %f7, %f24, %f570;
  1077. setp.lt.f32 %p38, %f25, 0f00000000;
  1078. mov.pred %p243, -1;
  1079. @%p38 bra BB1_17;
  1080. bra.uni BB1_18;
  1081.  
  1082. BB1_17:
  1083. mov.f32 %f1243, %f1242;
  1084. bra.uni BB1_20;
  1085.  
  1086. BB1_18:
  1087. mul.wide.u32 %rd282, %r130, 80;
  1088. add.s64 %rd281, %rd1, %rd282;
  1089. add.s64 %rd280, %rd281, 8;
  1090. mov.f32 %f1242, 0f4CBEBC20;
  1091. mul.f32 %f573, %f23, %f23;
  1092. fma.rn.f32 %f574, %f22, %f22, %f573;
  1093. fma.rn.f32 %f575, %f24, %f24, %f574;
  1094. mul.f32 %f576, %f25, %f25;
  1095. sub.f32 %f26, %f575, %f576;
  1096. ld.global.f32 %f27, [%rd280+16];
  1097. setp.gt.f32 %p40, %f26, %f27;
  1098. mov.f32 %f1243, %f1242;
  1099. @%p40 bra BB1_20;
  1100.  
  1101. sub.f32 %f577, %f27, %f26;
  1102. sqrt.rn.f32 %f578, %f577;
  1103. sub.f32 %f1243, %f25, %f578;
  1104. add.f32 %f1242, %f25, %f578;
  1105. mov.pred %p243, 0;
  1106.  
  1107. BB1_20:
  1108. @%p243 bra BB1_23;
  1109.  
  1110. setp.lt.f32 %p42, %f1243, 0f00000000;
  1111. selp.f32 %f32, %f1242, %f1243, %p42;
  1112. setp.geu.f32 %p43, %f32, %f1263;
  1113. @%p43 bra BB1_23;
  1114.  
  1115. cvt.u64.u32 %rd283, %r130;
  1116. mul.lo.s64 %rd173, %rd283, 80;
  1117. add.s64 %rd327, %rd1, %rd173;
  1118. add.s64 %rd328, %rd158, %rd173;
  1119. mov.f32 %f1263, %f32;
  1120.  
  1121. BB1_23:
  1122. cvt.u64.u32 %rd284, %r130;
  1123. cvt.u32.u64 %r52, %rd284;
  1124. add.s32 %r131, %r52, 1;
  1125.  
  1126. BB1_24:
  1127. mov.f32 %f1246, 0f4CBEBC20;
  1128. mul.wide.u32 %rd174, %r131, 80;
  1129. add.s64 %rd175, %rd1, %rd174;
  1130. ld.global.f32 %f35, [%rd175+8];
  1131. ld.global.f32 %f36, [%rd175+12];
  1132. mul.f32 %f581, %f8, %f36;
  1133. fma.rn.f32 %f582, %f9, %f35, %f581;
  1134. ld.global.f32 %f37, [%rd175+16];
  1135. fma.rn.f32 %f38, %f7, %f37, %f582;
  1136. setp.lt.f32 %p45, %f38, 0f00000000;
  1137. mov.pred %p244, -1;
  1138. @%p45 bra BB1_25;
  1139. bra.uni BB1_26;
  1140.  
  1141. BB1_25:
  1142. mov.f32 %f1247, %f1246;
  1143. bra.uni BB1_28;
  1144.  
  1145. BB1_26:
  1146. mul.wide.u32 %rd291, %r131, 80;
  1147. add.s64 %rd290, %rd1, %rd291;
  1148. add.s64 %rd289, %rd290, 8;
  1149. mov.f32 %f1246, 0f4CBEBC20;
  1150. mul.f32 %f585, %f36, %f36;
  1151. fma.rn.f32 %f586, %f35, %f35, %f585;
  1152. fma.rn.f32 %f587, %f37, %f37, %f586;
  1153. mul.f32 %f588, %f38, %f38;
  1154. sub.f32 %f39, %f587, %f588;
  1155. ld.global.f32 %f40, [%rd289+16];
  1156. setp.gt.f32 %p47, %f39, %f40;
  1157. mov.f32 %f1247, %f1246;
  1158. @%p47 bra BB1_28;
  1159.  
  1160. sub.f32 %f589, %f40, %f39;
  1161. sqrt.rn.f32 %f590, %f589;
  1162. sub.f32 %f1247, %f38, %f590;
  1163. add.f32 %f1246, %f38, %f590;
  1164. mov.pred %p244, 0;
  1165.  
  1166. BB1_28:
  1167. @%p244 bra BB1_31;
  1168.  
  1169. setp.lt.f32 %p49, %f1247, 0f00000000;
  1170. selp.f32 %f45, %f1246, %f1247, %p49;
  1171. setp.geu.f32 %p50, %f45, %f1263;
  1172. @%p50 bra BB1_31;
  1173.  
  1174. cvt.u64.u32 %rd292, %r131;
  1175. mul.lo.s64 %rd176, %rd292, 80;
  1176. add.s64 %rd327, %rd1, %rd176;
  1177. add.s64 %rd328, %rd158, %rd176;
  1178. mov.f32 %f1263, %f45;
  1179.  
  1180. BB1_31:
  1181. cvt.u64.u32 %rd293, %r131;
  1182. cvt.u32.u64 %r53, %rd293;
  1183. add.s32 %r132, %r53, 1;
  1184. mov.f32 %f1250, %f1263;
  1185.  
  1186. BB1_32:
  1187. setp.lt.u32 %p51, %r37, 4;
  1188. @%p51 bra BB1_33;
  1189. bra.uni BB1_34;
  1190.  
  1191. BB1_33:
  1192. mov.f32 %f1263, %f1250;
  1193. bra.uni BB1_63;
  1194.  
  1195. BB1_34:
  1196. mul.wide.u32 %rd177, %r132, 80;
  1197. add.s64 %rd178, %rd1, %rd177;
  1198. ld.global.f32 %f50, [%rd178+8];
  1199. ld.global.f32 %f51, [%rd178+12];
  1200. mul.f32 %f593, %f8, %f51;
  1201. fma.rn.f32 %f594, %f9, %f50, %f593;
  1202. ld.global.f32 %f52, [%rd178+16];
  1203. fma.rn.f32 %f53, %f7, %f52, %f594;
  1204. setp.lt.f32 %p53, %f53, 0f00000000;
  1205. mov.pred %p245, -1;
  1206. mov.f32 %f1252, 0f4CBEBC20;
  1207. @%p53 bra BB1_35;
  1208. bra.uni BB1_36;
  1209.  
  1210. BB1_35:
  1211. mov.f32 %f1253, %f1252;
  1212. bra.uni BB1_38;
  1213.  
  1214. BB1_36:
  1215. mul.wide.u32 %rd296, %r132, 80;
  1216. add.s64 %rd295, %rd1, %rd296;
  1217. add.s64 %rd294, %rd295, 8;
  1218. mul.f32 %f597, %f51, %f51;
  1219. fma.rn.f32 %f598, %f50, %f50, %f597;
  1220. fma.rn.f32 %f599, %f52, %f52, %f598;
  1221. mul.f32 %f600, %f53, %f53;
  1222. sub.f32 %f54, %f599, %f600;
  1223. ld.global.f32 %f55, [%rd294+16];
  1224. setp.gt.f32 %p55, %f54, %f55;
  1225. mov.f32 %f1253, %f1252;
  1226. @%p55 bra BB1_38;
  1227.  
  1228. sub.f32 %f601, %f55, %f54;
  1229. sqrt.rn.f32 %f602, %f601;
  1230. sub.f32 %f1253, %f53, %f602;
  1231. add.f32 %f1252, %f53, %f602;
  1232. mov.pred %p245, 0;
  1233.  
  1234. BB1_38:
  1235. @%p245 bra BB1_41;
  1236.  
  1237. setp.lt.f32 %p57, %f1253, 0f00000000;
  1238. selp.f32 %f60, %f1252, %f1253, %p57;
  1239. setp.geu.f32 %p58, %f60, %f1263;
  1240. @%p58 bra BB1_41;
  1241.  
  1242. cvt.u64.u32 %rd297, %r132;
  1243. mul.lo.s64 %rd179, %rd297, 80;
  1244. add.s64 %rd327, %rd1, %rd179;
  1245. add.s64 %rd328, %rd158, %rd179;
  1246. mov.f32 %f1263, %f60;
  1247.  
  1248. BB1_41:
  1249. add.s32 %r54, %r132, 1;
  1250. mul.wide.u32 %rd180, %r54, 80;
  1251. add.s64 %rd181, %rd1, %rd180;
  1252. ld.global.f32 %f62, [%rd181+8];
  1253. ld.global.f32 %f63, [%rd181+12];
  1254. mul.f32 %f605, %f8, %f63;
  1255. fma.rn.f32 %f606, %f9, %f62, %f605;
  1256. ld.global.f32 %f64, [%rd181+16];
  1257. fma.rn.f32 %f65, %f7, %f64, %f606;
  1258. setp.lt.f32 %p60, %f65, 0f00000000;
  1259. mov.pred %p246, -1;
  1260. mov.f32 %f1255, 0f4CBEBC20;
  1261. @%p60 bra BB1_42;
  1262. bra.uni BB1_43;
  1263.  
  1264. BB1_42:
  1265. mov.f32 %f1256, %f1255;
  1266. bra.uni BB1_45;
  1267.  
  1268. BB1_43:
  1269. add.s32 %r118, %r132, 1;
  1270. mul.wide.u32 %rd301, %r118, 80;
  1271. add.s64 %rd300, %rd1, %rd301;
  1272. add.s64 %rd299, %rd300, 8;
  1273. mul.f32 %f609, %f63, %f63;
  1274. fma.rn.f32 %f610, %f62, %f62, %f609;
  1275. fma.rn.f32 %f611, %f64, %f64, %f610;
  1276. mul.f32 %f612, %f65, %f65;
  1277. sub.f32 %f66, %f611, %f612;
  1278. ld.global.f32 %f67, [%rd299+16];
  1279. setp.gt.f32 %p62, %f66, %f67;
  1280. mov.f32 %f1256, %f1255;
  1281. @%p62 bra BB1_45;
  1282.  
  1283. sub.f32 %f613, %f67, %f66;
  1284. sqrt.rn.f32 %f614, %f613;
  1285. sub.f32 %f1256, %f65, %f614;
  1286. add.f32 %f1255, %f65, %f614;
  1287. mov.pred %p246, 0;
  1288.  
  1289. BB1_45:
  1290. @%p246 bra BB1_48;
  1291.  
  1292. setp.lt.f32 %p64, %f1256, 0f00000000;
  1293. selp.f32 %f72, %f1255, %f1256, %p64;
  1294. setp.geu.f32 %p65, %f72, %f1263;
  1295. @%p65 bra BB1_48;
  1296.  
  1297. add.s32 %r119, %r132, 1;
  1298. cvt.u64.u32 %rd302, %r119;
  1299. mul.lo.s64 %rd182, %rd302, 80;
  1300. add.s64 %rd327, %rd1, %rd182;
  1301. add.s64 %rd328, %rd158, %rd182;
  1302. mov.f32 %f1263, %f72;
  1303.  
  1304. BB1_48:
  1305. add.s32 %r55, %r132, 2;
  1306. mul.wide.u32 %rd183, %r55, 80;
  1307. add.s64 %rd184, %rd1, %rd183;
  1308. ld.global.f32 %f74, [%rd184+8];
  1309. ld.global.f32 %f75, [%rd184+12];
  1310. mul.f32 %f617, %f8, %f75;
  1311. fma.rn.f32 %f618, %f9, %f74, %f617;
  1312. ld.global.f32 %f76, [%rd184+16];
  1313. fma.rn.f32 %f77, %f7, %f76, %f618;
  1314. setp.lt.f32 %p67, %f77, 0f00000000;
  1315. mov.pred %p247, -1;
  1316. mov.f32 %f1258, 0f4CBEBC20;
  1317. @%p67 bra BB1_49;
  1318. bra.uni BB1_50;
  1319.  
  1320. BB1_49:
  1321. mov.f32 %f1259, %f1258;
  1322. bra.uni BB1_52;
  1323.  
  1324. BB1_50:
  1325. add.s32 %r120, %r132, 2;
  1326. mul.wide.u32 %rd305, %r120, 80;
  1327. add.s64 %rd304, %rd1, %rd305;
  1328. add.s64 %rd303, %rd304, 8;
  1329. mul.f32 %f621, %f75, %f75;
  1330. fma.rn.f32 %f622, %f74, %f74, %f621;
  1331. fma.rn.f32 %f623, %f76, %f76, %f622;
  1332. mul.f32 %f624, %f77, %f77;
  1333. sub.f32 %f78, %f623, %f624;
  1334. ld.global.f32 %f79, [%rd303+16];
  1335. setp.gt.f32 %p69, %f78, %f79;
  1336. mov.f32 %f1259, %f1258;
  1337. @%p69 bra BB1_52;
  1338.  
  1339. sub.f32 %f625, %f79, %f78;
  1340. sqrt.rn.f32 %f626, %f625;
  1341. sub.f32 %f1259, %f77, %f626;
  1342. add.f32 %f1258, %f77, %f626;
  1343. mov.pred %p247, 0;
  1344.  
  1345. BB1_52:
  1346. @%p247 bra BB1_55;
  1347.  
  1348. setp.lt.f32 %p71, %f1259, 0f00000000;
  1349. selp.f32 %f84, %f1258, %f1259, %p71;
  1350. setp.geu.f32 %p72, %f84, %f1263;
  1351. @%p72 bra BB1_55;
  1352.  
  1353. add.s32 %r121, %r132, 2;
  1354. cvt.u64.u32 %rd306, %r121;
  1355. mul.lo.s64 %rd185, %rd306, 80;
  1356. add.s64 %rd327, %rd1, %rd185;
  1357. add.s64 %rd328, %rd158, %rd185;
  1358. mov.f32 %f1263, %f84;
  1359.  
  1360. BB1_55:
  1361. add.s32 %r56, %r132, 3;
  1362. mul.wide.u32 %rd186, %r56, 80;
  1363. add.s64 %rd187, %rd1, %rd186;
  1364. ld.global.f32 %f86, [%rd187+8];
  1365. ld.global.f32 %f87, [%rd187+12];
  1366. mul.f32 %f629, %f8, %f87;
  1367. fma.rn.f32 %f630, %f9, %f86, %f629;
  1368. ld.global.f32 %f88, [%rd187+16];
  1369. fma.rn.f32 %f89, %f7, %f88, %f630;
  1370. setp.lt.f32 %p74, %f89, 0f00000000;
  1371. mov.pred %p248, -1;
  1372. mov.f32 %f1261, 0f4CBEBC20;
  1373. @%p74 bra BB1_56;
  1374. bra.uni BB1_57;
  1375.  
  1376. BB1_56:
  1377. mov.f32 %f1262, %f1261;
  1378. bra.uni BB1_59;
  1379.  
  1380. BB1_57:
  1381. add.s32 %r122, %r132, 3;
  1382. mul.wide.u32 %rd309, %r122, 80;
  1383. add.s64 %rd308, %rd1, %rd309;
  1384. add.s64 %rd307, %rd308, 8;
  1385. mul.f32 %f633, %f87, %f87;
  1386. fma.rn.f32 %f634, %f86, %f86, %f633;
  1387. fma.rn.f32 %f635, %f88, %f88, %f634;
  1388. mul.f32 %f636, %f89, %f89;
  1389. sub.f32 %f90, %f635, %f636;
  1390. ld.global.f32 %f91, [%rd307+16];
  1391. setp.gt.f32 %p76, %f90, %f91;
  1392. mov.f32 %f1262, %f1261;
  1393. @%p76 bra BB1_59;
  1394.  
  1395. sub.f32 %f637, %f91, %f90;
  1396. sqrt.rn.f32 %f638, %f637;
  1397. sub.f32 %f1262, %f89, %f638;
  1398. add.f32 %f1261, %f89, %f638;
  1399. mov.pred %p248, 0;
  1400.  
  1401. BB1_59:
  1402. @%p248 bra BB1_62;
  1403.  
  1404. setp.lt.f32 %p78, %f1262, 0f00000000;
  1405. selp.f32 %f96, %f1261, %f1262, %p78;
  1406. setp.geu.f32 %p79, %f96, %f1263;
  1407. @%p79 bra BB1_62;
  1408.  
  1409. add.s32 %r123, %r132, 3;
  1410. cvt.u64.u32 %rd310, %r123;
  1411. mul.lo.s64 %rd188, %rd310, 80;
  1412. add.s64 %rd327, %rd1, %rd188;
  1413. add.s64 %rd328, %rd158, %rd188;
  1414. mov.f32 %f1263, %f96;
  1415.  
  1416. BB1_62:
  1417. cvt.u64.u32 %rd298, %r132;
  1418. cvt.u32.u64 %r57, %rd298;
  1419. add.s32 %r132, %r57, 4;
  1420. setp.lt.u32 %p80, %r132, %r37;
  1421. @%p80 bra BB1_34;
  1422.  
  1423. BB1_63:
  1424. setp.eq.s64 %p81, %rd328, 0;
  1425. mov.f32 %f1396, 0f40000000;
  1426. mov.f32 %f1397, %f1396;
  1427. mov.f32 %f1398, %f1396;
  1428. @%p81 bra BB1_270;
  1429.  
  1430. fma.rn.f32 %f99, %f9, %f1263, 0f00000000;
  1431. fma.rn.f32 %f100, %f8, %f1263, 0f00000000;
  1432. fma.rn.f32 %f101, %f7, %f1263, 0f00000000;
  1433. add.s64 %rd49, %rd327, 8;
  1434. ld.global.f32 %f642, [%rd327+8];
  1435. sub.f32 %f1268, %f99, %f642;
  1436. ld.global.f32 %f643, [%rd327+12];
  1437. sub.f32 %f1269, %f100, %f643;
  1438. ld.global.f32 %f644, [%rd327+16];
  1439. sub.f32 %f1270, %f101, %f644;
  1440. mul.f32 %f645, %f1269, %f1269;
  1441. fma.rn.f32 %f646, %f1268, %f1268, %f645;
  1442. fma.rn.f32 %f105, %f1270, %f1270, %f646;
  1443. setp.leu.f32 %p82, %f105, 0f00000000;
  1444. @%p82 bra BB1_66;
  1445.  
  1446. sqrt.rn.f32 %f647, %f105;
  1447. rcp.rn.f32 %f648, %f647;
  1448. mul.f32 %f1268, %f1268, %f648;
  1449. mul.f32 %f1269, %f1269, %f648;
  1450. mul.f32 %f1270, %f1270, %f648;
  1451.  
  1452. BB1_66:
  1453. mul.f32 %f649, %f8, %f1269;
  1454. fma.rn.f32 %f650, %f9, %f1268, %f649;
  1455. fma.rn.f32 %f651, %f7, %f1270, %f650;
  1456. mov.u16 %rs1, 0;
  1457. setp.leu.f32 %p83, %f651, 0f00000000;
  1458. @%p83 bra BB1_68;
  1459.  
  1460. neg.f32 %f1268, %f1268;
  1461. neg.f32 %f1269, %f1269;
  1462. neg.f32 %f1270, %f1270;
  1463. mov.u16 %rs1, 1;
  1464.  
  1465. BB1_68:
  1466. ld.global.f32 %f652, [%rd49+44];
  1467. setp.gt.f32 %p84, %f652, 0f00000000;
  1468. @%p84 bra BB1_82;
  1469.  
  1470. ld.global.f32 %f653, [%rd49+48];
  1471. setp.gt.f32 %p85, %f653, 0f00000000;
  1472. @%p85 bra BB1_82;
  1473. bra.uni BB1_70;
  1474.  
  1475. BB1_82:
  1476. mul.f32 %f696, %f8, %f1269;
  1477. fma.rn.f32 %f697, %f9, %f1268, %f696;
  1478. fma.rn.f32 %f146, %f7, %f1270, %f697;
  1479. add.f32 %f698, %f146, 0f3F800000;
  1480. mov.u32 %r61, 3;
  1481. // Callseq Start 3
  1482. {
  1483. .reg .b32 temp_param_reg;
  1484. // <end>}
  1485. .param .b32 param0;
  1486. st.param.f32 [param0+0], %f698;
  1487. .param .b32 param1;
  1488. st.param.b32 [param1+0], %r61;
  1489. .param .b32 retval0;
  1490. call.uni (retval0),
  1491. _Z5powiffi,
  1492. (
  1493. param0,
  1494. param1
  1495. );
  1496. ld.param.f32 %f699, [retval0+0];
  1497.  
  1498. //{
  1499. }// Callseq End 3
  1500. fma.rn.f32 %f147, %f699, 0f3F666666, 0f3DCCCCCD;
  1501. add.f32 %f700, %f1268, %f1268;
  1502. mul.f32 %f701, %f700, %f146;
  1503. add.f32 %f702, %f1269, %f1269;
  1504. mul.f32 %f703, %f702, %f146;
  1505. add.f32 %f704, %f1270, %f1270;
  1506. mul.f32 %f705, %f704, %f146;
  1507. sub.f32 %f1281, %f9, %f701;
  1508. sub.f32 %f1282, %f8, %f703;
  1509. sub.f32 %f1283, %f7, %f705;
  1510. mul.f32 %f706, %f1282, %f1282;
  1511. fma.rn.f32 %f707, %f1281, %f1281, %f706;
  1512. fma.rn.f32 %f151, %f1283, %f1283, %f707;
  1513. setp.leu.f32 %p94, %f151, 0f00000000;
  1514. @%p94 bra BB1_84;
  1515.  
  1516. sqrt.rn.f32 %f708, %f151;
  1517. rcp.rn.f32 %f709, %f708;
  1518. mul.f32 %f1281, %f1281, %f709;
  1519. mul.f32 %f1282, %f1282, %f709;
  1520. mul.f32 %f1283, %f1283, %f709;
  1521.  
  1522. BB1_84:
  1523. mul.f32 %f158, %f1268, 0f38D1B717;
  1524. add.f32 %f159, %f99, %f158;
  1525. mul.f32 %f160, %f1269, 0f38D1B717;
  1526. add.f32 %f161, %f100, %f160;
  1527. mul.f32 %f162, %f1270, 0f38D1B717;
  1528. add.f32 %f163, %f101, %f162;
  1529. mov.f32 %f1308, 0f4CBEBC20;
  1530. mov.u64 %rd365, 0;
  1531. mov.u64 %rd350, %rd365;
  1532. @%p26 bra BB1_146;
  1533.  
  1534. and.b32 %r65, %r37, 3;
  1535. mov.f32 %f1308, 0f4CBEBC20;
  1536. mov.u64 %rd365, 0;
  1537. mov.u32 %r138, 0;
  1538. mov.f32 %f1295, 0f00000000;
  1539. setp.eq.s32 %p96, %r65, 0;
  1540. @%p96 bra BB1_86;
  1541.  
  1542. setp.eq.s32 %p97, %r65, 1;
  1543. @%p97 bra BB1_88;
  1544. bra.uni BB1_89;
  1545.  
  1546. BB1_88:
  1547. mov.u64 %rd350, %rd365;
  1548. bra.uni BB1_107;
  1549.  
  1550. BB1_70:
  1551. mov.f32 %f1393, 0f00000000;
  1552. mov.f32 %f1394, %f1393;
  1553. mov.f32 %f1395, %f1393;
  1554. @%p26 bra BB1_269;
  1555.  
  1556. fma.rn.f32 %f118, %f1268, 0f38D1B717, %f99;
  1557. fma.rn.f32 %f119, %f1269, 0f38D1B717, %f100;
  1558. fma.rn.f32 %f120, %f1270, 0f38D1B717, %f101;
  1559. mov.f32 %f1395, 0f00000000;
  1560. mov.u32 %r134, 0;
  1561. mov.f32 %f1394, %f1395;
  1562. mov.f32 %f1393, %f1395;
  1563.  
  1564. BB1_72:
  1565. cvt.u64.u32 %rd50, %r134;
  1566. mul.wide.u32 %rd189, %r134, 80;
  1567. add.s64 %rd190, %rd1, %rd189;
  1568. add.s64 %rd51, %rd190, 40;
  1569. ld.global.f32 %f124, [%rd190+40];
  1570. setp.leu.f32 %p87, %f124, 0f00000000;
  1571. @%p87 bra BB1_81;
  1572.  
  1573. ld.global.f32 %f660, [%rd51+-32];
  1574. sub.f32 %f1276, %f660, %f99;
  1575. ld.global.f32 %f661, [%rd51+-28];
  1576. sub.f32 %f1275, %f661, %f100;
  1577. ld.global.f32 %f662, [%rd51+-24];
  1578. sub.f32 %f1274, %f662, %f101;
  1579. mul.f32 %f663, %f1275, %f1275;
  1580. fma.rn.f32 %f664, %f1276, %f1276, %f663;
  1581. fma.rn.f32 %f128, %f1274, %f1274, %f664;
  1582. setp.leu.f32 %p88, %f128, 0f00000000;
  1583. @%p88 bra BB1_75;
  1584.  
  1585. sqrt.rn.f32 %f665, %f128;
  1586. rcp.rn.f32 %f666, %f665;
  1587. mul.f32 %f1276, %f1276, %f666;
  1588. mul.f32 %f1275, %f1275, %f666;
  1589. mul.f32 %f1274, %f1274, %f666;
  1590.  
  1591. BB1_75:
  1592. mov.u32 %r135, 0;
  1593.  
  1594. BB1_76:
  1595. setp.eq.s32 %p89, %r134, %r135;
  1596. @%p89 bra BB1_79;
  1597.  
  1598. mul.wide.u32 %rd191, %r135, 80;
  1599. add.s64 %rd192, %rd1, %rd191;
  1600. add.s64 %rd52, %rd192, 8;
  1601. ld.global.f32 %f667, [%rd192+8];
  1602. sub.f32 %f135, %f667, %f118;
  1603. ld.global.f32 %f668, [%rd192+12];
  1604. sub.f32 %f136, %f668, %f119;
  1605. ld.global.f32 %f669, [%rd192+16];
  1606. sub.f32 %f137, %f669, %f120;
  1607. mul.f32 %f670, %f1275, %f136;
  1608. fma.rn.f32 %f671, %f1276, %f135, %f670;
  1609. fma.rn.f32 %f138, %f1274, %f137, %f671;
  1610. setp.lt.f32 %p90, %f138, 0f00000000;
  1611. @%p90 bra BB1_79;
  1612.  
  1613. mul.f32 %f673, %f136, %f136;
  1614. fma.rn.f32 %f674, %f135, %f135, %f673;
  1615. fma.rn.f32 %f675, %f137, %f137, %f674;
  1616. mul.f32 %f676, %f138, %f138;
  1617. sub.f32 %f677, %f675, %f676;
  1618. ld.global.f32 %f678, [%rd52+16];
  1619. mov.f32 %f1277, 0f00000000;
  1620. setp.leu.f32 %p91, %f677, %f678;
  1621. @%p91 bra BB1_80;
  1622.  
  1623. BB1_79:
  1624. add.s32 %r135, %r135, 1;
  1625. setp.lt.u32 %p92, %r135, %r37;
  1626. mov.f32 %f1277, 0f3F800000;
  1627. @%p92 bra BB1_76;
  1628.  
  1629. BB1_80:
  1630. ld.global.f32 %f680, [%rd49+20];
  1631. mul.f32 %f681, %f1277, %f680;
  1632. ld.global.f32 %f682, [%rd49+24];
  1633. mul.f32 %f683, %f1277, %f682;
  1634. ld.global.f32 %f684, [%rd49+28];
  1635. mul.f32 %f685, %f1277, %f684;
  1636. mul.f32 %f686, %f1268, %f1276;
  1637. fma.rn.f32 %f687, %f1269, %f1275, %f686;
  1638. fma.rn.f32 %f688, %f1270, %f1274, %f687;
  1639. mov.f32 %f689, 0f00000000;
  1640. max.f32 %f690, %f689, %f688;
  1641. mul.f32 %f691, %f681, %f690;
  1642. mul.f32 %f692, %f683, %f690;
  1643. mul.f32 %f693, %f685, %f690;
  1644. ld.global.f32 %f694, [%rd51+4];
  1645. ld.global.f32 %f695, [%rd51+8];
  1646. fma.rn.f32 %f1393, %f691, %f124, %f1393;
  1647. fma.rn.f32 %f1394, %f692, %f694, %f1394;
  1648. fma.rn.f32 %f1395, %f693, %f695, %f1395;
  1649.  
  1650. BB1_81:
  1651. cvt.u32.u64 %r60, %rd50;
  1652. add.s32 %r134, %r60, 1;
  1653. setp.lt.u32 %p93, %r134, %r37;
  1654. @%p93 bra BB1_72;
  1655. bra.uni BB1_269;
  1656.  
  1657. BB1_86:
  1658. mov.u64 %rd350, %rd365;
  1659. bra.uni BB1_115;
  1660.  
  1661. BB1_89:
  1662. setp.eq.s32 %p98, %r65, 2;
  1663. mov.u64 %rd350, %rd365;
  1664. @%p98 bra BB1_99;
  1665.  
  1666. ld.global.f32 %f717, [%rd1+8];
  1667. sub.f32 %f164, %f717, %f159;
  1668. ld.global.f32 %f718, [%rd1+12];
  1669. sub.f32 %f165, %f718, %f161;
  1670. ld.global.f32 %f719, [%rd1+16];
  1671. sub.f32 %f166, %f719, %f163;
  1672. mul.f32 %f720, %f1282, %f165;
  1673. fma.rn.f32 %f721, %f1281, %f164, %f720;
  1674. fma.rn.f32 %f167, %f1283, %f166, %f721;
  1675. setp.lt.f32 %p100, %f167, 0f00000000;
  1676. mov.pred %p249, -1;
  1677. mov.f32 %f1284, 0f4CBEBC20;
  1678. @%p100 bra BB1_91;
  1679. bra.uni BB1_92;
  1680.  
  1681. BB1_91:
  1682. mov.f32 %f1285, %f1284;
  1683. bra.uni BB1_94;
  1684.  
  1685. BB1_92:
  1686. mov.f32 %f1284, 0f4CBEBC20;
  1687. mul.f32 %f724, %f165, %f165;
  1688. fma.rn.f32 %f725, %f164, %f164, %f724;
  1689. fma.rn.f32 %f726, %f166, %f166, %f725;
  1690. mul.f32 %f727, %f167, %f167;
  1691. sub.f32 %f168, %f726, %f727;
  1692. ld.global.f32 %f169, [%rd1+24];
  1693. setp.gt.f32 %p102, %f168, %f169;
  1694. mov.f32 %f1285, %f1284;
  1695. @%p102 bra BB1_94;
  1696.  
  1697. sub.f32 %f728, %f169, %f168;
  1698. sqrt.rn.f32 %f729, %f728;
  1699. sub.f32 %f1285, %f167, %f729;
  1700. add.f32 %f1284, %f167, %f729;
  1701. mov.pred %p249, 0;
  1702.  
  1703. BB1_94:
  1704. mov.u64 %rd365, 0;
  1705. mov.f32 %f1308, 0f4CBEBC20;
  1706. mov.u32 %r138, 1;
  1707. @%p249 bra BB1_95;
  1708. bra.uni BB1_96;
  1709.  
  1710. BB1_95:
  1711. mov.u64 %rd350, %rd365;
  1712. bra.uni BB1_99;
  1713.  
  1714. BB1_96:
  1715. mov.u64 %rd365, 0;
  1716. mov.f32 %f1234, 0f4CBEBC20;
  1717. setp.lt.f32 %p104, %f1285, 0f00000000;
  1718. selp.f32 %f1308, %f1284, %f1285, %p104;
  1719. setp.geu.f32 %p105, %f1308, 0f4CBEBC20;
  1720. @%p105 bra BB1_97;
  1721.  
  1722. mov.u64 %rd365, %rd1;
  1723. mov.u64 %rd350, %rd158;
  1724. bra.uni BB1_99;
  1725.  
  1726. BB1_97:
  1727. mov.u64 %rd350, %rd365;
  1728. mov.f32 %f1308, %f1234;
  1729.  
  1730. BB1_99:
  1731. mov.f32 %f1287, 0f4CBEBC20;
  1732. mul.wide.u32 %rd205, %r138, 80;
  1733. add.s64 %rd206, %rd1, %rd205;
  1734. ld.global.f32 %f734, [%rd206+8];
  1735. sub.f32 %f176, %f734, %f159;
  1736. ld.global.f32 %f735, [%rd206+12];
  1737. sub.f32 %f177, %f735, %f161;
  1738. ld.global.f32 %f736, [%rd206+16];
  1739. sub.f32 %f178, %f736, %f163;
  1740. mul.f32 %f737, %f1282, %f177;
  1741. fma.rn.f32 %f738, %f1281, %f176, %f737;
  1742. fma.rn.f32 %f179, %f1283, %f178, %f738;
  1743. setp.lt.f32 %p107, %f179, 0f00000000;
  1744. mov.pred %p250, -1;
  1745. @%p107 bra BB1_100;
  1746. bra.uni BB1_101;
  1747.  
  1748. BB1_100:
  1749. mov.f32 %f1288, %f1287;
  1750. bra.uni BB1_103;
  1751.  
  1752. BB1_101:
  1753. mul.wide.u32 %rd313, %r138, 80;
  1754. add.s64 %rd312, %rd1, %rd313;
  1755. add.s64 %rd311, %rd312, 8;
  1756. mov.f32 %f1287, 0f4CBEBC20;
  1757. mul.f32 %f741, %f177, %f177;
  1758. fma.rn.f32 %f742, %f176, %f176, %f741;
  1759. fma.rn.f32 %f743, %f178, %f178, %f742;
  1760. mul.f32 %f744, %f179, %f179;
  1761. sub.f32 %f180, %f743, %f744;
  1762. ld.global.f32 %f181, [%rd311+16];
  1763. setp.gt.f32 %p109, %f180, %f181;
  1764. mov.f32 %f1288, %f1287;
  1765. @%p109 bra BB1_103;
  1766.  
  1767. sub.f32 %f745, %f181, %f180;
  1768. sqrt.rn.f32 %f746, %f745;
  1769. sub.f32 %f1288, %f179, %f746;
  1770. add.f32 %f1287, %f179, %f746;
  1771. mov.pred %p250, 0;
  1772.  
  1773. BB1_103:
  1774. @%p250 bra BB1_106;
  1775.  
  1776. setp.lt.f32 %p111, %f1288, 0f00000000;
  1777. selp.f32 %f186, %f1287, %f1288, %p111;
  1778. setp.geu.f32 %p112, %f186, %f1308;
  1779. @%p112 bra BB1_106;
  1780.  
  1781. cvt.u64.u32 %rd314, %r138;
  1782. mul.lo.s64 %rd207, %rd314, 80;
  1783. add.s64 %rd365, %rd1, %rd207;
  1784. add.s64 %rd350, %rd158, %rd207;
  1785. mov.f32 %f1308, %f186;
  1786.  
  1787. BB1_106:
  1788. cvt.u64.u32 %rd315, %r138;
  1789. cvt.u32.u64 %r69, %rd315;
  1790. add.s32 %r138, %r69, 1;
  1791.  
  1792. BB1_107:
  1793. mov.f32 %f1291, 0f4CBEBC20;
  1794. mul.wide.u32 %rd208, %r138, 80;
  1795. add.s64 %rd209, %rd1, %rd208;
  1796. ld.global.f32 %f749, [%rd209+8];
  1797. sub.f32 %f189, %f749, %f159;
  1798. ld.global.f32 %f750, [%rd209+12];
  1799. sub.f32 %f190, %f750, %f161;
  1800. ld.global.f32 %f751, [%rd209+16];
  1801. sub.f32 %f191, %f751, %f163;
  1802. mul.f32 %f752, %f1282, %f190;
  1803. fma.rn.f32 %f753, %f1281, %f189, %f752;
  1804. fma.rn.f32 %f192, %f1283, %f191, %f753;
  1805. setp.lt.f32 %p114, %f192, 0f00000000;
  1806. mov.pred %p251, -1;
  1807. @%p114 bra BB1_108;
  1808. bra.uni BB1_109;
  1809.  
  1810. BB1_108:
  1811. mov.f32 %f1292, %f1291;
  1812. bra.uni BB1_111;
  1813.  
  1814. BB1_109:
  1815. mul.wide.u32 %rd318, %r138, 80;
  1816. add.s64 %rd317, %rd1, %rd318;
  1817. add.s64 %rd316, %rd317, 8;
  1818. mov.f32 %f1291, 0f4CBEBC20;
  1819. mul.f32 %f756, %f190, %f190;
  1820. fma.rn.f32 %f757, %f189, %f189, %f756;
  1821. fma.rn.f32 %f758, %f191, %f191, %f757;
  1822. mul.f32 %f759, %f192, %f192;
  1823. sub.f32 %f193, %f758, %f759;
  1824. ld.global.f32 %f194, [%rd316+16];
  1825. setp.gt.f32 %p116, %f193, %f194;
  1826. mov.f32 %f1292, %f1291;
  1827. @%p116 bra BB1_111;
  1828.  
  1829. sub.f32 %f760, %f194, %f193;
  1830. sqrt.rn.f32 %f761, %f760;
  1831. sub.f32 %f1292, %f192, %f761;
  1832. add.f32 %f1291, %f192, %f761;
  1833. mov.pred %p251, 0;
  1834.  
  1835. BB1_111:
  1836. @%p251 bra BB1_114;
  1837.  
  1838. setp.lt.f32 %p118, %f1292, 0f00000000;
  1839. selp.f32 %f199, %f1291, %f1292, %p118;
  1840. setp.geu.f32 %p119, %f199, %f1308;
  1841. @%p119 bra BB1_114;
  1842.  
  1843. cvt.u64.u32 %rd319, %r138;
  1844. mul.lo.s64 %rd210, %rd319, 80;
  1845. add.s64 %rd365, %rd1, %rd210;
  1846. add.s64 %rd350, %rd158, %rd210;
  1847. mov.f32 %f1308, %f199;
  1848.  
  1849. BB1_114:
  1850. cvt.u64.u32 %rd320, %r138;
  1851. cvt.u32.u64 %r70, %rd320;
  1852. add.s32 %r138, %r70, 1;
  1853. mov.f32 %f1295, %f1308;
  1854.  
  1855. BB1_115:
  1856. setp.lt.u32 %p120, %r37, 4;
  1857. @%p120 bra BB1_116;
  1858. bra.uni BB1_117;
  1859.  
  1860. BB1_116:
  1861. mov.f32 %f1308, %f1295;
  1862. bra.uni BB1_146;
  1863.  
  1864. BB1_117:
  1865. cvt.u64.u32 %rd73, %r138;
  1866. mul.wide.u32 %rd211, %r138, 80;
  1867. add.s64 %rd212, %rd1, %rd211;
  1868. add.s64 %rd74, %rd212, 8;
  1869. ld.global.f32 %f764, [%rd212+8];
  1870. sub.f32 %f204, %f764, %f159;
  1871. ld.global.f32 %f765, [%rd212+12];
  1872. sub.f32 %f205, %f765, %f161;
  1873. ld.global.f32 %f766, [%rd212+16];
  1874. sub.f32 %f206, %f766, %f163;
  1875. mul.f32 %f767, %f1282, %f205;
  1876. fma.rn.f32 %f768, %f1281, %f204, %f767;
  1877. fma.rn.f32 %f207, %f1283, %f206, %f768;
  1878. setp.lt.f32 %p122, %f207, 0f00000000;
  1879. mov.pred %p252, -1;
  1880. mov.f32 %f1297, 0f4CBEBC20;
  1881. @%p122 bra BB1_118;
  1882. bra.uni BB1_119;
  1883.  
  1884. BB1_118:
  1885. mov.f32 %f1298, %f1297;
  1886. bra.uni BB1_121;
  1887.  
  1888. BB1_119:
  1889. mul.f32 %f771, %f205, %f205;
  1890. fma.rn.f32 %f772, %f204, %f204, %f771;
  1891. fma.rn.f32 %f773, %f206, %f206, %f772;
  1892. mul.f32 %f774, %f207, %f207;
  1893. sub.f32 %f208, %f773, %f774;
  1894. ld.global.f32 %f209, [%rd74+16];
  1895. setp.gt.f32 %p124, %f208, %f209;
  1896. mov.f32 %f1298, %f1297;
  1897. @%p124 bra BB1_121;
  1898.  
  1899. sub.f32 %f775, %f209, %f208;
  1900. sqrt.rn.f32 %f776, %f775;
  1901. sub.f32 %f1298, %f207, %f776;
  1902. add.f32 %f1297, %f207, %f776;
  1903. mov.pred %p252, 0;
  1904.  
  1905. BB1_121:
  1906. @%p252 bra BB1_124;
  1907.  
  1908. setp.lt.f32 %p126, %f1298, 0f00000000;
  1909. selp.f32 %f214, %f1297, %f1298, %p126;
  1910. setp.geu.f32 %p127, %f214, %f1308;
  1911. @%p127 bra BB1_124;
  1912.  
  1913. mul.lo.s64 %rd213, %rd73, 80;
  1914. add.s64 %rd365, %rd1, %rd213;
  1915. add.s64 %rd350, %rd158, %rd213;
  1916. mov.f32 %f1308, %f214;
  1917.  
  1918. BB1_124:
  1919. add.s32 %r71, %r138, 1;
  1920. cvt.u64.u32 %rd79, %r71;
  1921. mul.wide.u32 %rd214, %r71, 80;
  1922. add.s64 %rd215, %rd1, %rd214;
  1923. add.s64 %rd80, %rd215, 8;
  1924. ld.global.f32 %f779, [%rd215+8];
  1925. sub.f32 %f216, %f779, %f159;
  1926. ld.global.f32 %f780, [%rd215+12];
  1927. sub.f32 %f217, %f780, %f161;
  1928. ld.global.f32 %f781, [%rd215+16];
  1929. sub.f32 %f218, %f781, %f163;
  1930. mul.f32 %f782, %f1282, %f217;
  1931. fma.rn.f32 %f783, %f1281, %f216, %f782;
  1932. fma.rn.f32 %f219, %f1283, %f218, %f783;
  1933. setp.lt.f32 %p129, %f219, 0f00000000;
  1934. mov.pred %p253, -1;
  1935. mov.f32 %f1300, 0f4CBEBC20;
  1936. @%p129 bra BB1_125;
  1937. bra.uni BB1_126;
  1938.  
  1939. BB1_125:
  1940. mov.f32 %f1301, %f1300;
  1941. bra.uni BB1_128;
  1942.  
  1943. BB1_126:
  1944. mul.f32 %f786, %f217, %f217;
  1945. fma.rn.f32 %f787, %f216, %f216, %f786;
  1946. fma.rn.f32 %f788, %f218, %f218, %f787;
  1947. mul.f32 %f789, %f219, %f219;
  1948. sub.f32 %f220, %f788, %f789;
  1949. ld.global.f32 %f221, [%rd80+16];
  1950. setp.gt.f32 %p131, %f220, %f221;
  1951. mov.f32 %f1301, %f1300;
  1952. @%p131 bra BB1_128;
  1953.  
  1954. sub.f32 %f790, %f221, %f220;
  1955. sqrt.rn.f32 %f791, %f790;
  1956. sub.f32 %f1301, %f219, %f791;
  1957. add.f32 %f1300, %f219, %f791;
  1958. mov.pred %p253, 0;
  1959.  
  1960. BB1_128:
  1961. @%p253 bra BB1_131;
  1962.  
  1963. setp.lt.f32 %p133, %f1301, 0f00000000;
  1964. selp.f32 %f226, %f1300, %f1301, %p133;
  1965. setp.geu.f32 %p134, %f226, %f1308;
  1966. @%p134 bra BB1_131;
  1967.  
  1968. mul.lo.s64 %rd216, %rd79, 80;
  1969. add.s64 %rd365, %rd1, %rd216;
  1970. add.s64 %rd350, %rd158, %rd216;
  1971. mov.f32 %f1308, %f226;
  1972.  
  1973. BB1_131:
  1974. add.s32 %r72, %r138, 2;
  1975. cvt.u64.u32 %rd85, %r72;
  1976. mul.wide.u32 %rd217, %r72, 80;
  1977. add.s64 %rd218, %rd1, %rd217;
  1978. add.s64 %rd86, %rd218, 8;
  1979. ld.global.f32 %f794, [%rd218+8];
  1980. sub.f32 %f228, %f794, %f159;
  1981. ld.global.f32 %f795, [%rd218+12];
  1982. sub.f32 %f229, %f795, %f161;
  1983. ld.global.f32 %f796, [%rd218+16];
  1984. sub.f32 %f230, %f796, %f163;
  1985. mul.f32 %f797, %f1282, %f229;
  1986. fma.rn.f32 %f798, %f1281, %f228, %f797;
  1987. fma.rn.f32 %f231, %f1283, %f230, %f798;
  1988. setp.lt.f32 %p136, %f231, 0f00000000;
  1989. mov.pred %p254, -1;
  1990. mov.f32 %f1303, 0f4CBEBC20;
  1991. @%p136 bra BB1_132;
  1992. bra.uni BB1_133;
  1993.  
  1994. BB1_132:
  1995. mov.f32 %f1304, %f1303;
  1996. bra.uni BB1_135;
  1997.  
  1998. BB1_133:
  1999. mul.f32 %f801, %f229, %f229;
  2000. fma.rn.f32 %f802, %f228, %f228, %f801;
  2001. fma.rn.f32 %f803, %f230, %f230, %f802;
  2002. mul.f32 %f804, %f231, %f231;
  2003. sub.f32 %f232, %f803, %f804;
  2004. ld.global.f32 %f233, [%rd86+16];
  2005. setp.gt.f32 %p138, %f232, %f233;
  2006. mov.f32 %f1304, %f1303;
  2007. @%p138 bra BB1_135;
  2008.  
  2009. sub.f32 %f805, %f233, %f232;
  2010. sqrt.rn.f32 %f806, %f805;
  2011. sub.f32 %f1304, %f231, %f806;
  2012. add.f32 %f1303, %f231, %f806;
  2013. mov.pred %p254, 0;
  2014.  
  2015. BB1_135:
  2016. @%p254 bra BB1_138;
  2017.  
  2018. setp.lt.f32 %p140, %f1304, 0f00000000;
  2019. selp.f32 %f238, %f1303, %f1304, %p140;
  2020. setp.geu.f32 %p141, %f238, %f1308;
  2021. @%p141 bra BB1_138;
  2022.  
  2023. mul.lo.s64 %rd219, %rd85, 80;
  2024. add.s64 %rd365, %rd1, %rd219;
  2025. add.s64 %rd350, %rd158, %rd219;
  2026. mov.f32 %f1308, %f238;
  2027.  
  2028. BB1_138:
  2029. add.s32 %r73, %r138, 3;
  2030. cvt.u64.u32 %rd91, %r73;
  2031. mul.wide.u32 %rd220, %r73, 80;
  2032. add.s64 %rd221, %rd1, %rd220;
  2033. add.s64 %rd92, %rd221, 8;
  2034. ld.global.f32 %f809, [%rd221+8];
  2035. sub.f32 %f240, %f809, %f159;
  2036. ld.global.f32 %f810, [%rd221+12];
  2037. sub.f32 %f241, %f810, %f161;
  2038. ld.global.f32 %f811, [%rd221+16];
  2039. sub.f32 %f242, %f811, %f163;
  2040. mul.f32 %f812, %f1282, %f241;
  2041. fma.rn.f32 %f813, %f1281, %f240, %f812;
  2042. fma.rn.f32 %f243, %f1283, %f242, %f813;
  2043. setp.lt.f32 %p143, %f243, 0f00000000;
  2044. mov.pred %p255, -1;
  2045. mov.f32 %f1306, 0f4CBEBC20;
  2046. @%p143 bra BB1_139;
  2047. bra.uni BB1_140;
  2048.  
  2049. BB1_139:
  2050. mov.f32 %f1307, %f1306;
  2051. bra.uni BB1_142;
  2052.  
  2053. BB1_140:
  2054. mul.f32 %f816, %f241, %f241;
  2055. fma.rn.f32 %f817, %f240, %f240, %f816;
  2056. fma.rn.f32 %f818, %f242, %f242, %f817;
  2057. mul.f32 %f819, %f243, %f243;
  2058. sub.f32 %f244, %f818, %f819;
  2059. ld.global.f32 %f245, [%rd92+16];
  2060. setp.gt.f32 %p145, %f244, %f245;
  2061. mov.f32 %f1307, %f1306;
  2062. @%p145 bra BB1_142;
  2063.  
  2064. sub.f32 %f820, %f245, %f244;
  2065. sqrt.rn.f32 %f821, %f820;
  2066. sub.f32 %f1307, %f243, %f821;
  2067. add.f32 %f1306, %f243, %f821;
  2068. mov.pred %p255, 0;
  2069.  
  2070. BB1_142:
  2071. @%p255 bra BB1_145;
  2072.  
  2073. setp.lt.f32 %p147, %f1307, 0f00000000;
  2074. selp.f32 %f250, %f1306, %f1307, %p147;
  2075. setp.geu.f32 %p148, %f250, %f1308;
  2076. @%p148 bra BB1_145;
  2077.  
  2078. mul.lo.s64 %rd222, %rd91, 80;
  2079. add.s64 %rd365, %rd1, %rd222;
  2080. add.s64 %rd350, %rd158, %rd222;
  2081. mov.f32 %f1308, %f250;
  2082.  
  2083. BB1_145:
  2084. cvt.u32.u64 %r74, %rd73;
  2085. add.s32 %r138, %r74, 4;
  2086. setp.lt.u32 %p149, %r138, %r37;
  2087. @%p149 bra BB1_117;
  2088.  
  2089. BB1_146:
  2090. setp.eq.s64 %p150, %rd350, 0;
  2091. mov.f32 %f1333, 0f40000000;
  2092. mov.f32 %f1334, %f1333;
  2093. mov.f32 %f1335, %f1333;
  2094. @%p150 bra BB1_173;
  2095.  
  2096. fma.rn.f32 %f253, %f1281, %f1308, %f159;
  2097. fma.rn.f32 %f254, %f1282, %f1308, %f161;
  2098. fma.rn.f32 %f255, %f1283, %f1308, %f163;
  2099. add.s64 %rd99, %rd365, 8;
  2100. ld.global.f32 %f825, [%rd365+8];
  2101. sub.f32 %f1310, %f253, %f825;
  2102. ld.global.f32 %f826, [%rd365+12];
  2103. sub.f32 %f1311, %f254, %f826;
  2104. ld.global.f32 %f827, [%rd365+16];
  2105. sub.f32 %f1312, %f255, %f827;
  2106. mul.f32 %f828, %f1311, %f1311;
  2107. fma.rn.f32 %f829, %f1310, %f1310, %f828;
  2108. fma.rn.f32 %f259, %f1312, %f1312, %f829;
  2109. setp.leu.f32 %p151, %f259, 0f00000000;
  2110. @%p151 bra BB1_149;
  2111.  
  2112. sqrt.rn.f32 %f830, %f259;
  2113. rcp.rn.f32 %f831, %f830;
  2114. mul.f32 %f1310, %f1310, %f831;
  2115. mul.f32 %f1311, %f1311, %f831;
  2116. mul.f32 %f1312, %f1312, %f831;
  2117.  
  2118. BB1_149:
  2119. mul.f32 %f832, %f1282, %f1311;
  2120. fma.rn.f32 %f833, %f1281, %f1310, %f832;
  2121. fma.rn.f32 %f834, %f1283, %f1312, %f833;
  2122. mov.u16 %rs11, 0;
  2123. setp.leu.f32 %p152, %f834, 0f00000000;
  2124. @%p152 bra BB1_151;
  2125.  
  2126. neg.f32 %f1310, %f1310;
  2127. neg.f32 %f1311, %f1311;
  2128. neg.f32 %f1312, %f1312;
  2129. mov.u16 %rs11, 1;
  2130.  
  2131. BB1_151:
  2132. ld.global.f32 %f835, [%rd99+44];
  2133. setp.gt.f32 %p153, %f835, 0f00000000;
  2134. @%p153 bra BB1_165;
  2135.  
  2136. ld.global.f32 %f836, [%rd99+48];
  2137. setp.gt.f32 %p154, %f836, 0f00000000;
  2138. @%p154 bra BB1_165;
  2139. bra.uni BB1_153;
  2140.  
  2141. BB1_165:
  2142. mov.u32 %r124, 3;
  2143. mul.f32 %f879, %f1282, %f1311;
  2144. fma.rn.f32 %f880, %f1281, %f1310, %f879;
  2145. fma.rn.f32 %f300, %f1283, %f1312, %f880;
  2146. add.f32 %f881, %f300, 0f3F800000;
  2147. // Callseq Start 4
  2148. {
  2149. .reg .b32 temp_param_reg;
  2150. // <end>}
  2151. .param .b32 param0;
  2152. st.param.f32 [param0+0], %f881;
  2153. .param .b32 param1;
  2154. st.param.b32 [param1+0], %r124;
  2155. .param .b32 retval0;
  2156. call.uni (retval0),
  2157. _Z5powiffi,
  2158. (
  2159. param0,
  2160. param1
  2161. );
  2162. ld.param.f32 %f882, [retval0+0];
  2163.  
  2164. //{
  2165. }// Callseq End 4
  2166. fma.rn.f32 %f301, %f882, 0f3F666666, 0f3DCCCCCD;
  2167. add.f32 %f883, %f1310, %f1310;
  2168. mul.f32 %f884, %f883, %f300;
  2169. add.f32 %f885, %f1311, %f1311;
  2170. mul.f32 %f886, %f885, %f300;
  2171. add.f32 %f887, %f1312, %f1312;
  2172. mul.f32 %f888, %f887, %f300;
  2173. sub.f32 %f302, %f1281, %f884;
  2174. sub.f32 %f303, %f1282, %f886;
  2175. sub.f32 %f304, %f1283, %f888;
  2176. add.u64 %rd227, %SP, 36;
  2177. add.u64 %rd103, %SPL, 36;
  2178. st.local.f32 [%rd103], %f302;
  2179. st.local.f32 [%rd103+4], %f303;
  2180. st.local.f32 [%rd103+8], %f304;
  2181. mul.f32 %f889, %f303, %f303;
  2182. fma.rn.f32 %f890, %f302, %f302, %f889;
  2183. fma.rn.f32 %f305, %f304, %f304, %f890;
  2184. setp.leu.f32 %p163, %f305, 0f00000000;
  2185. @%p163 bra BB1_167;
  2186.  
  2187. sqrt.rn.f32 %f891, %f305;
  2188. rcp.rn.f32 %f892, %f891;
  2189. mul.f32 %f893, %f302, %f892;
  2190. st.local.f32 [%rd103], %f893;
  2191. mul.f32 %f894, %f303, %f892;
  2192. st.local.f32 [%rd103+4], %f894;
  2193. mul.f32 %f895, %f304, %f892;
  2194. st.local.f32 [%rd103+8], %f895;
  2195.  
  2196. BB1_167:
  2197. mul.f32 %f306, %f1310, 0f38D1B717;
  2198. add.f32 %f899, %f253, %f306;
  2199. mul.f32 %f307, %f1311, 0f38D1B717;
  2200. add.f32 %f900, %f254, %f307;
  2201. mul.f32 %f308, %f1312, 0f38D1B717;
  2202. add.f32 %f901, %f255, %f308;
  2203. add.u64 %rd228, %SP, 52;
  2204. add.u64 %rd229, %SPL, 52;
  2205. mov.u32 %r79, 2;
  2206. st.local.u32 [%rd229], %r79;
  2207. // Callseq Start 5
  2208. {
  2209. .reg .b32 temp_param_reg;
  2210. // <end>}
  2211. .param .b32 param0;
  2212. st.param.f32 [param0+0], %f899;
  2213. .param .b32 param1;
  2214. st.param.f32 [param1+0], %f900;
  2215. .param .b32 param2;
  2216. st.param.f32 [param2+0], %f901;
  2217. .param .b64 param3;
  2218. st.param.b64 [param3+0], %rd227;
  2219. .param .b64 param4;
  2220. st.param.b64 [param4+0], %rd158;
  2221. .param .b32 param5;
  2222. st.param.b32 [param5+0], %r37;
  2223. .param .b64 param6;
  2224. st.param.b64 [param6+0], %rd228;
  2225. .param .align 4 .b8 retval0[12];
  2226. call.uni (retval0),
  2227. _Z5traceRK4Vec3IfES2_PK6SpherejRKi,
  2228. (
  2229. param0,
  2230. param1,
  2231. param2,
  2232. param3,
  2233. param4,
  2234. param5,
  2235. param6
  2236. );
  2237. ld.param.f32 %f309, [retval0+0];
  2238. ld.param.f32 %f310, [retval0+4];
  2239. ld.param.f32 %f311, [retval0+8];
  2240.  
  2241. //{
  2242. }// Callseq End 5
  2243. ld.global.f32 %f1326, [%rd99+44];
  2244. mov.f32 %f1327, 0f00000000;
  2245. setp.eq.f32 %p164, %f1326, 0f00000000;
  2246. mov.f32 %f1328, %f1327;
  2247. mov.f32 %f1329, %f1327;
  2248. @%p164 bra BB1_171;
  2249.  
  2250. setp.eq.s16 %p165, %rs11, 0;
  2251. selp.f32 %f902, 0f3F68BA2E, 0f3F8CCCCD, %p165;
  2252. mul.f32 %f903, %f902, %f902;
  2253. mul.f32 %f904, %f300, %f300;
  2254. mov.f32 %f905, 0f3F800000;
  2255. sub.f32 %f906, %f905, %f904;
  2256. mul.f32 %f907, %f903, %f906;
  2257. sub.f32 %f908, %f905, %f907;
  2258. mul.f32 %f909, %f300, %f902;
  2259. neg.f32 %f910, %f909;
  2260. sqrt.rn.f32 %f911, %f908;
  2261. sub.f32 %f912, %f910, %f911;
  2262. mul.f32 %f913, %f1310, %f912;
  2263. mul.f32 %f914, %f1311, %f912;
  2264. mul.f32 %f915, %f1312, %f912;
  2265. fma.rn.f32 %f313, %f1281, %f902, %f913;
  2266. fma.rn.f32 %f314, %f1282, %f902, %f914;
  2267. fma.rn.f32 %f315, %f1283, %f902, %f915;
  2268. add.u64 %rd231, %SP, 24;
  2269. add.u64 %rd104, %SPL, 24;
  2270. st.local.f32 [%rd104], %f313;
  2271. st.local.f32 [%rd104+4], %f314;
  2272. st.local.f32 [%rd104+8], %f315;
  2273. mul.f32 %f916, %f314, %f314;
  2274. fma.rn.f32 %f917, %f313, %f313, %f916;
  2275. fma.rn.f32 %f316, %f315, %f315, %f917;
  2276. setp.leu.f32 %p166, %f316, 0f00000000;
  2277. @%p166 bra BB1_170;
  2278.  
  2279. sqrt.rn.f32 %f918, %f316;
  2280. rcp.rn.f32 %f919, %f918;
  2281. mul.f32 %f920, %f313, %f919;
  2282. st.local.f32 [%rd104], %f920;
  2283. mul.f32 %f921, %f314, %f919;
  2284. st.local.f32 [%rd104+4], %f921;
  2285. mul.f32 %f922, %f315, %f919;
  2286. st.local.f32 [%rd104+8], %f922;
  2287.  
  2288. BB1_170:
  2289. sub.f32 %f923, %f253, %f306;
  2290. add.u64 %rd232, %SP, 48;
  2291. add.u64 %rd233, %SPL, 48;
  2292. st.local.u32 [%rd233], %r79;
  2293. sub.f32 %f924, %f255, %f308;
  2294. sub.f32 %f925, %f254, %f307;
  2295. // Callseq Start 6
  2296. {
  2297. .reg .b32 temp_param_reg;
  2298. // <end>}
  2299. .param .b32 param0;
  2300. st.param.f32 [param0+0], %f923;
  2301. .param .b32 param1;
  2302. st.param.f32 [param1+0], %f925;
  2303. .param .b32 param2;
  2304. st.param.f32 [param2+0], %f924;
  2305. .param .b64 param3;
  2306. st.param.b64 [param3+0], %rd231;
  2307. .param .b64 param4;
  2308. st.param.b64 [param4+0], %rd158;
  2309. .param .b32 param5;
  2310. st.param.b32 [param5+0], %r37;
  2311. .param .b64 param6;
  2312. st.param.b64 [param6+0], %rd232;
  2313. .param .align 4 .b8 retval0[12];
  2314. call.uni (retval0),
  2315. _Z5traceRK4Vec3IfES2_PK6SpherejRKi,
  2316. (
  2317. param0,
  2318. param1,
  2319. param2,
  2320. param3,
  2321. param4,
  2322. param5,
  2323. param6
  2324. );
  2325. ld.param.f32 %f1327, [retval0+0];
  2326. ld.param.f32 %f1328, [retval0+4];
  2327. ld.param.f32 %f1329, [retval0+8];
  2328.  
  2329. //{
  2330. }// Callseq End 6
  2331. ld.global.f32 %f1326, [%rd99+44];
  2332.  
  2333. BB1_171:
  2334. mov.f32 %f926, 0f3F800000;
  2335. sub.f32 %f927, %f926, %f301;
  2336. mul.f32 %f928, %f927, %f1327;
  2337. mul.f32 %f929, %f927, %f1328;
  2338. mul.f32 %f930, %f927, %f1329;
  2339. mul.f32 %f931, %f1326, %f928;
  2340. mul.f32 %f932, %f1326, %f929;
  2341. mul.f32 %f933, %f1326, %f930;
  2342. fma.rn.f32 %f934, %f301, %f309, %f931;
  2343. fma.rn.f32 %f935, %f301, %f310, %f932;
  2344. fma.rn.f32 %f936, %f301, %f311, %f933;
  2345. ld.global.f32 %f937, [%rd99+20];
  2346. mul.f32 %f1330, %f937, %f934;
  2347. ld.global.f32 %f938, [%rd99+24];
  2348. mul.f32 %f1331, %f935, %f938;
  2349. ld.global.f32 %f939, [%rd99+28];
  2350. mul.f32 %f1332, %f936, %f939;
  2351. bra.uni BB1_172;
  2352.  
  2353. BB1_153:
  2354. mov.f32 %f1330, 0f00000000;
  2355. mov.f32 %f1331, %f1330;
  2356. mov.f32 %f1332, %f1330;
  2357. @%p26 bra BB1_172;
  2358.  
  2359. fma.rn.f32 %f272, %f1310, 0f38D1B717, %f253;
  2360. fma.rn.f32 %f273, %f1311, 0f38D1B717, %f254;
  2361. fma.rn.f32 %f274, %f1312, 0f38D1B717, %f255;
  2362. mov.f32 %f1332, 0f00000000;
  2363. mov.u32 %r140, 0;
  2364. mov.f32 %f1331, %f1332;
  2365. mov.f32 %f1330, %f1332;
  2366.  
  2367. BB1_155:
  2368. cvt.u64.u32 %rd100, %r140;
  2369. mul.wide.u32 %rd223, %r140, 80;
  2370. add.s64 %rd224, %rd1, %rd223;
  2371. add.s64 %rd101, %rd224, 40;
  2372. ld.global.f32 %f278, [%rd224+40];
  2373. setp.leu.f32 %p156, %f278, 0f00000000;
  2374. @%p156 bra BB1_164;
  2375.  
  2376. ld.global.f32 %f843, [%rd101+-32];
  2377. sub.f32 %f1321, %f843, %f253;
  2378. ld.global.f32 %f844, [%rd101+-28];
  2379. sub.f32 %f1320, %f844, %f254;
  2380. ld.global.f32 %f845, [%rd101+-24];
  2381. sub.f32 %f1319, %f845, %f255;
  2382. mul.f32 %f846, %f1320, %f1320;
  2383. fma.rn.f32 %f847, %f1321, %f1321, %f846;
  2384. fma.rn.f32 %f282, %f1319, %f1319, %f847;
  2385. setp.leu.f32 %p157, %f282, 0f00000000;
  2386. @%p157 bra BB1_158;
  2387.  
  2388. sqrt.rn.f32 %f848, %f282;
  2389. rcp.rn.f32 %f849, %f848;
  2390. mul.f32 %f1321, %f1321, %f849;
  2391. mul.f32 %f1320, %f1320, %f849;
  2392. mul.f32 %f1319, %f1319, %f849;
  2393.  
  2394. BB1_158:
  2395. mov.u32 %r141, 0;
  2396.  
  2397. BB1_159:
  2398. cvt.u32.u64 %r77, %rd100;
  2399. setp.eq.s32 %p158, %r77, %r141;
  2400. @%p158 bra BB1_162;
  2401.  
  2402. mul.wide.u32 %rd225, %r141, 80;
  2403. add.s64 %rd226, %rd1, %rd225;
  2404. add.s64 %rd102, %rd226, 8;
  2405. ld.global.f32 %f850, [%rd226+8];
  2406. sub.f32 %f289, %f850, %f272;
  2407. ld.global.f32 %f851, [%rd226+12];
  2408. sub.f32 %f290, %f851, %f273;
  2409. ld.global.f32 %f852, [%rd226+16];
  2410. sub.f32 %f291, %f852, %f274;
  2411. mul.f32 %f853, %f1320, %f290;
  2412. fma.rn.f32 %f854, %f1321, %f289, %f853;
  2413. fma.rn.f32 %f292, %f1319, %f291, %f854;
  2414. setp.lt.f32 %p159, %f292, 0f00000000;
  2415. @%p159 bra BB1_162;
  2416.  
  2417. mul.f32 %f856, %f290, %f290;
  2418. fma.rn.f32 %f857, %f289, %f289, %f856;
  2419. fma.rn.f32 %f858, %f291, %f291, %f857;
  2420. mul.f32 %f859, %f292, %f292;
  2421. sub.f32 %f860, %f858, %f859;
  2422. ld.global.f32 %f861, [%rd102+16];
  2423. mov.f32 %f1322, 0f00000000;
  2424. setp.leu.f32 %p160, %f860, %f861;
  2425. @%p160 bra BB1_163;
  2426.  
  2427. BB1_162:
  2428. add.s32 %r141, %r141, 1;
  2429. setp.lt.u32 %p161, %r141, %r37;
  2430. mov.f32 %f1322, 0f3F800000;
  2431. @%p161 bra BB1_159;
  2432.  
  2433. BB1_163:
  2434. ld.global.f32 %f863, [%rd99+20];
  2435. mul.f32 %f864, %f1322, %f863;
  2436. ld.global.f32 %f865, [%rd99+24];
  2437. mul.f32 %f866, %f1322, %f865;
  2438. ld.global.f32 %f867, [%rd99+28];
  2439. mul.f32 %f868, %f1322, %f867;
  2440. mul.f32 %f869, %f1310, %f1321;
  2441. fma.rn.f32 %f870, %f1311, %f1320, %f869;
  2442. fma.rn.f32 %f871, %f1312, %f1319, %f870;
  2443. mov.f32 %f872, 0f00000000;
  2444. max.f32 %f873, %f872, %f871;
  2445. mul.f32 %f874, %f864, %f873;
  2446. mul.f32 %f875, %f866, %f873;
  2447. mul.f32 %f876, %f868, %f873;
  2448. ld.global.f32 %f877, [%rd101+4];
  2449. ld.global.f32 %f878, [%rd101+8];
  2450. fma.rn.f32 %f1330, %f278, %f874, %f1330;
  2451. fma.rn.f32 %f1331, %f875, %f877, %f1331;
  2452. fma.rn.f32 %f1332, %f876, %f878, %f1332;
  2453.  
  2454. BB1_164:
  2455. add.s32 %r140, %r140, 1;
  2456. setp.lt.u32 %p162, %r140, %r37;
  2457. @%p162 bra BB1_155;
  2458.  
  2459. BB1_172:
  2460. ld.global.f32 %f940, [%rd99+32];
  2461. add.f32 %f1333, %f1330, %f940;
  2462. ld.global.f32 %f941, [%rd99+36];
  2463. add.f32 %f1334, %f1331, %f941;
  2464. ld.global.f32 %f942, [%rd99+40];
  2465. add.f32 %f1335, %f1332, %f942;
  2466.  
  2467. BB1_173:
  2468. ld.global.f32 %f1389, [%rd49+44];
  2469. mov.f32 %f1390, 0f00000000;
  2470. setp.eq.f32 %p167, %f1389, 0f00000000;
  2471. @%p167 bra BB1_174;
  2472.  
  2473. setp.eq.s16 %p168, %rs1, 0;
  2474. selp.f32 %f946, 0f3F68BA2E, 0f3F8CCCCD, %p168;
  2475. mul.f32 %f947, %f946, %f946;
  2476. mul.f32 %f948, %f146, %f146;
  2477. mov.f32 %f949, 0f3F800000;
  2478. sub.f32 %f950, %f949, %f948;
  2479. mul.f32 %f951, %f947, %f950;
  2480. sub.f32 %f952, %f949, %f951;
  2481. mul.f32 %f953, %f146, %f946;
  2482. neg.f32 %f954, %f953;
  2483. sqrt.rn.f32 %f955, %f952;
  2484. sub.f32 %f956, %f954, %f955;
  2485. mul.f32 %f957, %f1268, %f956;
  2486. mul.f32 %f958, %f1269, %f956;
  2487. mul.f32 %f959, %f1270, %f956;
  2488. fma.rn.f32 %f347, %f9, %f946, %f957;
  2489. fma.rn.f32 %f346, %f8, %f946, %f958;
  2490. fma.rn.f32 %f345, %f7, %f946, %f959;
  2491. mul.f32 %f960, %f346, %f346;
  2492. fma.rn.f32 %f961, %f347, %f347, %f960;
  2493. fma.rn.f32 %f341, %f345, %f345, %f961;
  2494. setp.leu.f32 %p169, %f341, 0f00000000;
  2495. @%p169 bra BB1_177;
  2496.  
  2497. sqrt.rn.f32 %f962, %f341;
  2498. rcp.rn.f32 %f963, %f962;
  2499. mul.f32 %f347, %f347, %f963;
  2500. mul.f32 %f346, %f346, %f963;
  2501. mul.f32 %f345, %f345, %f963;
  2502.  
  2503. BB1_177:
  2504. fma.rn.f32 %f1229, %f7, %f1263, 0f00000000;
  2505. fma.rn.f32 %f1228, %f8, %f1263, 0f00000000;
  2506. fma.rn.f32 %f1227, %f9, %f1263, 0f00000000;
  2507. mul.f32 %f1226, %f1270, 0f38D1B717;
  2508. mul.f32 %f1225, %f1269, 0f38D1B717;
  2509. mul.f32 %f1224, %f1268, 0f38D1B717;
  2510. sub.f32 %f348, %f1227, %f1224;
  2511. sub.f32 %f349, %f1228, %f1225;
  2512. sub.f32 %f350, %f1229, %f1226;
  2513. mov.f32 %f1344, 0f4CBEBC20;
  2514. mov.u64 %rd387, 0;
  2515. mov.u64 %rd372, %rd387;
  2516. @%p26 bra BB1_239;
  2517.  
  2518. and.b32 %r84, %r37, 3;
  2519. mov.f32 %f1346, 0f4CBEBC20;
  2520. mov.u64 %rd387, 0;
  2521. mov.u32 %r144, 0;
  2522. mov.f32 %f1350, 0f00000000;
  2523. setp.eq.s32 %p171, %r84, 0;
  2524. @%p171 bra BB1_179;
  2525.  
  2526. setp.eq.s32 %p172, %r84, 1;
  2527. @%p172 bra BB1_181;
  2528. bra.uni BB1_182;
  2529.  
  2530. BB1_181:
  2531. mov.u64 %rd372, %rd387;
  2532. mov.f32 %f1344, %f1346;
  2533. bra.uni BB1_200;
  2534.  
  2535. BB1_179:
  2536. mov.u64 %rd372, %rd387;
  2537. mov.f32 %f1344, %f1346;
  2538. bra.uni BB1_208;
  2539.  
  2540. BB1_182:
  2541. setp.eq.s32 %p173, %r84, 2;
  2542. @%p173 bra BB1_183;
  2543. bra.uni BB1_184;
  2544.  
  2545. BB1_183:
  2546. mov.u64 %rd372, %rd387;
  2547. mov.f32 %f1344, %f1346;
  2548. bra.uni BB1_192;
  2549.  
  2550. BB1_184:
  2551. ld.global.f32 %f971, [%rd1+8];
  2552. sub.f32 %f351, %f971, %f348;
  2553. ld.global.f32 %f972, [%rd1+12];
  2554. sub.f32 %f352, %f972, %f349;
  2555. ld.global.f32 %f973, [%rd1+16];
  2556. sub.f32 %f353, %f973, %f350;
  2557. mul.f32 %f974, %f346, %f352;
  2558. fma.rn.f32 %f975, %f347, %f351, %f974;
  2559. fma.rn.f32 %f354, %f345, %f353, %f975;
  2560. setp.lt.f32 %p175, %f354, 0f00000000;
  2561. mov.pred %p256, -1;
  2562. mov.f32 %f970, 0f4CBEBC20;
  2563. @%p175 bra BB1_185;
  2564. bra.uni BB1_186;
  2565.  
  2566. BB1_185:
  2567. mov.f32 %f1339, %f970;
  2568. mov.f32 %f1340, %f970;
  2569. bra.uni BB1_188;
  2570.  
  2571. BB1_186:
  2572. mul.f32 %f978, %f352, %f352;
  2573. fma.rn.f32 %f979, %f351, %f351, %f978;
  2574. fma.rn.f32 %f980, %f353, %f353, %f979;
  2575. mul.f32 %f981, %f354, %f354;
  2576. sub.f32 %f355, %f980, %f981;
  2577. ld.global.f32 %f356, [%rd1+24];
  2578. setp.gt.f32 %p177, %f355, %f356;
  2579. mov.f32 %f1339, %f970;
  2580. mov.f32 %f1340, %f970;
  2581. @%p177 bra BB1_188;
  2582.  
  2583. sub.f32 %f982, %f356, %f355;
  2584. sqrt.rn.f32 %f983, %f982;
  2585. sub.f32 %f1340, %f354, %f983;
  2586. add.f32 %f1339, %f354, %f983;
  2587. mov.pred %p256, 0;
  2588.  
  2589. BB1_188:
  2590. mov.u32 %r144, 1;
  2591. @%p256 bra BB1_189;
  2592.  
  2593. setp.lt.f32 %p179, %f1340, 0f00000000;
  2594. selp.f32 %f1344, %f1339, %f1340, %p179;
  2595. setp.geu.f32 %p180, %f1344, 0f4CBEBC20;
  2596. @%p180 bra BB1_189;
  2597.  
  2598. mov.u64 %rd387, %rd1;
  2599. mov.u64 %rd372, %rd158;
  2600. bra.uni BB1_192;
  2601.  
  2602. BB1_189:
  2603. mov.u64 %rd372, %rd387;
  2604. mov.f32 %f1344, %f970;
  2605.  
  2606. BB1_192:
  2607. cvt.u64.u32 %rd107, %r144;
  2608. mul.wide.u32 %rd247, %r144, 80;
  2609. add.s64 %rd248, %rd1, %rd247;
  2610. add.s64 %rd108, %rd248, 8;
  2611. ld.global.f32 %f988, [%rd248+8];
  2612. sub.f32 %f363, %f988, %f348;
  2613. ld.global.f32 %f989, [%rd248+12];
  2614. sub.f32 %f364, %f989, %f349;
  2615. ld.global.f32 %f990, [%rd248+16];
  2616. sub.f32 %f365, %f990, %f350;
  2617. mul.f32 %f991, %f346, %f364;
  2618. fma.rn.f32 %f992, %f347, %f363, %f991;
  2619. fma.rn.f32 %f366, %f345, %f365, %f992;
  2620. setp.lt.f32 %p182, %f366, 0f00000000;
  2621. mov.pred %p257, -1;
  2622. @%p182 bra BB1_193;
  2623. bra.uni BB1_194;
  2624.  
  2625. BB1_193:
  2626. mov.f32 %f1342, %f1346;
  2627. mov.f32 %f1343, %f1346;
  2628. bra.uni BB1_196;
  2629.  
  2630. BB1_194:
  2631. mul.f32 %f995, %f364, %f364;
  2632. fma.rn.f32 %f996, %f363, %f363, %f995;
  2633. fma.rn.f32 %f997, %f365, %f365, %f996;
  2634. mul.f32 %f998, %f366, %f366;
  2635. sub.f32 %f367, %f997, %f998;
  2636. ld.global.f32 %f368, [%rd108+16];
  2637. setp.gt.f32 %p184, %f367, %f368;
  2638. mov.f32 %f1342, %f1346;
  2639. mov.f32 %f1343, %f1346;
  2640. @%p184 bra BB1_196;
  2641.  
  2642. sub.f32 %f999, %f368, %f367;
  2643. sqrt.rn.f32 %f1000, %f999;
  2644. sub.f32 %f1343, %f366, %f1000;
  2645. add.f32 %f1342, %f366, %f1000;
  2646. mov.pred %p257, 0;
  2647.  
  2648. BB1_196:
  2649. @%p257 bra BB1_199;
  2650.  
  2651. setp.lt.f32 %p186, %f1343, 0f00000000;
  2652. selp.f32 %f373, %f1342, %f1343, %p186;
  2653. setp.geu.f32 %p187, %f373, %f1344;
  2654. @%p187 bra BB1_199;
  2655.  
  2656. mul.lo.s64 %rd249, %rd107, 80;
  2657. add.s64 %rd387, %rd1, %rd249;
  2658. add.s64 %rd372, %rd158, %rd249;
  2659. mov.f32 %f1344, %f373;
  2660.  
  2661. BB1_199:
  2662. cvt.u32.u64 %r88, %rd107;
  2663. add.s32 %r144, %r88, 1;
  2664.  
  2665. BB1_200:
  2666. cvt.u64.u32 %rd115, %r144;
  2667. mul.wide.u32 %rd250, %r144, 80;
  2668. add.s64 %rd251, %rd1, %rd250;
  2669. add.s64 %rd116, %rd251, 8;
  2670. ld.global.f32 %f1003, [%rd251+8];
  2671. sub.f32 %f376, %f1003, %f348;
  2672. ld.global.f32 %f1004, [%rd251+12];
  2673. sub.f32 %f377, %f1004, %f349;
  2674. ld.global.f32 %f1005, [%rd251+16];
  2675. sub.f32 %f378, %f1005, %f350;
  2676. mul.f32 %f1006, %f346, %f377;
  2677. fma.rn.f32 %f1007, %f347, %f376, %f1006;
  2678. fma.rn.f32 %f379, %f345, %f378, %f1007;
  2679. setp.lt.f32 %p189, %f379, 0f00000000;
  2680. mov.pred %p258, -1;
  2681. @%p189 bra BB1_201;
  2682. bra.uni BB1_202;
  2683.  
  2684. BB1_201:
  2685. mov.f32 %f1347, %f1346;
  2686. bra.uni BB1_204;
  2687.  
  2688. BB1_202:
  2689. mul.f32 %f1010, %f377, %f377;
  2690. fma.rn.f32 %f1011, %f376, %f376, %f1010;
  2691. fma.rn.f32 %f1012, %f378, %f378, %f1011;
  2692. mul.f32 %f1013, %f379, %f379;
  2693. sub.f32 %f380, %f1012, %f1013;
  2694. ld.global.f32 %f381, [%rd116+16];
  2695. setp.gt.f32 %p191, %f380, %f381;
  2696. mov.f32 %f1347, %f1346;
  2697. @%p191 bra BB1_204;
  2698.  
  2699. sub.f32 %f1014, %f381, %f380;
  2700. sqrt.rn.f32 %f1015, %f1014;
  2701. sub.f32 %f1347, %f379, %f1015;
  2702. add.f32 %f1346, %f379, %f1015;
  2703. mov.pred %p258, 0;
  2704.  
  2705. BB1_204:
  2706. @%p258 bra BB1_207;
  2707.  
  2708. setp.lt.f32 %p193, %f1347, 0f00000000;
  2709. selp.f32 %f386, %f1346, %f1347, %p193;
  2710. setp.geu.f32 %p194, %f386, %f1344;
  2711. @%p194 bra BB1_207;
  2712.  
  2713. mul.lo.s64 %rd252, %rd115, 80;
  2714. add.s64 %rd387, %rd1, %rd252;
  2715. add.s64 %rd372, %rd158, %rd252;
  2716. mov.f32 %f1344, %f386;
  2717.  
  2718. BB1_207:
  2719. cvt.u32.u64 %r89, %rd115;
  2720. add.s32 %r144, %r89, 1;
  2721. mov.f32 %f1350, %f1344;
  2722.  
  2723. BB1_208:
  2724. setp.lt.u32 %p195, %r37, 4;
  2725. @%p195 bra BB1_209;
  2726. bra.uni BB1_210;
  2727.  
  2728. BB1_209:
  2729. mov.f32 %f1344, %f1350;
  2730. bra.uni BB1_239;
  2731.  
  2732. BB1_210:
  2733. cvt.u64.u32 %rd125, %r144;
  2734. mul.wide.u32 %rd253, %r144, 80;
  2735. add.s64 %rd254, %rd1, %rd253;
  2736. add.s64 %rd126, %rd254, 8;
  2737. ld.global.f32 %f1018, [%rd254+8];
  2738. sub.f32 %f391, %f1018, %f348;
  2739. ld.global.f32 %f1019, [%rd254+12];
  2740. sub.f32 %f392, %f1019, %f349;
  2741. ld.global.f32 %f1020, [%rd254+16];
  2742. sub.f32 %f393, %f1020, %f350;
  2743. mul.f32 %f1021, %f346, %f392;
  2744. fma.rn.f32 %f1022, %f347, %f391, %f1021;
  2745. fma.rn.f32 %f394, %f345, %f393, %f1022;
  2746. setp.lt.f32 %p197, %f394, 0f00000000;
  2747. mov.pred %p259, -1;
  2748. mov.f32 %f1352, 0f4CBEBC20;
  2749. @%p197 bra BB1_211;
  2750. bra.uni BB1_212;
  2751.  
  2752. BB1_211:
  2753. mov.f32 %f1353, %f1352;
  2754. bra.uni BB1_214;
  2755.  
  2756. BB1_212:
  2757. mul.f32 %f1025, %f392, %f392;
  2758. fma.rn.f32 %f1026, %f391, %f391, %f1025;
  2759. fma.rn.f32 %f1027, %f393, %f393, %f1026;
  2760. mul.f32 %f1028, %f394, %f394;
  2761. sub.f32 %f395, %f1027, %f1028;
  2762. ld.global.f32 %f396, [%rd126+16];
  2763. setp.gt.f32 %p199, %f395, %f396;
  2764. mov.f32 %f1353, %f1352;
  2765. @%p199 bra BB1_214;
  2766.  
  2767. sub.f32 %f1029, %f396, %f395;
  2768. sqrt.rn.f32 %f1030, %f1029;
  2769. sub.f32 %f1353, %f394, %f1030;
  2770. add.f32 %f1352, %f394, %f1030;
  2771. mov.pred %p259, 0;
  2772.  
  2773. BB1_214:
  2774. @%p259 bra BB1_217;
  2775.  
  2776. setp.lt.f32 %p201, %f1353, 0f00000000;
  2777. selp.f32 %f401, %f1352, %f1353, %p201;
  2778. setp.geu.f32 %p202, %f401, %f1344;
  2779. @%p202 bra BB1_217;
  2780.  
  2781. mul.lo.s64 %rd255, %rd125, 80;
  2782. add.s64 %rd387, %rd1, %rd255;
  2783. add.s64 %rd372, %rd158, %rd255;
  2784. mov.f32 %f1344, %f401;
  2785.  
  2786. BB1_217:
  2787. add.s32 %r90, %r144, 1;
  2788. cvt.u64.u32 %rd131, %r90;
  2789. mul.wide.u32 %rd256, %r90, 80;
  2790. add.s64 %rd257, %rd1, %rd256;
  2791. add.s64 %rd132, %rd257, 8;
  2792. ld.global.f32 %f1033, [%rd257+8];
  2793. sub.f32 %f403, %f1033, %f348;
  2794. ld.global.f32 %f1034, [%rd257+12];
  2795. sub.f32 %f404, %f1034, %f349;
  2796. ld.global.f32 %f1035, [%rd257+16];
  2797. sub.f32 %f405, %f1035, %f350;
  2798. mul.f32 %f1036, %f346, %f404;
  2799. fma.rn.f32 %f1037, %f347, %f403, %f1036;
  2800. fma.rn.f32 %f406, %f345, %f405, %f1037;
  2801. setp.lt.f32 %p204, %f406, 0f00000000;
  2802. mov.pred %p260, -1;
  2803. mov.f32 %f1355, 0f4CBEBC20;
  2804. @%p204 bra BB1_218;
  2805. bra.uni BB1_219;
  2806.  
  2807. BB1_218:
  2808. mov.f32 %f1356, %f1355;
  2809. bra.uni BB1_221;
  2810.  
  2811. BB1_219:
  2812. mul.f32 %f1040, %f404, %f404;
  2813. fma.rn.f32 %f1041, %f403, %f403, %f1040;
  2814. fma.rn.f32 %f1042, %f405, %f405, %f1041;
  2815. mul.f32 %f1043, %f406, %f406;
  2816. sub.f32 %f407, %f1042, %f1043;
  2817. ld.global.f32 %f408, [%rd132+16];
  2818. setp.gt.f32 %p206, %f407, %f408;
  2819. mov.f32 %f1356, %f1355;
  2820. @%p206 bra BB1_221;
  2821.  
  2822. sub.f32 %f1044, %f408, %f407;
  2823. sqrt.rn.f32 %f1045, %f1044;
  2824. sub.f32 %f1356, %f406, %f1045;
  2825. add.f32 %f1355, %f406, %f1045;
  2826. mov.pred %p260, 0;
  2827.  
  2828. BB1_221:
  2829. @%p260 bra BB1_224;
  2830.  
  2831. setp.lt.f32 %p208, %f1356, 0f00000000;
  2832. selp.f32 %f413, %f1355, %f1356, %p208;
  2833. setp.geu.f32 %p209, %f413, %f1344;
  2834. @%p209 bra BB1_224;
  2835.  
  2836. mul.lo.s64 %rd258, %rd131, 80;
  2837. add.s64 %rd387, %rd1, %rd258;
  2838. add.s64 %rd372, %rd158, %rd258;
  2839. mov.f32 %f1344, %f413;
  2840.  
  2841. BB1_224:
  2842. add.s32 %r91, %r144, 2;
  2843. cvt.u64.u32 %rd137, %r91;
  2844. mul.wide.u32 %rd259, %r91, 80;
  2845. add.s64 %rd260, %rd1, %rd259;
  2846. add.s64 %rd138, %rd260, 8;
  2847. ld.global.f32 %f1048, [%rd260+8];
  2848. sub.f32 %f415, %f1048, %f348;
  2849. ld.global.f32 %f1049, [%rd260+12];
  2850. sub.f32 %f416, %f1049, %f349;
  2851. ld.global.f32 %f1050, [%rd260+16];
  2852. sub.f32 %f417, %f1050, %f350;
  2853. mul.f32 %f1051, %f346, %f416;
  2854. fma.rn.f32 %f1052, %f347, %f415, %f1051;
  2855. fma.rn.f32 %f418, %f345, %f417, %f1052;
  2856. setp.lt.f32 %p211, %f418, 0f00000000;
  2857. mov.pred %p261, -1;
  2858. mov.f32 %f1358, 0f4CBEBC20;
  2859. @%p211 bra BB1_225;
  2860. bra.uni BB1_226;
  2861.  
  2862. BB1_225:
  2863. mov.f32 %f1359, %f1358;
  2864. bra.uni BB1_228;
  2865.  
  2866. BB1_226:
  2867. mul.f32 %f1055, %f416, %f416;
  2868. fma.rn.f32 %f1056, %f415, %f415, %f1055;
  2869. fma.rn.f32 %f1057, %f417, %f417, %f1056;
  2870. mul.f32 %f1058, %f418, %f418;
  2871. sub.f32 %f419, %f1057, %f1058;
  2872. ld.global.f32 %f420, [%rd138+16];
  2873. setp.gt.f32 %p213, %f419, %f420;
  2874. mov.f32 %f1359, %f1358;
  2875. @%p213 bra BB1_228;
  2876.  
  2877. sub.f32 %f1059, %f420, %f419;
  2878. sqrt.rn.f32 %f1060, %f1059;
  2879. sub.f32 %f1359, %f418, %f1060;
  2880. add.f32 %f1358, %f418, %f1060;
  2881. mov.pred %p261, 0;
  2882.  
  2883. BB1_228:
  2884. @%p261 bra BB1_231;
  2885.  
  2886. setp.lt.f32 %p215, %f1359, 0f00000000;
  2887. selp.f32 %f425, %f1358, %f1359, %p215;
  2888. setp.geu.f32 %p216, %f425, %f1344;
  2889. @%p216 bra BB1_231;
  2890.  
  2891. mul.lo.s64 %rd261, %rd137, 80;
  2892. add.s64 %rd387, %rd1, %rd261;
  2893. add.s64 %rd372, %rd158, %rd261;
  2894. mov.f32 %f1344, %f425;
  2895.  
  2896. BB1_231:
  2897. add.s32 %r92, %r144, 3;
  2898. cvt.u64.u32 %rd143, %r92;
  2899. mul.wide.u32 %rd262, %r92, 80;
  2900. add.s64 %rd263, %rd1, %rd262;
  2901. add.s64 %rd144, %rd263, 8;
  2902. ld.global.f32 %f1063, [%rd263+8];
  2903. sub.f32 %f427, %f1063, %f348;
  2904. ld.global.f32 %f1064, [%rd263+12];
  2905. sub.f32 %f428, %f1064, %f349;
  2906. ld.global.f32 %f1065, [%rd263+16];
  2907. sub.f32 %f429, %f1065, %f350;
  2908. mul.f32 %f1066, %f346, %f428;
  2909. fma.rn.f32 %f1067, %f347, %f427, %f1066;
  2910. fma.rn.f32 %f430, %f345, %f429, %f1067;
  2911. setp.lt.f32 %p218, %f430, 0f00000000;
  2912. mov.pred %p262, -1;
  2913. mov.f32 %f1361, 0f4CBEBC20;
  2914. @%p218 bra BB1_232;
  2915. bra.uni BB1_233;
  2916.  
  2917. BB1_232:
  2918. mov.f32 %f1362, %f1361;
  2919. bra.uni BB1_235;
  2920.  
  2921. BB1_233:
  2922. mul.f32 %f1070, %f428, %f428;
  2923. fma.rn.f32 %f1071, %f427, %f427, %f1070;
  2924. fma.rn.f32 %f1072, %f429, %f429, %f1071;
  2925. mul.f32 %f1073, %f430, %f430;
  2926. sub.f32 %f431, %f1072, %f1073;
  2927. ld.global.f32 %f432, [%rd144+16];
  2928. setp.gt.f32 %p220, %f431, %f432;
  2929. mov.f32 %f1362, %f1361;
  2930. @%p220 bra BB1_235;
  2931.  
  2932. sub.f32 %f1074, %f432, %f431;
  2933. sqrt.rn.f32 %f1075, %f1074;
  2934. sub.f32 %f1362, %f430, %f1075;
  2935. add.f32 %f1361, %f430, %f1075;
  2936. mov.pred %p262, 0;
  2937.  
  2938. BB1_235:
  2939. @%p262 bra BB1_238;
  2940.  
  2941. setp.lt.f32 %p222, %f1362, 0f00000000;
  2942. selp.f32 %f437, %f1361, %f1362, %p222;
  2943. setp.geu.f32 %p223, %f437, %f1344;
  2944. @%p223 bra BB1_238;
  2945.  
  2946. mul.lo.s64 %rd264, %rd143, 80;
  2947. add.s64 %rd387, %rd1, %rd264;
  2948. add.s64 %rd372, %rd158, %rd264;
  2949. mov.f32 %f1344, %f437;
  2950.  
  2951. BB1_238:
  2952. cvt.u32.u64 %r93, %rd125;
  2953. add.s32 %r144, %r93, 4;
  2954. setp.lt.u32 %p224, %r144, %r37;
  2955. @%p224 bra BB1_210;
  2956.  
  2957. BB1_239:
  2958. setp.eq.s64 %p225, %rd372, 0;
  2959. mov.f32 %f1390, 0f40000000;
  2960. @%p225 bra BB1_174;
  2961.  
  2962. fma.rn.f32 %f440, %f347, %f1344, %f348;
  2963. fma.rn.f32 %f441, %f346, %f1344, %f349;
  2964. fma.rn.f32 %f442, %f345, %f1344, %f350;
  2965. add.s64 %rd151, %rd387, 8;
  2966. ld.global.f32 %f1079, [%rd387+8];
  2967. sub.f32 %f456, %f440, %f1079;
  2968. ld.global.f32 %f1080, [%rd387+12];
  2969. sub.f32 %f457, %f441, %f1080;
  2970. ld.global.f32 %f1081, [%rd387+16];
  2971. sub.f32 %f458, %f442, %f1081;
  2972. mul.f32 %f1082, %f457, %f457;
  2973. fma.rn.f32 %f1083, %f456, %f456, %f1082;
  2974. fma.rn.f32 %f446, %f458, %f458, %f1083;
  2975. setp.leu.f32 %p226, %f446, 0f00000000;
  2976. @%p226 bra BB1_242;
  2977.  
  2978. sqrt.rn.f32 %f1084, %f446;
  2979. rcp.rn.f32 %f1085, %f1084;
  2980. mul.f32 %f456, %f456, %f1085;
  2981. mul.f32 %f457, %f457, %f1085;
  2982. mul.f32 %f458, %f458, %f1085;
  2983.  
  2984. BB1_242:
  2985. mul.f32 %f1086, %f346, %f457;
  2986. fma.rn.f32 %f1087, %f347, %f456, %f1086;
  2987. fma.rn.f32 %f1088, %f345, %f458, %f1087;
  2988. mov.u16 %rs12, 0;
  2989. setp.leu.f32 %p227, %f1088, 0f00000000;
  2990. @%p227 bra BB1_244;
  2991.  
  2992. neg.f32 %f456, %f456;
  2993. neg.f32 %f457, %f457;
  2994. neg.f32 %f458, %f458;
  2995. mov.u16 %rs12, 1;
  2996.  
  2997. BB1_244:
  2998. ld.global.f32 %f1089, [%rd151+44];
  2999. setp.gt.f32 %p228, %f1089, 0f00000000;
  3000. @%p228 bra BB1_260;
  3001.  
  3002. ld.global.f32 %f1090, [%rd151+48];
  3003. setp.gt.f32 %p229, %f1090, 0f00000000;
  3004. @%p229 bra BB1_260;
  3005. bra.uni BB1_246;
  3006.  
  3007. BB1_260:
  3008. mov.u32 %r125, 3;
  3009. mul.f32 %f1133, %f346, %f457;
  3010. fma.rn.f32 %f1134, %f347, %f456, %f1133;
  3011. fma.rn.f32 %f487, %f345, %f458, %f1134;
  3012. add.f32 %f1135, %f487, 0f3F800000;
  3013. // Callseq Start 7
  3014. {
  3015. .reg .b32 temp_param_reg;
  3016. // <end>}
  3017. .param .b32 param0;
  3018. st.param.f32 [param0+0], %f1135;
  3019. .param .b32 param1;
  3020. st.param.b32 [param1+0], %r125;
  3021. .param .b32 retval0;
  3022. call.uni (retval0),
  3023. _Z5powiffi,
  3024. (
  3025. param0,
  3026. param1
  3027. );
  3028. ld.param.f32 %f1136, [retval0+0];
  3029.  
  3030. //{
  3031. }// Callseq End 7
  3032. fma.rn.f32 %f488, %f1136, 0f3F666666, 0f3DCCCCCD;
  3033. add.f32 %f1137, %f456, %f456;
  3034. mul.f32 %f1138, %f1137, %f487;
  3035. add.f32 %f1139, %f457, %f457;
  3036. mul.f32 %f1140, %f1139, %f487;
  3037. add.f32 %f1141, %f458, %f458;
  3038. mul.f32 %f1142, %f1141, %f487;
  3039. sub.f32 %f489, %f347, %f1138;
  3040. sub.f32 %f490, %f346, %f1140;
  3041. sub.f32 %f491, %f345, %f1142;
  3042. add.u64 %rd269, %SP, 12;
  3043. add.u64 %rd155, %SPL, 12;
  3044. st.local.f32 [%rd155], %f489;
  3045. st.local.f32 [%rd155+4], %f490;
  3046. st.local.f32 [%rd155+8], %f491;
  3047. mul.f32 %f1143, %f490, %f490;
  3048. fma.rn.f32 %f1144, %f489, %f489, %f1143;
  3049. fma.rn.f32 %f492, %f491, %f491, %f1144;
  3050. setp.leu.f32 %p238, %f492, 0f00000000;
  3051. @%p238 bra BB1_262;
  3052.  
  3053. sqrt.rn.f32 %f1145, %f492;
  3054. rcp.rn.f32 %f1146, %f1145;
  3055. mul.f32 %f1147, %f489, %f1146;
  3056. st.local.f32 [%rd155], %f1147;
  3057. mul.f32 %f1148, %f490, %f1146;
  3058. st.local.f32 [%rd155+4], %f1148;
  3059. mul.f32 %f1149, %f491, %f1146;
  3060. st.local.f32 [%rd155+8], %f1149;
  3061.  
  3062. BB1_262:
  3063. mul.f32 %f493, %f456, 0f38D1B717;
  3064. add.f32 %f1153, %f440, %f493;
  3065. mul.f32 %f494, %f457, 0f38D1B717;
  3066. add.f32 %f1154, %f441, %f494;
  3067. mul.f32 %f495, %f458, 0f38D1B717;
  3068. add.f32 %f1155, %f442, %f495;
  3069. add.u64 %rd270, %SP, 52;
  3070. add.u64 %rd271, %SPL, 52;
  3071. mov.u32 %r98, 2;
  3072. st.local.u32 [%rd271], %r98;
  3073. // Callseq Start 8
  3074. {
  3075. .reg .b32 temp_param_reg;
  3076. // <end>}
  3077. .param .b32 param0;
  3078. st.param.f32 [param0+0], %f1153;
  3079. .param .b32 param1;
  3080. st.param.f32 [param1+0], %f1154;
  3081. .param .b32 param2;
  3082. st.param.f32 [param2+0], %f1155;
  3083. .param .b64 param3;
  3084. st.param.b64 [param3+0], %rd269;
  3085. .param .b64 param4;
  3086. st.param.b64 [param4+0], %rd158;
  3087. .param .b32 param5;
  3088. st.param.b32 [param5+0], %r37;
  3089. .param .b64 param6;
  3090. st.param.b64 [param6+0], %rd270;
  3091. .param .align 4 .b8 retval0[12];
  3092. call.uni (retval0),
  3093. _Z5traceRK4Vec3IfES2_PK6SpherejRKi,
  3094. (
  3095. param0,
  3096. param1,
  3097. param2,
  3098. param3,
  3099. param4,
  3100. param5,
  3101. param6
  3102. );
  3103. ld.param.f32 %f496, [retval0+0];
  3104. ld.param.f32 %f497, [retval0+4];
  3105. ld.param.f32 %f498, [retval0+8];
  3106.  
  3107. //{
  3108. }// Callseq End 8
  3109. ld.global.f32 %f1381, [%rd151+44];
  3110. mov.f32 %f1382, 0f00000000;
  3111. setp.eq.f32 %p239, %f1381, 0f00000000;
  3112. mov.f32 %f1383, %f1382;
  3113. mov.f32 %f1384, %f1382;
  3114. @%p239 bra BB1_266;
  3115.  
  3116. setp.eq.s16 %p240, %rs12, 0;
  3117. selp.f32 %f1156, 0f3F68BA2E, 0f3F8CCCCD, %p240;
  3118. mul.f32 %f1157, %f1156, %f1156;
  3119. mul.f32 %f1158, %f487, %f487;
  3120. mov.f32 %f1159, 0f3F800000;
  3121. sub.f32 %f1160, %f1159, %f1158;
  3122. mul.f32 %f1161, %f1157, %f1160;
  3123. sub.f32 %f1162, %f1159, %f1161;
  3124. mul.f32 %f1163, %f487, %f1156;
  3125. neg.f32 %f1164, %f1163;
  3126. sqrt.rn.f32 %f1165, %f1162;
  3127. sub.f32 %f1166, %f1164, %f1165;
  3128. mul.f32 %f1167, %f456, %f1166;
  3129. mul.f32 %f1168, %f457, %f1166;
  3130. mul.f32 %f1169, %f458, %f1166;
  3131. fma.rn.f32 %f500, %f347, %f1156, %f1167;
  3132. fma.rn.f32 %f501, %f346, %f1156, %f1168;
  3133. fma.rn.f32 %f502, %f345, %f1156, %f1169;
  3134. add.u64 %rd273, %SP, 0;
  3135. add.u64 %rd156, %SPL, 0;
  3136. st.local.f32 [%rd156], %f500;
  3137. st.local.f32 [%rd156+4], %f501;
  3138. st.local.f32 [%rd156+8], %f502;
  3139. mul.f32 %f1170, %f501, %f501;
  3140. fma.rn.f32 %f1171, %f500, %f500, %f1170;
  3141. fma.rn.f32 %f503, %f502, %f502, %f1171;
  3142. setp.leu.f32 %p241, %f503, 0f00000000;
  3143. @%p241 bra BB1_265;
  3144.  
  3145. sqrt.rn.f32 %f1172, %f503;
  3146. rcp.rn.f32 %f1173, %f1172;
  3147. mul.f32 %f1174, %f500, %f1173;
  3148. st.local.f32 [%rd156], %f1174;
  3149. mul.f32 %f1175, %f501, %f1173;
  3150. st.local.f32 [%rd156+4], %f1175;
  3151. mul.f32 %f1176, %f502, %f1173;
  3152. st.local.f32 [%rd156+8], %f1176;
  3153.  
  3154. BB1_265:
  3155. sub.f32 %f1177, %f440, %f493;
  3156. add.u64 %rd274, %SP, 48;
  3157. add.u64 %rd275, %SPL, 48;
  3158. st.local.u32 [%rd275], %r98;
  3159. sub.f32 %f1178, %f442, %f495;
  3160. sub.f32 %f1179, %f441, %f494;
  3161. // Callseq Start 9
  3162. {
  3163. .reg .b32 temp_param_reg;
  3164. // <end>}
  3165. .param .b32 param0;
  3166. st.param.f32 [param0+0], %f1177;
  3167. .param .b32 param1;
  3168. st.param.f32 [param1+0], %f1179;
  3169. .param .b32 param2;
  3170. st.param.f32 [param2+0], %f1178;
  3171. .param .b64 param3;
  3172. st.param.b64 [param3+0], %rd273;
  3173. .param .b64 param4;
  3174. st.param.b64 [param4+0], %rd158;
  3175. .param .b32 param5;
  3176. st.param.b32 [param5+0], %r37;
  3177. .param .b64 param6;
  3178. st.param.b64 [param6+0], %rd274;
  3179. .param .align 4 .b8 retval0[12];
  3180. call.uni (retval0),
  3181. _Z5traceRK4Vec3IfES2_PK6SpherejRKi,
  3182. (
  3183. param0,
  3184. param1,
  3185. param2,
  3186. param3,
  3187. param4,
  3188. param5,
  3189. param6
  3190. );
  3191. ld.param.f32 %f1382, [retval0+0];
  3192. ld.param.f32 %f1383, [retval0+4];
  3193. ld.param.f32 %f1384, [retval0+8];
  3194.  
  3195. //{
  3196. }// Callseq End 9
  3197. ld.global.f32 %f1381, [%rd151+44];
  3198.  
  3199. BB1_266:
  3200. mov.f32 %f1180, 0f3F800000;
  3201. sub.f32 %f1181, %f1180, %f488;
  3202. mul.f32 %f1182, %f1181, %f1382;
  3203. mul.f32 %f1183, %f1181, %f1383;
  3204. mul.f32 %f1184, %f1181, %f1384;
  3205. mul.f32 %f1185, %f1381, %f1182;
  3206. mul.f32 %f1186, %f1381, %f1183;
  3207. mul.f32 %f1187, %f1381, %f1184;
  3208. fma.rn.f32 %f1188, %f488, %f496, %f1185;
  3209. fma.rn.f32 %f1189, %f488, %f497, %f1186;
  3210. fma.rn.f32 %f1190, %f488, %f498, %f1187;
  3211. ld.global.f32 %f1191, [%rd151+20];
  3212. mul.f32 %f1386, %f1191, %f1188;
  3213. ld.global.f32 %f1192, [%rd151+24];
  3214. mul.f32 %f1387, %f1189, %f1192;
  3215. ld.global.f32 %f1193, [%rd151+28];
  3216. mul.f32 %f1388, %f1190, %f1193;
  3217. ld.global.f32 %f1389, [%rd49+44];
  3218. bra.uni BB1_267;
  3219.  
  3220. BB1_174:
  3221. mov.f32 %f1391, %f1390;
  3222. mov.f32 %f1392, %f1390;
  3223. bra.uni BB1_268;
  3224.  
  3225. BB1_246:
  3226. mov.f32 %f1386, 0f00000000;
  3227. @%p26 bra BB1_247;
  3228.  
  3229. fma.rn.f32 %f459, %f456, 0f38D1B717, %f440;
  3230. fma.rn.f32 %f460, %f457, 0f38D1B717, %f441;
  3231. fma.rn.f32 %f461, %f458, 0f38D1B717, %f442;
  3232. mov.f32 %f1388, 0f00000000;
  3233. mov.u32 %r146, 0;
  3234. mov.f32 %f1387, %f1388;
  3235. mov.f32 %f1386, %f1388;
  3236.  
  3237. BB1_249:
  3238. cvt.u64.u32 %rd152, %r146;
  3239. mul.wide.u32 %rd265, %r146, 80;
  3240. add.s64 %rd266, %rd1, %rd265;
  3241. add.s64 %rd153, %rd266, 40;
  3242. ld.global.f32 %f465, [%rd266+40];
  3243. setp.leu.f32 %p231, %f465, 0f00000000;
  3244. @%p231 bra BB1_259;
  3245.  
  3246. ld.global.f32 %f1097, [%rd153+-32];
  3247. sub.f32 %f1376, %f1097, %f440;
  3248. ld.global.f32 %f1098, [%rd153+-28];
  3249. sub.f32 %f1375, %f1098, %f441;
  3250. ld.global.f32 %f1099, [%rd153+-24];
  3251. sub.f32 %f1374, %f1099, %f442;
  3252. mul.f32 %f1100, %f1375, %f1375;
  3253. fma.rn.f32 %f1101, %f1376, %f1376, %f1100;
  3254. fma.rn.f32 %f469, %f1374, %f1374, %f1101;
  3255. setp.leu.f32 %p232, %f469, 0f00000000;
  3256. @%p232 bra BB1_252;
  3257.  
  3258. sqrt.rn.f32 %f1102, %f469;
  3259. rcp.rn.f32 %f1103, %f1102;
  3260. mul.f32 %f1376, %f1376, %f1103;
  3261. mul.f32 %f1375, %f1375, %f1103;
  3262. mul.f32 %f1374, %f1374, %f1103;
  3263.  
  3264. BB1_252:
  3265. mov.u32 %r147, 0;
  3266.  
  3267. BB1_253:
  3268. cvt.u32.u64 %r96, %rd152;
  3269. setp.eq.s32 %p233, %r96, %r147;
  3270. @%p233 bra BB1_256;
  3271.  
  3272. mul.wide.u32 %rd267, %r147, 80;
  3273. add.s64 %rd268, %rd1, %rd267;
  3274. add.s64 %rd154, %rd268, 8;
  3275. ld.global.f32 %f1104, [%rd268+8];
  3276. sub.f32 %f476, %f1104, %f459;
  3277. ld.global.f32 %f1105, [%rd268+12];
  3278. sub.f32 %f477, %f1105, %f460;
  3279. ld.global.f32 %f1106, [%rd268+16];
  3280. sub.f32 %f478, %f1106, %f461;
  3281. mul.f32 %f1107, %f1375, %f477;
  3282. fma.rn.f32 %f1108, %f1376, %f476, %f1107;
  3283. fma.rn.f32 %f479, %f1374, %f478, %f1108;
  3284. setp.lt.f32 %p234, %f479, 0f00000000;
  3285. @%p234 bra BB1_256;
  3286.  
  3287. mul.f32 %f1110, %f477, %f477;
  3288. fma.rn.f32 %f1111, %f476, %f476, %f1110;
  3289. fma.rn.f32 %f1112, %f478, %f478, %f1111;
  3290. mul.f32 %f1113, %f479, %f479;
  3291. sub.f32 %f1114, %f1112, %f1113;
  3292. ld.global.f32 %f1115, [%rd154+16];
  3293. mov.f32 %f1377, 0f00000000;
  3294. setp.leu.f32 %p235, %f1114, %f1115;
  3295. @%p235 bra BB1_258;
  3296.  
  3297. BB1_256:
  3298. add.s32 %r147, %r147, 1;
  3299. setp.lt.u32 %p236, %r147, %r37;
  3300. @%p236 bra BB1_253;
  3301.  
  3302. mov.f32 %f1377, %f949;
  3303.  
  3304. BB1_258:
  3305. ld.global.f32 %f1117, [%rd151+20];
  3306. mul.f32 %f1118, %f1377, %f1117;
  3307. ld.global.f32 %f1119, [%rd151+24];
  3308. mul.f32 %f1120, %f1377, %f1119;
  3309. ld.global.f32 %f1121, [%rd151+28];
  3310. mul.f32 %f1122, %f1377, %f1121;
  3311. mul.f32 %f1123, %f456, %f1376;
  3312. fma.rn.f32 %f1124, %f457, %f1375, %f1123;
  3313. fma.rn.f32 %f1125, %f458, %f1374, %f1124;
  3314. mov.f32 %f1126, 0f00000000;
  3315. max.f32 %f1127, %f1126, %f1125;
  3316. mul.f32 %f1128, %f1118, %f1127;
  3317. mul.f32 %f1129, %f1120, %f1127;
  3318. mul.f32 %f1130, %f1122, %f1127;
  3319. ld.global.f32 %f1131, [%rd153+4];
  3320. ld.global.f32 %f1132, [%rd153+8];
  3321. fma.rn.f32 %f1386, %f465, %f1128, %f1386;
  3322. fma.rn.f32 %f1387, %f1129, %f1131, %f1387;
  3323. fma.rn.f32 %f1388, %f1130, %f1132, %f1388;
  3324.  
  3325. BB1_259:
  3326. add.s32 %r146, %r146, 1;
  3327. setp.lt.u32 %p237, %r146, %r37;
  3328. @%p237 bra BB1_249;
  3329. bra.uni BB1_267;
  3330.  
  3331. BB1_247:
  3332. mov.f32 %f1387, %f1386;
  3333. mov.f32 %f1388, %f1386;
  3334.  
  3335. BB1_267:
  3336. ld.global.f32 %f1194, [%rd151+32];
  3337. add.f32 %f1390, %f1386, %f1194;
  3338. ld.global.f32 %f1195, [%rd151+36];
  3339. add.f32 %f1391, %f1387, %f1195;
  3340. ld.global.f32 %f1196, [%rd151+40];
  3341. add.f32 %f1392, %f1388, %f1196;
  3342.  
  3343. BB1_268:
  3344. mov.f32 %f1197, 0f3F800000;
  3345. sub.f32 %f1198, %f1197, %f147;
  3346. mul.f32 %f1199, %f1198, %f1390;
  3347. mul.f32 %f1200, %f1198, %f1391;
  3348. mul.f32 %f1201, %f1198, %f1392;
  3349. mul.f32 %f1202, %f1199, %f1389;
  3350. mul.f32 %f1203, %f1200, %f1389;
  3351. mul.f32 %f1204, %f1201, %f1389;
  3352. fma.rn.f32 %f1205, %f147, %f1333, %f1202;
  3353. fma.rn.f32 %f1206, %f147, %f1334, %f1203;
  3354. fma.rn.f32 %f1207, %f147, %f1335, %f1204;
  3355. ld.global.f32 %f1208, [%rd49+20];
  3356. mul.f32 %f1393, %f1208, %f1205;
  3357. ld.global.f32 %f1209, [%rd49+24];
  3358. mul.f32 %f1394, %f1206, %f1209;
  3359. ld.global.f32 %f1210, [%rd49+28];
  3360. mul.f32 %f1395, %f1207, %f1210;
  3361.  
  3362. BB1_269:
  3363. ld.global.f32 %f1211, [%rd49+32];
  3364. add.f32 %f1396, %f1393, %f1211;
  3365. ld.global.f32 %f1212, [%rd49+36];
  3366. add.f32 %f1397, %f1394, %f1212;
  3367. ld.global.f32 %f1213, [%rd49+40];
  3368. add.f32 %f1398, %f1395, %f1213;
  3369.  
  3370. BB1_270:
  3371. ld.param.u64 %rd285, [raytrace_kernel_param_2];
  3372. ld.param.u32 %r117, [raytrace_kernel_param_0];
  3373. mov.u32 %r116, %tid.y;
  3374. mov.u32 %r115, %ctaid.y;
  3375. mov.u32 %r114, %ntid.y;
  3376. mad.lo.s32 %r113, %r114, %r115, %r116;
  3377. mov.u32 %r112, %tid.x;
  3378. mov.u32 %r111, %ctaid.x;
  3379. mov.u32 %r110, %ntid.x;
  3380. mad.lo.s32 %r109, %r110, %r111, %r112;
  3381. mad.lo.s32 %r108, %r113, %r117, %r109;
  3382. cvta.to.global.u64 %rd277, %rd285;
  3383. mul.wide.u32 %rd278, %r108, 12;
  3384. add.s64 %rd279, %rd277, %rd278;
  3385. st.global.f32 [%rd279], %f1396;
  3386. st.global.f32 [%rd279+4], %f1397;
  3387. st.global.f32 [%rd279+8], %f1398;
  3388.  
  3389. BB1_271:
  3390. ret;
  3391. }
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement