Advertisement
dineiar

raytracer_kernel_ptx_nvcc

Mar 30th, 2020
461
0
Never
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
text 102.36 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. .func (.param .b64 func_retval0) __internal_accurate_pow
  14. (
  15. .param .b64 __internal_accurate_pow_param_0
  16. )
  17. ;
  18.  
  19. .func (.param .align 4 .b8 func_retval0[12]) _Z5traceRK4Vec3IfES2_PK6SpherejRKi(
  20. .param .b32 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_0,
  21. .param .b32 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_1,
  22. .param .b32 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_2,
  23. .param .b32 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_3,
  24. .param .b32 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_4,
  25. .param .b32 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_5,
  26. .param .b64 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_6,
  27. .param .b32 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_7,
  28. .param .b32 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_8
  29. )
  30. {
  31. .reg .pred %p<111>;
  32. .reg .b16 %rs<5>;
  33. .reg .f32 %f<474>;
  34. .reg .b32 %r<66>;
  35. .reg .f64 %fd<20>;
  36. .reg .b64 %rd<78>;
  37.  
  38.  
  39. ld.param.f32 %f185, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_0];
  40. ld.param.f32 %f186, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_1];
  41. ld.param.f32 %f187, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_2];
  42. ld.param.f32 %f188, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_3];
  43. ld.param.f32 %f189, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_4];
  44. ld.param.f32 %f190, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_5];
  45. ld.param.u64 %rd35, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_6];
  46. ld.param.u32 %r16, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_7];
  47. ld.param.u32 %r17, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_8];
  48. setp.eq.s32 %p9, %r16, 0;
  49. mov.f32 %f422, 0f7F800000;
  50. mov.u64 %rd68, 0;
  51. @%p9 bra BB0_61;
  52.  
  53. and.b32 %r1, %r16, 3;
  54. setp.eq.s32 %p10, %r1, 0;
  55. mov.f32 %f428, 0f00000000;
  56. mov.f32 %f422, 0f7F800000;
  57. mov.u64 %rd68, 0;
  58. mov.u32 %r62, 0;
  59. @%p10 bra BB0_30;
  60.  
  61. setp.eq.s32 %p11, %r1, 1;
  62. mov.f32 %f424, 0f7F800000;
  63. mov.u64 %rd68, 0;
  64. mov.u32 %r61, 0;
  65. mov.f32 %f422, %f424;
  66. @%p11 bra BB0_22;
  67.  
  68. setp.eq.s32 %p12, %r1, 2;
  69. mov.f32 %f420, 0f7F800000;
  70. mov.u64 %rd68, 0;
  71. mov.u32 %r60, 0;
  72. @%p12 bra BB0_4;
  73. bra.uni BB0_5;
  74.  
  75. BB0_4:
  76. mov.f32 %f422, %f420;
  77. bra.uni BB0_14;
  78.  
  79. BB0_5:
  80. add.s64 %rd1, %rd35, 8;
  81. ld.f32 %f198, [%rd35+8];
  82. sub.f32 %f1, %f198, %f185;
  83. ld.f32 %f199, [%rd35+12];
  84. sub.f32 %f2, %f199, %f186;
  85. ld.f32 %f200, [%rd35+16];
  86. sub.f32 %f3, %f200, %f187;
  87. mul.f32 %f201, %f2, %f189;
  88. fma.rn.f32 %f202, %f1, %f188, %f201;
  89. fma.rn.f32 %f4, %f3, %f190, %f202;
  90. setp.lt.f32 %p14, %f4, 0f00000000;
  91. mov.pred %p104, -1;
  92. mov.f32 %f197, 0f7F800000;
  93. @%p14 bra BB0_6;
  94. bra.uni BB0_7;
  95.  
  96. BB0_6:
  97. mov.f32 %f417, %f197;
  98. mov.f32 %f418, %f197;
  99. bra.uni BB0_9;
  100.  
  101. BB0_7:
  102. mul.f32 %f205, %f2, %f2;
  103. fma.rn.f32 %f206, %f1, %f1, %f205;
  104. fma.rn.f32 %f207, %f3, %f3, %f206;
  105. mul.f32 %f208, %f4, %f4;
  106. sub.f32 %f5, %f207, %f208;
  107. ld.f32 %f6, [%rd1+16];
  108. setp.gt.f32 %p16, %f5, %f6;
  109. mov.f32 %f417, %f197;
  110. mov.f32 %f418, %f197;
  111. @%p16 bra BB0_9;
  112.  
  113. sub.f32 %f209, %f6, %f5;
  114. sqrt.rn.f32 %f210, %f209;
  115. sub.f32 %f418, %f4, %f210;
  116. add.f32 %f417, %f4, %f210;
  117. mov.pred %p104, 0;
  118.  
  119. BB0_9:
  120. mov.u32 %r60, 1;
  121. @%p104 bra BB0_10;
  122. bra.uni BB0_11;
  123.  
  124. BB0_10:
  125. mov.f32 %f422, %f197;
  126. bra.uni BB0_14;
  127.  
  128. BB0_11:
  129. setp.lt.f32 %p18, %f418, 0f00000000;
  130. selp.f32 %f422, %f417, %f418, %p18;
  131. setp.geu.f32 %p19, %f422, 0f7F800000;
  132. @%p19 bra BB0_12;
  133.  
  134. mov.u64 %rd68, %rd35;
  135. bra.uni BB0_14;
  136.  
  137. BB0_12:
  138. mov.f32 %f422, %f197;
  139.  
  140. BB0_14:
  141. cvt.u64.u32 %rd3, %r60;
  142. mul.wide.u32 %rd42, %r60, 80;
  143. add.s64 %rd43, %rd35, %rd42;
  144. add.s64 %rd4, %rd43, 8;
  145. ld.f32 %f215, [%rd43+8];
  146. sub.f32 %f13, %f215, %f185;
  147. ld.f32 %f216, [%rd43+12];
  148. sub.f32 %f14, %f216, %f186;
  149. ld.f32 %f217, [%rd43+16];
  150. sub.f32 %f15, %f217, %f187;
  151. mul.f32 %f218, %f14, %f189;
  152. fma.rn.f32 %f219, %f13, %f188, %f218;
  153. fma.rn.f32 %f16, %f15, %f190, %f219;
  154. setp.lt.f32 %p21, %f16, 0f00000000;
  155. mov.pred %p105, -1;
  156. @%p21 bra BB0_15;
  157. bra.uni BB0_16;
  158.  
  159. BB0_15:
  160. mov.f32 %f421, %f420;
  161. bra.uni BB0_18;
  162.  
  163. BB0_16:
  164. mul.f32 %f222, %f14, %f14;
  165. fma.rn.f32 %f223, %f13, %f13, %f222;
  166. fma.rn.f32 %f224, %f15, %f15, %f223;
  167. mul.f32 %f225, %f16, %f16;
  168. sub.f32 %f17, %f224, %f225;
  169. ld.f32 %f18, [%rd4+16];
  170. setp.gt.f32 %p23, %f17, %f18;
  171. mov.f32 %f421, %f420;
  172. @%p23 bra BB0_18;
  173.  
  174. sub.f32 %f226, %f18, %f17;
  175. sqrt.rn.f32 %f227, %f226;
  176. sub.f32 %f421, %f16, %f227;
  177. add.f32 %f420, %f16, %f227;
  178. mov.pred %p105, 0;
  179.  
  180. BB0_18:
  181. @%p105 bra BB0_21;
  182.  
  183. setp.lt.f32 %p25, %f421, 0f00000000;
  184. selp.f32 %f23, %f420, %f421, %p25;
  185. setp.geu.f32 %p26, %f23, %f422;
  186. @%p26 bra BB0_21;
  187.  
  188. mul.lo.s64 %rd44, %rd3, 80;
  189. add.s64 %rd68, %rd35, %rd44;
  190. mov.f32 %f422, %f23;
  191.  
  192. BB0_21:
  193. cvt.u32.u64 %r24, %rd3;
  194. add.s32 %r61, %r24, 1;
  195.  
  196. BB0_22:
  197. cvt.u64.u32 %rd8, %r61;
  198. mul.wide.u32 %rd45, %r61, 80;
  199. add.s64 %rd46, %rd35, %rd45;
  200. add.s64 %rd9, %rd46, 8;
  201. ld.f32 %f230, [%rd46+8];
  202. sub.f32 %f26, %f230, %f185;
  203. ld.f32 %f231, [%rd46+12];
  204. sub.f32 %f27, %f231, %f186;
  205. ld.f32 %f232, [%rd46+16];
  206. sub.f32 %f28, %f232, %f187;
  207. mul.f32 %f233, %f27, %f189;
  208. fma.rn.f32 %f234, %f26, %f188, %f233;
  209. fma.rn.f32 %f29, %f28, %f190, %f234;
  210. setp.lt.f32 %p28, %f29, 0f00000000;
  211. mov.pred %p106, -1;
  212. @%p28 bra BB0_23;
  213. bra.uni BB0_24;
  214.  
  215. BB0_23:
  216. mov.f32 %f425, %f424;
  217. bra.uni BB0_26;
  218.  
  219. BB0_24:
  220. mul.f32 %f237, %f27, %f27;
  221. fma.rn.f32 %f238, %f26, %f26, %f237;
  222. fma.rn.f32 %f239, %f28, %f28, %f238;
  223. mul.f32 %f240, %f29, %f29;
  224. sub.f32 %f30, %f239, %f240;
  225. ld.f32 %f31, [%rd9+16];
  226. setp.gt.f32 %p30, %f30, %f31;
  227. mov.f32 %f425, %f424;
  228. @%p30 bra BB0_26;
  229.  
  230. sub.f32 %f241, %f31, %f30;
  231. sqrt.rn.f32 %f242, %f241;
  232. sub.f32 %f425, %f29, %f242;
  233. add.f32 %f424, %f29, %f242;
  234. mov.pred %p106, 0;
  235.  
  236. BB0_26:
  237. @%p106 bra BB0_29;
  238.  
  239. setp.lt.f32 %p32, %f425, 0f00000000;
  240. selp.f32 %f36, %f424, %f425, %p32;
  241. setp.geu.f32 %p33, %f36, %f422;
  242. @%p33 bra BB0_29;
  243.  
  244. mul.lo.s64 %rd47, %rd8, 80;
  245. add.s64 %rd68, %rd35, %rd47;
  246. mov.f32 %f422, %f36;
  247.  
  248. BB0_29:
  249. cvt.u32.u64 %r25, %rd8;
  250. add.s32 %r62, %r25, 1;
  251. mov.f32 %f428, %f422;
  252.  
  253. BB0_30:
  254. setp.lt.u32 %p34, %r16, 4;
  255. @%p34 bra BB0_31;
  256. bra.uni BB0_32;
  257.  
  258. BB0_31:
  259. mov.f32 %f422, %f428;
  260. bra.uni BB0_61;
  261.  
  262. BB0_32:
  263. cvt.u64.u32 %rd14, %r62;
  264. mul.wide.u32 %rd48, %r62, 80;
  265. add.s64 %rd49, %rd35, %rd48;
  266. add.s64 %rd15, %rd49, 8;
  267. ld.f32 %f245, [%rd49+8];
  268. sub.f32 %f41, %f245, %f185;
  269. ld.f32 %f246, [%rd49+12];
  270. sub.f32 %f42, %f246, %f186;
  271. ld.f32 %f247, [%rd49+16];
  272. sub.f32 %f43, %f247, %f187;
  273. mul.f32 %f248, %f42, %f189;
  274. fma.rn.f32 %f249, %f41, %f188, %f248;
  275. fma.rn.f32 %f44, %f43, %f190, %f249;
  276. setp.lt.f32 %p36, %f44, 0f00000000;
  277. mov.pred %p107, -1;
  278. mov.f32 %f430, 0f7F800000;
  279. @%p36 bra BB0_33;
  280. bra.uni BB0_34;
  281.  
  282. BB0_33:
  283. mov.f32 %f431, %f430;
  284. bra.uni BB0_36;
  285.  
  286. BB0_34:
  287. mul.f32 %f252, %f42, %f42;
  288. fma.rn.f32 %f253, %f41, %f41, %f252;
  289. fma.rn.f32 %f254, %f43, %f43, %f253;
  290. mul.f32 %f255, %f44, %f44;
  291. sub.f32 %f45, %f254, %f255;
  292. ld.f32 %f46, [%rd15+16];
  293. setp.gt.f32 %p38, %f45, %f46;
  294. mov.f32 %f431, %f430;
  295. @%p38 bra BB0_36;
  296.  
  297. sub.f32 %f256, %f46, %f45;
  298. sqrt.rn.f32 %f257, %f256;
  299. sub.f32 %f431, %f44, %f257;
  300. add.f32 %f430, %f44, %f257;
  301. mov.pred %p107, 0;
  302.  
  303. BB0_36:
  304. @%p107 bra BB0_39;
  305.  
  306. setp.lt.f32 %p40, %f431, 0f00000000;
  307. selp.f32 %f51, %f430, %f431, %p40;
  308. setp.geu.f32 %p41, %f51, %f422;
  309. @%p41 bra BB0_39;
  310.  
  311. mul.lo.s64 %rd50, %rd14, 80;
  312. add.s64 %rd68, %rd35, %rd50;
  313. mov.f32 %f422, %f51;
  314.  
  315. BB0_39:
  316. add.s32 %r26, %r62, 1;
  317. cvt.u64.u32 %rd18, %r26;
  318. mul.wide.u32 %rd51, %r26, 80;
  319. add.s64 %rd52, %rd35, %rd51;
  320. add.s64 %rd19, %rd52, 8;
  321. ld.f32 %f260, [%rd52+8];
  322. sub.f32 %f53, %f260, %f185;
  323. ld.f32 %f261, [%rd52+12];
  324. sub.f32 %f54, %f261, %f186;
  325. ld.f32 %f262, [%rd52+16];
  326. sub.f32 %f55, %f262, %f187;
  327. mul.f32 %f263, %f54, %f189;
  328. fma.rn.f32 %f264, %f53, %f188, %f263;
  329. fma.rn.f32 %f56, %f55, %f190, %f264;
  330. setp.lt.f32 %p43, %f56, 0f00000000;
  331. mov.pred %p108, -1;
  332. mov.f32 %f433, 0f7F800000;
  333. @%p43 bra BB0_40;
  334. bra.uni BB0_41;
  335.  
  336. BB0_40:
  337. mov.f32 %f434, %f433;
  338. bra.uni BB0_43;
  339.  
  340. BB0_41:
  341. mul.f32 %f267, %f54, %f54;
  342. fma.rn.f32 %f268, %f53, %f53, %f267;
  343. fma.rn.f32 %f269, %f55, %f55, %f268;
  344. mul.f32 %f270, %f56, %f56;
  345. sub.f32 %f57, %f269, %f270;
  346. ld.f32 %f58, [%rd19+16];
  347. setp.gt.f32 %p45, %f57, %f58;
  348. mov.f32 %f434, %f433;
  349. @%p45 bra BB0_43;
  350.  
  351. sub.f32 %f271, %f58, %f57;
  352. sqrt.rn.f32 %f272, %f271;
  353. sub.f32 %f434, %f56, %f272;
  354. add.f32 %f433, %f56, %f272;
  355. mov.pred %p108, 0;
  356.  
  357. BB0_43:
  358. @%p108 bra BB0_46;
  359.  
  360. setp.lt.f32 %p47, %f434, 0f00000000;
  361. selp.f32 %f63, %f433, %f434, %p47;
  362. setp.geu.f32 %p48, %f63, %f422;
  363. @%p48 bra BB0_46;
  364.  
  365. mul.lo.s64 %rd53, %rd18, 80;
  366. add.s64 %rd68, %rd35, %rd53;
  367. mov.f32 %f422, %f63;
  368.  
  369. BB0_46:
  370. add.s32 %r27, %r62, 2;
  371. cvt.u64.u32 %rd22, %r27;
  372. mul.wide.u32 %rd54, %r27, 80;
  373. add.s64 %rd55, %rd35, %rd54;
  374. add.s64 %rd23, %rd55, 8;
  375. ld.f32 %f275, [%rd55+8];
  376. sub.f32 %f65, %f275, %f185;
  377. ld.f32 %f276, [%rd55+12];
  378. sub.f32 %f66, %f276, %f186;
  379. ld.f32 %f277, [%rd55+16];
  380. sub.f32 %f67, %f277, %f187;
  381. mul.f32 %f278, %f66, %f189;
  382. fma.rn.f32 %f279, %f65, %f188, %f278;
  383. fma.rn.f32 %f68, %f67, %f190, %f279;
  384. setp.lt.f32 %p50, %f68, 0f00000000;
  385. mov.pred %p109, -1;
  386. mov.f32 %f436, 0f7F800000;
  387. @%p50 bra BB0_47;
  388. bra.uni BB0_48;
  389.  
  390. BB0_47:
  391. mov.f32 %f437, %f436;
  392. bra.uni BB0_50;
  393.  
  394. BB0_48:
  395. mul.f32 %f282, %f66, %f66;
  396. fma.rn.f32 %f283, %f65, %f65, %f282;
  397. fma.rn.f32 %f284, %f67, %f67, %f283;
  398. mul.f32 %f285, %f68, %f68;
  399. sub.f32 %f69, %f284, %f285;
  400. ld.f32 %f70, [%rd23+16];
  401. setp.gt.f32 %p52, %f69, %f70;
  402. mov.f32 %f437, %f436;
  403. @%p52 bra BB0_50;
  404.  
  405. sub.f32 %f286, %f70, %f69;
  406. sqrt.rn.f32 %f287, %f286;
  407. sub.f32 %f437, %f68, %f287;
  408. add.f32 %f436, %f68, %f287;
  409. mov.pred %p109, 0;
  410.  
  411. BB0_50:
  412. @%p109 bra BB0_53;
  413.  
  414. setp.lt.f32 %p54, %f437, 0f00000000;
  415. selp.f32 %f75, %f436, %f437, %p54;
  416. setp.geu.f32 %p55, %f75, %f422;
  417. @%p55 bra BB0_53;
  418.  
  419. mul.lo.s64 %rd56, %rd22, 80;
  420. add.s64 %rd68, %rd35, %rd56;
  421. mov.f32 %f422, %f75;
  422.  
  423. BB0_53:
  424. add.s32 %r28, %r62, 3;
  425. cvt.u64.u32 %rd26, %r28;
  426. mul.wide.u32 %rd57, %r28, 80;
  427. add.s64 %rd58, %rd35, %rd57;
  428. add.s64 %rd27, %rd58, 8;
  429. ld.f32 %f290, [%rd58+8];
  430. sub.f32 %f77, %f290, %f185;
  431. ld.f32 %f291, [%rd58+12];
  432. sub.f32 %f78, %f291, %f186;
  433. ld.f32 %f292, [%rd58+16];
  434. sub.f32 %f79, %f292, %f187;
  435. mul.f32 %f293, %f78, %f189;
  436. fma.rn.f32 %f294, %f77, %f188, %f293;
  437. fma.rn.f32 %f80, %f79, %f190, %f294;
  438. setp.lt.f32 %p57, %f80, 0f00000000;
  439. mov.pred %p110, -1;
  440. mov.f32 %f439, 0f7F800000;
  441. @%p57 bra BB0_54;
  442. bra.uni BB0_55;
  443.  
  444. BB0_54:
  445. mov.f32 %f440, %f439;
  446. bra.uni BB0_57;
  447.  
  448. BB0_55:
  449. mul.f32 %f297, %f78, %f78;
  450. fma.rn.f32 %f298, %f77, %f77, %f297;
  451. fma.rn.f32 %f299, %f79, %f79, %f298;
  452. mul.f32 %f300, %f80, %f80;
  453. sub.f32 %f81, %f299, %f300;
  454. ld.f32 %f82, [%rd27+16];
  455. setp.gt.f32 %p59, %f81, %f82;
  456. mov.f32 %f440, %f439;
  457. @%p59 bra BB0_57;
  458.  
  459. sub.f32 %f301, %f82, %f81;
  460. sqrt.rn.f32 %f302, %f301;
  461. sub.f32 %f440, %f80, %f302;
  462. add.f32 %f439, %f80, %f302;
  463. mov.pred %p110, 0;
  464.  
  465. BB0_57:
  466. @%p110 bra BB0_60;
  467.  
  468. setp.lt.f32 %p61, %f440, 0f00000000;
  469. selp.f32 %f87, %f439, %f440, %p61;
  470. setp.geu.f32 %p62, %f87, %f422;
  471. @%p62 bra BB0_60;
  472.  
  473. mul.lo.s64 %rd59, %rd26, 80;
  474. add.s64 %rd68, %rd35, %rd59;
  475. mov.f32 %f422, %f87;
  476.  
  477. BB0_60:
  478. cvt.u32.u64 %r29, %rd14;
  479. add.s32 %r62, %r29, 4;
  480. setp.lt.u32 %p63, %r62, %r16;
  481. @%p63 bra BB0_32;
  482.  
  483. BB0_61:
  484. setp.eq.s64 %p64, %rd68, 0;
  485. mov.f32 %f471, 0f40000000;
  486. mov.f32 %f472, %f471;
  487. mov.f32 %f473, %f471;
  488. @%p64 bra BB0_105;
  489.  
  490. fma.rn.f32 %f90, %f422, %f188, %f185;
  491. fma.rn.f32 %f91, %f422, %f189, %f186;
  492. fma.rn.f32 %f92, %f422, %f190, %f187;
  493. add.s64 %rd31, %rd68, 8;
  494. ld.f32 %f306, [%rd68+8];
  495. sub.f32 %f443, %f90, %f306;
  496. ld.f32 %f307, [%rd68+12];
  497. sub.f32 %f444, %f91, %f307;
  498. ld.f32 %f308, [%rd68+16];
  499. sub.f32 %f445, %f92, %f308;
  500. mul.f32 %f309, %f444, %f444;
  501. fma.rn.f32 %f310, %f443, %f443, %f309;
  502. fma.rn.f32 %f96, %f445, %f445, %f310;
  503. setp.leu.f32 %p65, %f96, 0f00000000;
  504. @%p65 bra BB0_64;
  505.  
  506. sqrt.rn.f32 %f311, %f96;
  507. rcp.rn.f32 %f312, %f311;
  508. mul.f32 %f443, %f443, %f312;
  509. mul.f32 %f444, %f444, %f312;
  510. mul.f32 %f445, %f445, %f312;
  511.  
  512. BB0_64:
  513. mul.f32 %f313, %f444, %f189;
  514. fma.rn.f32 %f314, %f443, %f188, %f313;
  515. fma.rn.f32 %f315, %f445, %f190, %f314;
  516. mov.u16 %rs4, 0;
  517. setp.leu.f32 %p66, %f315, 0f00000000;
  518. @%p66 bra BB0_66;
  519.  
  520. neg.f32 %f443, %f443;
  521. neg.f32 %f444, %f444;
  522. neg.f32 %f445, %f445;
  523. mov.u16 %rs4, 1;
  524.  
  525. BB0_66:
  526. ld.f32 %f109, [%rd31+44];
  527. setp.gt.f32 %p67, %f109, 0f00000000;
  528. @%p67 bra BB0_73;
  529. bra.uni BB0_67;
  530.  
  531. BB0_73:
  532. setp.lt.s32 %p71, %r17, 5;
  533. @%p71 bra BB0_68;
  534. bra.uni BB0_74;
  535.  
  536. BB0_67:
  537. setp.lt.s32 %p68, %r17, 5;
  538. ld.f32 %f316, [%rd31+48];
  539. setp.gt.f32 %p69, %f316, 0f00000000;
  540. and.pred %p70, %p69, %p68;
  541. @%p70 bra BB0_68;
  542. bra.uni BB0_74;
  543.  
  544. BB0_68:
  545. mul.f32 %f359, %f444, %f189;
  546. fma.rn.f32 %f360, %f443, %f188, %f359;
  547. fma.rn.f32 %f138, %f445, %f190, %f360;
  548. add.f32 %f139, %f138, 0f3F800000;
  549. cvt.f64.f32 %fd1, %f139;
  550. {
  551. .reg .b32 %temp;
  552. mov.b64 {%temp, %r13}, %fd1;
  553. }
  554. mov.f64 %fd13, 0d4008000000000000;
  555. {
  556. .reg .b32 %temp;
  557. mov.b64 {%temp, %r14}, %fd13;
  558. }
  559. bfe.u32 %r33, %r14, 20, 11;
  560. add.s32 %r34, %r33, -1012;
  561. mov.u64 %rd66, 4613937818241073152;
  562. shl.b64 %rd34, %rd66, %r34;
  563. setp.ne.s64 %p80, %rd34, -9223372036854775808;
  564. setp.eq.s64 %p81, %rd34, -9223372036854775808;
  565. abs.f64 %fd2, %fd1;
  566. // Callseq Start 0
  567. {
  568. .reg .b32 temp_param_reg;
  569. // <end>}
  570. .param .b64 param0;
  571. st.param.f64 [param0+0], %fd2;
  572. .param .b64 retval0;
  573. call.uni (retval0),
  574. __internal_accurate_pow,
  575. (
  576. param0
  577. );
  578. ld.param.f64 %fd18, [retval0+0];
  579.  
  580. //{
  581. }// Callseq End 0
  582. setp.gt.s32 %p82, %r13, -1;
  583. setp.lt.s32 %p83, %r13, 0;
  584. and.pred %p8, %p83, %p81;
  585. or.pred %p84, %p82, %p80;
  586. @%p84 bra BB0_70;
  587.  
  588. {
  589. .reg .b32 %temp;
  590. mov.b64 {%temp, %r35}, %fd18;
  591. }
  592. xor.b32 %r36, %r35, -2147483648;
  593. {
  594. .reg .b32 %temp;
  595. mov.b64 {%r37, %temp}, %fd18;
  596. }
  597. mov.b64 %fd18, {%r37, %r36};
  598.  
  599. BB0_70:
  600. setp.eq.f32 %p85, %f139, 0f00000000;
  601. @%p85 bra BB0_86;
  602. bra.uni BB0_71;
  603.  
  604. BB0_86:
  605. selp.b32 %r38, %r13, 0, %p81;
  606. mov.u32 %r39, 0;
  607. or.b32 %r40, %r38, 2146435072;
  608. setp.lt.s32 %p89, %r14, 0;
  609. selp.b32 %r41, %r40, %r38, %p89;
  610. mov.b64 %fd18, {%r39, %r41};
  611. bra.uni BB0_87;
  612.  
  613. BB0_74:
  614. mov.f32 %f468, 0f00000000;
  615. mov.f32 %f469, %f468;
  616. mov.f32 %f470, %f468;
  617. @%p9 bra BB0_104;
  618.  
  619. fma.rn.f32 %f110, %f443, 0f38D1B717, %f90;
  620. fma.rn.f32 %f111, %f444, 0f38D1B717, %f91;
  621. fma.rn.f32 %f112, %f445, 0f38D1B717, %f92;
  622. mov.f32 %f470, 0f00000000;
  623. mov.u32 %r64, 0;
  624. mov.f32 %f469, %f470;
  625. mov.f32 %f468, %f470;
  626.  
  627. BB0_76:
  628. cvt.u64.u32 %rd32, %r64;
  629. mul.wide.u32 %rd60, %r64, 80;
  630. add.s64 %rd61, %rd35, %rd60;
  631. add.s64 %rd33, %rd61, 40;
  632. ld.f32 %f116, [%rd61+40];
  633. setp.leu.f32 %p73, %f116, 0f00000000;
  634. @%p73 bra BB0_85;
  635.  
  636. ld.f32 %f323, [%rd33+-32];
  637. sub.f32 %f454, %f323, %f90;
  638. ld.f32 %f324, [%rd33+-28];
  639. sub.f32 %f453, %f324, %f91;
  640. ld.f32 %f325, [%rd33+-24];
  641. sub.f32 %f452, %f325, %f92;
  642. mul.f32 %f326, %f453, %f453;
  643. fma.rn.f32 %f327, %f454, %f454, %f326;
  644. fma.rn.f32 %f120, %f452, %f452, %f327;
  645. setp.leu.f32 %p74, %f120, 0f00000000;
  646. @%p74 bra BB0_79;
  647.  
  648. sqrt.rn.f32 %f328, %f120;
  649. rcp.rn.f32 %f329, %f328;
  650. mul.f32 %f454, %f454, %f329;
  651. mul.f32 %f453, %f453, %f329;
  652. mul.f32 %f452, %f452, %f329;
  653.  
  654. BB0_79:
  655. mov.u32 %r65, 0;
  656.  
  657. BB0_80:
  658. cvt.u32.u64 %r32, %rd32;
  659. setp.eq.s32 %p75, %r32, %r65;
  660. @%p75 bra BB0_83;
  661.  
  662. mul.wide.u32 %rd62, %r65, 80;
  663. add.s64 %rd63, %rd35, %rd62;
  664. ld.f32 %f330, [%rd63+8];
  665. sub.f32 %f127, %f330, %f110;
  666. ld.f32 %f331, [%rd63+12];
  667. sub.f32 %f128, %f331, %f111;
  668. ld.f32 %f332, [%rd63+16];
  669. sub.f32 %f129, %f332, %f112;
  670. mul.f32 %f333, %f453, %f128;
  671. fma.rn.f32 %f334, %f454, %f127, %f333;
  672. fma.rn.f32 %f130, %f452, %f129, %f334;
  673. setp.lt.f32 %p76, %f130, 0f00000000;
  674. @%p76 bra BB0_83;
  675.  
  676. mul.f32 %f336, %f128, %f128;
  677. fma.rn.f32 %f337, %f127, %f127, %f336;
  678. fma.rn.f32 %f338, %f129, %f129, %f337;
  679. mul.f32 %f339, %f130, %f130;
  680. sub.f32 %f340, %f338, %f339;
  681. ld.f32 %f341, [%rd63+24];
  682. mov.f32 %f455, 0f00000000;
  683. setp.leu.f32 %p77, %f340, %f341;
  684. @%p77 bra BB0_84;
  685.  
  686. BB0_83:
  687. add.s32 %r65, %r65, 1;
  688. setp.lt.u32 %p78, %r65, %r16;
  689. mov.f32 %f455, 0f3F800000;
  690. @%p78 bra BB0_80;
  691.  
  692. BB0_84:
  693. ld.f32 %f343, [%rd31+20];
  694. mul.f32 %f344, %f455, %f343;
  695. ld.f32 %f345, [%rd31+24];
  696. mul.f32 %f346, %f455, %f345;
  697. ld.f32 %f347, [%rd31+28];
  698. mul.f32 %f348, %f455, %f347;
  699. mul.f32 %f349, %f443, %f454;
  700. fma.rn.f32 %f350, %f444, %f453, %f349;
  701. fma.rn.f32 %f351, %f445, %f452, %f350;
  702. mov.f32 %f352, 0f00000000;
  703. max.f32 %f353, %f352, %f351;
  704. mul.f32 %f354, %f344, %f353;
  705. mul.f32 %f355, %f346, %f353;
  706. mul.f32 %f356, %f348, %f353;
  707. ld.f32 %f357, [%rd33+4];
  708. ld.f32 %f358, [%rd33+8];
  709. fma.rn.f32 %f468, %f354, %f116, %f468;
  710. fma.rn.f32 %f469, %f355, %f357, %f469;
  711. fma.rn.f32 %f470, %f356, %f358, %f470;
  712.  
  713. BB0_85:
  714. add.s32 %r64, %r64, 1;
  715. setp.lt.u32 %p79, %r64, %r16;
  716. @%p79 bra BB0_76;
  717. bra.uni BB0_104;
  718.  
  719. BB0_71:
  720. @%p82 bra BB0_87;
  721.  
  722. cvt.rzi.f64.f64 %fd15, %fd13;
  723. setp.neu.f64 %p87, %fd15, 0d4008000000000000;
  724. selp.f64 %fd18, 0dFFF8000000000000, %fd18, %p87;
  725.  
  726. BB0_87:
  727. add.f64 %fd19, %fd1, 0d4008000000000000;
  728. {
  729. .reg .b32 %temp;
  730. mov.b64 {%temp, %r42}, %fd19;
  731. }
  732. and.b32 %r43, %r42, 2146435072;
  733. setp.ne.s32 %p90, %r43, 2146435072;
  734. @%p90 bra BB0_88;
  735.  
  736. setp.gtu.f64 %p91, %fd2, 0d7FF0000000000000;
  737. @%p91 bra BB0_97;
  738.  
  739. and.b32 %r44, %r14, 2147483647;
  740. setp.ne.s32 %p92, %r44, 2146435072;
  741. @%p92 bra BB0_92;
  742.  
  743. {
  744. .reg .b32 %temp;
  745. mov.b64 {%r45, %temp}, %fd13;
  746. }
  747. setp.eq.s32 %p93, %r45, 0;
  748. @%p93 bra BB0_96;
  749.  
  750. BB0_92:
  751. and.b32 %r46, %r13, 2147483647;
  752. setp.ne.s32 %p94, %r46, 2146435072;
  753. @%p94 bra BB0_93;
  754.  
  755. {
  756. .reg .b32 %temp;
  757. mov.b64 {%r47, %temp}, %fd1;
  758. }
  759. setp.ne.s32 %p95, %r47, 0;
  760. mov.f64 %fd19, %fd18;
  761. @%p95 bra BB0_97;
  762.  
  763. shr.s32 %r48, %r14, 31;
  764. and.b32 %r49, %r48, -2146435072;
  765. add.s32 %r50, %r49, 2146435072;
  766. or.b32 %r51, %r50, -2147483648;
  767. selp.b32 %r52, %r51, %r50, %p8;
  768. mov.u32 %r53, 0;
  769. mov.b64 %fd19, {%r53, %r52};
  770. bra.uni BB0_97;
  771.  
  772. BB0_88:
  773. mov.f64 %fd19, %fd18;
  774. bra.uni BB0_97;
  775.  
  776. BB0_93:
  777. mov.f64 %fd19, %fd18;
  778. bra.uni BB0_97;
  779.  
  780. BB0_96:
  781. setp.gt.f64 %p96, %fd2, 0d3FF0000000000000;
  782. selp.b32 %r54, 2146435072, 0, %p96;
  783. mov.u32 %r55, 0;
  784. xor.b32 %r56, %r54, 2146435072;
  785. setp.lt.s32 %p97, %r14, 0;
  786. selp.b32 %r57, %r56, %r54, %p97;
  787. setp.eq.f32 %p98, %f139, 0fBF800000;
  788. selp.b32 %r58, 1072693248, %r57, %p98;
  789. mov.b64 %fd19, {%r55, %r58};
  790.  
  791. BB0_97:
  792. cvt.rn.f32.f64 %f361, %fd19;
  793. fma.rn.f32 %f362, %f361, 0f3F666666, 0f3DCCCCCD;
  794. setp.eq.f32 %p99, %f139, 0f3F800000;
  795. selp.f32 %f140, 0f3F800000, %f362, %p99;
  796. add.f32 %f363, %f443, %f443;
  797. mul.f32 %f364, %f363, %f138;
  798. add.f32 %f365, %f444, %f444;
  799. mul.f32 %f366, %f365, %f138;
  800. add.f32 %f367, %f445, %f445;
  801. mul.f32 %f368, %f367, %f138;
  802. sub.f32 %f459, %f188, %f364;
  803. sub.f32 %f460, %f189, %f366;
  804. sub.f32 %f461, %f190, %f368;
  805. mul.f32 %f369, %f460, %f460;
  806. fma.rn.f32 %f370, %f459, %f459, %f369;
  807. fma.rn.f32 %f144, %f461, %f461, %f370;
  808. setp.leu.f32 %p100, %f144, 0f00000000;
  809. @%p100 bra BB0_99;
  810.  
  811. sqrt.rn.f32 %f371, %f144;
  812. rcp.rn.f32 %f372, %f371;
  813. mul.f32 %f459, %f459, %f372;
  814. mul.f32 %f460, %f460, %f372;
  815. mul.f32 %f461, %f461, %f372;
  816.  
  817. BB0_99:
  818. mul.f32 %f151, %f443, 0f38D1B717;
  819. add.f32 %f376, %f90, %f151;
  820. mul.f32 %f152, %f444, 0f38D1B717;
  821. add.f32 %f377, %f91, %f152;
  822. mul.f32 %f153, %f445, 0f38D1B717;
  823. add.f32 %f378, %f92, %f153;
  824. add.s32 %r15, %r17, 1;
  825. // Callseq Start 1
  826. {
  827. .reg .b32 temp_param_reg;
  828. // <end>}
  829. .param .b32 param0;
  830. st.param.f32 [param0+0], %f376;
  831. .param .b32 param1;
  832. st.param.f32 [param1+0], %f377;
  833. .param .b32 param2;
  834. st.param.f32 [param2+0], %f378;
  835. .param .b32 param3;
  836. st.param.f32 [param3+0], %f459;
  837. .param .b32 param4;
  838. st.param.f32 [param4+0], %f460;
  839. .param .b32 param5;
  840. st.param.f32 [param5+0], %f461;
  841. .param .b64 param6;
  842. st.param.b64 [param6+0], %rd35;
  843. .param .b32 param7;
  844. st.param.b32 [param7+0], %r16;
  845. .param .b32 param8;
  846. st.param.b32 [param8+0], %r15;
  847. .param .align 4 .b8 retval0[12];
  848. call.uni (retval0),
  849. _Z5traceRK4Vec3IfES2_PK6SpherejRKi,
  850. (
  851. param0,
  852. param1,
  853. param2,
  854. param3,
  855. param4,
  856. param5,
  857. param6,
  858. param7,
  859. param8
  860. );
  861. ld.param.f32 %f154, [retval0+0];
  862. ld.param.f32 %f155, [retval0+4];
  863. ld.param.f32 %f156, [retval0+8];
  864.  
  865. //{
  866. }// Callseq End 1
  867. mov.f32 %f465, 0f00000000;
  868. setp.eq.f32 %p101, %f109, 0f00000000;
  869. mov.f32 %f466, %f465;
  870. mov.f32 %f467, %f465;
  871. @%p101 bra BB0_103;
  872.  
  873. setp.eq.s16 %p102, %rs4, 0;
  874. selp.f32 %f379, 0f3F68BA2E, 0f3F8CCCCD, %p102;
  875. mul.f32 %f380, %f379, %f379;
  876. mul.f32 %f381, %f138, %f138;
  877. mov.f32 %f382, 0f3F800000;
  878. sub.f32 %f383, %f382, %f381;
  879. mul.f32 %f384, %f380, %f383;
  880. sub.f32 %f385, %f382, %f384;
  881. mul.f32 %f386, %f138, %f379;
  882. neg.f32 %f387, %f386;
  883. sqrt.rn.f32 %f388, %f385;
  884. sub.f32 %f389, %f387, %f388;
  885. mul.f32 %f390, %f443, %f389;
  886. mul.f32 %f391, %f444, %f389;
  887. mul.f32 %f392, %f445, %f389;
  888. fma.rn.f32 %f464, %f379, %f188, %f390;
  889. fma.rn.f32 %f463, %f379, %f189, %f391;
  890. fma.rn.f32 %f462, %f379, %f190, %f392;
  891. mul.f32 %f393, %f463, %f463;
  892. fma.rn.f32 %f394, %f464, %f464, %f393;
  893. fma.rn.f32 %f160, %f462, %f462, %f394;
  894. setp.leu.f32 %p103, %f160, 0f00000000;
  895. @%p103 bra BB0_102;
  896.  
  897. sqrt.rn.f32 %f395, %f160;
  898. rcp.rn.f32 %f396, %f395;
  899. mul.f32 %f464, %f464, %f396;
  900. mul.f32 %f463, %f463, %f396;
  901. mul.f32 %f462, %f462, %f396;
  902.  
  903. BB0_102:
  904. sub.f32 %f397, %f92, %f153;
  905. sub.f32 %f398, %f91, %f152;
  906. sub.f32 %f399, %f90, %f151;
  907. // Callseq Start 2
  908. {
  909. .reg .b32 temp_param_reg;
  910. // <end>}
  911. .param .b32 param0;
  912. st.param.f32 [param0+0], %f399;
  913. .param .b32 param1;
  914. st.param.f32 [param1+0], %f398;
  915. .param .b32 param2;
  916. st.param.f32 [param2+0], %f397;
  917. .param .b32 param3;
  918. st.param.f32 [param3+0], %f464;
  919. .param .b32 param4;
  920. st.param.f32 [param4+0], %f463;
  921. .param .b32 param5;
  922. st.param.f32 [param5+0], %f462;
  923. .param .b64 param6;
  924. st.param.b64 [param6+0], %rd35;
  925. .param .b32 param7;
  926. st.param.b32 [param7+0], %r16;
  927. .param .b32 param8;
  928. add.s32 %r59, %r17, 1;
  929. st.param.b32 [param8+0], %r59;
  930. .param .align 4 .b8 retval0[12];
  931. call.uni (retval0),
  932. _Z5traceRK4Vec3IfES2_PK6SpherejRKi,
  933. (
  934. param0,
  935. param1,
  936. param2,
  937. param3,
  938. param4,
  939. param5,
  940. param6,
  941. param7,
  942. param8
  943. );
  944. ld.param.f32 %f465, [retval0+0];
  945. ld.param.f32 %f466, [retval0+4];
  946. ld.param.f32 %f467, [retval0+8];
  947.  
  948. //{
  949. }// Callseq End 2
  950.  
  951. BB0_103:
  952. mov.f32 %f400, 0f3F800000;
  953. sub.f32 %f401, %f400, %f140;
  954. mul.f32 %f402, %f401, %f465;
  955. mul.f32 %f403, %f401, %f466;
  956. mul.f32 %f404, %f401, %f467;
  957. mul.f32 %f405, %f402, %f109;
  958. mul.f32 %f406, %f403, %f109;
  959. mul.f32 %f407, %f404, %f109;
  960. fma.rn.f32 %f408, %f140, %f154, %f405;
  961. fma.rn.f32 %f409, %f140, %f155, %f406;
  962. fma.rn.f32 %f410, %f140, %f156, %f407;
  963. ld.f32 %f411, [%rd31+20];
  964. mul.f32 %f468, %f411, %f408;
  965. ld.f32 %f412, [%rd31+24];
  966. mul.f32 %f469, %f409, %f412;
  967. ld.f32 %f413, [%rd31+28];
  968. mul.f32 %f470, %f410, %f413;
  969.  
  970. BB0_104:
  971. ld.f32 %f414, [%rd31+32];
  972. add.f32 %f471, %f468, %f414;
  973. ld.f32 %f415, [%rd31+36];
  974. add.f32 %f472, %f469, %f415;
  975. ld.f32 %f416, [%rd31+40];
  976. add.f32 %f473, %f470, %f416;
  977.  
  978. BB0_105:
  979. st.param.f32 [func_retval0+0], %f471;
  980. st.param.f32 [func_retval0+4], %f472;
  981. st.param.f32 [func_retval0+8], %f473;
  982. ret;
  983. }
  984.  
  985. // .globl raytrace_kernel
  986. .visible .entry raytrace_kernel(
  987. .param .u32 raytrace_kernel_param_0,
  988. .param .u32 raytrace_kernel_param_1,
  989. .param .u64 raytrace_kernel_param_2,
  990. .param .u64 raytrace_kernel_param_3,
  991. .param .u32 raytrace_kernel_param_4,
  992. .param .f32 raytrace_kernel_param_5,
  993. .param .f32 raytrace_kernel_param_6,
  994. .param .f32 raytrace_kernel_param_7,
  995. .param .f32 raytrace_kernel_param_8
  996. )
  997. {
  998. .reg .pred %p<326>;
  999. .reg .b16 %rs<13>;
  1000. .reg .f32 %f<1415>;
  1001. .reg .b32 %r<220>;
  1002. .reg .f64 %fd<79>;
  1003. .reg .b64 %rd<373>;
  1004.  
  1005.  
  1006. ld.param.u32 %r40, [raytrace_kernel_param_0];
  1007. ld.param.u32 %r42, [raytrace_kernel_param_1];
  1008. ld.param.u64 %rd153, [raytrace_kernel_param_3];
  1009. ld.param.u32 %r41, [raytrace_kernel_param_4];
  1010. ld.param.f32 %f559, [raytrace_kernel_param_5];
  1011. ld.param.f32 %f560, [raytrace_kernel_param_6];
  1012. ld.param.f32 %f561, [raytrace_kernel_param_7];
  1013. ld.param.f32 %f562, [raytrace_kernel_param_8];
  1014. cvta.to.global.u64 %rd1, %rd153;
  1015. mov.u32 %r43, %ntid.x;
  1016. mov.u32 %r44, %ctaid.x;
  1017. mov.u32 %r45, %tid.x;
  1018. mad.lo.s32 %r1, %r43, %r44, %r45;
  1019. mov.u32 %r46, %ntid.y;
  1020. mov.u32 %r47, %ctaid.y;
  1021. mov.u32 %r48, %tid.y;
  1022. mad.lo.s32 %r2, %r46, %r47, %r48;
  1023. setp.ge.u32 %p25, %r2, %r42;
  1024. setp.ge.u32 %p26, %r1, %r40;
  1025. or.pred %p27, %p25, %p26;
  1026. @%p27 bra BB1_317;
  1027.  
  1028. cvt.rn.f64.s32 %fd37, %r1;
  1029. add.f64 %fd38, %fd37, 0d3FE0000000000000;
  1030. cvt.f64.f32 %fd39, %f559;
  1031. mul.f64 %fd40, %fd39, %fd38;
  1032. fma.rn.f64 %fd41, %fd40, 0d4000000000000000, 0dBFF0000000000000;
  1033. cvt.f64.f32 %fd42, %f562;
  1034. mul.f64 %fd43, %fd42, %fd41;
  1035. cvt.f64.f32 %fd44, %f561;
  1036. mul.f64 %fd45, %fd44, %fd43;
  1037. cvt.rn.f32.f64 %f1246, %fd45;
  1038. cvt.rn.f64.s32 %fd46, %r2;
  1039. add.f64 %fd47, %fd46, 0d3FE0000000000000;
  1040. cvt.f64.f32 %fd48, %f560;
  1041. mul.f64 %fd49, %fd48, %fd47;
  1042. fma.rn.f64 %fd50, %fd48, %fd47, %fd49;
  1043. mov.f64 %fd51, 0d3FF0000000000000;
  1044. sub.f64 %fd52, %fd51, %fd50;
  1045. mul.f64 %fd53, %fd42, %fd52;
  1046. cvt.rn.f32.f64 %f1245, %fd53;
  1047. mul.f32 %f564, %f1245, %f1245;
  1048. fma.rn.f32 %f565, %f1246, %f1246, %f564;
  1049. add.f32 %f3, %f565, 0f3F800000;
  1050. mov.f32 %f1244, 0fBF800000;
  1051. setp.leu.f32 %p28, %f3, 0f00000000;
  1052. @%p28 bra BB1_3;
  1053.  
  1054. sqrt.rn.f32 %f566, %f3;
  1055. rcp.rn.f32 %f567, %f566;
  1056. mul.f32 %f1246, %f1246, %f567;
  1057. mul.f32 %f1245, %f1245, %f567;
  1058. neg.f32 %f1244, %f567;
  1059.  
  1060. BB1_3:
  1061. setp.eq.s32 %p29, %r41, 0;
  1062. mov.f32 %f1252, 0f7F800000;
  1063. mov.u64 %rd310, 0;
  1064. mov.u64 %rd309, %rd310;
  1065. @%p29 bra BB1_63;
  1066.  
  1067. and.b32 %r52, %r41, 3;
  1068. mov.f32 %f1252, 0f7F800000;
  1069. mov.u64 %rd310, 0;
  1070. mov.u32 %r204, 0;
  1071. mov.f32 %f1258, 0f00000000;
  1072. setp.eq.s32 %p30, %r52, 0;
  1073. mov.u64 %rd309, %rd310;
  1074. @%p30 bra BB1_32;
  1075.  
  1076. mov.f32 %f1252, 0f7F800000;
  1077. mov.u64 %rd310, 0;
  1078. mov.u32 %r203, 0;
  1079. and.b32 %r192, %r41, 3;
  1080. setp.eq.s32 %p31, %r192, 1;
  1081. mov.u64 %rd309, %rd310;
  1082. @%p31 bra BB1_24;
  1083.  
  1084. mov.f32 %f1252, 0f7F800000;
  1085. mov.u64 %rd310, 0;
  1086. mov.u32 %r202, 0;
  1087. and.b32 %r194, %r41, 3;
  1088. setp.eq.s32 %p32, %r194, 2;
  1089. mov.u64 %rd309, %rd310;
  1090. @%p32 bra BB1_16;
  1091.  
  1092. ld.global.f32 %f10, [%rd1+8];
  1093. ld.global.f32 %f11, [%rd1+12];
  1094. mul.f32 %f575, %f1245, %f11;
  1095. fma.rn.f32 %f576, %f1246, %f10, %f575;
  1096. ld.global.f32 %f12, [%rd1+16];
  1097. fma.rn.f32 %f13, %f1244, %f12, %f576;
  1098. setp.lt.f32 %p34, %f13, 0f00000000;
  1099. mov.pred %p305, -1;
  1100. mov.f32 %f1247, 0f7F800000;
  1101. @%p34 bra BB1_8;
  1102. bra.uni BB1_9;
  1103.  
  1104. BB1_8:
  1105. mov.f32 %f1248, %f1247;
  1106. bra.uni BB1_11;
  1107.  
  1108. BB1_9:
  1109. mov.f32 %f1247, 0f7F800000;
  1110. add.s64 %rd268, %rd1, 8;
  1111. mul.f32 %f579, %f11, %f11;
  1112. fma.rn.f32 %f580, %f10, %f10, %f579;
  1113. fma.rn.f32 %f581, %f12, %f12, %f580;
  1114. mul.f32 %f582, %f13, %f13;
  1115. sub.f32 %f14, %f581, %f582;
  1116. ld.global.f32 %f15, [%rd268+16];
  1117. setp.gt.f32 %p36, %f14, %f15;
  1118. mov.f32 %f1248, %f1247;
  1119. @%p36 bra BB1_11;
  1120.  
  1121. sub.f32 %f583, %f15, %f14;
  1122. sqrt.rn.f32 %f584, %f583;
  1123. sub.f32 %f1248, %f13, %f584;
  1124. add.f32 %f1247, %f13, %f584;
  1125. mov.pred %p305, 0;
  1126.  
  1127. BB1_11:
  1128. mov.u64 %rd310, 0;
  1129. mov.f32 %f1252, 0f7F800000;
  1130. mov.u32 %r202, 1;
  1131. @%p305 bra BB1_12;
  1132. bra.uni BB1_13;
  1133.  
  1134. BB1_12:
  1135. mov.u64 %rd309, %rd310;
  1136. bra.uni BB1_16;
  1137.  
  1138. BB1_13:
  1139. mov.u64 %rd310, 0;
  1140. mov.f32 %f1224, 0f7F800000;
  1141. setp.lt.f32 %p38, %f1248, 0f00000000;
  1142. selp.f32 %f1252, %f1247, %f1248, %p38;
  1143. setp.geu.f32 %p39, %f1252, 0f7F800000;
  1144. @%p39 bra BB1_14;
  1145.  
  1146. mov.u64 %rd310, %rd153;
  1147. mov.u64 %rd309, %rd1;
  1148. bra.uni BB1_16;
  1149.  
  1150. BB1_14:
  1151. mov.u64 %rd309, %rd310;
  1152. mov.f32 %f1252, %f1224;
  1153.  
  1154. BB1_16:
  1155. mov.f32 %f1250, 0f7F800000;
  1156. mul.wide.u32 %rd166, %r202, 80;
  1157. add.s64 %rd167, %rd1, %rd166;
  1158. ld.global.f32 %f22, [%rd167+8];
  1159. ld.global.f32 %f23, [%rd167+12];
  1160. mul.f32 %f589, %f1245, %f23;
  1161. fma.rn.f32 %f590, %f1246, %f22, %f589;
  1162. ld.global.f32 %f24, [%rd167+16];
  1163. fma.rn.f32 %f25, %f1244, %f24, %f590;
  1164. setp.lt.f32 %p41, %f25, 0f00000000;
  1165. mov.pred %p306, -1;
  1166. @%p41 bra BB1_17;
  1167. bra.uni BB1_18;
  1168.  
  1169. BB1_17:
  1170. mov.f32 %f1251, %f1250;
  1171. bra.uni BB1_20;
  1172.  
  1173. BB1_18:
  1174. mul.wide.u32 %rd263, %r202, 80;
  1175. add.s64 %rd262, %rd1, %rd263;
  1176. mov.f32 %f1250, 0f7F800000;
  1177. mul.f32 %f593, %f23, %f23;
  1178. fma.rn.f32 %f594, %f22, %f22, %f593;
  1179. fma.rn.f32 %f595, %f24, %f24, %f594;
  1180. mul.f32 %f596, %f25, %f25;
  1181. sub.f32 %f26, %f595, %f596;
  1182. ld.global.f32 %f27, [%rd262+24];
  1183. setp.gt.f32 %p43, %f26, %f27;
  1184. mov.f32 %f1251, %f1250;
  1185. @%p43 bra BB1_20;
  1186.  
  1187. sub.f32 %f597, %f27, %f26;
  1188. sqrt.rn.f32 %f598, %f597;
  1189. sub.f32 %f1251, %f25, %f598;
  1190. add.f32 %f1250, %f25, %f598;
  1191. mov.pred %p306, 0;
  1192.  
  1193. BB1_20:
  1194. @%p306 bra BB1_23;
  1195.  
  1196. setp.lt.f32 %p45, %f1251, 0f00000000;
  1197. selp.f32 %f32, %f1250, %f1251, %p45;
  1198. setp.geu.f32 %p46, %f32, %f1252;
  1199. @%p46 bra BB1_23;
  1200.  
  1201. mul.wide.u32 %rd265, %r202, 80;
  1202. add.s64 %rd309, %rd1, %rd265;
  1203. add.s64 %rd310, %rd153, %rd265;
  1204. mov.f32 %f1252, %f32;
  1205.  
  1206. BB1_23:
  1207. add.s32 %r203, %r202, 1;
  1208.  
  1209. BB1_24:
  1210. mov.f32 %f1254, 0f7F800000;
  1211. mul.wide.u32 %rd171, %r203, 80;
  1212. add.s64 %rd172, %rd1, %rd171;
  1213. ld.global.f32 %f35, [%rd172+8];
  1214. ld.global.f32 %f36, [%rd172+12];
  1215. mul.f32 %f601, %f1245, %f36;
  1216. fma.rn.f32 %f602, %f1246, %f35, %f601;
  1217. ld.global.f32 %f37, [%rd172+16];
  1218. fma.rn.f32 %f38, %f1244, %f37, %f602;
  1219. setp.lt.f32 %p48, %f38, 0f00000000;
  1220. mov.pred %p307, -1;
  1221. @%p48 bra BB1_25;
  1222. bra.uni BB1_26;
  1223.  
  1224. BB1_25:
  1225. mov.f32 %f1255, %f1254;
  1226. bra.uni BB1_28;
  1227.  
  1228. BB1_26:
  1229. mul.wide.u32 %rd273, %r203, 80;
  1230. add.s64 %rd272, %rd1, %rd273;
  1231. add.s64 %rd271, %rd272, 8;
  1232. mov.f32 %f1254, 0f7F800000;
  1233. mul.f32 %f605, %f36, %f36;
  1234. fma.rn.f32 %f606, %f35, %f35, %f605;
  1235. fma.rn.f32 %f607, %f37, %f37, %f606;
  1236. mul.f32 %f608, %f38, %f38;
  1237. sub.f32 %f39, %f607, %f608;
  1238. ld.global.f32 %f40, [%rd271+16];
  1239. setp.gt.f32 %p50, %f39, %f40;
  1240. mov.f32 %f1255, %f1254;
  1241. @%p50 bra BB1_28;
  1242.  
  1243. sub.f32 %f609, %f40, %f39;
  1244. sqrt.rn.f32 %f610, %f609;
  1245. sub.f32 %f1255, %f38, %f610;
  1246. add.f32 %f1254, %f38, %f610;
  1247. mov.pred %p307, 0;
  1248.  
  1249. BB1_28:
  1250. @%p307 bra BB1_31;
  1251.  
  1252. setp.lt.f32 %p52, %f1255, 0f00000000;
  1253. selp.f32 %f45, %f1254, %f1255, %p52;
  1254. setp.geu.f32 %p53, %f45, %f1252;
  1255. @%p53 bra BB1_31;
  1256.  
  1257. cvt.u64.u32 %rd274, %r203;
  1258. mul.lo.s64 %rd173, %rd274, 80;
  1259. add.s64 %rd309, %rd1, %rd173;
  1260. add.s64 %rd310, %rd153, %rd173;
  1261. mov.f32 %f1252, %f45;
  1262.  
  1263. BB1_31:
  1264. cvt.u64.u32 %rd275, %r203;
  1265. cvt.u32.u64 %r56, %rd275;
  1266. add.s32 %r204, %r56, 1;
  1267. mov.f32 %f1258, %f1252;
  1268.  
  1269. BB1_32:
  1270. setp.lt.u32 %p54, %r41, 4;
  1271. @%p54 bra BB1_33;
  1272. bra.uni BB1_34;
  1273.  
  1274. BB1_33:
  1275. mov.f32 %f1252, %f1258;
  1276. bra.uni BB1_63;
  1277.  
  1278. BB1_34:
  1279. mul.wide.u32 %rd174, %r204, 80;
  1280. add.s64 %rd175, %rd1, %rd174;
  1281. ld.global.f32 %f50, [%rd175+8];
  1282. ld.global.f32 %f51, [%rd175+12];
  1283. mul.f32 %f613, %f1245, %f51;
  1284. fma.rn.f32 %f614, %f1246, %f50, %f613;
  1285. ld.global.f32 %f52, [%rd175+16];
  1286. fma.rn.f32 %f53, %f1244, %f52, %f614;
  1287. setp.lt.f32 %p56, %f53, 0f00000000;
  1288. mov.pred %p308, -1;
  1289. mov.f32 %f1260, 0f7F800000;
  1290. @%p56 bra BB1_35;
  1291. bra.uni BB1_36;
  1292.  
  1293. BB1_35:
  1294. mov.f32 %f1261, %f1260;
  1295. bra.uni BB1_38;
  1296.  
  1297. BB1_36:
  1298. mul.wide.u32 %rd290, %r204, 80;
  1299. add.s64 %rd289, %rd1, %rd290;
  1300. add.s64 %rd288, %rd289, 8;
  1301. mul.f32 %f617, %f51, %f51;
  1302. fma.rn.f32 %f618, %f50, %f50, %f617;
  1303. fma.rn.f32 %f619, %f52, %f52, %f618;
  1304. mul.f32 %f620, %f53, %f53;
  1305. sub.f32 %f54, %f619, %f620;
  1306. ld.global.f32 %f55, [%rd288+16];
  1307. setp.gt.f32 %p58, %f54, %f55;
  1308. mov.f32 %f1261, %f1260;
  1309. @%p58 bra BB1_38;
  1310.  
  1311. sub.f32 %f621, %f55, %f54;
  1312. sqrt.rn.f32 %f622, %f621;
  1313. sub.f32 %f1261, %f53, %f622;
  1314. add.f32 %f1260, %f53, %f622;
  1315. mov.pred %p308, 0;
  1316.  
  1317. BB1_38:
  1318. @%p308 bra BB1_41;
  1319.  
  1320. setp.lt.f32 %p60, %f1261, 0f00000000;
  1321. selp.f32 %f60, %f1260, %f1261, %p60;
  1322. setp.geu.f32 %p61, %f60, %f1252;
  1323. @%p61 bra BB1_41;
  1324.  
  1325. cvt.u64.u32 %rd291, %r204;
  1326. mul.lo.s64 %rd176, %rd291, 80;
  1327. add.s64 %rd309, %rd1, %rd176;
  1328. add.s64 %rd310, %rd153, %rd176;
  1329. mov.f32 %f1252, %f60;
  1330.  
  1331. BB1_41:
  1332. add.s32 %r57, %r204, 1;
  1333. mul.wide.u32 %rd177, %r57, 80;
  1334. add.s64 %rd178, %rd1, %rd177;
  1335. ld.global.f32 %f62, [%rd178+8];
  1336. ld.global.f32 %f63, [%rd178+12];
  1337. mul.f32 %f625, %f1245, %f63;
  1338. fma.rn.f32 %f626, %f1246, %f62, %f625;
  1339. ld.global.f32 %f64, [%rd178+16];
  1340. fma.rn.f32 %f65, %f1244, %f64, %f626;
  1341. setp.lt.f32 %p63, %f65, 0f00000000;
  1342. mov.pred %p309, -1;
  1343. mov.f32 %f1263, 0f7F800000;
  1344. @%p63 bra BB1_42;
  1345. bra.uni BB1_43;
  1346.  
  1347. BB1_42:
  1348. mov.f32 %f1264, %f1263;
  1349. bra.uni BB1_45;
  1350.  
  1351. BB1_43:
  1352. add.s32 %r196, %r204, 1;
  1353. mul.wide.u32 %rd295, %r196, 80;
  1354. add.s64 %rd294, %rd1, %rd295;
  1355. add.s64 %rd293, %rd294, 8;
  1356. mul.f32 %f629, %f63, %f63;
  1357. fma.rn.f32 %f630, %f62, %f62, %f629;
  1358. fma.rn.f32 %f631, %f64, %f64, %f630;
  1359. mul.f32 %f632, %f65, %f65;
  1360. sub.f32 %f66, %f631, %f632;
  1361. ld.global.f32 %f67, [%rd293+16];
  1362. setp.gt.f32 %p65, %f66, %f67;
  1363. mov.f32 %f1264, %f1263;
  1364. @%p65 bra BB1_45;
  1365.  
  1366. sub.f32 %f633, %f67, %f66;
  1367. sqrt.rn.f32 %f634, %f633;
  1368. sub.f32 %f1264, %f65, %f634;
  1369. add.f32 %f1263, %f65, %f634;
  1370. mov.pred %p309, 0;
  1371.  
  1372. BB1_45:
  1373. @%p309 bra BB1_48;
  1374.  
  1375. setp.lt.f32 %p67, %f1264, 0f00000000;
  1376. selp.f32 %f72, %f1263, %f1264, %p67;
  1377. setp.geu.f32 %p68, %f72, %f1252;
  1378. @%p68 bra BB1_48;
  1379.  
  1380. add.s32 %r197, %r204, 1;
  1381. cvt.u64.u32 %rd296, %r197;
  1382. mul.lo.s64 %rd179, %rd296, 80;
  1383. add.s64 %rd309, %rd1, %rd179;
  1384. add.s64 %rd310, %rd153, %rd179;
  1385. mov.f32 %f1252, %f72;
  1386.  
  1387. BB1_48:
  1388. add.s32 %r58, %r204, 2;
  1389. mul.wide.u32 %rd180, %r58, 80;
  1390. add.s64 %rd181, %rd1, %rd180;
  1391. ld.global.f32 %f74, [%rd181+8];
  1392. ld.global.f32 %f75, [%rd181+12];
  1393. mul.f32 %f637, %f1245, %f75;
  1394. fma.rn.f32 %f638, %f1246, %f74, %f637;
  1395. ld.global.f32 %f76, [%rd181+16];
  1396. fma.rn.f32 %f77, %f1244, %f76, %f638;
  1397. setp.lt.f32 %p70, %f77, 0f00000000;
  1398. mov.pred %p310, -1;
  1399. mov.f32 %f1266, 0f7F800000;
  1400. @%p70 bra BB1_49;
  1401. bra.uni BB1_50;
  1402.  
  1403. BB1_49:
  1404. mov.f32 %f1267, %f1266;
  1405. bra.uni BB1_52;
  1406.  
  1407. BB1_50:
  1408. add.s32 %r198, %r204, 2;
  1409. mul.wide.u32 %rd299, %r198, 80;
  1410. add.s64 %rd298, %rd1, %rd299;
  1411. add.s64 %rd297, %rd298, 8;
  1412. mul.f32 %f641, %f75, %f75;
  1413. fma.rn.f32 %f642, %f74, %f74, %f641;
  1414. fma.rn.f32 %f643, %f76, %f76, %f642;
  1415. mul.f32 %f644, %f77, %f77;
  1416. sub.f32 %f78, %f643, %f644;
  1417. ld.global.f32 %f79, [%rd297+16];
  1418. setp.gt.f32 %p72, %f78, %f79;
  1419. mov.f32 %f1267, %f1266;
  1420. @%p72 bra BB1_52;
  1421.  
  1422. sub.f32 %f645, %f79, %f78;
  1423. sqrt.rn.f32 %f646, %f645;
  1424. sub.f32 %f1267, %f77, %f646;
  1425. add.f32 %f1266, %f77, %f646;
  1426. mov.pred %p310, 0;
  1427.  
  1428. BB1_52:
  1429. @%p310 bra BB1_55;
  1430.  
  1431. setp.lt.f32 %p74, %f1267, 0f00000000;
  1432. selp.f32 %f84, %f1266, %f1267, %p74;
  1433. setp.geu.f32 %p75, %f84, %f1252;
  1434. @%p75 bra BB1_55;
  1435.  
  1436. add.s32 %r199, %r204, 2;
  1437. cvt.u64.u32 %rd300, %r199;
  1438. mul.lo.s64 %rd182, %rd300, 80;
  1439. add.s64 %rd309, %rd1, %rd182;
  1440. add.s64 %rd310, %rd153, %rd182;
  1441. mov.f32 %f1252, %f84;
  1442.  
  1443. BB1_55:
  1444. add.s32 %r59, %r204, 3;
  1445. mul.wide.u32 %rd183, %r59, 80;
  1446. add.s64 %rd184, %rd1, %rd183;
  1447. ld.global.f32 %f86, [%rd184+8];
  1448. ld.global.f32 %f87, [%rd184+12];
  1449. mul.f32 %f649, %f1245, %f87;
  1450. fma.rn.f32 %f650, %f1246, %f86, %f649;
  1451. ld.global.f32 %f88, [%rd184+16];
  1452. fma.rn.f32 %f89, %f1244, %f88, %f650;
  1453. setp.lt.f32 %p77, %f89, 0f00000000;
  1454. mov.pred %p311, -1;
  1455. mov.f32 %f1269, 0f7F800000;
  1456. @%p77 bra BB1_56;
  1457. bra.uni BB1_57;
  1458.  
  1459. BB1_56:
  1460. mov.f32 %f1270, %f1269;
  1461. bra.uni BB1_59;
  1462.  
  1463. BB1_57:
  1464. add.s32 %r200, %r204, 3;
  1465. mul.wide.u32 %rd303, %r200, 80;
  1466. add.s64 %rd302, %rd1, %rd303;
  1467. add.s64 %rd301, %rd302, 8;
  1468. mul.f32 %f653, %f87, %f87;
  1469. fma.rn.f32 %f654, %f86, %f86, %f653;
  1470. fma.rn.f32 %f655, %f88, %f88, %f654;
  1471. mul.f32 %f656, %f89, %f89;
  1472. sub.f32 %f90, %f655, %f656;
  1473. ld.global.f32 %f91, [%rd301+16];
  1474. setp.gt.f32 %p79, %f90, %f91;
  1475. mov.f32 %f1270, %f1269;
  1476. @%p79 bra BB1_59;
  1477.  
  1478. sub.f32 %f657, %f91, %f90;
  1479. sqrt.rn.f32 %f658, %f657;
  1480. sub.f32 %f1270, %f89, %f658;
  1481. add.f32 %f1269, %f89, %f658;
  1482. mov.pred %p311, 0;
  1483.  
  1484. BB1_59:
  1485. @%p311 bra BB1_62;
  1486.  
  1487. setp.lt.f32 %p81, %f1270, 0f00000000;
  1488. selp.f32 %f96, %f1269, %f1270, %p81;
  1489. setp.geu.f32 %p82, %f96, %f1252;
  1490. @%p82 bra BB1_62;
  1491.  
  1492. add.s32 %r201, %r204, 3;
  1493. cvt.u64.u32 %rd304, %r201;
  1494. mul.lo.s64 %rd185, %rd304, 80;
  1495. add.s64 %rd309, %rd1, %rd185;
  1496. add.s64 %rd310, %rd153, %rd185;
  1497. mov.f32 %f1252, %f96;
  1498.  
  1499. BB1_62:
  1500. cvt.u64.u32 %rd292, %r204;
  1501. cvt.u32.u64 %r60, %rd292;
  1502. add.s32 %r204, %r60, 4;
  1503. setp.lt.u32 %p83, %r204, %r41;
  1504. @%p83 bra BB1_34;
  1505.  
  1506. BB1_63:
  1507. setp.eq.s64 %p84, %rd310, 0;
  1508. mov.f32 %f1412, 0f40000000;
  1509. mov.f32 %f1413, %f1412;
  1510. mov.f32 %f1414, %f1412;
  1511. @%p84 bra BB1_316;
  1512.  
  1513. fma.rn.f32 %f99, %f1246, %f1252, 0f00000000;
  1514. fma.rn.f32 %f100, %f1245, %f1252, 0f00000000;
  1515. fma.rn.f32 %f101, %f1244, %f1252, 0f00000000;
  1516. add.s64 %rd47, %rd309, 8;
  1517. ld.global.f32 %f662, [%rd309+8];
  1518. sub.f32 %f1273, %f99, %f662;
  1519. ld.global.f32 %f663, [%rd309+12];
  1520. sub.f32 %f1274, %f100, %f663;
  1521. ld.global.f32 %f664, [%rd309+16];
  1522. sub.f32 %f1275, %f101, %f664;
  1523. mul.f32 %f665, %f1274, %f1274;
  1524. fma.rn.f32 %f666, %f1273, %f1273, %f665;
  1525. fma.rn.f32 %f105, %f1275, %f1275, %f666;
  1526. setp.leu.f32 %p85, %f105, 0f00000000;
  1527. @%p85 bra BB1_66;
  1528.  
  1529. sqrt.rn.f32 %f667, %f105;
  1530. rcp.rn.f32 %f668, %f667;
  1531. mul.f32 %f1273, %f1273, %f668;
  1532. mul.f32 %f1274, %f1274, %f668;
  1533. mul.f32 %f1275, %f1275, %f668;
  1534.  
  1535. BB1_66:
  1536. mul.f32 %f669, %f1245, %f1274;
  1537. fma.rn.f32 %f670, %f1246, %f1273, %f669;
  1538. fma.rn.f32 %f671, %f1244, %f1275, %f670;
  1539. mov.u16 %rs10, 0;
  1540. setp.leu.f32 %p86, %f671, 0f00000000;
  1541. @%p86 bra BB1_68;
  1542.  
  1543. neg.f32 %f1273, %f1273;
  1544. neg.f32 %f1274, %f1274;
  1545. neg.f32 %f1275, %f1275;
  1546. mov.u16 %rs10, 1;
  1547.  
  1548. BB1_68:
  1549. ld.global.f32 %f118, [%rd47+44];
  1550. setp.gt.f32 %p87, %f118, 0f00000000;
  1551. @%p87 bra BB1_82;
  1552.  
  1553. ld.global.f32 %f672, [%rd47+48];
  1554. setp.gt.f32 %p88, %f672, 0f00000000;
  1555. @%p88 bra BB1_82;
  1556. bra.uni BB1_70;
  1557.  
  1558. BB1_82:
  1559. mul.f32 %f715, %f1245, %f1274;
  1560. fma.rn.f32 %f716, %f1246, %f1273, %f715;
  1561. fma.rn.f32 %f147, %f1244, %f1275, %f716;
  1562. add.f32 %f148, %f147, 0f3F800000;
  1563. cvt.f64.f32 %fd1, %f148;
  1564. {
  1565. .reg .b32 %temp;
  1566. mov.b64 {%temp, %r14}, %fd1;
  1567. }
  1568. mov.f64 %fd54, 0d4008000000000000;
  1569. {
  1570. .reg .b32 %temp;
  1571. mov.b64 {%temp, %r15}, %fd54;
  1572. }
  1573. bfe.u32 %r64, %r15, 20, 11;
  1574. add.s32 %r65, %r64, -1012;
  1575. mov.u64 %rd190, 4613937818241073152;
  1576. shl.b64 %rd51, %rd190, %r65;
  1577. setp.ne.s64 %p97, %rd51, -9223372036854775808;
  1578. setp.eq.s64 %p98, %rd51, -9223372036854775808;
  1579. abs.f64 %fd2, %fd1;
  1580. // Callseq Start 3
  1581. {
  1582. .reg .b32 temp_param_reg;
  1583. // <end>}
  1584. .param .b64 param0;
  1585. st.param.f64 [param0+0], %fd2;
  1586. .param .b64 retval0;
  1587. call.uni (retval0),
  1588. __internal_accurate_pow,
  1589. (
  1590. param0
  1591. );
  1592. ld.param.f64 %fd71, [retval0+0];
  1593.  
  1594. //{
  1595. }// Callseq End 3
  1596. setp.gt.s32 %p99, %r14, -1;
  1597. setp.lt.s32 %p100, %r14, 0;
  1598. and.pred %p8, %p100, %p98;
  1599. or.pred %p101, %p99, %p97;
  1600. @%p101 bra BB1_84;
  1601.  
  1602. {
  1603. .reg .b32 %temp;
  1604. mov.b64 {%temp, %r66}, %fd71;
  1605. }
  1606. xor.b32 %r67, %r66, -2147483648;
  1607. {
  1608. .reg .b32 %temp;
  1609. mov.b64 {%r68, %temp}, %fd71;
  1610. }
  1611. mov.b64 %fd71, {%r68, %r67};
  1612.  
  1613. BB1_84:
  1614. setp.eq.f32 %p102, %f148, 0f00000000;
  1615. @%p102 bra BB1_87;
  1616. bra.uni BB1_85;
  1617.  
  1618. BB1_87:
  1619. selp.b32 %r69, %r14, 0, %p98;
  1620. mov.u32 %r70, 0;
  1621. or.b32 %r71, %r69, 2146435072;
  1622. setp.lt.s32 %p106, %r15, 0;
  1623. selp.b32 %r72, %r71, %r69, %p106;
  1624. mov.b64 %fd71, {%r70, %r72};
  1625. bra.uni BB1_88;
  1626.  
  1627. BB1_85:
  1628. @%p99 bra BB1_88;
  1629.  
  1630. mov.f64 %fd69, 0d4008000000000000;
  1631. cvt.rzi.f64.f64 %fd56, %fd69;
  1632. setp.neu.f64 %p104, %fd56, 0d4008000000000000;
  1633. selp.f64 %fd71, 0dFFF8000000000000, %fd71, %p104;
  1634.  
  1635. BB1_88:
  1636. add.f64 %fd72, %fd1, 0d4008000000000000;
  1637. {
  1638. .reg .b32 %temp;
  1639. mov.b64 {%temp, %r73}, %fd72;
  1640. }
  1641. and.b32 %r74, %r73, 2146435072;
  1642. setp.ne.s32 %p107, %r74, 2146435072;
  1643. @%p107 bra BB1_89;
  1644.  
  1645. setp.gtu.f64 %p108, %fd2, 0d7FF0000000000000;
  1646. @%p108 bra BB1_98;
  1647.  
  1648. and.b32 %r75, %r15, 2147483647;
  1649. setp.ne.s32 %p109, %r75, 2146435072;
  1650. @%p109 bra BB1_93;
  1651.  
  1652. mov.f64 %fd68, 0d4008000000000000;
  1653. {
  1654. .reg .b32 %temp;
  1655. mov.b64 {%r76, %temp}, %fd68;
  1656. }
  1657. setp.eq.s32 %p110, %r76, 0;
  1658. @%p110 bra BB1_97;
  1659.  
  1660. BB1_93:
  1661. and.b32 %r77, %r14, 2147483647;
  1662. setp.ne.s32 %p111, %r77, 2146435072;
  1663. @%p111 bra BB1_94;
  1664.  
  1665. {
  1666. .reg .b32 %temp;
  1667. mov.b64 {%r78, %temp}, %fd1;
  1668. }
  1669. setp.ne.s32 %p112, %r78, 0;
  1670. mov.f64 %fd72, %fd71;
  1671. @%p112 bra BB1_98;
  1672.  
  1673. shr.s32 %r79, %r15, 31;
  1674. and.b32 %r80, %r79, -2146435072;
  1675. add.s32 %r81, %r80, 2146435072;
  1676. or.b32 %r82, %r81, -2147483648;
  1677. selp.b32 %r83, %r82, %r81, %p8;
  1678. mov.u32 %r84, 0;
  1679. mov.b64 %fd72, {%r84, %r83};
  1680. bra.uni BB1_98;
  1681.  
  1682. BB1_89:
  1683. mov.f64 %fd72, %fd71;
  1684. bra.uni BB1_98;
  1685.  
  1686. BB1_70:
  1687. mov.f32 %f1409, 0f00000000;
  1688. mov.f32 %f1410, %f1409;
  1689. mov.f32 %f1411, %f1409;
  1690. @%p29 bra BB1_315;
  1691.  
  1692. fma.rn.f32 %f119, %f1273, 0f38D1B717, %f99;
  1693. fma.rn.f32 %f120, %f1274, 0f38D1B717, %f100;
  1694. fma.rn.f32 %f121, %f1275, 0f38D1B717, %f101;
  1695. mov.f32 %f1411, 0f00000000;
  1696. mov.u32 %r206, 0;
  1697. mov.f32 %f1410, %f1411;
  1698. mov.f32 %f1409, %f1411;
  1699.  
  1700. BB1_72:
  1701. cvt.u64.u32 %rd48, %r206;
  1702. mul.wide.u32 %rd186, %r206, 80;
  1703. add.s64 %rd187, %rd1, %rd186;
  1704. add.s64 %rd49, %rd187, 40;
  1705. ld.global.f32 %f125, [%rd187+40];
  1706. setp.leu.f32 %p90, %f125, 0f00000000;
  1707. @%p90 bra BB1_81;
  1708.  
  1709. ld.global.f32 %f679, [%rd49+-32];
  1710. sub.f32 %f1284, %f679, %f99;
  1711. ld.global.f32 %f680, [%rd49+-28];
  1712. sub.f32 %f1283, %f680, %f100;
  1713. ld.global.f32 %f681, [%rd49+-24];
  1714. sub.f32 %f1282, %f681, %f101;
  1715. mul.f32 %f682, %f1283, %f1283;
  1716. fma.rn.f32 %f683, %f1284, %f1284, %f682;
  1717. fma.rn.f32 %f129, %f1282, %f1282, %f683;
  1718. setp.leu.f32 %p91, %f129, 0f00000000;
  1719. @%p91 bra BB1_75;
  1720.  
  1721. sqrt.rn.f32 %f684, %f129;
  1722. rcp.rn.f32 %f685, %f684;
  1723. mul.f32 %f1284, %f1284, %f685;
  1724. mul.f32 %f1283, %f1283, %f685;
  1725. mul.f32 %f1282, %f1282, %f685;
  1726.  
  1727. BB1_75:
  1728. mov.u32 %r207, 0;
  1729.  
  1730. BB1_76:
  1731. setp.eq.s32 %p92, %r206, %r207;
  1732. @%p92 bra BB1_79;
  1733.  
  1734. mul.wide.u32 %rd188, %r207, 80;
  1735. add.s64 %rd189, %rd1, %rd188;
  1736. add.s64 %rd50, %rd189, 8;
  1737. ld.global.f32 %f686, [%rd189+8];
  1738. sub.f32 %f136, %f686, %f119;
  1739. ld.global.f32 %f687, [%rd189+12];
  1740. sub.f32 %f137, %f687, %f120;
  1741. ld.global.f32 %f688, [%rd189+16];
  1742. sub.f32 %f138, %f688, %f121;
  1743. mul.f32 %f689, %f1283, %f137;
  1744. fma.rn.f32 %f690, %f1284, %f136, %f689;
  1745. fma.rn.f32 %f139, %f1282, %f138, %f690;
  1746. setp.lt.f32 %p93, %f139, 0f00000000;
  1747. @%p93 bra BB1_79;
  1748.  
  1749. mul.f32 %f692, %f137, %f137;
  1750. fma.rn.f32 %f693, %f136, %f136, %f692;
  1751. fma.rn.f32 %f694, %f138, %f138, %f693;
  1752. mul.f32 %f695, %f139, %f139;
  1753. sub.f32 %f696, %f694, %f695;
  1754. ld.global.f32 %f697, [%rd50+16];
  1755. mov.f32 %f1285, 0f00000000;
  1756. setp.leu.f32 %p94, %f696, %f697;
  1757. @%p94 bra BB1_80;
  1758.  
  1759. BB1_79:
  1760. add.s32 %r207, %r207, 1;
  1761. setp.lt.u32 %p95, %r207, %r41;
  1762. mov.f32 %f1285, 0f3F800000;
  1763. @%p95 bra BB1_76;
  1764.  
  1765. BB1_80:
  1766. ld.global.f32 %f699, [%rd47+20];
  1767. mul.f32 %f700, %f1285, %f699;
  1768. ld.global.f32 %f701, [%rd47+24];
  1769. mul.f32 %f702, %f1285, %f701;
  1770. ld.global.f32 %f703, [%rd47+28];
  1771. mul.f32 %f704, %f1285, %f703;
  1772. mul.f32 %f705, %f1273, %f1284;
  1773. fma.rn.f32 %f706, %f1274, %f1283, %f705;
  1774. fma.rn.f32 %f707, %f1275, %f1282, %f706;
  1775. mov.f32 %f708, 0f00000000;
  1776. max.f32 %f709, %f708, %f707;
  1777. mul.f32 %f710, %f700, %f709;
  1778. mul.f32 %f711, %f702, %f709;
  1779. mul.f32 %f712, %f704, %f709;
  1780. ld.global.f32 %f713, [%rd49+4];
  1781. ld.global.f32 %f714, [%rd49+8];
  1782. fma.rn.f32 %f1409, %f710, %f125, %f1409;
  1783. fma.rn.f32 %f1410, %f711, %f713, %f1410;
  1784. fma.rn.f32 %f1411, %f712, %f714, %f1411;
  1785.  
  1786. BB1_81:
  1787. cvt.u32.u64 %r63, %rd48;
  1788. add.s32 %r206, %r63, 1;
  1789. setp.lt.u32 %p96, %r206, %r41;
  1790. @%p96 bra BB1_72;
  1791. bra.uni BB1_315;
  1792.  
  1793. BB1_94:
  1794. mov.f64 %fd72, %fd71;
  1795. bra.uni BB1_98;
  1796.  
  1797. BB1_97:
  1798. setp.gt.f64 %p113, %fd2, 0d3FF0000000000000;
  1799. selp.b32 %r85, 2146435072, 0, %p113;
  1800. mov.u32 %r86, 0;
  1801. xor.b32 %r87, %r85, 2146435072;
  1802. setp.lt.s32 %p114, %r15, 0;
  1803. selp.b32 %r88, %r87, %r85, %p114;
  1804. setp.eq.f32 %p115, %f148, 0fBF800000;
  1805. selp.b32 %r89, 1072693248, %r88, %p115;
  1806. mov.b64 %fd72, {%r86, %r89};
  1807.  
  1808. BB1_98:
  1809. cvt.rn.f32.f64 %f717, %fd72;
  1810. fma.rn.f32 %f718, %f717, 0f3F666666, 0f3DCCCCCD;
  1811. setp.eq.f32 %p116, %f148, 0f3F800000;
  1812. selp.f32 %f149, 0f3F800000, %f718, %p116;
  1813. add.f32 %f719, %f1273, %f1273;
  1814. mul.f32 %f720, %f719, %f147;
  1815. add.f32 %f721, %f1274, %f1274;
  1816. mul.f32 %f722, %f721, %f147;
  1817. add.f32 %f723, %f1275, %f1275;
  1818. mul.f32 %f724, %f723, %f147;
  1819. sub.f32 %f1289, %f1246, %f720;
  1820. sub.f32 %f1290, %f1245, %f722;
  1821. sub.f32 %f1291, %f1244, %f724;
  1822. mul.f32 %f725, %f1290, %f1290;
  1823. fma.rn.f32 %f726, %f1289, %f1289, %f725;
  1824. fma.rn.f32 %f153, %f1291, %f1291, %f726;
  1825. setp.leu.f32 %p117, %f153, 0f00000000;
  1826. @%p117 bra BB1_100;
  1827.  
  1828. sqrt.rn.f32 %f727, %f153;
  1829. rcp.rn.f32 %f728, %f727;
  1830. mul.f32 %f1289, %f1289, %f728;
  1831. mul.f32 %f1290, %f1290, %f728;
  1832. mul.f32 %f1291, %f1291, %f728;
  1833.  
  1834. BB1_100:
  1835. fma.rn.f32 %f1239, %f1244, %f1252, 0f00000000;
  1836. fma.rn.f32 %f1238, %f1245, %f1252, 0f00000000;
  1837. fma.rn.f32 %f1237, %f1246, %f1252, 0f00000000;
  1838. mul.f32 %f160, %f1273, 0f38D1B717;
  1839. add.f32 %f161, %f1237, %f160;
  1840. mul.f32 %f162, %f1274, 0f38D1B717;
  1841. add.f32 %f163, %f1238, %f162;
  1842. mul.f32 %f164, %f1275, 0f38D1B717;
  1843. add.f32 %f165, %f1239, %f164;
  1844. mov.f32 %f1316, 0f7F800000;
  1845. mov.u64 %rd347, 0;
  1846. mov.u64 %rd332, %rd347;
  1847. @%p29 bra BB1_162;
  1848.  
  1849. and.b32 %r93, %r41, 3;
  1850. mov.f32 %f1316, 0f7F800000;
  1851. mov.u64 %rd347, 0;
  1852. mov.u32 %r210, 0;
  1853. mov.f32 %f1303, 0f00000000;
  1854. setp.eq.s32 %p119, %r93, 0;
  1855. @%p119 bra BB1_102;
  1856.  
  1857. setp.eq.s32 %p120, %r93, 1;
  1858. @%p120 bra BB1_104;
  1859. bra.uni BB1_105;
  1860.  
  1861. BB1_104:
  1862. mov.u64 %rd332, %rd347;
  1863. bra.uni BB1_123;
  1864.  
  1865. BB1_102:
  1866. mov.u64 %rd332, %rd347;
  1867. bra.uni BB1_131;
  1868.  
  1869. BB1_105:
  1870. setp.eq.s32 %p121, %r93, 2;
  1871. mov.u64 %rd332, %rd347;
  1872. @%p121 bra BB1_115;
  1873.  
  1874. ld.global.f32 %f736, [%rd1+8];
  1875. sub.f32 %f166, %f736, %f161;
  1876. ld.global.f32 %f737, [%rd1+12];
  1877. sub.f32 %f167, %f737, %f163;
  1878. ld.global.f32 %f738, [%rd1+16];
  1879. sub.f32 %f168, %f738, %f165;
  1880. mul.f32 %f739, %f1290, %f167;
  1881. fma.rn.f32 %f740, %f1289, %f166, %f739;
  1882. fma.rn.f32 %f169, %f1291, %f168, %f740;
  1883. setp.lt.f32 %p123, %f169, 0f00000000;
  1884. mov.pred %p312, -1;
  1885. mov.f32 %f1292, 0f7F800000;
  1886. @%p123 bra BB1_107;
  1887. bra.uni BB1_108;
  1888.  
  1889. BB1_107:
  1890. mov.f32 %f1293, %f1292;
  1891. bra.uni BB1_110;
  1892.  
  1893. BB1_108:
  1894. mov.f32 %f1292, 0f7F800000;
  1895. mul.f32 %f743, %f167, %f167;
  1896. fma.rn.f32 %f744, %f166, %f166, %f743;
  1897. fma.rn.f32 %f745, %f168, %f168, %f744;
  1898. mul.f32 %f746, %f169, %f169;
  1899. sub.f32 %f170, %f745, %f746;
  1900. ld.global.f32 %f171, [%rd1+24];
  1901. setp.gt.f32 %p125, %f170, %f171;
  1902. mov.f32 %f1293, %f1292;
  1903. @%p125 bra BB1_110;
  1904.  
  1905. sub.f32 %f747, %f171, %f170;
  1906. sqrt.rn.f32 %f748, %f747;
  1907. sub.f32 %f1293, %f169, %f748;
  1908. add.f32 %f1292, %f169, %f748;
  1909. mov.pred %p312, 0;
  1910.  
  1911. BB1_110:
  1912. mov.u64 %rd347, 0;
  1913. mov.f32 %f1316, 0f7F800000;
  1914. mov.u32 %r210, 1;
  1915. @%p312 bra BB1_111;
  1916. bra.uni BB1_112;
  1917.  
  1918. BB1_111:
  1919. mov.u64 %rd332, %rd347;
  1920. bra.uni BB1_115;
  1921.  
  1922. BB1_112:
  1923. mov.u64 %rd347, 0;
  1924. mov.f32 %f1242, 0f7F800000;
  1925. setp.lt.f32 %p127, %f1293, 0f00000000;
  1926. selp.f32 %f1316, %f1292, %f1293, %p127;
  1927. setp.geu.f32 %p128, %f1316, 0f7F800000;
  1928. @%p128 bra BB1_113;
  1929.  
  1930. mov.u64 %rd347, %rd1;
  1931. mov.u64 %rd332, %rd153;
  1932. bra.uni BB1_115;
  1933.  
  1934. BB1_113:
  1935. mov.u64 %rd332, %rd347;
  1936. mov.f32 %f1316, %f1242;
  1937.  
  1938. BB1_115:
  1939. mov.f32 %f1295, 0f7F800000;
  1940. mul.wide.u32 %rd203, %r210, 80;
  1941. add.s64 %rd204, %rd1, %rd203;
  1942. ld.global.f32 %f753, [%rd204+8];
  1943. sub.f32 %f178, %f753, %f161;
  1944. ld.global.f32 %f754, [%rd204+12];
  1945. sub.f32 %f179, %f754, %f163;
  1946. ld.global.f32 %f755, [%rd204+16];
  1947. sub.f32 %f180, %f755, %f165;
  1948. mul.f32 %f756, %f1290, %f179;
  1949. fma.rn.f32 %f757, %f1289, %f178, %f756;
  1950. fma.rn.f32 %f181, %f1291, %f180, %f757;
  1951. setp.lt.f32 %p130, %f181, 0f00000000;
  1952. mov.pred %p313, -1;
  1953. @%p130 bra BB1_116;
  1954. bra.uni BB1_117;
  1955.  
  1956. BB1_116:
  1957. mov.f32 %f1296, %f1295;
  1958. bra.uni BB1_119;
  1959.  
  1960. BB1_117:
  1961. mul.wide.u32 %rd278, %r210, 80;
  1962. add.s64 %rd277, %rd1, %rd278;
  1963. add.s64 %rd276, %rd277, 8;
  1964. mov.f32 %f1295, 0f7F800000;
  1965. mul.f32 %f760, %f179, %f179;
  1966. fma.rn.f32 %f761, %f178, %f178, %f760;
  1967. fma.rn.f32 %f762, %f180, %f180, %f761;
  1968. mul.f32 %f763, %f181, %f181;
  1969. sub.f32 %f182, %f762, %f763;
  1970. ld.global.f32 %f183, [%rd276+16];
  1971. setp.gt.f32 %p132, %f182, %f183;
  1972. mov.f32 %f1296, %f1295;
  1973. @%p132 bra BB1_119;
  1974.  
  1975. sub.f32 %f764, %f183, %f182;
  1976. sqrt.rn.f32 %f765, %f764;
  1977. sub.f32 %f1296, %f181, %f765;
  1978. add.f32 %f1295, %f181, %f765;
  1979. mov.pred %p313, 0;
  1980.  
  1981. BB1_119:
  1982. @%p313 bra BB1_122;
  1983.  
  1984. setp.lt.f32 %p134, %f1296, 0f00000000;
  1985. selp.f32 %f188, %f1295, %f1296, %p134;
  1986. setp.geu.f32 %p135, %f188, %f1316;
  1987. @%p135 bra BB1_122;
  1988.  
  1989. cvt.u64.u32 %rd279, %r210;
  1990. mul.lo.s64 %rd205, %rd279, 80;
  1991. add.s64 %rd347, %rd1, %rd205;
  1992. add.s64 %rd332, %rd153, %rd205;
  1993. mov.f32 %f1316, %f188;
  1994.  
  1995. BB1_122:
  1996. cvt.u64.u32 %rd280, %r210;
  1997. cvt.u32.u64 %r97, %rd280;
  1998. add.s32 %r210, %r97, 1;
  1999.  
  2000. BB1_123:
  2001. mov.f32 %f1299, 0f7F800000;
  2002. mul.wide.u32 %rd206, %r210, 80;
  2003. add.s64 %rd207, %rd1, %rd206;
  2004. ld.global.f32 %f768, [%rd207+8];
  2005. sub.f32 %f191, %f768, %f161;
  2006. ld.global.f32 %f769, [%rd207+12];
  2007. sub.f32 %f192, %f769, %f163;
  2008. ld.global.f32 %f770, [%rd207+16];
  2009. sub.f32 %f193, %f770, %f165;
  2010. mul.f32 %f771, %f1290, %f192;
  2011. fma.rn.f32 %f772, %f1289, %f191, %f771;
  2012. fma.rn.f32 %f194, %f1291, %f193, %f772;
  2013. setp.lt.f32 %p137, %f194, 0f00000000;
  2014. mov.pred %p314, -1;
  2015. @%p137 bra BB1_124;
  2016. bra.uni BB1_125;
  2017.  
  2018. BB1_124:
  2019. mov.f32 %f1300, %f1299;
  2020. bra.uni BB1_127;
  2021.  
  2022. BB1_125:
  2023. mul.wide.u32 %rd283, %r210, 80;
  2024. add.s64 %rd282, %rd1, %rd283;
  2025. add.s64 %rd281, %rd282, 8;
  2026. mov.f32 %f1299, 0f7F800000;
  2027. mul.f32 %f775, %f192, %f192;
  2028. fma.rn.f32 %f776, %f191, %f191, %f775;
  2029. fma.rn.f32 %f777, %f193, %f193, %f776;
  2030. mul.f32 %f778, %f194, %f194;
  2031. sub.f32 %f195, %f777, %f778;
  2032. ld.global.f32 %f196, [%rd281+16];
  2033. setp.gt.f32 %p139, %f195, %f196;
  2034. mov.f32 %f1300, %f1299;
  2035. @%p139 bra BB1_127;
  2036.  
  2037. sub.f32 %f779, %f196, %f195;
  2038. sqrt.rn.f32 %f780, %f779;
  2039. sub.f32 %f1300, %f194, %f780;
  2040. add.f32 %f1299, %f194, %f780;
  2041. mov.pred %p314, 0;
  2042.  
  2043. BB1_127:
  2044. @%p314 bra BB1_130;
  2045.  
  2046. setp.lt.f32 %p141, %f1300, 0f00000000;
  2047. selp.f32 %f201, %f1299, %f1300, %p141;
  2048. setp.geu.f32 %p142, %f201, %f1316;
  2049. @%p142 bra BB1_130;
  2050.  
  2051. cvt.u64.u32 %rd284, %r210;
  2052. mul.lo.s64 %rd208, %rd284, 80;
  2053. add.s64 %rd347, %rd1, %rd208;
  2054. add.s64 %rd332, %rd153, %rd208;
  2055. mov.f32 %f1316, %f201;
  2056.  
  2057. BB1_130:
  2058. cvt.u64.u32 %rd285, %r210;
  2059. cvt.u32.u64 %r98, %rd285;
  2060. add.s32 %r210, %r98, 1;
  2061. mov.f32 %f1303, %f1316;
  2062.  
  2063. BB1_131:
  2064. setp.lt.u32 %p143, %r41, 4;
  2065. @%p143 bra BB1_132;
  2066. bra.uni BB1_133;
  2067.  
  2068. BB1_132:
  2069. mov.f32 %f1316, %f1303;
  2070. bra.uni BB1_162;
  2071.  
  2072. BB1_133:
  2073. cvt.u64.u32 %rd72, %r210;
  2074. mul.wide.u32 %rd209, %r210, 80;
  2075. add.s64 %rd210, %rd1, %rd209;
  2076. add.s64 %rd73, %rd210, 8;
  2077. ld.global.f32 %f783, [%rd210+8];
  2078. sub.f32 %f206, %f783, %f161;
  2079. ld.global.f32 %f784, [%rd210+12];
  2080. sub.f32 %f207, %f784, %f163;
  2081. ld.global.f32 %f785, [%rd210+16];
  2082. sub.f32 %f208, %f785, %f165;
  2083. mul.f32 %f786, %f1290, %f207;
  2084. fma.rn.f32 %f787, %f1289, %f206, %f786;
  2085. fma.rn.f32 %f209, %f1291, %f208, %f787;
  2086. setp.lt.f32 %p145, %f209, 0f00000000;
  2087. mov.pred %p315, -1;
  2088. mov.f32 %f1305, 0f7F800000;
  2089. @%p145 bra BB1_134;
  2090. bra.uni BB1_135;
  2091.  
  2092. BB1_134:
  2093. mov.f32 %f1306, %f1305;
  2094. bra.uni BB1_137;
  2095.  
  2096. BB1_135:
  2097. mul.f32 %f790, %f207, %f207;
  2098. fma.rn.f32 %f791, %f206, %f206, %f790;
  2099. fma.rn.f32 %f792, %f208, %f208, %f791;
  2100. mul.f32 %f793, %f209, %f209;
  2101. sub.f32 %f210, %f792, %f793;
  2102. ld.global.f32 %f211, [%rd73+16];
  2103. setp.gt.f32 %p147, %f210, %f211;
  2104. mov.f32 %f1306, %f1305;
  2105. @%p147 bra BB1_137;
  2106.  
  2107. sub.f32 %f794, %f211, %f210;
  2108. sqrt.rn.f32 %f795, %f794;
  2109. sub.f32 %f1306, %f209, %f795;
  2110. add.f32 %f1305, %f209, %f795;
  2111. mov.pred %p315, 0;
  2112.  
  2113. BB1_137:
  2114. @%p315 bra BB1_140;
  2115.  
  2116. setp.lt.f32 %p149, %f1306, 0f00000000;
  2117. selp.f32 %f216, %f1305, %f1306, %p149;
  2118. setp.geu.f32 %p150, %f216, %f1316;
  2119. @%p150 bra BB1_140;
  2120.  
  2121. mul.lo.s64 %rd211, %rd72, 80;
  2122. add.s64 %rd347, %rd1, %rd211;
  2123. add.s64 %rd332, %rd153, %rd211;
  2124. mov.f32 %f1316, %f216;
  2125.  
  2126. BB1_140:
  2127. add.s32 %r99, %r210, 1;
  2128. cvt.u64.u32 %rd78, %r99;
  2129. mul.wide.u32 %rd212, %r99, 80;
  2130. add.s64 %rd213, %rd1, %rd212;
  2131. add.s64 %rd79, %rd213, 8;
  2132. ld.global.f32 %f798, [%rd213+8];
  2133. sub.f32 %f218, %f798, %f161;
  2134. ld.global.f32 %f799, [%rd213+12];
  2135. sub.f32 %f219, %f799, %f163;
  2136. ld.global.f32 %f800, [%rd213+16];
  2137. sub.f32 %f220, %f800, %f165;
  2138. mul.f32 %f801, %f1290, %f219;
  2139. fma.rn.f32 %f802, %f1289, %f218, %f801;
  2140. fma.rn.f32 %f221, %f1291, %f220, %f802;
  2141. setp.lt.f32 %p152, %f221, 0f00000000;
  2142. mov.pred %p316, -1;
  2143. mov.f32 %f1308, 0f7F800000;
  2144. @%p152 bra BB1_141;
  2145. bra.uni BB1_142;
  2146.  
  2147. BB1_141:
  2148. mov.f32 %f1309, %f1308;
  2149. bra.uni BB1_144;
  2150.  
  2151. BB1_142:
  2152. mul.f32 %f805, %f219, %f219;
  2153. fma.rn.f32 %f806, %f218, %f218, %f805;
  2154. fma.rn.f32 %f807, %f220, %f220, %f806;
  2155. mul.f32 %f808, %f221, %f221;
  2156. sub.f32 %f222, %f807, %f808;
  2157. ld.global.f32 %f223, [%rd79+16];
  2158. setp.gt.f32 %p154, %f222, %f223;
  2159. mov.f32 %f1309, %f1308;
  2160. @%p154 bra BB1_144;
  2161.  
  2162. sub.f32 %f809, %f223, %f222;
  2163. sqrt.rn.f32 %f810, %f809;
  2164. sub.f32 %f1309, %f221, %f810;
  2165. add.f32 %f1308, %f221, %f810;
  2166. mov.pred %p316, 0;
  2167.  
  2168. BB1_144:
  2169. @%p316 bra BB1_147;
  2170.  
  2171. setp.lt.f32 %p156, %f1309, 0f00000000;
  2172. selp.f32 %f228, %f1308, %f1309, %p156;
  2173. setp.geu.f32 %p157, %f228, %f1316;
  2174. @%p157 bra BB1_147;
  2175.  
  2176. mul.lo.s64 %rd214, %rd78, 80;
  2177. add.s64 %rd347, %rd1, %rd214;
  2178. add.s64 %rd332, %rd153, %rd214;
  2179. mov.f32 %f1316, %f228;
  2180.  
  2181. BB1_147:
  2182. add.s32 %r100, %r210, 2;
  2183. cvt.u64.u32 %rd84, %r100;
  2184. mul.wide.u32 %rd215, %r100, 80;
  2185. add.s64 %rd216, %rd1, %rd215;
  2186. add.s64 %rd85, %rd216, 8;
  2187. ld.global.f32 %f813, [%rd216+8];
  2188. sub.f32 %f230, %f813, %f161;
  2189. ld.global.f32 %f814, [%rd216+12];
  2190. sub.f32 %f231, %f814, %f163;
  2191. ld.global.f32 %f815, [%rd216+16];
  2192. sub.f32 %f232, %f815, %f165;
  2193. mul.f32 %f816, %f1290, %f231;
  2194. fma.rn.f32 %f817, %f1289, %f230, %f816;
  2195. fma.rn.f32 %f233, %f1291, %f232, %f817;
  2196. setp.lt.f32 %p159, %f233, 0f00000000;
  2197. mov.pred %p317, -1;
  2198. mov.f32 %f1311, 0f7F800000;
  2199. @%p159 bra BB1_148;
  2200. bra.uni BB1_149;
  2201.  
  2202. BB1_148:
  2203. mov.f32 %f1312, %f1311;
  2204. bra.uni BB1_151;
  2205.  
  2206. BB1_149:
  2207. mul.f32 %f820, %f231, %f231;
  2208. fma.rn.f32 %f821, %f230, %f230, %f820;
  2209. fma.rn.f32 %f822, %f232, %f232, %f821;
  2210. mul.f32 %f823, %f233, %f233;
  2211. sub.f32 %f234, %f822, %f823;
  2212. ld.global.f32 %f235, [%rd85+16];
  2213. setp.gt.f32 %p161, %f234, %f235;
  2214. mov.f32 %f1312, %f1311;
  2215. @%p161 bra BB1_151;
  2216.  
  2217. sub.f32 %f824, %f235, %f234;
  2218. sqrt.rn.f32 %f825, %f824;
  2219. sub.f32 %f1312, %f233, %f825;
  2220. add.f32 %f1311, %f233, %f825;
  2221. mov.pred %p317, 0;
  2222.  
  2223. BB1_151:
  2224. @%p317 bra BB1_154;
  2225.  
  2226. setp.lt.f32 %p163, %f1312, 0f00000000;
  2227. selp.f32 %f240, %f1311, %f1312, %p163;
  2228. setp.geu.f32 %p164, %f240, %f1316;
  2229. @%p164 bra BB1_154;
  2230.  
  2231. mul.lo.s64 %rd217, %rd84, 80;
  2232. add.s64 %rd347, %rd1, %rd217;
  2233. add.s64 %rd332, %rd153, %rd217;
  2234. mov.f32 %f1316, %f240;
  2235.  
  2236. BB1_154:
  2237. add.s32 %r101, %r210, 3;
  2238. cvt.u64.u32 %rd90, %r101;
  2239. mul.wide.u32 %rd218, %r101, 80;
  2240. add.s64 %rd219, %rd1, %rd218;
  2241. add.s64 %rd91, %rd219, 8;
  2242. ld.global.f32 %f828, [%rd219+8];
  2243. sub.f32 %f242, %f828, %f161;
  2244. ld.global.f32 %f829, [%rd219+12];
  2245. sub.f32 %f243, %f829, %f163;
  2246. ld.global.f32 %f830, [%rd219+16];
  2247. sub.f32 %f244, %f830, %f165;
  2248. mul.f32 %f831, %f1290, %f243;
  2249. fma.rn.f32 %f832, %f1289, %f242, %f831;
  2250. fma.rn.f32 %f245, %f1291, %f244, %f832;
  2251. setp.lt.f32 %p166, %f245, 0f00000000;
  2252. mov.pred %p318, -1;
  2253. mov.f32 %f1314, 0f7F800000;
  2254. @%p166 bra BB1_155;
  2255. bra.uni BB1_156;
  2256.  
  2257. BB1_155:
  2258. mov.f32 %f1315, %f1314;
  2259. bra.uni BB1_158;
  2260.  
  2261. BB1_156:
  2262. mul.f32 %f835, %f243, %f243;
  2263. fma.rn.f32 %f836, %f242, %f242, %f835;
  2264. fma.rn.f32 %f837, %f244, %f244, %f836;
  2265. mul.f32 %f838, %f245, %f245;
  2266. sub.f32 %f246, %f837, %f838;
  2267. ld.global.f32 %f247, [%rd91+16];
  2268. setp.gt.f32 %p168, %f246, %f247;
  2269. mov.f32 %f1315, %f1314;
  2270. @%p168 bra BB1_158;
  2271.  
  2272. sub.f32 %f839, %f247, %f246;
  2273. sqrt.rn.f32 %f840, %f839;
  2274. sub.f32 %f1315, %f245, %f840;
  2275. add.f32 %f1314, %f245, %f840;
  2276. mov.pred %p318, 0;
  2277.  
  2278. BB1_158:
  2279. @%p318 bra BB1_161;
  2280.  
  2281. setp.lt.f32 %p170, %f1315, 0f00000000;
  2282. selp.f32 %f252, %f1314, %f1315, %p170;
  2283. setp.geu.f32 %p171, %f252, %f1316;
  2284. @%p171 bra BB1_161;
  2285.  
  2286. mul.lo.s64 %rd220, %rd90, 80;
  2287. add.s64 %rd347, %rd1, %rd220;
  2288. add.s64 %rd332, %rd153, %rd220;
  2289. mov.f32 %f1316, %f252;
  2290.  
  2291. BB1_161:
  2292. cvt.u32.u64 %r102, %rd72;
  2293. add.s32 %r210, %r102, 4;
  2294. setp.lt.u32 %p172, %r210, %r41;
  2295. @%p172 bra BB1_133;
  2296.  
  2297. BB1_162:
  2298. setp.eq.s64 %p173, %rd332, 0;
  2299. mov.f32 %f1346, 0f40000000;
  2300. mov.f32 %f1347, %f1346;
  2301. mov.f32 %f1348, %f1346;
  2302. @%p173 bra BB1_205;
  2303.  
  2304. fma.rn.f32 %f255, %f1289, %f1316, %f161;
  2305. fma.rn.f32 %f256, %f1290, %f1316, %f163;
  2306. fma.rn.f32 %f257, %f1291, %f1316, %f165;
  2307. add.s64 %rd98, %rd347, 8;
  2308. ld.global.f32 %f844, [%rd347+8];
  2309. sub.f32 %f1318, %f255, %f844;
  2310. ld.global.f32 %f845, [%rd347+12];
  2311. sub.f32 %f1319, %f256, %f845;
  2312. ld.global.f32 %f846, [%rd347+16];
  2313. sub.f32 %f1320, %f257, %f846;
  2314. mul.f32 %f847, %f1319, %f1319;
  2315. fma.rn.f32 %f848, %f1318, %f1318, %f847;
  2316. fma.rn.f32 %f261, %f1320, %f1320, %f848;
  2317. setp.leu.f32 %p174, %f261, 0f00000000;
  2318. @%p174 bra BB1_165;
  2319.  
  2320. sqrt.rn.f32 %f849, %f261;
  2321. rcp.rn.f32 %f850, %f849;
  2322. mul.f32 %f1318, %f1318, %f850;
  2323. mul.f32 %f1319, %f1319, %f850;
  2324. mul.f32 %f1320, %f1320, %f850;
  2325.  
  2326. BB1_165:
  2327. mul.f32 %f851, %f1290, %f1319;
  2328. fma.rn.f32 %f852, %f1289, %f1318, %f851;
  2329. fma.rn.f32 %f853, %f1291, %f1320, %f852;
  2330. mov.u16 %rs11, 0;
  2331. setp.leu.f32 %p175, %f853, 0f00000000;
  2332. @%p175 bra BB1_167;
  2333.  
  2334. neg.f32 %f1318, %f1318;
  2335. neg.f32 %f1319, %f1319;
  2336. neg.f32 %f1320, %f1320;
  2337. mov.u16 %rs11, 1;
  2338.  
  2339. BB1_167:
  2340. ld.global.f32 %f274, [%rd98+44];
  2341. setp.gt.f32 %p176, %f274, 0f00000000;
  2342. @%p176 bra BB1_181;
  2343.  
  2344. ld.global.f32 %f854, [%rd98+48];
  2345. setp.gt.f32 %p177, %f854, 0f00000000;
  2346. @%p177 bra BB1_181;
  2347. bra.uni BB1_169;
  2348.  
  2349. BB1_181:
  2350. mul.f32 %f897, %f1290, %f1319;
  2351. fma.rn.f32 %f898, %f1289, %f1318, %f897;
  2352. fma.rn.f32 %f303, %f1291, %f1320, %f898;
  2353. add.f32 %f304, %f303, 0f3F800000;
  2354. cvt.f64.f32 %fd13, %f304;
  2355. {
  2356. .reg .b32 %temp;
  2357. mov.b64 {%temp, %r27}, %fd13;
  2358. }
  2359. abs.f64 %fd14, %fd13;
  2360. // Callseq Start 4
  2361. {
  2362. .reg .b32 temp_param_reg;
  2363. // <end>}
  2364. .param .b64 param0;
  2365. st.param.f64 [param0+0], %fd14;
  2366. .param .b64 retval0;
  2367. call.uni (retval0),
  2368. __internal_accurate_pow,
  2369. (
  2370. param0
  2371. );
  2372. ld.param.f64 %fd74, [retval0+0];
  2373.  
  2374. //{
  2375. }// Callseq End 4
  2376. setp.gt.s32 %p188, %r27, -1;
  2377. setp.lt.s32 %p189, %r27, 0;
  2378. and.pred %p16, %p189, %p98;
  2379. or.pred %p190, %p188, %p97;
  2380. @%p190 bra BB1_183;
  2381.  
  2382. {
  2383. .reg .b32 %temp;
  2384. mov.b64 {%temp, %r106}, %fd74;
  2385. }
  2386. xor.b32 %r107, %r106, -2147483648;
  2387. {
  2388. .reg .b32 %temp;
  2389. mov.b64 {%r108, %temp}, %fd74;
  2390. }
  2391. mov.b64 %fd74, {%r108, %r107};
  2392.  
  2393. BB1_183:
  2394. setp.eq.f32 %p191, %f304, 0f00000000;
  2395. @%p191 bra BB1_186;
  2396. bra.uni BB1_184;
  2397.  
  2398. BB1_186:
  2399. selp.b32 %r109, %r27, 0, %p98;
  2400. mov.u32 %r110, 0;
  2401. or.b32 %r111, %r109, 2146435072;
  2402. setp.lt.s32 %p195, %r15, 0;
  2403. selp.b32 %r112, %r111, %r109, %p195;
  2404. mov.b64 %fd74, {%r110, %r112};
  2405. bra.uni BB1_187;
  2406.  
  2407. BB1_184:
  2408. @%p188 bra BB1_187;
  2409.  
  2410. mov.f64 %fd65, 0d4008000000000000;
  2411. cvt.rzi.f64.f64 %fd59, %fd65;
  2412. setp.neu.f64 %p193, %fd59, 0d4008000000000000;
  2413. selp.f64 %fd74, 0dFFF8000000000000, %fd74, %p193;
  2414.  
  2415. BB1_187:
  2416. add.f64 %fd75, %fd13, 0d4008000000000000;
  2417. {
  2418. .reg .b32 %temp;
  2419. mov.b64 {%temp, %r113}, %fd75;
  2420. }
  2421. and.b32 %r114, %r113, 2146435072;
  2422. setp.ne.s32 %p196, %r114, 2146435072;
  2423. @%p196 bra BB1_188;
  2424.  
  2425. setp.gtu.f64 %p197, %fd14, 0d7FF0000000000000;
  2426. @%p197 bra BB1_197;
  2427.  
  2428. and.b32 %r115, %r15, 2147483647;
  2429. setp.ne.s32 %p198, %r115, 2146435072;
  2430. @%p198 bra BB1_192;
  2431.  
  2432. mov.f64 %fd64, 0d4008000000000000;
  2433. {
  2434. .reg .b32 %temp;
  2435. mov.b64 {%r116, %temp}, %fd64;
  2436. }
  2437. setp.eq.s32 %p199, %r116, 0;
  2438. @%p199 bra BB1_196;
  2439.  
  2440. BB1_192:
  2441. and.b32 %r117, %r27, 2147483647;
  2442. setp.ne.s32 %p200, %r117, 2146435072;
  2443. @%p200 bra BB1_193;
  2444.  
  2445. {
  2446. .reg .b32 %temp;
  2447. mov.b64 {%r118, %temp}, %fd13;
  2448. }
  2449. setp.ne.s32 %p201, %r118, 0;
  2450. mov.f64 %fd75, %fd74;
  2451. @%p201 bra BB1_197;
  2452.  
  2453. shr.s32 %r119, %r15, 31;
  2454. and.b32 %r120, %r119, -2146435072;
  2455. add.s32 %r121, %r120, 2146435072;
  2456. or.b32 %r122, %r121, -2147483648;
  2457. selp.b32 %r123, %r122, %r121, %p16;
  2458. mov.u32 %r124, 0;
  2459. mov.b64 %fd75, {%r124, %r123};
  2460. bra.uni BB1_197;
  2461.  
  2462. BB1_188:
  2463. mov.f64 %fd75, %fd74;
  2464. bra.uni BB1_197;
  2465.  
  2466. BB1_169:
  2467. mov.f32 %f1343, 0f00000000;
  2468. mov.f32 %f1344, %f1343;
  2469. mov.f32 %f1345, %f1343;
  2470. @%p29 bra BB1_204;
  2471.  
  2472. fma.rn.f32 %f275, %f1318, 0f38D1B717, %f255;
  2473. fma.rn.f32 %f276, %f1319, 0f38D1B717, %f256;
  2474. fma.rn.f32 %f277, %f1320, 0f38D1B717, %f257;
  2475. mov.f32 %f1345, 0f00000000;
  2476. mov.u32 %r212, 0;
  2477. mov.f32 %f1344, %f1345;
  2478. mov.f32 %f1343, %f1345;
  2479.  
  2480. BB1_171:
  2481. cvt.u64.u32 %rd99, %r212;
  2482. mul.wide.u32 %rd221, %r212, 80;
  2483. add.s64 %rd222, %rd1, %rd221;
  2484. add.s64 %rd100, %rd222, 40;
  2485. ld.global.f32 %f281, [%rd222+40];
  2486. setp.leu.f32 %p179, %f281, 0f00000000;
  2487. @%p179 bra BB1_180;
  2488.  
  2489. ld.global.f32 %f861, [%rd100+-32];
  2490. sub.f32 %f1329, %f861, %f255;
  2491. ld.global.f32 %f862, [%rd100+-28];
  2492. sub.f32 %f1328, %f862, %f256;
  2493. ld.global.f32 %f863, [%rd100+-24];
  2494. sub.f32 %f1327, %f863, %f257;
  2495. mul.f32 %f864, %f1328, %f1328;
  2496. fma.rn.f32 %f865, %f1329, %f1329, %f864;
  2497. fma.rn.f32 %f285, %f1327, %f1327, %f865;
  2498. setp.leu.f32 %p180, %f285, 0f00000000;
  2499. @%p180 bra BB1_174;
  2500.  
  2501. sqrt.rn.f32 %f866, %f285;
  2502. rcp.rn.f32 %f867, %f866;
  2503. mul.f32 %f1329, %f1329, %f867;
  2504. mul.f32 %f1328, %f1328, %f867;
  2505. mul.f32 %f1327, %f1327, %f867;
  2506.  
  2507. BB1_174:
  2508. mov.u32 %r213, 0;
  2509.  
  2510. BB1_175:
  2511. cvt.u32.u64 %r105, %rd99;
  2512. setp.eq.s32 %p181, %r105, %r213;
  2513. @%p181 bra BB1_178;
  2514.  
  2515. mul.wide.u32 %rd223, %r213, 80;
  2516. add.s64 %rd224, %rd1, %rd223;
  2517. add.s64 %rd101, %rd224, 8;
  2518. ld.global.f32 %f868, [%rd224+8];
  2519. sub.f32 %f292, %f868, %f275;
  2520. ld.global.f32 %f869, [%rd224+12];
  2521. sub.f32 %f293, %f869, %f276;
  2522. ld.global.f32 %f870, [%rd224+16];
  2523. sub.f32 %f294, %f870, %f277;
  2524. mul.f32 %f871, %f1328, %f293;
  2525. fma.rn.f32 %f872, %f1329, %f292, %f871;
  2526. fma.rn.f32 %f295, %f1327, %f294, %f872;
  2527. setp.lt.f32 %p182, %f295, 0f00000000;
  2528. @%p182 bra BB1_178;
  2529.  
  2530. mul.f32 %f874, %f293, %f293;
  2531. fma.rn.f32 %f875, %f292, %f292, %f874;
  2532. fma.rn.f32 %f876, %f294, %f294, %f875;
  2533. mul.f32 %f877, %f295, %f295;
  2534. sub.f32 %f878, %f876, %f877;
  2535. ld.global.f32 %f879, [%rd101+16];
  2536. mov.f32 %f1330, 0f00000000;
  2537. setp.leu.f32 %p183, %f878, %f879;
  2538. @%p183 bra BB1_179;
  2539.  
  2540. BB1_178:
  2541. add.s32 %r213, %r213, 1;
  2542. setp.lt.u32 %p184, %r213, %r41;
  2543. mov.f32 %f1330, 0f3F800000;
  2544. @%p184 bra BB1_175;
  2545.  
  2546. BB1_179:
  2547. ld.global.f32 %f881, [%rd98+20];
  2548. mul.f32 %f882, %f1330, %f881;
  2549. ld.global.f32 %f883, [%rd98+24];
  2550. mul.f32 %f884, %f1330, %f883;
  2551. ld.global.f32 %f885, [%rd98+28];
  2552. mul.f32 %f886, %f1330, %f885;
  2553. mul.f32 %f887, %f1318, %f1329;
  2554. fma.rn.f32 %f888, %f1319, %f1328, %f887;
  2555. fma.rn.f32 %f889, %f1320, %f1327, %f888;
  2556. mov.f32 %f890, 0f00000000;
  2557. max.f32 %f891, %f890, %f889;
  2558. mul.f32 %f892, %f882, %f891;
  2559. mul.f32 %f893, %f884, %f891;
  2560. mul.f32 %f894, %f886, %f891;
  2561. ld.global.f32 %f895, [%rd100+4];
  2562. ld.global.f32 %f896, [%rd100+8];
  2563. fma.rn.f32 %f1343, %f281, %f892, %f1343;
  2564. fma.rn.f32 %f1344, %f893, %f895, %f1344;
  2565. fma.rn.f32 %f1345, %f894, %f896, %f1345;
  2566.  
  2567. BB1_180:
  2568. add.s32 %r212, %r212, 1;
  2569. setp.lt.u32 %p185, %r212, %r41;
  2570. @%p185 bra BB1_171;
  2571. bra.uni BB1_204;
  2572.  
  2573. BB1_193:
  2574. mov.f64 %fd75, %fd74;
  2575. bra.uni BB1_197;
  2576.  
  2577. BB1_196:
  2578. setp.gt.f64 %p202, %fd14, 0d3FF0000000000000;
  2579. selp.b32 %r125, 2146435072, 0, %p202;
  2580. mov.u32 %r126, 0;
  2581. xor.b32 %r127, %r125, 2146435072;
  2582. setp.lt.s32 %p203, %r15, 0;
  2583. selp.b32 %r128, %r127, %r125, %p203;
  2584. setp.eq.f32 %p204, %f304, 0fBF800000;
  2585. selp.b32 %r129, 1072693248, %r128, %p204;
  2586. mov.b64 %fd75, {%r126, %r129};
  2587.  
  2588. BB1_197:
  2589. cvt.rn.f32.f64 %f899, %fd75;
  2590. fma.rn.f32 %f900, %f899, 0f3F666666, 0f3DCCCCCD;
  2591. setp.eq.f32 %p205, %f304, 0f3F800000;
  2592. selp.f32 %f305, 0f3F800000, %f900, %p205;
  2593. add.f32 %f901, %f1318, %f1318;
  2594. mul.f32 %f902, %f901, %f303;
  2595. add.f32 %f903, %f1319, %f1319;
  2596. mul.f32 %f904, %f903, %f303;
  2597. add.f32 %f905, %f1320, %f1320;
  2598. mul.f32 %f906, %f905, %f303;
  2599. sub.f32 %f1334, %f1289, %f902;
  2600. sub.f32 %f1335, %f1290, %f904;
  2601. sub.f32 %f1336, %f1291, %f906;
  2602. mul.f32 %f907, %f1335, %f1335;
  2603. fma.rn.f32 %f908, %f1334, %f1334, %f907;
  2604. fma.rn.f32 %f309, %f1336, %f1336, %f908;
  2605. setp.leu.f32 %p206, %f309, 0f00000000;
  2606. @%p206 bra BB1_199;
  2607.  
  2608. sqrt.rn.f32 %f909, %f309;
  2609. rcp.rn.f32 %f910, %f909;
  2610. mul.f32 %f1334, %f1334, %f910;
  2611. mul.f32 %f1335, %f1335, %f910;
  2612. mul.f32 %f1336, %f1336, %f910;
  2613.  
  2614. BB1_199:
  2615. mul.f32 %f316, %f1318, 0f38D1B717;
  2616. add.f32 %f914, %f255, %f316;
  2617. mul.f32 %f317, %f1319, 0f38D1B717;
  2618. add.f32 %f915, %f256, %f317;
  2619. mul.f32 %f318, %f1320, 0f38D1B717;
  2620. add.f32 %f916, %f257, %f318;
  2621. mov.u32 %r130, 2;
  2622. // Callseq Start 5
  2623. {
  2624. .reg .b32 temp_param_reg;
  2625. // <end>}
  2626. .param .b32 param0;
  2627. st.param.f32 [param0+0], %f914;
  2628. .param .b32 param1;
  2629. st.param.f32 [param1+0], %f915;
  2630. .param .b32 param2;
  2631. st.param.f32 [param2+0], %f916;
  2632. .param .b32 param3;
  2633. st.param.f32 [param3+0], %f1334;
  2634. .param .b32 param4;
  2635. st.param.f32 [param4+0], %f1335;
  2636. .param .b32 param5;
  2637. st.param.f32 [param5+0], %f1336;
  2638. .param .b64 param6;
  2639. st.param.b64 [param6+0], %rd153;
  2640. .param .b32 param7;
  2641. st.param.b32 [param7+0], %r41;
  2642. .param .b32 param8;
  2643. st.param.b32 [param8+0], %r130;
  2644. .param .align 4 .b8 retval0[12];
  2645. call.uni (retval0),
  2646. _Z5traceRK4Vec3IfES2_PK6SpherejRKi,
  2647. (
  2648. param0,
  2649. param1,
  2650. param2,
  2651. param3,
  2652. param4,
  2653. param5,
  2654. param6,
  2655. param7,
  2656. param8
  2657. );
  2658. ld.param.f32 %f319, [retval0+0];
  2659. ld.param.f32 %f320, [retval0+4];
  2660. ld.param.f32 %f321, [retval0+8];
  2661.  
  2662. //{
  2663. }// Callseq End 5
  2664. mov.f32 %f1340, 0f00000000;
  2665. setp.eq.f32 %p207, %f274, 0f00000000;
  2666. mov.f32 %f1341, %f1340;
  2667. mov.f32 %f1342, %f1340;
  2668. @%p207 bra BB1_203;
  2669.  
  2670. setp.eq.s16 %p208, %rs11, 0;
  2671. selp.f32 %f917, 0f3F68BA2E, 0f3F8CCCCD, %p208;
  2672. mul.f32 %f918, %f917, %f917;
  2673. mul.f32 %f919, %f303, %f303;
  2674. mov.f32 %f920, 0f3F800000;
  2675. sub.f32 %f921, %f920, %f919;
  2676. mul.f32 %f922, %f918, %f921;
  2677. sub.f32 %f923, %f920, %f922;
  2678. mul.f32 %f924, %f303, %f917;
  2679. neg.f32 %f925, %f924;
  2680. sqrt.rn.f32 %f926, %f923;
  2681. sub.f32 %f927, %f925, %f926;
  2682. mul.f32 %f928, %f1318, %f927;
  2683. mul.f32 %f929, %f1319, %f927;
  2684. mul.f32 %f930, %f1320, %f927;
  2685. fma.rn.f32 %f1339, %f1289, %f917, %f928;
  2686. fma.rn.f32 %f1338, %f1290, %f917, %f929;
  2687. fma.rn.f32 %f1337, %f1291, %f917, %f930;
  2688. mul.f32 %f931, %f1338, %f1338;
  2689. fma.rn.f32 %f932, %f1339, %f1339, %f931;
  2690. fma.rn.f32 %f325, %f1337, %f1337, %f932;
  2691. setp.leu.f32 %p209, %f325, 0f00000000;
  2692. @%p209 bra BB1_202;
  2693.  
  2694. sqrt.rn.f32 %f933, %f325;
  2695. rcp.rn.f32 %f934, %f933;
  2696. mul.f32 %f1339, %f1339, %f934;
  2697. mul.f32 %f1338, %f1338, %f934;
  2698. mul.f32 %f1337, %f1337, %f934;
  2699.  
  2700. BB1_202:
  2701. sub.f32 %f935, %f257, %f318;
  2702. sub.f32 %f936, %f256, %f317;
  2703. sub.f32 %f937, %f255, %f316;
  2704. // Callseq Start 6
  2705. {
  2706. .reg .b32 temp_param_reg;
  2707. // <end>}
  2708. .param .b32 param0;
  2709. st.param.f32 [param0+0], %f937;
  2710. .param .b32 param1;
  2711. st.param.f32 [param1+0], %f936;
  2712. .param .b32 param2;
  2713. st.param.f32 [param2+0], %f935;
  2714. .param .b32 param3;
  2715. st.param.f32 [param3+0], %f1339;
  2716. .param .b32 param4;
  2717. st.param.f32 [param4+0], %f1338;
  2718. .param .b32 param5;
  2719. st.param.f32 [param5+0], %f1337;
  2720. .param .b64 param6;
  2721. st.param.b64 [param6+0], %rd153;
  2722. .param .b32 param7;
  2723. st.param.b32 [param7+0], %r41;
  2724. .param .b32 param8;
  2725. st.param.b32 [param8+0], %r130;
  2726. .param .align 4 .b8 retval0[12];
  2727. call.uni (retval0),
  2728. _Z5traceRK4Vec3IfES2_PK6SpherejRKi,
  2729. (
  2730. param0,
  2731. param1,
  2732. param2,
  2733. param3,
  2734. param4,
  2735. param5,
  2736. param6,
  2737. param7,
  2738. param8
  2739. );
  2740. ld.param.f32 %f1340, [retval0+0];
  2741. ld.param.f32 %f1341, [retval0+4];
  2742. ld.param.f32 %f1342, [retval0+8];
  2743.  
  2744. //{
  2745. }// Callseq End 6
  2746.  
  2747. BB1_203:
  2748. mov.f32 %f938, 0f3F800000;
  2749. sub.f32 %f939, %f938, %f305;
  2750. mul.f32 %f940, %f939, %f1340;
  2751. mul.f32 %f941, %f939, %f1341;
  2752. mul.f32 %f942, %f939, %f1342;
  2753. mul.f32 %f943, %f274, %f940;
  2754. mul.f32 %f944, %f274, %f941;
  2755. mul.f32 %f945, %f274, %f942;
  2756. fma.rn.f32 %f946, %f305, %f319, %f943;
  2757. fma.rn.f32 %f947, %f305, %f320, %f944;
  2758. fma.rn.f32 %f948, %f305, %f321, %f945;
  2759. ld.global.f32 %f949, [%rd98+20];
  2760. mul.f32 %f1343, %f949, %f946;
  2761. ld.global.f32 %f950, [%rd98+24];
  2762. mul.f32 %f1344, %f947, %f950;
  2763. ld.global.f32 %f951, [%rd98+28];
  2764. mul.f32 %f1345, %f948, %f951;
  2765.  
  2766. BB1_204:
  2767. ld.global.f32 %f952, [%rd98+32];
  2768. add.f32 %f1346, %f1343, %f952;
  2769. ld.global.f32 %f953, [%rd98+36];
  2770. add.f32 %f1347, %f1344, %f953;
  2771. ld.global.f32 %f954, [%rd98+40];
  2772. add.f32 %f1348, %f1345, %f954;
  2773.  
  2774. BB1_205:
  2775. mov.f32 %f1406, 0f00000000;
  2776. setp.eq.f32 %p210, %f118, 0f00000000;
  2777. mov.f32 %f1407, %f1406;
  2778. mov.f32 %f1408, %f1406;
  2779. @%p210 bra BB1_314;
  2780.  
  2781. setp.eq.s16 %p211, %rs10, 0;
  2782. selp.f32 %f958, 0f3F68BA2E, 0f3F8CCCCD, %p211;
  2783. mul.f32 %f959, %f958, %f958;
  2784. mul.f32 %f960, %f147, %f147;
  2785. mov.f32 %f961, 0f3F800000;
  2786. sub.f32 %f962, %f961, %f960;
  2787. mul.f32 %f963, %f959, %f962;
  2788. sub.f32 %f964, %f961, %f963;
  2789. mul.f32 %f965, %f147, %f958;
  2790. neg.f32 %f966, %f965;
  2791. sqrt.rn.f32 %f967, %f964;
  2792. sub.f32 %f968, %f966, %f967;
  2793. mul.f32 %f969, %f1273, %f968;
  2794. mul.f32 %f970, %f1274, %f968;
  2795. mul.f32 %f971, %f1275, %f968;
  2796. fma.rn.f32 %f1351, %f1246, %f958, %f969;
  2797. fma.rn.f32 %f1350, %f1245, %f958, %f970;
  2798. fma.rn.f32 %f1349, %f1244, %f958, %f971;
  2799. mul.f32 %f972, %f1350, %f1350;
  2800. fma.rn.f32 %f973, %f1351, %f1351, %f972;
  2801. fma.rn.f32 %f353, %f1349, %f1349, %f973;
  2802. setp.leu.f32 %p212, %f353, 0f00000000;
  2803. @%p212 bra BB1_208;
  2804.  
  2805. sqrt.rn.f32 %f974, %f353;
  2806. rcp.rn.f32 %f975, %f974;
  2807. mul.f32 %f1351, %f1351, %f975;
  2808. mul.f32 %f1350, %f1350, %f975;
  2809. mul.f32 %f1349, %f1349, %f975;
  2810.  
  2811. BB1_208:
  2812. fma.rn.f32 %f1234, %f1244, %f1252, 0f00000000;
  2813. fma.rn.f32 %f1233, %f1245, %f1252, 0f00000000;
  2814. fma.rn.f32 %f1232, %f1246, %f1252, 0f00000000;
  2815. mul.f32 %f1231, %f1275, 0f38D1B717;
  2816. mul.f32 %f1230, %f1274, 0f38D1B717;
  2817. mul.f32 %f1229, %f1273, 0f38D1B717;
  2818. sub.f32 %f360, %f1232, %f1229;
  2819. sub.f32 %f361, %f1233, %f1230;
  2820. sub.f32 %f362, %f1234, %f1231;
  2821. mov.f32 %f1357, 0f7F800000;
  2822. mov.u64 %rd353, 0;
  2823. mov.u64 %rd354, %rd353;
  2824. @%p29 bra BB1_270;
  2825.  
  2826. and.b32 %r135, %r41, 3;
  2827. mov.f32 %f1359, 0f7F800000;
  2828. mov.u64 %rd353, 0;
  2829. mov.u32 %r216, 0;
  2830. mov.f32 %f1363, 0f00000000;
  2831. setp.eq.s32 %p214, %r135, 0;
  2832. @%p214 bra BB1_210;
  2833.  
  2834. setp.eq.s32 %p215, %r135, 1;
  2835. @%p215 bra BB1_212;
  2836. bra.uni BB1_213;
  2837.  
  2838. BB1_212:
  2839. mov.u64 %rd354, %rd353;
  2840. mov.f32 %f1357, %f1359;
  2841. bra.uni BB1_231;
  2842.  
  2843. BB1_210:
  2844. mov.u64 %rd354, %rd353;
  2845. mov.f32 %f1357, %f1359;
  2846. bra.uni BB1_239;
  2847.  
  2848. BB1_213:
  2849. setp.eq.s32 %p216, %r135, 2;
  2850. @%p216 bra BB1_214;
  2851. bra.uni BB1_215;
  2852.  
  2853. BB1_214:
  2854. mov.u64 %rd354, %rd353;
  2855. mov.f32 %f1357, %f1359;
  2856. bra.uni BB1_223;
  2857.  
  2858. BB1_215:
  2859. ld.global.f32 %f983, [%rd1+8];
  2860. sub.f32 %f363, %f983, %f360;
  2861. ld.global.f32 %f984, [%rd1+12];
  2862. sub.f32 %f364, %f984, %f361;
  2863. ld.global.f32 %f985, [%rd1+16];
  2864. sub.f32 %f365, %f985, %f362;
  2865. mul.f32 %f986, %f1350, %f364;
  2866. fma.rn.f32 %f987, %f1351, %f363, %f986;
  2867. fma.rn.f32 %f366, %f1349, %f365, %f987;
  2868. setp.lt.f32 %p218, %f366, 0f00000000;
  2869. mov.pred %p319, -1;
  2870. mov.f32 %f982, 0f7F800000;
  2871. @%p218 bra BB1_216;
  2872. bra.uni BB1_217;
  2873.  
  2874. BB1_216:
  2875. mov.f32 %f1352, %f982;
  2876. mov.f32 %f1353, %f982;
  2877. bra.uni BB1_219;
  2878.  
  2879. BB1_217:
  2880. mul.f32 %f990, %f364, %f364;
  2881. fma.rn.f32 %f991, %f363, %f363, %f990;
  2882. fma.rn.f32 %f992, %f365, %f365, %f991;
  2883. mul.f32 %f993, %f366, %f366;
  2884. sub.f32 %f367, %f992, %f993;
  2885. ld.global.f32 %f368, [%rd1+24];
  2886. setp.gt.f32 %p220, %f367, %f368;
  2887. mov.f32 %f1352, %f982;
  2888. mov.f32 %f1353, %f982;
  2889. @%p220 bra BB1_219;
  2890.  
  2891. sub.f32 %f994, %f368, %f367;
  2892. sqrt.rn.f32 %f995, %f994;
  2893. sub.f32 %f1353, %f366, %f995;
  2894. add.f32 %f1352, %f366, %f995;
  2895. mov.pred %p319, 0;
  2896.  
  2897. BB1_219:
  2898. mov.u32 %r216, 1;
  2899. @%p319 bra BB1_220;
  2900.  
  2901. setp.lt.f32 %p222, %f1353, 0f00000000;
  2902. selp.f32 %f1357, %f1352, %f1353, %p222;
  2903. setp.geu.f32 %p223, %f1357, 0f7F800000;
  2904. @%p223 bra BB1_220;
  2905.  
  2906. mov.u64 %rd353, %rd1;
  2907. mov.u64 %rd354, %rd153;
  2908. bra.uni BB1_223;
  2909.  
  2910. BB1_220:
  2911. mov.u64 %rd354, %rd353;
  2912. mov.f32 %f1357, %f982;
  2913.  
  2914. BB1_223:
  2915. cvt.u64.u32 %rd104, %r216;
  2916. mul.wide.u32 %rd237, %r216, 80;
  2917. add.s64 %rd238, %rd1, %rd237;
  2918. add.s64 %rd105, %rd238, 8;
  2919. ld.global.f32 %f1000, [%rd238+8];
  2920. sub.f32 %f375, %f1000, %f360;
  2921. ld.global.f32 %f1001, [%rd238+12];
  2922. sub.f32 %f376, %f1001, %f361;
  2923. ld.global.f32 %f1002, [%rd238+16];
  2924. sub.f32 %f377, %f1002, %f362;
  2925. mul.f32 %f1003, %f1350, %f376;
  2926. fma.rn.f32 %f1004, %f1351, %f375, %f1003;
  2927. fma.rn.f32 %f378, %f1349, %f377, %f1004;
  2928. setp.lt.f32 %p225, %f378, 0f00000000;
  2929. mov.pred %p320, -1;
  2930. @%p225 bra BB1_224;
  2931. bra.uni BB1_225;
  2932.  
  2933. BB1_224:
  2934. mov.f32 %f1355, %f1359;
  2935. mov.f32 %f1356, %f1359;
  2936. bra.uni BB1_227;
  2937.  
  2938. BB1_225:
  2939. mul.f32 %f1007, %f376, %f376;
  2940. fma.rn.f32 %f1008, %f375, %f375, %f1007;
  2941. fma.rn.f32 %f1009, %f377, %f377, %f1008;
  2942. mul.f32 %f1010, %f378, %f378;
  2943. sub.f32 %f379, %f1009, %f1010;
  2944. ld.global.f32 %f380, [%rd105+16];
  2945. setp.gt.f32 %p227, %f379, %f380;
  2946. mov.f32 %f1355, %f1359;
  2947. mov.f32 %f1356, %f1359;
  2948. @%p227 bra BB1_227;
  2949.  
  2950. sub.f32 %f1011, %f380, %f379;
  2951. sqrt.rn.f32 %f1012, %f1011;
  2952. sub.f32 %f1356, %f378, %f1012;
  2953. add.f32 %f1355, %f378, %f1012;
  2954. mov.pred %p320, 0;
  2955.  
  2956. BB1_227:
  2957. @%p320 bra BB1_230;
  2958.  
  2959. setp.lt.f32 %p229, %f1356, 0f00000000;
  2960. selp.f32 %f385, %f1355, %f1356, %p229;
  2961. setp.geu.f32 %p230, %f385, %f1357;
  2962. @%p230 bra BB1_230;
  2963.  
  2964. mul.lo.s64 %rd239, %rd104, 80;
  2965. add.s64 %rd353, %rd1, %rd239;
  2966. add.s64 %rd354, %rd153, %rd239;
  2967. mov.f32 %f1357, %f385;
  2968.  
  2969. BB1_230:
  2970. cvt.u32.u64 %r139, %rd104;
  2971. add.s32 %r216, %r139, 1;
  2972.  
  2973. BB1_231:
  2974. cvt.u64.u32 %rd112, %r216;
  2975. mul.wide.u32 %rd240, %r216, 80;
  2976. add.s64 %rd241, %rd1, %rd240;
  2977. add.s64 %rd113, %rd241, 8;
  2978. ld.global.f32 %f1015, [%rd241+8];
  2979. sub.f32 %f388, %f1015, %f360;
  2980. ld.global.f32 %f1016, [%rd241+12];
  2981. sub.f32 %f389, %f1016, %f361;
  2982. ld.global.f32 %f1017, [%rd241+16];
  2983. sub.f32 %f390, %f1017, %f362;
  2984. mul.f32 %f1018, %f1350, %f389;
  2985. fma.rn.f32 %f1019, %f1351, %f388, %f1018;
  2986. fma.rn.f32 %f391, %f1349, %f390, %f1019;
  2987. setp.lt.f32 %p232, %f391, 0f00000000;
  2988. mov.pred %p321, -1;
  2989. @%p232 bra BB1_232;
  2990. bra.uni BB1_233;
  2991.  
  2992. BB1_232:
  2993. mov.f32 %f1360, %f1359;
  2994. bra.uni BB1_235;
  2995.  
  2996. BB1_233:
  2997. mul.f32 %f1022, %f389, %f389;
  2998. fma.rn.f32 %f1023, %f388, %f388, %f1022;
  2999. fma.rn.f32 %f1024, %f390, %f390, %f1023;
  3000. mul.f32 %f1025, %f391, %f391;
  3001. sub.f32 %f392, %f1024, %f1025;
  3002. ld.global.f32 %f393, [%rd113+16];
  3003. setp.gt.f32 %p234, %f392, %f393;
  3004. mov.f32 %f1360, %f1359;
  3005. @%p234 bra BB1_235;
  3006.  
  3007. sub.f32 %f1026, %f393, %f392;
  3008. sqrt.rn.f32 %f1027, %f1026;
  3009. sub.f32 %f1360, %f391, %f1027;
  3010. add.f32 %f1359, %f391, %f1027;
  3011. mov.pred %p321, 0;
  3012.  
  3013. BB1_235:
  3014. @%p321 bra BB1_238;
  3015.  
  3016. setp.lt.f32 %p236, %f1360, 0f00000000;
  3017. selp.f32 %f398, %f1359, %f1360, %p236;
  3018. setp.geu.f32 %p237, %f398, %f1357;
  3019. @%p237 bra BB1_238;
  3020.  
  3021. mul.lo.s64 %rd242, %rd112, 80;
  3022. add.s64 %rd353, %rd1, %rd242;
  3023. add.s64 %rd354, %rd153, %rd242;
  3024. mov.f32 %f1357, %f398;
  3025.  
  3026. BB1_238:
  3027. cvt.u32.u64 %r140, %rd112;
  3028. add.s32 %r216, %r140, 1;
  3029. mov.f32 %f1363, %f1357;
  3030.  
  3031. BB1_239:
  3032. setp.lt.u32 %p238, %r41, 4;
  3033. @%p238 bra BB1_240;
  3034. bra.uni BB1_241;
  3035.  
  3036. BB1_240:
  3037. mov.f32 %f1357, %f1363;
  3038. bra.uni BB1_270;
  3039.  
  3040. BB1_241:
  3041. cvt.u64.u32 %rd122, %r216;
  3042. mul.wide.u32 %rd243, %r216, 80;
  3043. add.s64 %rd244, %rd1, %rd243;
  3044. add.s64 %rd123, %rd244, 8;
  3045. ld.global.f32 %f1030, [%rd244+8];
  3046. sub.f32 %f403, %f1030, %f360;
  3047. ld.global.f32 %f1031, [%rd244+12];
  3048. sub.f32 %f404, %f1031, %f361;
  3049. ld.global.f32 %f1032, [%rd244+16];
  3050. sub.f32 %f405, %f1032, %f362;
  3051. mul.f32 %f1033, %f1350, %f404;
  3052. fma.rn.f32 %f1034, %f1351, %f403, %f1033;
  3053. fma.rn.f32 %f406, %f1349, %f405, %f1034;
  3054. setp.lt.f32 %p240, %f406, 0f00000000;
  3055. mov.pred %p322, -1;
  3056. mov.f32 %f1365, 0f7F800000;
  3057. @%p240 bra BB1_242;
  3058. bra.uni BB1_243;
  3059.  
  3060. BB1_242:
  3061. mov.f32 %f1366, %f1365;
  3062. bra.uni BB1_245;
  3063.  
  3064. BB1_243:
  3065. mul.f32 %f1037, %f404, %f404;
  3066. fma.rn.f32 %f1038, %f403, %f403, %f1037;
  3067. fma.rn.f32 %f1039, %f405, %f405, %f1038;
  3068. mul.f32 %f1040, %f406, %f406;
  3069. sub.f32 %f407, %f1039, %f1040;
  3070. ld.global.f32 %f408, [%rd123+16];
  3071. setp.gt.f32 %p242, %f407, %f408;
  3072. mov.f32 %f1366, %f1365;
  3073. @%p242 bra BB1_245;
  3074.  
  3075. sub.f32 %f1041, %f408, %f407;
  3076. sqrt.rn.f32 %f1042, %f1041;
  3077. sub.f32 %f1366, %f406, %f1042;
  3078. add.f32 %f1365, %f406, %f1042;
  3079. mov.pred %p322, 0;
  3080.  
  3081. BB1_245:
  3082. @%p322 bra BB1_248;
  3083.  
  3084. setp.lt.f32 %p244, %f1366, 0f00000000;
  3085. selp.f32 %f413, %f1365, %f1366, %p244;
  3086. setp.geu.f32 %p245, %f413, %f1357;
  3087. @%p245 bra BB1_248;
  3088.  
  3089. mul.lo.s64 %rd245, %rd122, 80;
  3090. add.s64 %rd353, %rd1, %rd245;
  3091. add.s64 %rd354, %rd153, %rd245;
  3092. mov.f32 %f1357, %f413;
  3093.  
  3094. BB1_248:
  3095. add.s32 %r141, %r216, 1;
  3096. cvt.u64.u32 %rd128, %r141;
  3097. mul.wide.u32 %rd246, %r141, 80;
  3098. add.s64 %rd247, %rd1, %rd246;
  3099. add.s64 %rd129, %rd247, 8;
  3100. ld.global.f32 %f1045, [%rd247+8];
  3101. sub.f32 %f415, %f1045, %f360;
  3102. ld.global.f32 %f1046, [%rd247+12];
  3103. sub.f32 %f416, %f1046, %f361;
  3104. ld.global.f32 %f1047, [%rd247+16];
  3105. sub.f32 %f417, %f1047, %f362;
  3106. mul.f32 %f1048, %f1350, %f416;
  3107. fma.rn.f32 %f1049, %f1351, %f415, %f1048;
  3108. fma.rn.f32 %f418, %f1349, %f417, %f1049;
  3109. setp.lt.f32 %p247, %f418, 0f00000000;
  3110. mov.pred %p323, -1;
  3111. mov.f32 %f1368, 0f7F800000;
  3112. @%p247 bra BB1_249;
  3113. bra.uni BB1_250;
  3114.  
  3115. BB1_249:
  3116. mov.f32 %f1369, %f1368;
  3117. bra.uni BB1_252;
  3118.  
  3119. BB1_250:
  3120. mul.f32 %f1052, %f416, %f416;
  3121. fma.rn.f32 %f1053, %f415, %f415, %f1052;
  3122. fma.rn.f32 %f1054, %f417, %f417, %f1053;
  3123. mul.f32 %f1055, %f418, %f418;
  3124. sub.f32 %f419, %f1054, %f1055;
  3125. ld.global.f32 %f420, [%rd129+16];
  3126. setp.gt.f32 %p249, %f419, %f420;
  3127. mov.f32 %f1369, %f1368;
  3128. @%p249 bra BB1_252;
  3129.  
  3130. sub.f32 %f1056, %f420, %f419;
  3131. sqrt.rn.f32 %f1057, %f1056;
  3132. sub.f32 %f1369, %f418, %f1057;
  3133. add.f32 %f1368, %f418, %f1057;
  3134. mov.pred %p323, 0;
  3135.  
  3136. BB1_252:
  3137. @%p323 bra BB1_255;
  3138.  
  3139. setp.lt.f32 %p251, %f1369, 0f00000000;
  3140. selp.f32 %f425, %f1368, %f1369, %p251;
  3141. setp.geu.f32 %p252, %f425, %f1357;
  3142. @%p252 bra BB1_255;
  3143.  
  3144. mul.lo.s64 %rd248, %rd128, 80;
  3145. add.s64 %rd353, %rd1, %rd248;
  3146. add.s64 %rd354, %rd153, %rd248;
  3147. mov.f32 %f1357, %f425;
  3148.  
  3149. BB1_255:
  3150. add.s32 %r142, %r216, 2;
  3151. cvt.u64.u32 %rd134, %r142;
  3152. mul.wide.u32 %rd249, %r142, 80;
  3153. add.s64 %rd250, %rd1, %rd249;
  3154. add.s64 %rd135, %rd250, 8;
  3155. ld.global.f32 %f1060, [%rd250+8];
  3156. sub.f32 %f427, %f1060, %f360;
  3157. ld.global.f32 %f1061, [%rd250+12];
  3158. sub.f32 %f428, %f1061, %f361;
  3159. ld.global.f32 %f1062, [%rd250+16];
  3160. sub.f32 %f429, %f1062, %f362;
  3161. mul.f32 %f1063, %f1350, %f428;
  3162. fma.rn.f32 %f1064, %f1351, %f427, %f1063;
  3163. fma.rn.f32 %f430, %f1349, %f429, %f1064;
  3164. setp.lt.f32 %p254, %f430, 0f00000000;
  3165. mov.pred %p324, -1;
  3166. mov.f32 %f1371, 0f7F800000;
  3167. @%p254 bra BB1_256;
  3168. bra.uni BB1_257;
  3169.  
  3170. BB1_256:
  3171. mov.f32 %f1372, %f1371;
  3172. bra.uni BB1_259;
  3173.  
  3174. BB1_257:
  3175. mul.f32 %f1067, %f428, %f428;
  3176. fma.rn.f32 %f1068, %f427, %f427, %f1067;
  3177. fma.rn.f32 %f1069, %f429, %f429, %f1068;
  3178. mul.f32 %f1070, %f430, %f430;
  3179. sub.f32 %f431, %f1069, %f1070;
  3180. ld.global.f32 %f432, [%rd135+16];
  3181. setp.gt.f32 %p256, %f431, %f432;
  3182. mov.f32 %f1372, %f1371;
  3183. @%p256 bra BB1_259;
  3184.  
  3185. sub.f32 %f1071, %f432, %f431;
  3186. sqrt.rn.f32 %f1072, %f1071;
  3187. sub.f32 %f1372, %f430, %f1072;
  3188. add.f32 %f1371, %f430, %f1072;
  3189. mov.pred %p324, 0;
  3190.  
  3191. BB1_259:
  3192. @%p324 bra BB1_262;
  3193.  
  3194. setp.lt.f32 %p258, %f1372, 0f00000000;
  3195. selp.f32 %f437, %f1371, %f1372, %p258;
  3196. setp.geu.f32 %p259, %f437, %f1357;
  3197. @%p259 bra BB1_262;
  3198.  
  3199. mul.lo.s64 %rd251, %rd134, 80;
  3200. add.s64 %rd353, %rd1, %rd251;
  3201. add.s64 %rd354, %rd153, %rd251;
  3202. mov.f32 %f1357, %f437;
  3203.  
  3204. BB1_262:
  3205. add.s32 %r143, %r216, 3;
  3206. cvt.u64.u32 %rd140, %r143;
  3207. mul.wide.u32 %rd252, %r143, 80;
  3208. add.s64 %rd253, %rd1, %rd252;
  3209. add.s64 %rd141, %rd253, 8;
  3210. ld.global.f32 %f1075, [%rd253+8];
  3211. sub.f32 %f439, %f1075, %f360;
  3212. ld.global.f32 %f1076, [%rd253+12];
  3213. sub.f32 %f440, %f1076, %f361;
  3214. ld.global.f32 %f1077, [%rd253+16];
  3215. sub.f32 %f441, %f1077, %f362;
  3216. mul.f32 %f1078, %f1350, %f440;
  3217. fma.rn.f32 %f1079, %f1351, %f439, %f1078;
  3218. fma.rn.f32 %f442, %f1349, %f441, %f1079;
  3219. setp.lt.f32 %p261, %f442, 0f00000000;
  3220. mov.pred %p325, -1;
  3221. mov.f32 %f1374, 0f7F800000;
  3222. @%p261 bra BB1_263;
  3223. bra.uni BB1_264;
  3224.  
  3225. BB1_263:
  3226. mov.f32 %f1375, %f1374;
  3227. bra.uni BB1_266;
  3228.  
  3229. BB1_264:
  3230. mul.f32 %f1082, %f440, %f440;
  3231. fma.rn.f32 %f1083, %f439, %f439, %f1082;
  3232. fma.rn.f32 %f1084, %f441, %f441, %f1083;
  3233. mul.f32 %f1085, %f442, %f442;
  3234. sub.f32 %f443, %f1084, %f1085;
  3235. ld.global.f32 %f444, [%rd141+16];
  3236. setp.gt.f32 %p263, %f443, %f444;
  3237. mov.f32 %f1375, %f1374;
  3238. @%p263 bra BB1_266;
  3239.  
  3240. sub.f32 %f1086, %f444, %f443;
  3241. sqrt.rn.f32 %f1087, %f1086;
  3242. sub.f32 %f1375, %f442, %f1087;
  3243. add.f32 %f1374, %f442, %f1087;
  3244. mov.pred %p325, 0;
  3245.  
  3246. BB1_266:
  3247. @%p325 bra BB1_269;
  3248.  
  3249. setp.lt.f32 %p265, %f1375, 0f00000000;
  3250. selp.f32 %f449, %f1374, %f1375, %p265;
  3251. setp.geu.f32 %p266, %f449, %f1357;
  3252. @%p266 bra BB1_269;
  3253.  
  3254. mul.lo.s64 %rd254, %rd140, 80;
  3255. add.s64 %rd353, %rd1, %rd254;
  3256. add.s64 %rd354, %rd153, %rd254;
  3257. mov.f32 %f1357, %f449;
  3258.  
  3259. BB1_269:
  3260. cvt.u32.u64 %r144, %rd122;
  3261. add.s32 %r216, %r144, 4;
  3262. setp.lt.u32 %p267, %r216, %r41;
  3263. @%p267 bra BB1_241;
  3264.  
  3265. BB1_270:
  3266. setp.eq.s64 %p268, %rd354, 0;
  3267. mov.f32 %f1406, 0f40000000;
  3268. mov.f32 %f1407, %f1406;
  3269. mov.f32 %f1408, %f1406;
  3270. @%p268 bra BB1_314;
  3271.  
  3272. fma.rn.f32 %f452, %f1351, %f1357, %f360;
  3273. fma.rn.f32 %f453, %f1350, %f1357, %f361;
  3274. fma.rn.f32 %f454, %f1349, %f1357, %f362;
  3275. add.s64 %rd148, %rd353, 8;
  3276. ld.global.f32 %f1091, [%rd353+8];
  3277. sub.f32 %f1378, %f452, %f1091;
  3278. ld.global.f32 %f1092, [%rd353+12];
  3279. sub.f32 %f1379, %f453, %f1092;
  3280. ld.global.f32 %f1093, [%rd353+16];
  3281. sub.f32 %f1380, %f454, %f1093;
  3282. mul.f32 %f1094, %f1379, %f1379;
  3283. fma.rn.f32 %f1095, %f1378, %f1378, %f1094;
  3284. fma.rn.f32 %f458, %f1380, %f1380, %f1095;
  3285. setp.leu.f32 %p269, %f458, 0f00000000;
  3286. @%p269 bra BB1_273;
  3287.  
  3288. sqrt.rn.f32 %f1096, %f458;
  3289. rcp.rn.f32 %f1097, %f1096;
  3290. mul.f32 %f1378, %f1378, %f1097;
  3291. mul.f32 %f1379, %f1379, %f1097;
  3292. mul.f32 %f1380, %f1380, %f1097;
  3293.  
  3294. BB1_273:
  3295. mul.f32 %f1098, %f1350, %f1379;
  3296. fma.rn.f32 %f1099, %f1351, %f1378, %f1098;
  3297. fma.rn.f32 %f1100, %f1349, %f1380, %f1099;
  3298. mov.u16 %rs12, 0;
  3299. setp.leu.f32 %p270, %f1100, 0f00000000;
  3300. @%p270 bra BB1_275;
  3301.  
  3302. neg.f32 %f1378, %f1378;
  3303. neg.f32 %f1379, %f1379;
  3304. neg.f32 %f1380, %f1380;
  3305. mov.u16 %rs12, 1;
  3306.  
  3307. BB1_275:
  3308. ld.global.f32 %f471, [%rd148+44];
  3309. setp.gt.f32 %p271, %f471, 0f00000000;
  3310. @%p271 bra BB1_290;
  3311.  
  3312. ld.global.f32 %f1101, [%rd148+48];
  3313. setp.gt.f32 %p272, %f1101, 0f00000000;
  3314. @%p272 bra BB1_290;
  3315. bra.uni BB1_277;
  3316.  
  3317. BB1_290:
  3318. mul.f32 %f1144, %f1350, %f1379;
  3319. fma.rn.f32 %f1145, %f1351, %f1378, %f1144;
  3320. fma.rn.f32 %f500, %f1349, %f1380, %f1145;
  3321. add.f32 %f501, %f500, 0f3F800000;
  3322. cvt.f64.f32 %fd25, %f501;
  3323. {
  3324. .reg .b32 %temp;
  3325. mov.b64 {%temp, %r39}, %fd25;
  3326. }
  3327. abs.f64 %fd26, %fd25;
  3328. // Callseq Start 7
  3329. {
  3330. .reg .b32 temp_param_reg;
  3331. // <end>}
  3332. .param .b64 param0;
  3333. st.param.f64 [param0+0], %fd26;
  3334. .param .b64 retval0;
  3335. call.uni (retval0),
  3336. __internal_accurate_pow,
  3337. (
  3338. param0
  3339. );
  3340. ld.param.f64 %fd77, [retval0+0];
  3341.  
  3342. //{
  3343. }// Callseq End 7
  3344. setp.gt.s32 %p283, %r39, -1;
  3345. setp.lt.s32 %p284, %r39, 0;
  3346. and.pred %p24, %p284, %p98;
  3347. or.pred %p285, %p283, %p97;
  3348. @%p285 bra BB1_292;
  3349.  
  3350. {
  3351. .reg .b32 %temp;
  3352. mov.b64 {%temp, %r148}, %fd77;
  3353. }
  3354. xor.b32 %r149, %r148, -2147483648;
  3355. {
  3356. .reg .b32 %temp;
  3357. mov.b64 {%r150, %temp}, %fd77;
  3358. }
  3359. mov.b64 %fd77, {%r150, %r149};
  3360.  
  3361. BB1_292:
  3362. setp.eq.f32 %p286, %f501, 0f00000000;
  3363. @%p286 bra BB1_295;
  3364. bra.uni BB1_293;
  3365.  
  3366. BB1_295:
  3367. selp.b32 %r151, %r39, 0, %p98;
  3368. mov.u32 %r152, 0;
  3369. or.b32 %r153, %r151, 2146435072;
  3370. setp.lt.s32 %p290, %r15, 0;
  3371. selp.b32 %r154, %r153, %r151, %p290;
  3372. mov.b64 %fd77, {%r152, %r154};
  3373. bra.uni BB1_296;
  3374.  
  3375. BB1_293:
  3376. @%p283 bra BB1_296;
  3377.  
  3378. mov.f64 %fd67, 0d4008000000000000;
  3379. cvt.rzi.f64.f64 %fd62, %fd67;
  3380. setp.neu.f64 %p288, %fd62, 0d4008000000000000;
  3381. selp.f64 %fd77, 0dFFF8000000000000, %fd77, %p288;
  3382.  
  3383. BB1_296:
  3384. add.f64 %fd78, %fd25, 0d4008000000000000;
  3385. {
  3386. .reg .b32 %temp;
  3387. mov.b64 {%temp, %r155}, %fd78;
  3388. }
  3389. and.b32 %r156, %r155, 2146435072;
  3390. setp.ne.s32 %p291, %r156, 2146435072;
  3391. @%p291 bra BB1_297;
  3392.  
  3393. setp.gtu.f64 %p292, %fd26, 0d7FF0000000000000;
  3394. @%p292 bra BB1_306;
  3395.  
  3396. and.b32 %r157, %r15, 2147483647;
  3397. setp.ne.s32 %p293, %r157, 2146435072;
  3398. @%p293 bra BB1_301;
  3399.  
  3400. mov.f64 %fd66, 0d4008000000000000;
  3401. {
  3402. .reg .b32 %temp;
  3403. mov.b64 {%r158, %temp}, %fd66;
  3404. }
  3405. setp.eq.s32 %p294, %r158, 0;
  3406. @%p294 bra BB1_305;
  3407.  
  3408. BB1_301:
  3409. and.b32 %r159, %r39, 2147483647;
  3410. setp.ne.s32 %p295, %r159, 2146435072;
  3411. @%p295 bra BB1_302;
  3412.  
  3413. {
  3414. .reg .b32 %temp;
  3415. mov.b64 {%r160, %temp}, %fd25;
  3416. }
  3417. setp.ne.s32 %p296, %r160, 0;
  3418. mov.f64 %fd78, %fd77;
  3419. @%p296 bra BB1_306;
  3420.  
  3421. shr.s32 %r161, %r15, 31;
  3422. and.b32 %r162, %r161, -2146435072;
  3423. add.s32 %r163, %r162, 2146435072;
  3424. or.b32 %r164, %r163, -2147483648;
  3425. selp.b32 %r165, %r164, %r163, %p24;
  3426. mov.u32 %r166, 0;
  3427. mov.b64 %fd78, {%r166, %r165};
  3428. bra.uni BB1_306;
  3429.  
  3430. BB1_297:
  3431. mov.f64 %fd78, %fd77;
  3432. bra.uni BB1_306;
  3433.  
  3434. BB1_277:
  3435. mov.f32 %f1403, 0f00000000;
  3436. mov.f32 %f1404, %f1403;
  3437. mov.f32 %f1405, %f1403;
  3438. @%p29 bra BB1_313;
  3439.  
  3440. fma.rn.f32 %f472, %f1378, 0f38D1B717, %f452;
  3441. fma.rn.f32 %f473, %f1379, 0f38D1B717, %f453;
  3442. fma.rn.f32 %f474, %f1380, 0f38D1B717, %f454;
  3443. mov.f32 %f1405, 0f00000000;
  3444. mov.u32 %r218, 0;
  3445. mov.f32 %f1404, %f1405;
  3446. mov.f32 %f1403, %f1405;
  3447.  
  3448. BB1_279:
  3449. cvt.u64.u32 %rd149, %r218;
  3450. mul.wide.u32 %rd255, %r218, 80;
  3451. add.s64 %rd256, %rd1, %rd255;
  3452. add.s64 %rd150, %rd256, 40;
  3453. ld.global.f32 %f478, [%rd256+40];
  3454. setp.leu.f32 %p274, %f478, 0f00000000;
  3455. @%p274 bra BB1_289;
  3456.  
  3457. ld.global.f32 %f1108, [%rd150+-32];
  3458. sub.f32 %f1389, %f1108, %f452;
  3459. ld.global.f32 %f1109, [%rd150+-28];
  3460. sub.f32 %f1388, %f1109, %f453;
  3461. ld.global.f32 %f1110, [%rd150+-24];
  3462. sub.f32 %f1387, %f1110, %f454;
  3463. mul.f32 %f1111, %f1388, %f1388;
  3464. fma.rn.f32 %f1112, %f1389, %f1389, %f1111;
  3465. fma.rn.f32 %f482, %f1387, %f1387, %f1112;
  3466. setp.leu.f32 %p275, %f482, 0f00000000;
  3467. @%p275 bra BB1_282;
  3468.  
  3469. sqrt.rn.f32 %f1113, %f482;
  3470. rcp.rn.f32 %f1114, %f1113;
  3471. mul.f32 %f1389, %f1389, %f1114;
  3472. mul.f32 %f1388, %f1388, %f1114;
  3473. mul.f32 %f1387, %f1387, %f1114;
  3474.  
  3475. BB1_282:
  3476. mov.u32 %r219, 0;
  3477.  
  3478. BB1_283:
  3479. cvt.u32.u64 %r147, %rd149;
  3480. setp.eq.s32 %p276, %r147, %r219;
  3481. @%p276 bra BB1_286;
  3482.  
  3483. mul.wide.u32 %rd257, %r219, 80;
  3484. add.s64 %rd258, %rd1, %rd257;
  3485. add.s64 %rd151, %rd258, 8;
  3486. ld.global.f32 %f1115, [%rd258+8];
  3487. sub.f32 %f489, %f1115, %f472;
  3488. ld.global.f32 %f1116, [%rd258+12];
  3489. sub.f32 %f490, %f1116, %f473;
  3490. ld.global.f32 %f1117, [%rd258+16];
  3491. sub.f32 %f491, %f1117, %f474;
  3492. mul.f32 %f1118, %f1388, %f490;
  3493. fma.rn.f32 %f1119, %f1389, %f489, %f1118;
  3494. fma.rn.f32 %f492, %f1387, %f491, %f1119;
  3495. setp.lt.f32 %p277, %f492, 0f00000000;
  3496. @%p277 bra BB1_286;
  3497.  
  3498. mul.f32 %f1121, %f490, %f490;
  3499. fma.rn.f32 %f1122, %f489, %f489, %f1121;
  3500. fma.rn.f32 %f1123, %f491, %f491, %f1122;
  3501. mul.f32 %f1124, %f492, %f492;
  3502. sub.f32 %f1125, %f1123, %f1124;
  3503. ld.global.f32 %f1126, [%rd151+16];
  3504. mov.f32 %f1390, 0f00000000;
  3505. setp.leu.f32 %p278, %f1125, %f1126;
  3506. @%p278 bra BB1_288;
  3507.  
  3508. BB1_286:
  3509. add.s32 %r219, %r219, 1;
  3510. setp.lt.u32 %p279, %r219, %r41;
  3511. @%p279 bra BB1_283;
  3512.  
  3513. mov.f32 %f1390, %f961;
  3514.  
  3515. BB1_288:
  3516. ld.global.f32 %f1128, [%rd148+20];
  3517. mul.f32 %f1129, %f1390, %f1128;
  3518. ld.global.f32 %f1130, [%rd148+24];
  3519. mul.f32 %f1131, %f1390, %f1130;
  3520. ld.global.f32 %f1132, [%rd148+28];
  3521. mul.f32 %f1133, %f1390, %f1132;
  3522. mul.f32 %f1134, %f1378, %f1389;
  3523. fma.rn.f32 %f1135, %f1379, %f1388, %f1134;
  3524. fma.rn.f32 %f1136, %f1380, %f1387, %f1135;
  3525. mov.f32 %f1137, 0f00000000;
  3526. max.f32 %f1138, %f1137, %f1136;
  3527. mul.f32 %f1139, %f1129, %f1138;
  3528. mul.f32 %f1140, %f1131, %f1138;
  3529. mul.f32 %f1141, %f1133, %f1138;
  3530. ld.global.f32 %f1142, [%rd150+4];
  3531. ld.global.f32 %f1143, [%rd150+8];
  3532. fma.rn.f32 %f1403, %f478, %f1139, %f1403;
  3533. fma.rn.f32 %f1404, %f1140, %f1142, %f1404;
  3534. fma.rn.f32 %f1405, %f1141, %f1143, %f1405;
  3535.  
  3536. BB1_289:
  3537. add.s32 %r218, %r218, 1;
  3538. setp.lt.u32 %p280, %r218, %r41;
  3539. @%p280 bra BB1_279;
  3540. bra.uni BB1_313;
  3541.  
  3542. BB1_302:
  3543. mov.f64 %fd78, %fd77;
  3544. bra.uni BB1_306;
  3545.  
  3546. BB1_305:
  3547. setp.gt.f64 %p297, %fd26, 0d3FF0000000000000;
  3548. selp.b32 %r167, 2146435072, 0, %p297;
  3549. mov.u32 %r168, 0;
  3550. xor.b32 %r169, %r167, 2146435072;
  3551. setp.lt.s32 %p298, %r15, 0;
  3552. selp.b32 %r170, %r169, %r167, %p298;
  3553. setp.eq.f32 %p299, %f501, 0fBF800000;
  3554. selp.b32 %r171, 1072693248, %r170, %p299;
  3555. mov.b64 %fd78, {%r168, %r171};
  3556.  
  3557. BB1_306:
  3558. cvt.rn.f32.f64 %f1146, %fd78;
  3559. fma.rn.f32 %f1147, %f1146, 0f3F666666, 0f3DCCCCCD;
  3560. setp.eq.f32 %p300, %f501, 0f3F800000;
  3561. selp.f32 %f502, 0f3F800000, %f1147, %p300;
  3562. add.f32 %f1148, %f1378, %f1378;
  3563. mul.f32 %f1149, %f1148, %f500;
  3564. add.f32 %f1150, %f1379, %f1379;
  3565. mul.f32 %f1151, %f1150, %f500;
  3566. add.f32 %f1152, %f1380, %f1380;
  3567. mul.f32 %f1153, %f1152, %f500;
  3568. sub.f32 %f1394, %f1351, %f1149;
  3569. sub.f32 %f1395, %f1350, %f1151;
  3570. sub.f32 %f1396, %f1349, %f1153;
  3571. mul.f32 %f1154, %f1395, %f1395;
  3572. fma.rn.f32 %f1155, %f1394, %f1394, %f1154;
  3573. fma.rn.f32 %f506, %f1396, %f1396, %f1155;
  3574. setp.leu.f32 %p301, %f506, 0f00000000;
  3575. @%p301 bra BB1_308;
  3576.  
  3577. sqrt.rn.f32 %f1156, %f506;
  3578. rcp.rn.f32 %f1157, %f1156;
  3579. mul.f32 %f1394, %f1394, %f1157;
  3580. mul.f32 %f1395, %f1395, %f1157;
  3581. mul.f32 %f1396, %f1396, %f1157;
  3582.  
  3583. BB1_308:
  3584. mul.f32 %f513, %f1378, 0f38D1B717;
  3585. add.f32 %f1161, %f452, %f513;
  3586. mul.f32 %f514, %f1379, 0f38D1B717;
  3587. add.f32 %f1162, %f453, %f514;
  3588. mul.f32 %f515, %f1380, 0f38D1B717;
  3589. add.f32 %f1163, %f454, %f515;
  3590. mov.u32 %r172, 2;
  3591. // Callseq Start 8
  3592. {
  3593. .reg .b32 temp_param_reg;
  3594. // <end>}
  3595. .param .b32 param0;
  3596. st.param.f32 [param0+0], %f1161;
  3597. .param .b32 param1;
  3598. st.param.f32 [param1+0], %f1162;
  3599. .param .b32 param2;
  3600. st.param.f32 [param2+0], %f1163;
  3601. .param .b32 param3;
  3602. st.param.f32 [param3+0], %f1394;
  3603. .param .b32 param4;
  3604. st.param.f32 [param4+0], %f1395;
  3605. .param .b32 param5;
  3606. st.param.f32 [param5+0], %f1396;
  3607. .param .b64 param6;
  3608. st.param.b64 [param6+0], %rd153;
  3609. .param .b32 param7;
  3610. st.param.b32 [param7+0], %r41;
  3611. .param .b32 param8;
  3612. st.param.b32 [param8+0], %r172;
  3613. .param .align 4 .b8 retval0[12];
  3614. call.uni (retval0),
  3615. _Z5traceRK4Vec3IfES2_PK6SpherejRKi,
  3616. (
  3617. param0,
  3618. param1,
  3619. param2,
  3620. param3,
  3621. param4,
  3622. param5,
  3623. param6,
  3624. param7,
  3625. param8
  3626. );
  3627. ld.param.f32 %f516, [retval0+0];
  3628. ld.param.f32 %f517, [retval0+4];
  3629. ld.param.f32 %f518, [retval0+8];
  3630.  
  3631. //{
  3632. }// Callseq End 8
  3633. mov.f32 %f1400, 0f00000000;
  3634. setp.eq.f32 %p302, %f471, 0f00000000;
  3635. mov.f32 %f1401, %f1400;
  3636. mov.f32 %f1402, %f1400;
  3637. @%p302 bra BB1_312;
  3638.  
  3639. setp.eq.s16 %p303, %rs12, 0;
  3640. selp.f32 %f1164, 0f3F68BA2E, 0f3F8CCCCD, %p303;
  3641. mul.f32 %f1165, %f1164, %f1164;
  3642. mul.f32 %f1166, %f500, %f500;
  3643. mov.f32 %f1167, 0f3F800000;
  3644. sub.f32 %f1168, %f1167, %f1166;
  3645. mul.f32 %f1169, %f1165, %f1168;
  3646. sub.f32 %f1170, %f1167, %f1169;
  3647. mul.f32 %f1171, %f500, %f1164;
  3648. neg.f32 %f1172, %f1171;
  3649. sqrt.rn.f32 %f1173, %f1170;
  3650. sub.f32 %f1174, %f1172, %f1173;
  3651. mul.f32 %f1175, %f1378, %f1174;
  3652. mul.f32 %f1176, %f1379, %f1174;
  3653. mul.f32 %f1177, %f1380, %f1174;
  3654. fma.rn.f32 %f1399, %f1351, %f1164, %f1175;
  3655. fma.rn.f32 %f1398, %f1350, %f1164, %f1176;
  3656. fma.rn.f32 %f1397, %f1349, %f1164, %f1177;
  3657. mul.f32 %f1178, %f1398, %f1398;
  3658. fma.rn.f32 %f1179, %f1399, %f1399, %f1178;
  3659. fma.rn.f32 %f522, %f1397, %f1397, %f1179;
  3660. setp.leu.f32 %p304, %f522, 0f00000000;
  3661. @%p304 bra BB1_311;
  3662.  
  3663. sqrt.rn.f32 %f1180, %f522;
  3664. rcp.rn.f32 %f1181, %f1180;
  3665. mul.f32 %f1399, %f1399, %f1181;
  3666. mul.f32 %f1398, %f1398, %f1181;
  3667. mul.f32 %f1397, %f1397, %f1181;
  3668.  
  3669. BB1_311:
  3670. sub.f32 %f1182, %f454, %f515;
  3671. sub.f32 %f1183, %f453, %f514;
  3672. sub.f32 %f1184, %f452, %f513;
  3673. // Callseq Start 9
  3674. {
  3675. .reg .b32 temp_param_reg;
  3676. // <end>}
  3677. .param .b32 param0;
  3678. st.param.f32 [param0+0], %f1184;
  3679. .param .b32 param1;
  3680. st.param.f32 [param1+0], %f1183;
  3681. .param .b32 param2;
  3682. st.param.f32 [param2+0], %f1182;
  3683. .param .b32 param3;
  3684. st.param.f32 [param3+0], %f1399;
  3685. .param .b32 param4;
  3686. st.param.f32 [param4+0], %f1398;
  3687. .param .b32 param5;
  3688. st.param.f32 [param5+0], %f1397;
  3689. .param .b64 param6;
  3690. st.param.b64 [param6+0], %rd153;
  3691. .param .b32 param7;
  3692. st.param.b32 [param7+0], %r41;
  3693. .param .b32 param8;
  3694. st.param.b32 [param8+0], %r172;
  3695. .param .align 4 .b8 retval0[12];
  3696. call.uni (retval0),
  3697. _Z5traceRK4Vec3IfES2_PK6SpherejRKi,
  3698. (
  3699. param0,
  3700. param1,
  3701. param2,
  3702. param3,
  3703. param4,
  3704. param5,
  3705. param6,
  3706. param7,
  3707. param8
  3708. );
  3709. ld.param.f32 %f1400, [retval0+0];
  3710. ld.param.f32 %f1401, [retval0+4];
  3711. ld.param.f32 %f1402, [retval0+8];
  3712.  
  3713. //{
  3714. }// Callseq End 9
  3715.  
  3716. BB1_312:
  3717. mov.f32 %f1185, 0f3F800000;
  3718. sub.f32 %f1186, %f1185, %f502;
  3719. mul.f32 %f1187, %f1186, %f1400;
  3720. mul.f32 %f1188, %f1186, %f1401;
  3721. mul.f32 %f1189, %f1186, %f1402;
  3722. mul.f32 %f1190, %f471, %f1187;
  3723. mul.f32 %f1191, %f471, %f1188;
  3724. mul.f32 %f1192, %f471, %f1189;
  3725. fma.rn.f32 %f1193, %f502, %f516, %f1190;
  3726. fma.rn.f32 %f1194, %f502, %f517, %f1191;
  3727. fma.rn.f32 %f1195, %f502, %f518, %f1192;
  3728. ld.global.f32 %f1196, [%rd148+20];
  3729. mul.f32 %f1403, %f1196, %f1193;
  3730. ld.global.f32 %f1197, [%rd148+24];
  3731. mul.f32 %f1404, %f1194, %f1197;
  3732. ld.global.f32 %f1198, [%rd148+28];
  3733. mul.f32 %f1405, %f1195, %f1198;
  3734.  
  3735. BB1_313:
  3736. ld.global.f32 %f1199, [%rd148+32];
  3737. add.f32 %f1406, %f1403, %f1199;
  3738. ld.global.f32 %f1200, [%rd148+36];
  3739. add.f32 %f1407, %f1404, %f1200;
  3740. ld.global.f32 %f1201, [%rd148+40];
  3741. add.f32 %f1408, %f1405, %f1201;
  3742.  
  3743. BB1_314:
  3744. mov.f32 %f1202, 0f3F800000;
  3745. sub.f32 %f1203, %f1202, %f149;
  3746. mul.f32 %f1204, %f1203, %f1406;
  3747. mul.f32 %f1205, %f1203, %f1407;
  3748. mul.f32 %f1206, %f1203, %f1408;
  3749. mul.f32 %f1207, %f1204, %f118;
  3750. mul.f32 %f1208, %f1205, %f118;
  3751. mul.f32 %f1209, %f1206, %f118;
  3752. fma.rn.f32 %f1210, %f149, %f1346, %f1207;
  3753. fma.rn.f32 %f1211, %f149, %f1347, %f1208;
  3754. fma.rn.f32 %f1212, %f149, %f1348, %f1209;
  3755. ld.global.f32 %f1213, [%rd47+20];
  3756. mul.f32 %f1409, %f1213, %f1210;
  3757. ld.global.f32 %f1214, [%rd47+24];
  3758. mul.f32 %f1410, %f1211, %f1214;
  3759. ld.global.f32 %f1215, [%rd47+28];
  3760. mul.f32 %f1411, %f1212, %f1215;
  3761.  
  3762. BB1_315:
  3763. ld.global.f32 %f1216, [%rd47+32];
  3764. add.f32 %f1412, %f1409, %f1216;
  3765. ld.global.f32 %f1217, [%rd47+36];
  3766. add.f32 %f1413, %f1410, %f1217;
  3767. ld.global.f32 %f1218, [%rd47+40];
  3768. add.f32 %f1414, %f1411, %f1218;
  3769.  
  3770. BB1_316:
  3771. ld.param.u64 %rd267, [raytrace_kernel_param_2];
  3772. ld.param.u32 %r191, [raytrace_kernel_param_0];
  3773. mov.u32 %r190, %tid.y;
  3774. mov.u32 %r189, %ctaid.y;
  3775. mov.u32 %r188, %ntid.y;
  3776. mad.lo.s32 %r187, %r188, %r189, %r190;
  3777. mov.u32 %r186, %tid.x;
  3778. mov.u32 %r185, %ctaid.x;
  3779. mov.u32 %r184, %ntid.x;
  3780. mad.lo.s32 %r183, %r184, %r185, %r186;
  3781. mad.lo.s32 %r182, %r187, %r191, %r183;
  3782. cvta.to.global.u64 %rd259, %rd267;
  3783. mul.wide.u32 %rd260, %r182, 12;
  3784. add.s64 %rd261, %rd259, %rd260;
  3785. st.global.f32 [%rd261], %f1412;
  3786. st.global.f32 [%rd261+4], %f1413;
  3787. st.global.f32 [%rd261+8], %f1414;
  3788.  
  3789. BB1_317:
  3790. ret;
  3791. }
  3792.  
  3793. .func (.param .b64 func_retval0) __internal_accurate_pow(
  3794. .param .b64 __internal_accurate_pow_param_0
  3795. )
  3796. {
  3797. .reg .pred %p<9>;
  3798. .reg .f32 %f<3>;
  3799. .reg .b32 %r<53>;
  3800. .reg .f64 %fd<138>;
  3801.  
  3802.  
  3803. ld.param.f64 %fd12, [__internal_accurate_pow_param_0];
  3804. {
  3805. .reg .b32 %temp;
  3806. mov.b64 {%temp, %r50}, %fd12;
  3807. }
  3808. {
  3809. .reg .b32 %temp;
  3810. mov.b64 {%r49, %temp}, %fd12;
  3811. }
  3812. shr.u32 %r51, %r50, 20;
  3813. setp.ne.s32 %p1, %r51, 0;
  3814. @%p1 bra BB2_2;
  3815.  
  3816. mul.f64 %fd13, %fd12, 0d4350000000000000;
  3817. {
  3818. .reg .b32 %temp;
  3819. mov.b64 {%temp, %r50}, %fd13;
  3820. }
  3821. {
  3822. .reg .b32 %temp;
  3823. mov.b64 {%r49, %temp}, %fd13;
  3824. }
  3825. shr.u32 %r16, %r50, 20;
  3826. add.s32 %r51, %r16, -54;
  3827.  
  3828. BB2_2:
  3829. add.s32 %r52, %r51, -1023;
  3830. and.b32 %r17, %r50, -2146435073;
  3831. or.b32 %r18, %r17, 1072693248;
  3832. mov.b64 %fd135, {%r49, %r18};
  3833. setp.lt.u32 %p2, %r18, 1073127583;
  3834. @%p2 bra BB2_4;
  3835.  
  3836. {
  3837. .reg .b32 %temp;
  3838. mov.b64 {%r19, %temp}, %fd135;
  3839. }
  3840. {
  3841. .reg .b32 %temp;
  3842. mov.b64 {%temp, %r20}, %fd135;
  3843. }
  3844. add.s32 %r21, %r20, -1048576;
  3845. mov.b64 %fd135, {%r19, %r21};
  3846. add.s32 %r52, %r51, -1022;
  3847.  
  3848. BB2_4:
  3849. add.f64 %fd14, %fd135, 0d3FF0000000000000;
  3850. rcp.approx.ftz.f64 %fd15, %fd14;
  3851. neg.f64 %fd16, %fd14;
  3852. mov.f64 %fd17, 0d3FF0000000000000;
  3853. fma.rn.f64 %fd18, %fd16, %fd15, %fd17;
  3854. fma.rn.f64 %fd19, %fd18, %fd18, %fd18;
  3855. fma.rn.f64 %fd20, %fd19, %fd15, %fd15;
  3856. add.f64 %fd21, %fd135, 0dBFF0000000000000;
  3857. mul.f64 %fd22, %fd21, %fd20;
  3858. fma.rn.f64 %fd23, %fd21, %fd20, %fd22;
  3859. mul.f64 %fd24, %fd23, %fd23;
  3860. mov.f64 %fd25, 0d3ED0F5D241AD3B5A;
  3861. mov.f64 %fd26, 0d3EB0F5FF7D2CAFE2;
  3862. fma.rn.f64 %fd27, %fd26, %fd24, %fd25;
  3863. mov.f64 %fd28, 0d3EF3B20A75488A3F;
  3864. fma.rn.f64 %fd29, %fd27, %fd24, %fd28;
  3865. mov.f64 %fd30, 0d3F1745CDE4FAECD5;
  3866. fma.rn.f64 %fd31, %fd29, %fd24, %fd30;
  3867. mov.f64 %fd32, 0d3F3C71C7258A578B;
  3868. fma.rn.f64 %fd33, %fd31, %fd24, %fd32;
  3869. mov.f64 %fd34, 0d3F6249249242B910;
  3870. fma.rn.f64 %fd35, %fd33, %fd24, %fd34;
  3871. mov.f64 %fd36, 0d3F89999999999DFB;
  3872. fma.rn.f64 %fd37, %fd35, %fd24, %fd36;
  3873. sub.f64 %fd38, %fd21, %fd23;
  3874. add.f64 %fd39, %fd38, %fd38;
  3875. neg.f64 %fd40, %fd23;
  3876. fma.rn.f64 %fd41, %fd40, %fd21, %fd39;
  3877. mul.f64 %fd42, %fd20, %fd41;
  3878. fma.rn.f64 %fd43, %fd24, %fd37, 0d3FB5555555555555;
  3879. mov.f64 %fd44, 0d3FB5555555555555;
  3880. sub.f64 %fd45, %fd44, %fd43;
  3881. fma.rn.f64 %fd46, %fd24, %fd37, %fd45;
  3882. add.f64 %fd47, %fd46, 0d0000000000000000;
  3883. add.f64 %fd48, %fd47, 0dBC46A4CB00B9E7B0;
  3884. add.f64 %fd49, %fd43, %fd48;
  3885. sub.f64 %fd50, %fd43, %fd49;
  3886. add.f64 %fd51, %fd48, %fd50;
  3887. mul.rn.f64 %fd52, %fd23, %fd23;
  3888. neg.f64 %fd53, %fd52;
  3889. fma.rn.f64 %fd54, %fd23, %fd23, %fd53;
  3890. {
  3891. .reg .b32 %temp;
  3892. mov.b64 {%r22, %temp}, %fd42;
  3893. }
  3894. {
  3895. .reg .b32 %temp;
  3896. mov.b64 {%temp, %r23}, %fd42;
  3897. }
  3898. add.s32 %r24, %r23, 1048576;
  3899. mov.b64 %fd55, {%r22, %r24};
  3900. fma.rn.f64 %fd56, %fd23, %fd55, %fd54;
  3901. mul.rn.f64 %fd57, %fd52, %fd23;
  3902. neg.f64 %fd58, %fd57;
  3903. fma.rn.f64 %fd59, %fd52, %fd23, %fd58;
  3904. fma.rn.f64 %fd60, %fd52, %fd42, %fd59;
  3905. fma.rn.f64 %fd61, %fd56, %fd23, %fd60;
  3906. mul.rn.f64 %fd62, %fd49, %fd57;
  3907. neg.f64 %fd63, %fd62;
  3908. fma.rn.f64 %fd64, %fd49, %fd57, %fd63;
  3909. fma.rn.f64 %fd65, %fd49, %fd61, %fd64;
  3910. fma.rn.f64 %fd66, %fd51, %fd57, %fd65;
  3911. add.f64 %fd67, %fd62, %fd66;
  3912. sub.f64 %fd68, %fd62, %fd67;
  3913. add.f64 %fd69, %fd66, %fd68;
  3914. add.f64 %fd70, %fd23, %fd67;
  3915. sub.f64 %fd71, %fd23, %fd70;
  3916. add.f64 %fd72, %fd67, %fd71;
  3917. add.f64 %fd73, %fd69, %fd72;
  3918. add.f64 %fd74, %fd42, %fd73;
  3919. add.f64 %fd75, %fd70, %fd74;
  3920. sub.f64 %fd76, %fd70, %fd75;
  3921. add.f64 %fd77, %fd74, %fd76;
  3922. xor.b32 %r25, %r52, -2147483648;
  3923. mov.u32 %r26, -2147483648;
  3924. mov.u32 %r27, 1127219200;
  3925. mov.b64 %fd78, {%r25, %r27};
  3926. mov.b64 %fd79, {%r26, %r27};
  3927. sub.f64 %fd80, %fd78, %fd79;
  3928. mov.f64 %fd81, 0d3FE62E42FEFA39EF;
  3929. fma.rn.f64 %fd82, %fd80, %fd81, %fd75;
  3930. neg.f64 %fd83, %fd80;
  3931. fma.rn.f64 %fd84, %fd83, %fd81, %fd82;
  3932. sub.f64 %fd85, %fd84, %fd75;
  3933. sub.f64 %fd86, %fd77, %fd85;
  3934. mov.f64 %fd87, 0d3C7ABC9E3B39803F;
  3935. fma.rn.f64 %fd88, %fd80, %fd87, %fd86;
  3936. add.f64 %fd89, %fd82, %fd88;
  3937. sub.f64 %fd90, %fd82, %fd89;
  3938. add.f64 %fd91, %fd88, %fd90;
  3939. mov.f64 %fd92, 0d4008000000000000;
  3940. {
  3941. .reg .b32 %temp;
  3942. mov.b64 {%temp, %r28}, %fd92;
  3943. }
  3944. add.s32 %r29, %r28, %r28;
  3945. setp.gt.u32 %p3, %r29, -33554433;
  3946. and.b32 %r30, %r28, -15728641;
  3947. selp.b32 %r31, %r30, %r28, %p3;
  3948. {
  3949. .reg .b32 %temp;
  3950. mov.b64 {%r32, %temp}, %fd92;
  3951. }
  3952. mov.b64 %fd93, {%r32, %r31};
  3953. mul.rn.f64 %fd94, %fd89, %fd93;
  3954. neg.f64 %fd95, %fd94;
  3955. fma.rn.f64 %fd96, %fd89, %fd93, %fd95;
  3956. fma.rn.f64 %fd97, %fd91, %fd93, %fd96;
  3957. add.f64 %fd4, %fd94, %fd97;
  3958. sub.f64 %fd98, %fd94, %fd4;
  3959. add.f64 %fd5, %fd97, %fd98;
  3960. mov.f64 %fd99, 0d4338000000000000;
  3961. mov.f64 %fd100, 0d3FF71547652B82FE;
  3962. fma.rn.f64 %fd101, %fd4, %fd100, %fd99;
  3963. {
  3964. .reg .b32 %temp;
  3965. mov.b64 {%r13, %temp}, %fd101;
  3966. }
  3967. mov.f64 %fd102, 0dC338000000000000;
  3968. add.rn.f64 %fd103, %fd101, %fd102;
  3969. mov.f64 %fd104, 0dBFE62E42FEFA39EF;
  3970. fma.rn.f64 %fd105, %fd103, %fd104, %fd4;
  3971. mov.f64 %fd106, 0dBC7ABC9E3B39803F;
  3972. fma.rn.f64 %fd107, %fd103, %fd106, %fd105;
  3973. mov.f64 %fd108, 0d3E928AF3FCA213EA;
  3974. mov.f64 %fd109, 0d3E5ADE1569CE2BDF;
  3975. fma.rn.f64 %fd110, %fd109, %fd107, %fd108;
  3976. mov.f64 %fd111, 0d3EC71DEE62401315;
  3977. fma.rn.f64 %fd112, %fd110, %fd107, %fd111;
  3978. mov.f64 %fd113, 0d3EFA01997C89EB71;
  3979. fma.rn.f64 %fd114, %fd112, %fd107, %fd113;
  3980. mov.f64 %fd115, 0d3F2A01A014761F65;
  3981. fma.rn.f64 %fd116, %fd114, %fd107, %fd115;
  3982. mov.f64 %fd117, 0d3F56C16C1852B7AF;
  3983. fma.rn.f64 %fd118, %fd116, %fd107, %fd117;
  3984. mov.f64 %fd119, 0d3F81111111122322;
  3985. fma.rn.f64 %fd120, %fd118, %fd107, %fd119;
  3986. mov.f64 %fd121, 0d3FA55555555502A1;
  3987. fma.rn.f64 %fd122, %fd120, %fd107, %fd121;
  3988. mov.f64 %fd123, 0d3FC5555555555511;
  3989. fma.rn.f64 %fd124, %fd122, %fd107, %fd123;
  3990. mov.f64 %fd125, 0d3FE000000000000B;
  3991. fma.rn.f64 %fd126, %fd124, %fd107, %fd125;
  3992. fma.rn.f64 %fd127, %fd126, %fd107, %fd17;
  3993. fma.rn.f64 %fd128, %fd127, %fd107, %fd17;
  3994. {
  3995. .reg .b32 %temp;
  3996. mov.b64 {%r14, %temp}, %fd128;
  3997. }
  3998. {
  3999. .reg .b32 %temp;
  4000. mov.b64 {%temp, %r15}, %fd128;
  4001. }
  4002. shl.b32 %r33, %r13, 20;
  4003. add.s32 %r34, %r15, %r33;
  4004. mov.b64 %fd136, {%r14, %r34};
  4005. {
  4006. .reg .b32 %temp;
  4007. mov.b64 {%temp, %r35}, %fd4;
  4008. }
  4009. mov.b32 %f2, %r35;
  4010. abs.f32 %f1, %f2;
  4011. setp.lt.f32 %p4, %f1, 0f4086232B;
  4012. @%p4 bra BB2_7;
  4013.  
  4014. setp.lt.f64 %p5, %fd4, 0d0000000000000000;
  4015. add.f64 %fd129, %fd4, 0d7FF0000000000000;
  4016. selp.f64 %fd136, 0d0000000000000000, %fd129, %p5;
  4017. setp.geu.f32 %p6, %f1, 0f40874800;
  4018. @%p6 bra BB2_7;
  4019.  
  4020. mov.f64 %fd134, 0d4338000000000000;
  4021. mov.f64 %fd133, 0d3FF71547652B82FE;
  4022. fma.rn.f64 %fd132, %fd4, %fd133, %fd134;
  4023. {
  4024. .reg .b32 %temp;
  4025. mov.b64 {%r48, %temp}, %fd132;
  4026. }
  4027. shr.u32 %r36, %r48, 31;
  4028. add.s32 %r37, %r48, %r36;
  4029. shr.s32 %r38, %r37, 1;
  4030. shl.b32 %r39, %r38, 20;
  4031. add.s32 %r40, %r39, %r15;
  4032. mov.b64 %fd130, {%r14, %r40};
  4033. sub.s32 %r41, %r48, %r38;
  4034. shl.b32 %r42, %r41, 20;
  4035. add.s32 %r43, %r42, 1072693248;
  4036. mov.u32 %r44, 0;
  4037. mov.b64 %fd131, {%r44, %r43};
  4038. mul.f64 %fd136, %fd130, %fd131;
  4039.  
  4040. BB2_7:
  4041. {
  4042. .reg .b32 %temp;
  4043. mov.b64 {%temp, %r45}, %fd136;
  4044. }
  4045. and.b32 %r46, %r45, 2147483647;
  4046. setp.ne.s32 %p7, %r46, 2146435072;
  4047. @%p7 bra BB2_9;
  4048.  
  4049. {
  4050. .reg .b32 %temp;
  4051. mov.b64 {%r47, %temp}, %fd136;
  4052. }
  4053. setp.eq.s32 %p8, %r47, 0;
  4054. @%p8 bra BB2_10;
  4055.  
  4056. BB2_9:
  4057. fma.rn.f64 %fd136, %fd136, %fd5, %fd136;
  4058.  
  4059. BB2_10:
  4060. st.param.f64 [func_retval0+0], %fd136;
  4061. ret;
  4062. }
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement