You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@systemml.apache.org by ni...@apache.org on 2018/09/18 23:22:58 UTC
[2/3] systemml git commit: [SYSTEMML-445] Added sparse scalar-matrix
arithmetic/relational operators
http://git-wip-us.apache.org/repos/asf/systemml/blob/61139e40/src/main/cpp/kernels/SystemML.ptx
----------------------------------------------------------------------
diff --git a/src/main/cpp/kernels/SystemML.ptx b/src/main/cpp/kernels/SystemML.ptx
index 1ab32f5..ac04967 100644
--- a/src/main/cpp/kernels/SystemML.ptx
+++ b/src/main/cpp/kernels/SystemML.ptx
@@ -4595,6 +4595,1739 @@ BB31_126:
ret;
}
+ // .globl sparse_dense_matrix_scalar_op_d
+.visible .entry sparse_dense_matrix_scalar_op_d(
+ .param .u64 sparse_dense_matrix_scalar_op_d_param_0,
+ .param .u64 sparse_dense_matrix_scalar_op_d_param_1,
+ .param .u64 sparse_dense_matrix_scalar_op_d_param_2,
+ .param .f64 sparse_dense_matrix_scalar_op_d_param_3,
+ .param .u64 sparse_dense_matrix_scalar_op_d_param_4,
+ .param .u32 sparse_dense_matrix_scalar_op_d_param_5,
+ .param .u32 sparse_dense_matrix_scalar_op_d_param_6,
+ .param .u32 sparse_dense_matrix_scalar_op_d_param_7,
+ .param .u32 sparse_dense_matrix_scalar_op_d_param_8
+)
+{
+ .reg .pred %p<133>;
+ .reg .b32 %r<92>;
+ .reg .f64 %fd<99>;
+ .reg .b64 %rd<28>;
+
+
+ ld.param.u64 %rd4, [sparse_dense_matrix_scalar_op_d_param_0];
+ ld.param.u64 %rd5, [sparse_dense_matrix_scalar_op_d_param_1];
+ ld.param.u64 %rd6, [sparse_dense_matrix_scalar_op_d_param_2];
+ ld.param.f64 %fd68, [sparse_dense_matrix_scalar_op_d_param_3];
+ ld.param.u64 %rd7, [sparse_dense_matrix_scalar_op_d_param_4];
+ ld.param.u32 %r9, [sparse_dense_matrix_scalar_op_d_param_5];
+ ld.param.u32 %r6, [sparse_dense_matrix_scalar_op_d_param_6];
+ ld.param.u32 %r7, [sparse_dense_matrix_scalar_op_d_param_7];
+ ld.param.u32 %r8, [sparse_dense_matrix_scalar_op_d_param_8];
+ mov.u32 %r10, %ntid.x;
+ mov.u32 %r11, %ctaid.x;
+ mov.u32 %r12, %tid.x;
+ mad.lo.s32 %r1, %r10, %r11, %r12;
+ setp.ge.s32 %p3, %r1, %r9;
+ @%p3 bra BB32_142;
+
+ cvta.to.global.u64 %rd8, %rd7;
+ cvta.to.global.u64 %rd9, %rd6;
+ mul.wide.s32 %rd10, %r1, 8;
+ add.s64 %rd11, %rd9, %rd10;
+ ld.global.f64 %fd1, [%rd11];
+ cvta.to.global.u64 %rd12, %rd4;
+ mul.wide.s32 %rd13, %r1, 4;
+ add.s64 %rd14, %rd12, %rd13;
+ ld.global.u32 %r13, [%rd14];
+ cvta.to.global.u64 %rd15, %rd5;
+ add.s64 %rd16, %rd15, %rd13;
+ ld.global.u32 %r14, [%rd16];
+ mad.lo.s32 %r15, %r13, %r6, %r14;
+ mul.wide.s32 %rd17, %r15, 8;
+ add.s64 %rd1, %rd8, %rd17;
+ setp.eq.s32 %p4, %r8, 0;
+ @%p4 bra BB32_72;
+
+ mov.f64 %fd94, 0d7FEFFFFFFFFFFFFF;
+ setp.gt.s32 %p5, %r7, 8;
+ @%p5 bra BB32_19;
+
+ setp.gt.s32 %p19, %r7, 3;
+ @%p19 bra BB32_11;
+
+ setp.gt.s32 %p26, %r7, 1;
+ @%p26 bra BB32_8;
+
+ setp.eq.s32 %p29, %r7, 0;
+ @%p29 bra BB32_70;
+ bra.uni BB32_6;
+
+BB32_70:
+ add.f64 %fd94, %fd1, %fd68;
+ bra.uni BB32_71;
+
+BB32_72:
+ mov.f64 %fd98, 0d7FEFFFFFFFFFFFFF;
+ setp.gt.s32 %p69, %r7, 8;
+ @%p69 bra BB32_89;
+
+ setp.gt.s32 %p83, %r7, 3;
+ @%p83 bra BB32_81;
+
+ setp.gt.s32 %p90, %r7, 1;
+ @%p90 bra BB32_78;
+
+ setp.eq.s32 %p93, %r7, 0;
+ @%p93 bra BB32_140;
+ bra.uni BB32_76;
+
+BB32_140:
+ add.f64 %fd98, %fd1, %fd68;
+ bra.uni BB32_141;
+
+BB32_19:
+ setp.gt.s32 %p6, %r7, 13;
+ @%p6 bra BB32_28;
+
+ setp.gt.s32 %p13, %r7, 10;
+ @%p13 bra BB32_24;
+
+ setp.eq.s32 %p17, %r7, 9;
+ @%p17 bra BB32_48;
+ bra.uni BB32_22;
+
+BB32_48:
+ setp.eq.f64 %p44, %fd1, %fd68;
+ selp.f64 %fd94, 0d3FF0000000000000, 0d0000000000000000, %p44;
+ bra.uni BB32_71;
+
+BB32_89:
+ setp.gt.s32 %p70, %r7, 13;
+ @%p70 bra BB32_98;
+
+ setp.gt.s32 %p77, %r7, 10;
+ @%p77 bra BB32_94;
+
+ setp.eq.s32 %p81, %r7, 9;
+ @%p81 bra BB32_118;
+ bra.uni BB32_92;
+
+BB32_118:
+ setp.eq.f64 %p108, %fd1, %fd68;
+ selp.f64 %fd98, 0d3FF0000000000000, 0d0000000000000000, %p108;
+ bra.uni BB32_141;
+
+BB32_11:
+ setp.gt.s32 %p20, %r7, 5;
+ @%p20 bra BB32_15;
+
+ setp.eq.s32 %p24, %r7, 4;
+ @%p24 bra BB32_51;
+ bra.uni BB32_13;
+
+BB32_51:
+ {
+ .reg .b32 %temp;
+ mov.b64 {%temp, %r2}, %fd68;
+ }
+ {
+ .reg .b32 %temp;
+ mov.b64 {%temp, %r3}, %fd1;
+ }
+ bfe.u32 %r28, %r3, 20, 11;
+ add.s32 %r29, %r28, -1012;
+ mov.b64 %rd22, %fd1;
+ shl.b64 %rd2, %rd22, %r29;
+ setp.eq.s64 %p49, %rd2, -9223372036854775808;
+ abs.f64 %fd18, %fd68;
+ // Callseq Start 3
+ {
+ .reg .b32 temp_param_reg;
+ // <end>}
+ .param .b64 param0;
+ st.param.f64 [param0+0], %fd18;
+ .param .b64 param1;
+ st.param.f64 [param1+0], %fd1;
+ .param .b64 retval0;
+ call.uni (retval0),
+ __internal_accurate_pow,
+ (
+ param0,
+ param1
+ );
+ ld.param.f64 %fd24, [retval0+0];
+
+ //{
+ }// Callseq End 3
+ setp.lt.s32 %p50, %r2, 0;
+ and.pred %p1, %p50, %p49;
+ @!%p1 bra BB32_53;
+ bra.uni BB32_52;
+
+BB32_52:
+ {
+ .reg .b32 %temp;
+ mov.b64 {%temp, %r30}, %fd24;
+ }
+ xor.b32 %r31, %r30, -2147483648;
+ {
+ .reg .b32 %temp;
+ mov.b64 {%r32, %temp}, %fd24;
+ }
+ mov.b64 %fd24, {%r32, %r31};
+
+BB32_53:
+ setp.eq.f64 %p51, %fd68, 0d0000000000000000;
+ @%p51 bra BB32_56;
+ bra.uni BB32_54;
+
+BB32_56:
+ selp.b32 %r33, %r2, 0, %p49;
+ or.b32 %r34, %r33, 2146435072;
+ setp.lt.s32 %p55, %r3, 0;
+ selp.b32 %r35, %r34, %r33, %p55;
+ mov.u32 %r36, 0;
+ mov.b64 %fd24, {%r36, %r35};
+ bra.uni BB32_57;
+
+BB32_28:
+ setp.gt.s32 %p7, %r7, 15;
+ @%p7 bra BB32_32;
+
+ setp.eq.s32 %p11, %r7, 14;
+ @%p11 bra BB32_45;
+ bra.uni BB32_30;
+
+BB32_45:
+ cvt.rni.s64.f64 %rd18, %fd68;
+ cvt.u32.u64 %r22, %rd18;
+ cvt.rni.s64.f64 %rd19, %fd1;
+ cvt.u32.u64 %r23, %rd19;
+ or.b32 %r24, %r23, %r22;
+ setp.eq.s32 %p41, %r24, 0;
+ selp.f64 %fd94, 0d0000000000000000, 0d3FF0000000000000, %p41;
+ bra.uni BB32_71;
+
+BB32_81:
+ setp.gt.s32 %p84, %r7, 5;
+ @%p84 bra BB32_85;
+
+ setp.eq.s32 %p88, %r7, 4;
+ @%p88 bra BB32_121;
+ bra.uni BB32_83;
+
+BB32_121:
+ {
+ .reg .b32 %temp;
+ mov.b64 {%temp, %r4}, %fd1;
+ }
+ {
+ .reg .b32 %temp;
+ mov.b64 {%temp, %r5}, %fd68;
+ }
+ bfe.u32 %r66, %r5, 20, 11;
+ add.s32 %r67, %r66, -1012;
+ mov.b64 %rd27, %fd68;
+ shl.b64 %rd3, %rd27, %r67;
+ setp.eq.s64 %p113, %rd3, -9223372036854775808;
+ abs.f64 %fd51, %fd1;
+ // Callseq Start 4
+ {
+ .reg .b32 temp_param_reg;
+ // <end>}
+ .param .b64 param0;
+ st.param.f64 [param0+0], %fd51;
+ .param .b64 param1;
+ st.param.f64 [param1+0], %fd68;
+ .param .b64 retval0;
+ call.uni (retval0),
+ __internal_accurate_pow,
+ (
+ param0,
+ param1
+ );
+ ld.param.f64 %fd57, [retval0+0];
+
+ //{
+ }// Callseq End 4
+ setp.lt.s32 %p114, %r4, 0;
+ and.pred %p2, %p114, %p113;
+ @!%p2 bra BB32_123;
+ bra.uni BB32_122;
+
+BB32_122:
+ {
+ .reg .b32 %temp;
+ mov.b64 {%temp, %r68}, %fd57;
+ }
+ xor.b32 %r69, %r68, -2147483648;
+ {
+ .reg .b32 %temp;
+ mov.b64 {%r70, %temp}, %fd57;
+ }
+ mov.b64 %fd57, {%r70, %r69};
+
+BB32_123:
+ setp.eq.f64 %p115, %fd1, 0d0000000000000000;
+ @%p115 bra BB32_126;
+ bra.uni BB32_124;
+
+BB32_126:
+ selp.b32 %r71, %r4, 0, %p113;
+ or.b32 %r72, %r71, 2146435072;
+ setp.lt.s32 %p119, %r5, 0;
+ selp.b32 %r73, %r72, %r71, %p119;
+ mov.u32 %r74, 0;
+ mov.b64 %fd57, {%r74, %r73};
+ bra.uni BB32_127;
+
+BB32_98:
+ setp.gt.s32 %p71, %r7, 15;
+ @%p71 bra BB32_102;
+
+ setp.eq.s32 %p75, %r7, 14;
+ @%p75 bra BB32_115;
+ bra.uni BB32_100;
+
+BB32_115:
+ cvt.rni.s64.f64 %rd23, %fd1;
+ cvt.u32.u64 %r60, %rd23;
+ cvt.rni.s64.f64 %rd24, %fd68;
+ cvt.u32.u64 %r61, %rd24;
+ or.b32 %r62, %r61, %r60;
+ setp.eq.s32 %p105, %r62, 0;
+ selp.f64 %fd98, 0d0000000000000000, 0d3FF0000000000000, %p105;
+ bra.uni BB32_141;
+
+BB32_8:
+ setp.eq.s32 %p27, %r7, 2;
+ @%p27 bra BB32_69;
+ bra.uni BB32_9;
+
+BB32_69:
+ mul.f64 %fd94, %fd1, %fd68;
+ bra.uni BB32_71;
+
+BB32_24:
+ setp.eq.s32 %p14, %r7, 11;
+ @%p14 bra BB32_47;
+
+ setp.eq.s32 %p15, %r7, 12;
+ @%p15 bra BB32_46;
+ bra.uni BB32_26;
+
+BB32_46:
+ max.f64 %fd94, %fd68, %fd1;
+ bra.uni BB32_71;
+
+BB32_15:
+ setp.eq.s32 %p21, %r7, 6;
+ @%p21 bra BB32_50;
+
+ setp.eq.s32 %p22, %r7, 7;
+ @%p22 bra BB32_49;
+ bra.uni BB32_17;
+
+BB32_49:
+ setp.lt.f64 %p46, %fd1, %fd68;
+ selp.f64 %fd94, 0d3FF0000000000000, 0d0000000000000000, %p46;
+ bra.uni BB32_71;
+
+BB32_32:
+ setp.eq.s32 %p8, %r7, 16;
+ @%p8 bra BB32_44;
+
+ setp.eq.s32 %p9, %r7, 17;
+ @%p9 bra BB32_39;
+ bra.uni BB32_34;
+
+BB32_39:
+ setp.eq.f64 %p34, %fd1, 0d0000000000000000;
+ setp.eq.f64 %p35, %fd1, 0d8000000000000000;
+ or.pred %p36, %p34, %p35;
+ mov.f64 %fd94, 0d7FF8000000000000;
+ @%p36 bra BB32_71;
+
+ div.rn.f64 %fd94, %fd68, %fd1;
+ abs.f64 %fd72, %fd94;
+ setp.gtu.f64 %p37, %fd72, 0d7FF0000000000000;
+ @%p37 bra BB32_71;
+
+ {
+ .reg .b32 %temp;
+ mov.b64 {%temp, %r19}, %fd94;
+ }
+ and.b32 %r20, %r19, 2147483647;
+ setp.ne.s32 %p38, %r20, 2146435072;
+ @%p38 bra BB32_43;
+
+ {
+ .reg .b32 %temp;
+ mov.b64 {%r21, %temp}, %fd94;
+ }
+ setp.eq.s32 %p39, %r21, 0;
+ @%p39 bra BB32_71;
+
+BB32_43:
+ cvt.rmi.f64.f64 %fd73, %fd94;
+ mul.f64 %fd74, %fd1, %fd73;
+ sub.f64 %fd94, %fd68, %fd74;
+ bra.uni BB32_71;
+
+BB32_78:
+ setp.eq.s32 %p91, %r7, 2;
+ @%p91 bra BB32_139;
+ bra.uni BB32_79;
+
+BB32_139:
+ mul.f64 %fd98, %fd1, %fd68;
+ bra.uni BB32_141;
+
+BB32_94:
+ setp.eq.s32 %p78, %r7, 11;
+ @%p78 bra BB32_117;
+
+ setp.eq.s32 %p79, %r7, 12;
+ @%p79 bra BB32_116;
+ bra.uni BB32_96;
+
+BB32_116:
+ max.f64 %fd98, %fd1, %fd68;
+ bra.uni BB32_141;
+
+BB32_85:
+ setp.eq.s32 %p85, %r7, 6;
+ @%p85 bra BB32_120;
+
+ setp.eq.s32 %p86, %r7, 7;
+ @%p86 bra BB32_119;
+ bra.uni BB32_87;
+
+BB32_119:
+ setp.gt.f64 %p110, %fd1, %fd68;
+ selp.f64 %fd98, 0d3FF0000000000000, 0d0000000000000000, %p110;
+ bra.uni BB32_141;
+
+BB32_102:
+ setp.eq.s32 %p72, %r7, 16;
+ @%p72 bra BB32_114;
+
+ setp.eq.s32 %p73, %r7, 17;
+ @%p73 bra BB32_109;
+ bra.uni BB32_104;
+
+BB32_109:
+ setp.eq.f64 %p98, %fd68, 0d0000000000000000;
+ setp.eq.f64 %p99, %fd68, 0d8000000000000000;
+ or.pred %p100, %p98, %p99;
+ mov.f64 %fd98, 0d7FF8000000000000;
+ @%p100 bra BB32_141;
+
+ div.rn.f64 %fd98, %fd1, %fd68;
+ abs.f64 %fd83, %fd98;
+ setp.gtu.f64 %p101, %fd83, 0d7FF0000000000000;
+ @%p101 bra BB32_141;
+
+ {
+ .reg .b32 %temp;
+ mov.b64 {%temp, %r57}, %fd98;
+ }
+ and.b32 %r58, %r57, 2147483647;
+ setp.ne.s32 %p102, %r58, 2146435072;
+ @%p102 bra BB32_113;
+
+ {
+ .reg .b32 %temp;
+ mov.b64 {%r59, %temp}, %fd98;
+ }
+ setp.eq.s32 %p103, %r59, 0;
+ @%p103 bra BB32_141;
+
+BB32_113:
+ cvt.rmi.f64.f64 %fd84, %fd98;
+ mul.f64 %fd85, %fd84, %fd68;
+ sub.f64 %fd98, %fd1, %fd85;
+ bra.uni BB32_141;
+
+BB32_6:
+ setp.eq.s32 %p30, %r7, 1;
+ @%p30 bra BB32_7;
+ bra.uni BB32_71;
+
+BB32_7:
+ sub.f64 %fd94, %fd68, %fd1;
+ bra.uni BB32_71;
+
+BB32_22:
+ setp.eq.s32 %p18, %r7, 10;
+ @%p18 bra BB32_23;
+ bra.uni BB32_71;
+
+BB32_23:
+ setp.neu.f64 %p43, %fd1, %fd68;
+ selp.f64 %fd94, 0d3FF0000000000000, 0d0000000000000000, %p43;
+ bra.uni BB32_71;
+
+BB32_13:
+ setp.eq.s32 %p25, %r7, 5;
+ @%p25 bra BB32_14;
+ bra.uni BB32_71;
+
+BB32_14:
+ setp.gt.f64 %p48, %fd1, %fd68;
+ selp.f64 %fd94, 0d3FF0000000000000, 0d0000000000000000, %p48;
+ bra.uni BB32_71;
+
+BB32_30:
+ setp.eq.s32 %p12, %r7, 15;
+ @%p12 bra BB32_31;
+ bra.uni BB32_71;
+
+BB32_31:
+ mul.f64 %fd76, %fd1, %fd68;
+ mov.f64 %fd77, 0d3FF0000000000000;
+ sub.f64 %fd94, %fd77, %fd76;
+ bra.uni BB32_71;
+
+BB32_9:
+ setp.eq.s32 %p28, %r7, 3;
+ @%p28 bra BB32_10;
+ bra.uni BB32_71;
+
+BB32_10:
+ div.rn.f64 %fd94, %fd68, %fd1;
+ bra.uni BB32_71;
+
+BB32_47:
+ min.f64 %fd94, %fd68, %fd1;
+ bra.uni BB32_71;
+
+BB32_26:
+ setp.eq.s32 %p16, %r7, 13;
+ @%p16 bra BB32_27;
+ bra.uni BB32_71;
+
+BB32_27:
+ cvt.rni.s64.f64 %rd20, %fd68;
+ cvt.u32.u64 %r25, %rd20;
+ cvt.rni.s64.f64 %rd21, %fd1;
+ cvt.u32.u64 %r26, %rd21;
+ and.b32 %r27, %r26, %r25;
+ setp.eq.s32 %p42, %r27, 0;
+ selp.f64 %fd94, 0d0000000000000000, 0d3FF0000000000000, %p42;
+ bra.uni BB32_71;
+
+BB32_50:
+ setp.ltu.f64 %p47, %fd1, %fd68;
+ selp.f64 %fd94, 0d0000000000000000, 0d3FF0000000000000, %p47;
+ bra.uni BB32_71;
+
+BB32_17:
+ setp.eq.s32 %p23, %r7, 8;
+ @%p23 bra BB32_18;
+ bra.uni BB32_71;
+
+BB32_18:
+ setp.gtu.f64 %p45, %fd1, %fd68;
+ selp.f64 %fd94, 0d0000000000000000, 0d3FF0000000000000, %p45;
+ bra.uni BB32_71;
+
+BB32_44:
+ setp.neu.f64 %p40, %fd68, 0d0000000000000000;
+ sub.f64 %fd75, %fd68, %fd1;
+ selp.f64 %fd94, %fd75, 0d0000000000000000, %p40;
+ bra.uni BB32_71;
+
+BB32_34:
+ setp.ne.s32 %p10, %r7, 18;
+ @%p10 bra BB32_71;
+
+ div.rn.f64 %fd94, %fd68, %fd1;
+ abs.f64 %fd70, %fd94;
+ setp.gtu.f64 %p31, %fd70, 0d7FF0000000000000;
+ @%p31 bra BB32_71;
+
+ {
+ .reg .b32 %temp;
+ mov.b64 {%temp, %r16}, %fd94;
+ }
+ and.b32 %r17, %r16, 2147483647;
+ setp.ne.s32 %p32, %r17, 2146435072;
+ @%p32 bra BB32_38;
+
+ {
+ .reg .b32 %temp;
+ mov.b64 {%r18, %temp}, %fd94;
+ }
+ setp.eq.s32 %p33, %r18, 0;
+ @%p33 bra BB32_71;
+
+BB32_38:
+ cvt.rmi.f64.f64 %fd94, %fd94;
+ bra.uni BB32_71;
+
+BB32_76:
+ setp.eq.s32 %p94, %r7, 1;
+ @%p94 bra BB32_77;
+ bra.uni BB32_141;
+
+BB32_77:
+ sub.f64 %fd98, %fd1, %fd68;
+ bra.uni BB32_141;
+
+BB32_92:
+ setp.eq.s32 %p82, %r7, 10;
+ @%p82 bra BB32_93;
+ bra.uni BB32_141;
+
+BB32_93:
+ setp.neu.f64 %p107, %fd1, %fd68;
+ selp.f64 %fd98, 0d3FF0000000000000, 0d0000000000000000, %p107;
+ bra.uni BB32_141;
+
+BB32_83:
+ setp.eq.s32 %p89, %r7, 5;
+ @%p89 bra BB32_84;
+ bra.uni BB32_141;
+
+BB32_84:
+ setp.lt.f64 %p112, %fd1, %fd68;
+ selp.f64 %fd98, 0d3FF0000000000000, 0d0000000000000000, %p112;
+ bra.uni BB32_141;
+
+BB32_100:
+ setp.eq.s32 %p76, %r7, 15;
+ @%p76 bra BB32_101;
+ bra.uni BB32_141;
+
+BB32_101:
+ mul.f64 %fd87, %fd1, %fd68;
+ mov.f64 %fd88, 0d3FF0000000000000;
+ sub.f64 %fd98, %fd88, %fd87;
+ bra.uni BB32_141;
+
+BB32_79:
+ setp.eq.s32 %p92, %r7, 3;
+ @%p92 bra BB32_80;
+ bra.uni BB32_141;
+
+BB32_80:
+ div.rn.f64 %fd98, %fd1, %fd68;
+ bra.uni BB32_141;
+
+BB32_117:
+ min.f64 %fd98, %fd1, %fd68;
+ bra.uni BB32_141;
+
+BB32_96:
+ setp.eq.s32 %p80, %r7, 13;
+ @%p80 bra BB32_97;
+ bra.uni BB32_141;
+
+BB32_97:
+ cvt.rni.s64.f64 %rd25, %fd1;
+ cvt.u32.u64 %r63, %rd25;
+ cvt.rni.s64.f64 %rd26, %fd68;
+ cvt.u32.u64 %r64, %rd26;
+ and.b32 %r65, %r64, %r63;
+ setp.eq.s32 %p106, %r65, 0;
+ selp.f64 %fd98, 0d0000000000000000, 0d3FF0000000000000, %p106;
+ bra.uni BB32_141;
+
+BB32_120:
+ setp.gtu.f64 %p111, %fd1, %fd68;
+ selp.f64 %fd98, 0d0000000000000000, 0d3FF0000000000000, %p111;
+ bra.uni BB32_141;
+
+BB32_87:
+ setp.eq.s32 %p87, %r7, 8;
+ @%p87 bra BB32_88;
+ bra.uni BB32_141;
+
+BB32_88:
+ setp.ltu.f64 %p109, %fd1, %fd68;
+ selp.f64 %fd98, 0d0000000000000000, 0d3FF0000000000000, %p109;
+ bra.uni BB32_141;
+
+BB32_114:
+ setp.neu.f64 %p104, %fd1, 0d0000000000000000;
+ sub.f64 %fd86, %fd1, %fd68;
+ selp.f64 %fd98, %fd86, 0d0000000000000000, %p104;
+ bra.uni BB32_141;
+
+BB32_104:
+ setp.ne.s32 %p74, %r7, 18;
+ @%p74 bra BB32_141;
+
+ div.rn.f64 %fd98, %fd1, %fd68;
+ abs.f64 %fd81, %fd98;
+ setp.gtu.f64 %p95, %fd81, 0d7FF0000000000000;
+ @%p95 bra BB32_141;
+
+ {
+ .reg .b32 %temp;
+ mov.b64 {%temp, %r54}, %fd98;
+ }
+ and.b32 %r55, %r54, 2147483647;
+ setp.ne.s32 %p96, %r55, 2146435072;
+ @%p96 bra BB32_108;
+
+ {
+ .reg .b32 %temp;
+ mov.b64 {%r56, %temp}, %fd98;
+ }
+ setp.eq.s32 %p97, %r56, 0;
+ @%p97 bra BB32_141;
+
+BB32_108:
+ cvt.rmi.f64.f64 %fd98, %fd98;
+ bra.uni BB32_141;
+
+BB32_54:
+ setp.gt.s32 %p52, %r2, -1;
+ @%p52 bra BB32_57;
+
+ cvt.rzi.f64.f64 %fd78, %fd1;
+ setp.neu.f64 %p53, %fd78, %fd1;
+ selp.f64 %fd24, 0dFFF8000000000000, %fd24, %p53;
+
+BB32_57:
+ add.f64 %fd93, %fd1, %fd68;
+ {
+ .reg .b32 %temp;
+ mov.b64 {%temp, %r37}, %fd93;
+ }
+ and.b32 %r38, %r37, 2146435072;
+ setp.ne.s32 %p56, %r38, 2146435072;
+ @%p56 bra BB32_58;
+
+ setp.gtu.f64 %p57, %fd18, 0d7FF0000000000000;
+ @%p57 bra BB32_68;
+
+ abs.f64 %fd79, %fd1;
+ setp.gtu.f64 %p58, %fd79, 0d7FF0000000000000;
+ @%p58 bra BB32_68;
+
+ and.b32 %r39, %r3, 2147483647;
+ setp.ne.s32 %p59, %r39, 2146435072;
+ @%p59 bra BB32_63;
+
+ {
+ .reg .b32 %temp;
+ mov.b64 {%r40, %temp}, %fd1;
+ }
+ setp.eq.s32 %p60, %r40, 0;
+ @%p60 bra BB32_67;
+
+BB32_63:
+ and.b32 %r41, %r2, 2147483647;
+ setp.ne.s32 %p61, %r41, 2146435072;
+ @%p61 bra BB32_64;
+
+ {
+ .reg .b32 %temp;
+ mov.b64 {%r42, %temp}, %fd68;
+ }
+ setp.ne.s32 %p62, %r42, 0;
+ mov.f64 %fd93, %fd24;
+ @%p62 bra BB32_68;
+
+ shr.s32 %r43, %r3, 31;
+ and.b32 %r44, %r43, -2146435072;
+ add.s32 %r45, %r44, 2146435072;
+ or.b32 %r46, %r45, -2147483648;
+ selp.b32 %r47, %r46, %r45, %p1;
+ mov.u32 %r48, 0;
+ mov.b64 %fd93, {%r48, %r47};
+ bra.uni BB32_68;
+
+BB32_58:
+ mov.f64 %fd93, %fd24;
+
+BB32_68:
+ setp.eq.f64 %p66, %fd1, 0d0000000000000000;
+ setp.eq.f64 %p67, %fd68, 0d3FF0000000000000;
+ or.pred %p68, %p67, %p66;
+ selp.f64 %fd94, 0d3FF0000000000000, %fd93, %p68;
+
+BB32_71:
+ st.global.f64 [%rd1], %fd94;
+ bra.uni BB32_142;
+
+BB32_124:
+ setp.gt.s32 %p116, %r4, -1;
+ @%p116 bra BB32_127;
+
+ cvt.rzi.f64.f64 %fd89, %fd68;
+ setp.neu.f64 %p117, %fd89, %fd68;
+ selp.f64 %fd57, 0dFFF8000000000000, %fd57, %p117;
+
+BB32_127:
+ add.f64 %fd97, %fd1, %fd68;
+ {
+ .reg .b32 %temp;
+ mov.b64 {%temp, %r75}, %fd97;
+ }
+ and.b32 %r76, %r75, 2146435072;
+ setp.ne.s32 %p120, %r76, 2146435072;
+ @%p120 bra BB32_128;
+
+ setp.gtu.f64 %p121, %fd51, 0d7FF0000000000000;
+ @%p121 bra BB32_138;
+
+ abs.f64 %fd90, %fd68;
+ setp.gtu.f64 %p122, %fd90, 0d7FF0000000000000;
+ @%p122 bra BB32_138;
+
+ and.b32 %r77, %r5, 2147483647;
+ setp.ne.s32 %p123, %r77, 2146435072;
+ @%p123 bra BB32_133;
+
+ {
+ .reg .b32 %temp;
+ mov.b64 {%r78, %temp}, %fd68;
+ }
+ setp.eq.s32 %p124, %r78, 0;
+ @%p124 bra BB32_137;
+
+BB32_133:
+ and.b32 %r79, %r4, 2147483647;
+ setp.ne.s32 %p125, %r79, 2146435072;
+ @%p125 bra BB32_134;
+
+ {
+ .reg .b32 %temp;
+ mov.b64 {%r80, %temp}, %fd1;
+ }
+ setp.ne.s32 %p126, %r80, 0;
+ mov.f64 %fd97, %fd57;
+ @%p126 bra BB32_138;
+
+ shr.s32 %r81, %r5, 31;
+ and.b32 %r82, %r81, -2146435072;
+ add.s32 %r83, %r82, 2146435072;
+ or.b32 %r84, %r83, -2147483648;
+ selp.b32 %r85, %r84, %r83, %p2;
+ mov.u32 %r86, 0;
+ mov.b64 %fd97, {%r86, %r85};
+ bra.uni BB32_138;
+
+BB32_128:
+ mov.f64 %fd97, %fd57;
+
+BB32_138:
+ setp.eq.f64 %p130, %fd68, 0d0000000000000000;
+ setp.eq.f64 %p131, %fd1, 0d3FF0000000000000;
+ or.pred %p132, %p131, %p130;
+ selp.f64 %fd98, 0d3FF0000000000000, %fd97, %p132;
+
+BB32_141:
+ st.global.f64 [%rd1], %fd98;
+
+BB32_142:
+ bar.sync 0;
+ ret;
+
+BB32_64:
+ mov.f64 %fd93, %fd24;
+ bra.uni BB32_68;
+
+BB32_134:
+ mov.f64 %fd97, %fd57;
+ bra.uni BB32_138;
+
+BB32_67:
+ setp.gt.f64 %p63, %fd18, 0d3FF0000000000000;
+ selp.b32 %r49, 2146435072, 0, %p63;
+ xor.b32 %r50, %r49, 2146435072;
+ setp.lt.s32 %p64, %r3, 0;
+ selp.b32 %r51, %r50, %r49, %p64;
+ setp.eq.f64 %p65, %fd68, 0dBFF0000000000000;
+ selp.b32 %r52, 1072693248, %r51, %p65;
+ mov.u32 %r53, 0;
+ mov.b64 %fd93, {%r53, %r52};
+ bra.uni BB32_68;
+
+BB32_137:
+ setp.gt.f64 %p127, %fd51, 0d3FF0000000000000;
+ selp.b32 %r87, 2146435072, 0, %p127;
+ xor.b32 %r88, %r87, 2146435072;
+ setp.lt.s32 %p128, %r5, 0;
+ selp.b32 %r89, %r88, %r87, %p128;
+ setp.eq.f64 %p129, %fd1, 0dBFF0000000000000;
+ selp.b32 %r90, 1072693248, %r89, %p129;
+ mov.u32 %r91, 0;
+ mov.b64 %fd97, {%r91, %r90};
+ bra.uni BB32_138;
+}
+
+ // .globl sparse_dense_matrix_scalar_op_f
+.visible .entry sparse_dense_matrix_scalar_op_f(
+ .param .u64 sparse_dense_matrix_scalar_op_f_param_0,
+ .param .u64 sparse_dense_matrix_scalar_op_f_param_1,
+ .param .u64 sparse_dense_matrix_scalar_op_f_param_2,
+ .param .f64 sparse_dense_matrix_scalar_op_f_param_3,
+ .param .u64 sparse_dense_matrix_scalar_op_f_param_4,
+ .param .u32 sparse_dense_matrix_scalar_op_f_param_5,
+ .param .u32 sparse_dense_matrix_scalar_op_f_param_6,
+ .param .u32 sparse_dense_matrix_scalar_op_f_param_7,
+ .param .u32 sparse_dense_matrix_scalar_op_f_param_8
+)
+{
+ .reg .pred %p<139>;
+ .reg .f32 %f<267>;
+ .reg .b32 %r<62>;
+ .reg .f64 %fd<2>;
+ .reg .b64 %rd<23>;
+
+
+ ld.param.u64 %rd2, [sparse_dense_matrix_scalar_op_f_param_0];
+ ld.param.u64 %rd3, [sparse_dense_matrix_scalar_op_f_param_1];
+ ld.param.u64 %rd4, [sparse_dense_matrix_scalar_op_f_param_2];
+ ld.param.f64 %fd1, [sparse_dense_matrix_scalar_op_f_param_3];
+ ld.param.u64 %rd5, [sparse_dense_matrix_scalar_op_f_param_4];
+ ld.param.u32 %r5, [sparse_dense_matrix_scalar_op_f_param_5];
+ ld.param.u32 %r2, [sparse_dense_matrix_scalar_op_f_param_6];
+ ld.param.u32 %r3, [sparse_dense_matrix_scalar_op_f_param_7];
+ ld.param.u32 %r4, [sparse_dense_matrix_scalar_op_f_param_8];
+ cvt.rn.f32.f64 %f1, %fd1;
+ mov.u32 %r6, %ntid.x;
+ mov.u32 %r7, %ctaid.x;
+ mov.u32 %r8, %tid.x;
+ mad.lo.s32 %r1, %r6, %r7, %r8;
+ setp.ge.s32 %p3, %r1, %r5;
+ @%p3 bra BB33_126;
+
+ cvta.to.global.u64 %rd6, %rd5;
+ cvta.to.global.u64 %rd7, %rd4;
+ mul.wide.s32 %rd8, %r1, 4;
+ add.s64 %rd9, %rd7, %rd8;
+ ld.global.f32 %f2, [%rd9];
+ cvta.to.global.u64 %rd10, %rd2;
+ add.s64 %rd11, %rd10, %rd8;
+ ld.global.u32 %r9, [%rd11];
+ cvta.to.global.u64 %rd12, %rd3;
+ add.s64 %rd13, %rd12, %rd8;
+ ld.global.u32 %r10, [%rd13];
+ mad.lo.s32 %r11, %r9, %r2, %r10;
+ mul.wide.s32 %rd14, %r11, 4;
+ add.s64 %rd1, %rd6, %rd14;
+ setp.eq.s32 %p4, %r4, 0;
+ @%p4 bra BB33_64;
+
+ mov.f32 %f262, 0f7F7FFFFF;
+ setp.gt.s32 %p5, %r3, 8;
+ @%p5 bra BB33_19;
+
+ setp.gt.s32 %p19, %r3, 3;
+ @%p19 bra BB33_11;
+
+ setp.gt.s32 %p26, %r3, 1;
+ @%p26 bra BB33_8;
+
+ setp.eq.s32 %p29, %r3, 0;
+ @%p29 bra BB33_62;
+ bra.uni BB33_6;
+
+BB33_62:
+ add.f32 %f262, %f1, %f2;
+ bra.uni BB33_63;
+
+BB33_64:
+ mov.f32 %f266, 0f7F7FFFFF;
+ setp.gt.s32 %p72, %r3, 8;
+ @%p72 bra BB33_81;
+
+ setp.gt.s32 %p86, %r3, 3;
+ @%p86 bra BB33_73;
+
+ setp.gt.s32 %p93, %r3, 1;
+ @%p93 bra BB33_70;
+
+ setp.eq.s32 %p96, %r3, 0;
+ @%p96 bra BB33_124;
+ bra.uni BB33_68;
+
+BB33_124:
+ add.f32 %f266, %f1, %f2;
+ bra.uni BB33_125;
+
+BB33_19:
+ setp.gt.s32 %p6, %r3, 13;
+ @%p6 bra BB33_28;
+
+ setp.gt.s32 %p13, %r3, 10;
+ @%p13 bra BB33_24;
+
+ setp.eq.s32 %p17, %r3, 9;
+ @%p17 bra BB33_44;
+ bra.uni BB33_22;
+
+BB33_44:
+ setp.eq.f32 %p40, %f1, %f2;
+ selp.f32 %f262, 0f3F800000, 0f00000000, %p40;
+ bra.uni BB33_63;
+
+BB33_81:
+ setp.gt.s32 %p73, %r3, 13;
+ @%p73 bra BB33_90;
+
+ setp.gt.s32 %p80, %r3, 10;
+ @%p80 bra BB33_86;
+
+ setp.eq.s32 %p84, %r3, 9;
+ @%p84 bra BB33_106;
+ bra.uni BB33_84;
+
+BB33_106:
+ setp.eq.f32 %p107, %f2, %f1;
+ selp.f32 %f266, 0f3F800000, 0f00000000, %p107;
+ bra.uni BB33_125;
+
+BB33_11:
+ setp.gt.s32 %p20, %r3, 5;
+ @%p20 bra BB33_15;
+
+ setp.eq.s32 %p24, %r3, 4;
+ @%p24 bra BB33_47;
+ bra.uni BB33_13;
+
+BB33_47:
+ mul.f32 %f88, %f2, 0f3F000000;
+ cvt.rzi.f32.f32 %f89, %f88;
+ fma.rn.f32 %f90, %f89, 0fC0000000, %f2;
+ abs.f32 %f19, %f90;
+ abs.f32 %f20, %f1;
+ setp.lt.f32 %p45, %f20, 0f00800000;
+ mul.f32 %f91, %f20, 0f4B800000;
+ selp.f32 %f92, 0fC3170000, 0fC2FE0000, %p45;
+ selp.f32 %f93, %f91, %f20, %p45;
+ mov.b32 %r18, %f93;
+ and.b32 %r19, %r18, 8388607;
+ or.b32 %r20, %r19, 1065353216;
+ mov.b32 %f94, %r20;
+ shr.u32 %r21, %r18, 23;
+ cvt.rn.f32.u32 %f95, %r21;
+ add.f32 %f96, %f92, %f95;
+ setp.gt.f32 %p46, %f94, 0f3FB504F3;
+ mul.f32 %f97, %f94, 0f3F000000;
+ add.f32 %f98, %f96, 0f3F800000;
+ selp.f32 %f99, %f97, %f94, %p46;
+ selp.f32 %f100, %f98, %f96, %p46;
+ add.f32 %f101, %f99, 0fBF800000;
+ add.f32 %f87, %f99, 0f3F800000;
+ // inline asm
+ rcp.approx.ftz.f32 %f86,%f87;
+ // inline asm
+ add.f32 %f102, %f101, %f101;
+ mul.f32 %f103, %f86, %f102;
+ mul.f32 %f104, %f103, %f103;
+ mov.f32 %f105, 0f3C4CAF63;
+ mov.f32 %f106, 0f3B18F0FE;
+ fma.rn.f32 %f107, %f106, %f104, %f105;
+ mov.f32 %f108, 0f3DAAAABD;
+ fma.rn.f32 %f109, %f107, %f104, %f108;
+ mul.rn.f32 %f110, %f109, %f104;
+ mul.rn.f32 %f111, %f110, %f103;
+ sub.f32 %f112, %f101, %f103;
+ neg.f32 %f113, %f103;
+ add.f32 %f114, %f112, %f112;
+ fma.rn.f32 %f115, %f113, %f101, %f114;
+ mul.rn.f32 %f116, %f86, %f115;
+ add.f32 %f117, %f111, %f103;
+ sub.f32 %f118, %f103, %f117;
+ add.f32 %f119, %f111, %f118;
+ add.f32 %f120, %f116, %f119;
+ add.f32 %f121, %f117, %f120;
+ sub.f32 %f122, %f117, %f121;
+ add.f32 %f123, %f120, %f122;
+ mov.f32 %f124, 0f3F317200;
+ mul.rn.f32 %f125, %f100, %f124;
+ mov.f32 %f126, 0f35BFBE8E;
+ mul.rn.f32 %f127, %f100, %f126;
+ add.f32 %f128, %f125, %f121;
+ sub.f32 %f129, %f125, %f128;
+ add.f32 %f130, %f121, %f129;
+ add.f32 %f131, %f123, %f130;
+ add.f32 %f132, %f127, %f131;
+ add.f32 %f133, %f128, %f132;
+ sub.f32 %f134, %f128, %f133;
+ add.f32 %f135, %f132, %f134;
+ abs.f32 %f21, %f2;
+ setp.gt.f32 %p47, %f21, 0f77F684DF;
+ mul.f32 %f136, %f2, 0f39000000;
+ selp.f32 %f137, %f136, %f2, %p47;
+ mul.rn.f32 %f138, %f137, %f133;
+ neg.f32 %f139, %f138;
+ fma.rn.f32 %f140, %f137, %f133, %f139;
+ fma.rn.f32 %f141, %f137, %f135, %f140;
+ mov.f32 %f142, 0f00000000;
+ fma.rn.f32 %f143, %f142, %f133, %f141;
+ add.rn.f32 %f144, %f138, %f143;
+ neg.f32 %f145, %f144;
+ add.rn.f32 %f146, %f138, %f145;
+ add.rn.f32 %f147, %f146, %f143;
+ mov.b32 %r22, %f144;
+ setp.eq.s32 %p48, %r22, 1118925336;
+ add.s32 %r23, %r22, -1;
+ mov.b32 %f148, %r23;
+ add.f32 %f149, %f147, 0f37000000;
+ selp.f32 %f150, %f148, %f144, %p48;
+ selp.f32 %f22, %f149, %f147, %p48;
+ mul.f32 %f151, %f150, 0f3FB8AA3B;
+ cvt.rzi.f32.f32 %f152, %f151;
+ mov.f32 %f153, 0fBF317200;
+ fma.rn.f32 %f154, %f152, %f153, %f150;
+ mov.f32 %f155, 0fB5BFBE8E;
+ fma.rn.f32 %f156, %f152, %f155, %f154;
+ mul.f32 %f157, %f156, 0f3FB8AA3B;
+ ex2.approx.ftz.f32 %f158, %f157;
+ add.f32 %f159, %f152, 0f00000000;
+ ex2.approx.f32 %f160, %f159;
+ mul.f32 %f161, %f158, %f160;
+ setp.lt.f32 %p49, %f150, 0fC2D20000;
+ selp.f32 %f162, 0f00000000, %f161, %p49;
+ setp.gt.f32 %p50, %f150, 0f42D20000;
+ selp.f32 %f259, 0f7F800000, %f162, %p50;
+ setp.eq.f32 %p51, %f259, 0f7F800000;
+ @%p51 bra BB33_49;
+
+ fma.rn.f32 %f259, %f259, %f22, %f259;
+
+BB33_49:
+ setp.lt.f32 %p52, %f1, 0f00000000;
+ setp.eq.f32 %p53, %f19, 0f3F800000;
+ and.pred %p1, %p52, %p53;
+ mov.b32 %r24, %f259;
+ xor.b32 %r25, %r24, -2147483648;
+ mov.b32 %f163, %r25;
+ selp.f32 %f261, %f163, %f259, %p1;
+ setp.eq.f32 %p54, %f1, 0f00000000;
+ @%p54 bra BB33_52;
+ bra.uni BB33_50;
+
+BB33_52:
+ add.f32 %f165, %f1, %f1;
+ mov.b32 %r26, %f165;
+ selp.b32 %r27, %r26, 0, %p53;
+ or.b32 %r28, %r27, 2139095040;
+ setp.lt.f32 %p58, %f2, 0f00000000;
+ selp.b32 %r29, %r28, %r27, %p58;
+ mov.b32 %f261, %r29;
+ bra.uni BB33_53;
+
+BB33_28:
+ setp.gt.s32 %p7, %r3, 15;
+ @%p7 bra BB33_32;
+
+ setp.eq.s32 %p11, %r3, 14;
+ @%p11 bra BB33_41;
+ bra.uni BB33_30;
+
+BB33_41:
+ cvt.rni.s64.f32 %rd15, %f1;
+ cvt.u32.u64 %r12, %rd15;
+ cvt.rni.s64.f32 %rd16, %f2;
+ cvt.u32.u64 %r13, %rd16;
+ or.b32 %r14, %r13, %r12;
+ setp.eq.s32 %p37, %r14, 0;
+ selp.f32 %f262, 0f00000000, 0f3F800000, %p37;
+ bra.uni BB33_63;
+
+BB33_73:
+ setp.gt.s32 %p87, %r3, 5;
+ @%p87 bra BB33_77;
+
+ setp.eq.s32 %p91, %r3, 4;
+ @%p91 bra BB33_109;
+ bra.uni BB33_75;
+
+BB33_109:
+ mul.f32 %f179, %f1, 0f3F000000;
+ cvt.rzi.f32.f32 %f180, %f179;
+ fma.rn.f32 %f181, %f180, 0fC0000000, %f1;
+ abs.f32 %f56, %f181;
+ abs.f32 %f57, %f2;
+ setp.lt.f32 %p112, %f57, 0f00800000;
+ mul.f32 %f182, %f57, 0f4B800000;
+ selp.f32 %f183, 0fC3170000, 0fC2FE0000, %p112;
+ selp.f32 %f184, %f182, %f57, %p112;
+ mov.b32 %r43, %f184;
+ and.b32 %r44, %r43, 8388607;
+ or.b32 %r45, %r44, 1065353216;
+ mov.b32 %f185, %r45;
+ shr.u32 %r46, %r43, 23;
+ cvt.rn.f32.u32 %f186, %r46;
+ add.f32 %f187, %f183, %f186;
+ setp.gt.f32 %p113, %f185, 0f3FB504F3;
+ mul.f32 %f188, %f185, 0f3F000000;
+ add.f32 %f189, %f187, 0f3F800000;
+ selp.f32 %f190, %f188, %f185, %p113;
+ selp.f32 %f191, %f189, %f187, %p113;
+ add.f32 %f192, %f190, 0fBF800000;
+ add.f32 %f178, %f190, 0f3F800000;
+ // inline asm
+ rcp.approx.ftz.f32 %f177,%f178;
+ // inline asm
+ add.f32 %f193, %f192, %f192;
+ mul.f32 %f194, %f177, %f193;
+ mul.f32 %f195, %f194, %f194;
+ mov.f32 %f196, 0f3C4CAF63;
+ mov.f32 %f197, 0f3B18F0FE;
+ fma.rn.f32 %f198, %f197, %f195, %f196;
+ mov.f32 %f199, 0f3DAAAABD;
+ fma.rn.f32 %f200, %f198, %f195, %f199;
+ mul.rn.f32 %f201, %f200, %f195;
+ mul.rn.f32 %f202, %f201, %f194;
+ sub.f32 %f203, %f192, %f194;
+ neg.f32 %f204, %f194;
+ add.f32 %f205, %f203, %f203;
+ fma.rn.f32 %f206, %f204, %f192, %f205;
+ mul.rn.f32 %f207, %f177, %f206;
+ add.f32 %f208, %f202, %f194;
+ sub.f32 %f209, %f194, %f208;
+ add.f32 %f210, %f202, %f209;
+ add.f32 %f211, %f207, %f210;
+ add.f32 %f212, %f208, %f211;
+ sub.f32 %f213, %f208, %f212;
+ add.f32 %f214, %f211, %f213;
+ mov.f32 %f215, 0f3F317200;
+ mul.rn.f32 %f216, %f191, %f215;
+ mov.f32 %f217, 0f35BFBE8E;
+ mul.rn.f32 %f218, %f191, %f217;
+ add.f32 %f219, %f216, %f212;
+ sub.f32 %f220, %f216, %f219;
+ add.f32 %f221, %f212, %f220;
+ add.f32 %f222, %f214, %f221;
+ add.f32 %f223, %f218, %f222;
+ add.f32 %f224, %f219, %f223;
+ sub.f32 %f225, %f219, %f224;
+ add.f32 %f226, %f223, %f225;
+ abs.f32 %f58, %f1;
+ setp.gt.f32 %p114, %f58, 0f77F684DF;
+ mul.f32 %f227, %f1, 0f39000000;
+ selp.f32 %f228, %f227, %f1, %p114;
+ mul.rn.f32 %f229, %f228, %f224;
+ neg.f32 %f230, %f229;
+ fma.rn.f32 %f231, %f228, %f224, %f230;
+ fma.rn.f32 %f232, %f228, %f226, %f231;
+ mov.f32 %f233, 0f00000000;
+ fma.rn.f32 %f234, %f233, %f224, %f232;
+ add.rn.f32 %f235, %f229, %f234;
+ neg.f32 %f236, %f235;
+ add.rn.f32 %f237, %f229, %f236;
+ add.rn.f32 %f238, %f237, %f234;
+ mov.b32 %r47, %f235;
+ setp.eq.s32 %p115, %r47, 1118925336;
+ add.s32 %r48, %r47, -1;
+ mov.b32 %f239, %r48;
+ add.f32 %f240, %f238, 0f37000000;
+ selp.f32 %f241, %f239, %f235, %p115;
+ selp.f32 %f59, %f240, %f238, %p115;
+ mul.f32 %f242, %f241, 0f3FB8AA3B;
+ cvt.rzi.f32.f32 %f243, %f242;
+ mov.f32 %f244, 0fBF317200;
+ fma.rn.f32 %f245, %f243, %f244, %f241;
+ mov.f32 %f246, 0fB5BFBE8E;
+ fma.rn.f32 %f247, %f243, %f246, %f245;
+ mul.f32 %f248, %f247, 0f3FB8AA3B;
+ ex2.approx.ftz.f32 %f249, %f248;
+ add.f32 %f250, %f243, 0f00000000;
+ ex2.approx.f32 %f251, %f250;
+ mul.f32 %f252, %f249, %f251;
+ setp.lt.f32 %p116, %f241, 0fC2D20000;
+ selp.f32 %f253, 0f00000000, %f252, %p116;
+ setp.gt.f32 %p117, %f241, 0f42D20000;
+ selp.f32 %f263, 0f7F800000, %f253, %p117;
+ setp.eq.f32 %p118, %f263, 0f7F800000;
+ @%p118 bra BB33_111;
+
+ fma.rn.f32 %f263, %f263, %f59, %f263;
+
+BB33_111:
+ setp.lt.f32 %p119, %f2, 0f00000000;
+ setp.eq.f32 %p120, %f56, 0f3F800000;
+ and.pred %p2, %p119, %p120;
+ mov.b32 %r49, %f263;
+ xor.b32 %r50, %r49, -2147483648;
+ mov.b32 %f254, %r50;
+ selp.f32 %f265, %f254, %f263, %p2;
+ setp.eq.f32 %p121, %f2, 0f00000000;
+ @%p121 bra BB33_114;
+ bra.uni BB33_112;
+
+BB33_114:
+ add.f32 %f256, %f2, %f2;
+ mov.b32 %r51, %f256;
+ selp.b32 %r52, %r51, 0, %p120;
+ or.b32 %r53, %r52, 2139095040;
+ setp.lt.f32 %p125, %f1, 0f00000000;
+ selp.b32 %r54, %r53, %r52, %p125;
+ mov.b32 %f265, %r54;
+ bra.uni BB33_115;
+
+BB33_90:
+ setp.gt.s32 %p74, %r3, 15;
+ @%p74 bra BB33_94;
+
+ setp.eq.s32 %p78, %r3, 14;
+ @%p78 bra BB33_103;
+ bra.uni BB33_92;
+
+BB33_103:
+ cvt.rni.s64.f32 %rd19, %f2;
+ cvt.u32.u64 %r37, %rd19;
+ cvt.rni.s64.f32 %rd20, %f1;
+ cvt.u32.u64 %r38, %rd20;
+ or.b32 %r39, %r38, %r37;
+ setp.eq.s32 %p104, %r39, 0;
+ selp.f32 %f266, 0f00000000, 0f3F800000, %p104;
+ bra.uni BB33_125;
+
+BB33_8:
+ setp.eq.s32 %p27, %r3, 2;
+ @%p27 bra BB33_61;
+ bra.uni BB33_9;
+
+BB33_61:
+ mul.f32 %f262, %f1, %f2;
+ bra.uni BB33_63;
+
+BB33_24:
+ setp.eq.s32 %p14, %r3, 11;
+ @%p14 bra BB33_43;
+
+ setp.eq.s32 %p15, %r3, 12;
+ @%p15 bra BB33_42;
+ bra.uni BB33_26;
+
+BB33_42:
+ max.f32 %f262, %f1, %f2;
+ bra.uni BB33_63;
+
+BB33_15:
+ setp.eq.s32 %p21, %r3, 6;
+ @%p21 bra BB33_46;
+
+ setp.eq.s32 %p22, %r3, 7;
+ @%p22 bra BB33_45;
+ bra.uni BB33_17;
+
+BB33_45:
+ setp.gt.f32 %p42, %f1, %f2;
+ selp.f32 %f262, 0f3F800000, 0f00000000, %p42;
+ bra.uni BB33_63;
+
+BB33_32:
+ setp.eq.s32 %p8, %r3, 16;
+ @%p8 bra BB33_40;
+
+ setp.eq.s32 %p9, %r3, 17;
+ @%p9 bra BB33_37;
+ bra.uni BB33_34;
+
+BB33_37:
+ setp.eq.f32 %p32, %f2, 0f00000000;
+ setp.eq.f32 %p33, %f2, 0f80000000;
+ or.pred %p34, %p32, %p33;
+ mov.f32 %f262, 0f7FC00000;
+ @%p34 bra BB33_63;
+
+ div.rn.f32 %f262, %f1, %f2;
+ abs.f32 %f80, %f262;
+ setp.geu.f32 %p35, %f80, 0f7F800000;
+ @%p35 bra BB33_63;
+
+ cvt.rmi.f32.f32 %f81, %f262;
+ mul.f32 %f82, %f2, %f81;
+ sub.f32 %f262, %f1, %f82;
+ bra.uni BB33_63;
+
+BB33_70:
+ setp.eq.s32 %p94, %r3, 2;
+ @%p94 bra BB33_123;
+ bra.uni BB33_71;
+
+BB33_123:
+ mul.f32 %f266, %f1, %f2;
+ bra.uni BB33_125;
+
+BB33_86:
+ setp.eq.s32 %p81, %r3, 11;
+ @%p81 bra BB33_105;
+
+ setp.eq.s32 %p82, %r3, 12;
+ @%p82 bra BB33_104;
+ bra.uni BB33_88;
+
+BB33_104:
+ max.f32 %f266, %f2, %f1;
+ bra.uni BB33_125;
+
+BB33_77:
+ setp.eq.s32 %p88, %r3, 6;
+ @%p88 bra BB33_108;
+
+ setp.eq.s32 %p89, %r3, 7;
+ @%p89 bra BB33_107;
+ bra.uni BB33_79;
+
+BB33_107:
+ setp.gt.f32 %p109, %f2, %f1;
+ selp.f32 %f266, 0f3F800000, 0f00000000, %p109;
+ bra.uni BB33_125;
+
+BB33_94:
+ setp.eq.s32 %p75, %r3, 16;
+ @%p75 bra BB33_102;
+
+ setp.eq.s32 %p76, %r3, 17;
+ @%p76 bra BB33_99;
+ bra.uni BB33_96;
+
+BB33_99:
+ setp.eq.f32 %p99, %f1, 0f00000000;
+ setp.eq.f32 %p100, %f1, 0f80000000;
+ or.pred %p101, %p99, %p100;
+ mov.f32 %f266, 0f7FC00000;
+ @%p101 bra BB33_125;
+
+ div.rn.f32 %f266, %f2, %f1;
+ abs.f32 %f171, %f266;
+ setp.geu.f32 %p102, %f171, 0f7F800000;
+ @%p102 bra BB33_125;
+
+ cvt.rmi.f32.f32 %f172, %f266;
+ mul.f32 %f173, %f1, %f172;
+ sub.f32 %f266, %f2, %f173;
+ bra.uni BB33_125;
+
+BB33_6:
+ setp.eq.s32 %p30, %r3, 1;
+ @%p30 bra BB33_7;
+ bra.uni BB33_63;
+
+BB33_7:
+ sub.f32 %f262, %f1, %f2;
+ bra.uni BB33_63;
+
+BB33_22:
+ setp.eq.s32 %p18, %r3, 10;
+ @%p18 bra BB33_23;
+ bra.uni BB33_63;
+
+BB33_23:
+ setp.neu.f32 %p39, %f1, %f2;
+ selp.f32 %f262, 0f3F800000, 0f00000000, %p39;
+ bra.uni BB33_63;
+
+BB33_13:
+ setp.eq.s32 %p25, %r3, 5;
+ @%p25 bra BB33_14;
+ bra.uni BB33_63;
+
+BB33_14:
+ setp.lt.f32 %p44, %f1, %f2;
+ selp.f32 %f262, 0f3F800000, 0f00000000, %p44;
+ bra.uni BB33_63;
+
+BB33_30:
+ setp.eq.s32 %p12, %r3, 15;
+ @%p12 bra BB33_31;
+ bra.uni BB33_63;
+
+BB33_31:
+ mul.f32 %f84, %f1, %f2;
+ mov.f32 %f85, 0f3F800000;
+ sub.f32 %f262, %f85, %f84;
+ bra.uni BB33_63;
+
+BB33_9:
+ setp.eq.s32 %p28, %r3, 3;
+ @%p28 bra BB33_10;
+ bra.uni BB33_63;
+
+BB33_10:
+ div.rn.f32 %f262, %f1, %f2;
+ bra.uni BB33_63;
+
+BB33_43:
+ min.f32 %f262, %f1, %f2;
+ bra.uni BB33_63;
+
+BB33_26:
+ setp.eq.s32 %p16, %r3, 13;
+ @%p16 bra BB33_27;
+ bra.uni BB33_63;
+
+BB33_27:
+ cvt.rni.s64.f32 %rd17, %f1;
+ cvt.u32.u64 %r15, %rd17;
+ cvt.rni.s64.f32 %rd18, %f2;
+ cvt.u32.u64 %r16, %rd18;
+ and.b32 %r17, %r16, %r15;
+ setp.eq.s32 %p38, %r17, 0;
+ selp.f32 %f262, 0f00000000, 0f3F800000, %p38;
+ bra.uni BB33_63;
+
+BB33_46:
+ setp.gtu.f32 %p43, %f1, %f2;
+ selp.f32 %f262, 0f00000000, 0f3F800000, %p43;
+ bra.uni BB33_63;
+
+BB33_17:
+ setp.eq.s32 %p23, %r3, 8;
+ @%p23 bra BB33_18;
+ bra.uni BB33_63;
+
+BB33_18:
+ setp.ltu.f32 %p41, %f1, %f2;
+ selp.f32 %f262, 0f00000000, 0f3F800000, %p41;
+ bra.uni BB33_63;
+
+BB33_40:
+ setp.neu.f32 %p36, %f1, 0f00000000;
+ sub.f32 %f83, %f1, %f2;
+ selp.f32 %f262, %f83, 0f00000000, %p36;
+ bra.uni BB33_63;
+
+BB33_34:
+ setp.ne.s32 %p10, %r3, 18;
+ @%p10 bra BB33_63;
+
+ div.rn.f32 %f262, %f1, %f2;
+ abs.f32 %f78, %f262;
+ setp.geu.f32 %p31, %f78, 0f7F800000;
+ @%p31 bra BB33_63;
+
+ cvt.rmi.f32.f32 %f262, %f262;
+ bra.uni BB33_63;
+
+BB33_68:
+ setp.eq.s32 %p97, %r3, 1;
+ @%p97 bra BB33_69;
+ bra.uni BB33_125;
+
+BB33_69:
+ sub.f32 %f266, %f2, %f1;
+ bra.uni BB33_125;
+
+BB33_84:
+ setp.eq.s32 %p85, %r3, 10;
+ @%p85 bra BB33_85;
+ bra.uni BB33_125;
+
+BB33_85:
+ setp.neu.f32 %p106, %f2, %f1;
+ selp.f32 %f266, 0f3F800000, 0f00000000, %p106;
+ bra.uni BB33_125;
+
+BB33_75:
+ setp.eq.s32 %p92, %r3, 5;
+ @%p92 bra BB33_76;
+ bra.uni BB33_125;
+
+BB33_76:
+ setp.lt.f32 %p111, %f2, %f1;
+ selp.f32 %f266, 0f3F800000, 0f00000000, %p111;
+ bra.uni BB33_125;
+
+BB33_92:
+ setp.eq.s32 %p79, %r3, 15;
+ @%p79 bra BB33_93;
+ bra.uni BB33_125;
+
+BB33_93:
+ mul.f32 %f175, %f1, %f2;
+ mov.f32 %f176, 0f3F800000;
+ sub.f32 %f266, %f176, %f175;
+ bra.uni BB33_125;
+
+BB33_71:
+ setp.eq.s32 %p95, %r3, 3;
+ @%p95 bra BB33_72;
+ bra.uni BB33_125;
+
+BB33_72:
+ div.rn.f32 %f266, %f2, %f1;
+ bra.uni BB33_125;
+
+BB33_105:
+ min.f32 %f266, %f2, %f1;
+ bra.uni BB33_125;
+
+BB33_88:
+ setp.eq.s32 %p83, %r3, 13;
+ @%p83 bra BB33_89;
+ bra.uni BB33_125;
+
+BB33_89:
+ cvt.rni.s64.f32 %rd21, %f2;
+ cvt.u32.u64 %r40, %rd21;
+ cvt.rni.s64.f32 %rd22, %f1;
+ cvt.u32.u64 %r41, %rd22;
+ and.b32 %r42, %r41, %r40;
+ setp.eq.s32 %p105, %r42, 0;
+ selp.f32 %f266, 0f00000000, 0f3F800000, %p105;
+ bra.uni BB33_125;
+
+BB33_108:
+ setp.gtu.f32 %p110, %f2, %f1;
+ selp.f32 %f266, 0f00000000, 0f3F800000, %p110;
+ bra.uni BB33_125;
+
+BB33_79:
+ setp.eq.s32 %p90, %r3, 8;
+ @%p90 bra BB33_80;
+ bra.uni BB33_125;
+
+BB33_80:
+ setp.ltu.f32 %p108, %f2, %f1;
+ selp.f32 %f266, 0f00000000, 0f3F800000, %p108;
+ bra.uni BB33_125;
+
+BB33_102:
+ setp.neu.f32 %p103, %f2, 0f00000000;
+ sub.f32 %f174, %f2, %f1;
+ selp.f32 %f266, %f174, 0f00000000, %p103;
+ bra.uni BB33_125;
+
+BB33_96:
+ setp.ne.s32 %p77, %r3, 18;
+ @%p77 bra BB33_125;
+
+ div.rn.f32 %f266, %f2, %f1;
+ abs.f32 %f169, %f266;
+ setp.geu.f32 %p98, %f169, 0f7F800000;
+ @%p98 bra BB33_125;
+
+ cvt.rmi.f32.f32 %f266, %f266;
+ bra.uni BB33_125;
+
+BB33_50:
+ setp.geu.f32 %p55, %f1, 0f00000000;
+ @%p55 bra BB33_53;
+
+ cvt.rzi.f32.f32 %f164, %f2;
+ setp.neu.f32 %p56, %f164, %f2;
+ selp.f32 %f261, 0f7FFFFFFF, %f261, %p56;
+
+BB33_53:
+ add.f32 %f166, %f20, %f21;
+ mov.b32 %r30, %f166;
+ setp.lt.s32 %p59, %r30, 2139095040;
+ @%p59 bra BB33_60;
+
+ setp.gtu.f32 %p60, %f20, 0f7F800000;
+ setp.gtu.f32 %p61, %f21, 0f7F800000;
+ or.pred %p62, %p60, %p61;
+ @%p62 bra BB33_59;
+ bra.uni BB33_55;
+
+BB33_59:
+ add.f32 %f261, %f1, %f2;
+ bra.uni BB33_60;
+
+BB33_55:
+ setp.eq.f32 %p63, %f21, 0f7F800000;
+ @%p63 bra BB33_58;
+ bra.uni BB33_56;
+
+BB33_58:
+ setp.gt.f32 %p66, %f20, 0f3F800000;
+ selp.b32 %r34, 2139095040, 0, %p66;
+ xor.b32 %r35, %r34, 2139095040;
+ setp.lt.f32 %p67, %f2, 0f00000000;
+ selp.b32 %r36, %r35, %r34, %p67;
+ mov.b32 %f167, %r36;
+ setp.eq.f32 %p68, %f1, 0fBF800000;
+ selp.f32 %f261, 0f3F800000, %f167, %p68;
+ bra.uni BB33_60;
+
+BB33_112:
+ setp.geu.f32 %p122, %f2, 0f00000000;
+ @%p122 bra BB33_115;
+
+ cvt.rzi.f32.f32 %f255, %f1;
+ setp.neu.f32 %p123, %f255, %f1;
+ selp.f32 %f265, 0f7FFFFFFF, %f265, %p123;
+
+BB33_115:
+ add.f32 %f257, %f57, %f58;
+ mov.b32 %r55, %f257;
+ setp.lt.s32 %p126, %r55, 2139095040;
+ @%p126 bra BB33_122;
+
+ setp.gtu.f32 %p127, %f57, 0f7F800000;
+ setp.gtu.f32 %p128, %f58, 0f7F800000;
+ or.pred %p129, %p127, %p128;
+ @%p129 bra BB33_121;
+ bra.uni BB33_117;
+
+BB33_121:
+ add.f32 %f265, %f1, %f2;
+ bra.uni BB33_122;
+
+BB33_117:
+ setp.eq.f32 %p130, %f58, 0f7F800000;
+ @%p130 bra BB33_120;
+ bra.uni BB33_118;
+
+BB33_120:
+ setp.gt.f32 %p133, %f57, 0f3F800000;
+ selp.b32 %r59, 2139095040, 0, %p133;
+ xor.b32 %r60, %r59, 2139095040;
+ setp.lt.f32 %p134, %f1, 0f00000000;
+ selp.b32 %r61, %r60, %r59, %p134;
+ mov.b32 %f258, %r61;
+ setp.eq.f32 %p135, %f2, 0fBF800000;
+ selp.f32 %f265, 0f3F800000, %f258, %p135;
+ bra.uni BB33_122;
+
+BB33_56:
+ setp.neu.f32 %p64, %f20, 0f7F800000;
+ @%p64 bra BB33_60;
+
+ setp.ltu.f32 %p65, %f2, 0f00000000;
+ selp.b32 %r31, 0, 2139095040, %p65;
+ or.b32 %r32, %r31, -2147483648;
+ selp.b32 %r33, %r32, %r31, %p1;
+ mov.b32 %f261, %r33;
+
+BB33_60:
+ setp.eq.f32 %p69, %f2, 0f00000000;
+ setp.eq.f32 %p70, %f1, 0f3F800000;
+ or.pred %p71, %p70, %p69;
+ selp.f32 %f262, 0f3F800000, %f261, %p71;
+
+BB33_63:
+ st.global.f32 [%rd1], %f262;
+ bra.uni BB33_126;
+
+BB33_118:
+ setp.neu.f32 %p131, %f57, 0f7F800000;
+ @%p131 bra BB33_122;
+
+ setp.ltu.f32 %p132, %f1, 0f00000000;
+ selp.b32 %r56, 0, 2139095040, %p132;
+ or.b32 %r57, %r56, -2147483648;
+ selp.b32 %r58, %r57, %r56, %p2;
+ mov.b32 %f265, %r58;
+
+BB33_122:
+ setp.eq.f32 %p136, %f1, 0f00000000;
+ setp.eq.f32 %p137, %f2, 0f3F800000;
+ or.pred %p138, %p137, %p136;
+ selp.f32 %f266, 0f3F800000, %f265, %p138;
+
+BB33_125:
+ st.global.f32 [%rd1], %f266;
+
+BB33_126:
+ bar.sync 0;
+ ret;
+}
+
// .globl fill_d
.visible .entry fill_d(
.param .u64 fill_d_param_0,
@@ -4616,14 +6349,14 @@ BB31_126:
mov.u32 %r5, %tid.x;
mad.lo.s32 %r1, %r4, %r3, %r5;
setp.ge.s32 %p1, %r1, %r2;
- @%p1 bra BB32_2;
+ @%p1 bra BB34_2;
cvta.to.global.u64 %rd2, %rd1;
mul.wide.s32 %rd3, %r1, 8;
add.s64 %rd4, %rd2, %rd3;
st.global.f64 [%rd4], %fd1;
-BB32_2:
+BB34_2:
ret;
}
@@ -4649,7 +6382,7 @@ BB32_2:
mov.u32 %r5, %tid.x;
mad.lo.s32 %r1, %r4, %r3, %r5;
setp.ge.s32 %p1, %r1, %r2;
- @%p1 bra BB33_2;
+ @%p1 bra BB35_2;
cvt.rn.f32.f64 %f1, %fd1;
cvta.to.global.u64 %rd2, %rd1;
@@ -4657,7 +6390,7 @@ BB32_2:
add.s64 %rd4, %rd2, %rd3;
st.global.f32 [%rd4], %f1;
-BB33_2:
+BB35_2:
ret;
}
@@ -4697,10 +6430,10 @@ BB33_2:
setp.lt.s32 %p1, %r1, %r7;
setp.lt.s32 %p2, %r2, %r4;
and.pred %p3, %p1, %p2;
- @!%p3 bra BB34_2;
- bra.uni BB34_1;
+ @!%p3 bra BB36_2;
+ bra.uni BB36_1;
-BB34_1:
+BB36_1:
cvta.to.global.u64 %rd5, %rd2;
mad.lo.s32 %r13, %r1, %r4, %r2;
mul.wide.s32 %rd6, %r13, 8;
@@ -4711,14 +6444,14 @@ BB34_1:
add.s64 %rd9, %rd1, %rd8;
st.global.f64 [%rd9], %fd1;
-BB34_2:
+BB36_2:
setp.lt.s32 %p4, %r1, %r5;
setp.lt.s32 %p5, %r2, %r6;
and.pred %p6, %p4, %p5;
- @!%p6 bra BB34_4;
- bra.uni BB34_3;
+ @!%p6 bra BB36_4;
+ bra.uni BB36_3;
-BB34_3:
+BB36_3:
cvta.to.global.u64 %rd10, %rd3;
mad.lo.s32 %r15, %r1, %r6, %r2;
mul.wide.s32 %rd11, %r15, 8;
@@ -4730,7 +6463,7 @@ BB34_3:
add.s64 %rd14, %rd1, %rd13;
st.global.f64 [%rd14], %fd2;
-BB34_4:
+BB36_4:
ret;
}
@@ -4770,10 +6503,10 @@ BB34_4:
setp.lt.s32 %p1, %r1, %r7;
setp.lt.s32 %p2, %r2, %r4;
and.pred %p3, %p1, %p2;
- @!%p3 bra BB35_2;
- bra.uni BB35_1;
+ @!%p3 bra BB37_2;
+ bra.uni BB37_1;
-BB35_1:
+BB37_1:
cvta.to.global.u64 %rd5, %rd2;
mad.lo.s32 %r13, %r1, %r4, %r2;
mul.wide.s32 %rd6, %r13, 4;
@@ -4784,14 +6517,14 @@ BB35_1:
add.s64 %rd9, %rd1, %rd8;
st.global.f32 [%rd9], %f1;
-BB35_2:
+BB37_2:
setp.lt.s32 %p4, %r1, %r5;
setp.lt.s32 %p5, %r2, %r6;
and.pred %p6, %p4, %p5;
- @!%p6 bra BB35_4;
- bra.uni BB35_3;
+ @!%p6 bra BB37_4;
+ bra.uni BB37_3;
-BB35_3:
+BB37_3:
cvta.to.global.u64 %rd10, %rd3;
mad.lo.s32 %r15, %r1, %r6, %r2;
mul.wide.s32 %rd11, %r15, 4;
@@ -4803,7 +6536,7 @@ BB35_3:
add.s64 %rd14, %rd1, %rd13;
st.global.f32 [%rd14], %f2;
-BB35_4:
+BB37_4:
ret;
}
@@ -4842,10 +6575,10 @@ BB35_4:
setp.lt.s32 %p1, %r1, %r3;
setp.lt.s32 %p2, %r2, %r4;
and.pred %p3, %p1, %p2;
- @!%p3 bra BB36_2;
- bra.uni BB36_1;
+ @!%p3 bra BB38_2;
+ bra.uni BB38_1;
-BB36_1:
+BB38_1:
cvta.to.global.u64 %rd5, %rd2;
mad.lo.s32 %r12, %r1, %r4, %r2;
mul.wide.s32 %rd6, %r12, 8;
@@ -4854,14 +6587,14 @@ BB36_1:
add.s64 %rd8, %rd1, %rd6;
st.global.f64 [%rd8], %fd1;
-BB36_2:
+BB38_2:
setp.lt.s32 %p4, %r1, %r5;
setp.lt.s32 %p5, %r2, %r6;
and.pred %p6, %p4, %p5;
- @!%p6 bra BB36_4;
- bra.uni BB36_3;
+ @!%p6 bra BB38_4;
+ bra.uni BB38_3;
-BB36_3:
+BB38_3:
cvta.to.global.u64 %rd9, %rd3;
mad.lo.s32 %r13, %r1, %r6, %r2;
mul.wide.s32 %rd10, %r13, 8;
@@ -4873,7 +6606,7 @@ BB36_3:
add.s64 %rd13, %rd1, %rd12;
st.global.f64 [%rd13], %fd2;
-BB36_4:
+BB38_4:
ret;
}
@@ -4912,10 +6645,10 @@ BB36_4:
setp.lt.s32 %p1, %r1, %r3;
setp.lt.s32 %p2, %r2, %r4;
and.pred %p3, %p1, %p2;
- @!%p3 bra BB37_2;
- bra.uni BB37_1;
+ @!%p3 bra BB39_2;
+ bra.uni BB39_1;
-BB37_1:
+BB39_1:
cvta.to.global.u64 %rd5, %rd2;
mad.lo.s32 %r12, %r1, %r4, %r2;
mul.wide.s32 %rd6, %r12, 4;
@@ -4924,14 +6657,14 @@ BB37_1:
add.s64 %rd8, %rd1, %rd6;
st.global.f32 [%rd8], %f1;
-BB37_2:
+BB39_2:
setp.lt.s32 %p4, %r1, %r5;
setp.lt.s32 %p5, %r2, %r6;
and.pred %p6, %p4, %p5;
- @!%p6 bra BB37_4;
- bra.uni BB37_3;
+ @!%p6 bra BB39_4;
+ bra.uni BB39_3;
-BB37_3:
+BB39_3:
cvta.to.global.u64 %rd9, %rd3;
mad.lo.s32 %r13, %r1, %r6, %r2;
mul.wide.s32 %rd10, %r13, 4;
@@ -4943,7 +6676,7 @@ BB37_3:
add.s64 %rd13, %rd1, %rd12;
st.global.f32 [%rd13], %f2;
-BB37_4:
+BB39_4:
ret;
}
@@ -4970,9 +6703,9 @@ BB37_4:
mad.lo.s32 %r35, %r9, %r10, %r7;
mov.f64 %fd44, 0d0000000000000000;
setp.ge.u32 %p1, %r35, %r6;
- @%p1 bra BB38_4;
+ @%p1 bra BB40_4;
-BB38_1:
+BB40_1:
cvta.to.global.u64 %rd3, %rd1;
mul.wide.u32 %rd4, %r35, 8;
add.s64 %rd5, %rd3, %rd4;
@@ -4980,135 +6713,135 @@ BB38_1:
add.f64 %fd44, %fd44, %fd30;
add.s32 %r3, %r35, %r10;
setp.ge.u32 %p2, %r3, %r6;
- @%p2 bra BB38_3;
+ @%p2 bra BB40_3;
mul.wide.u32 %rd7, %r3, 8;
add.s64 %rd8, %rd3, %rd7;
ld.global.f64 %fd31, [%rd8];
add.f64 %fd44, %fd44, %fd31;
-BB38_3:
+BB40_3:
shl.b32 %r13, %r10, 1;
mov.u32 %r14, %nctaid.x;
mad.lo.s32 %r35, %r13, %r14, %r35;
setp.lt.u32 %p3, %r35, %r6;
- @%p3 bra BB38_1;
+ @%p3 bra BB40_1;
-BB38_4:
+BB40_4:
shl.b32 %r16, %r7, 3;
mov.u32 %r17, my_sdata;
add.s32 %r5, %r17, %r16;
st.shared.f64 [%r5], %fd44;
bar.sync 0;
setp.lt.u32 %p4, %r10, 1024;
- @%p4 bra BB38_8;
+ @%p4 bra BB40_8;
setp.gt.u32 %p5, %r7, 511;
- @%p5 bra BB38_7;
+ @%p5 bra BB40_7;
ld.shared.f64 %fd32, [%r5+4096];
add.f64 %fd44, %fd44, %fd32;
st.shared.f64 [%r5], %fd44;
-BB38_7:
+BB40_7:
bar.sync 0;
-BB38_8:
+BB40_8:
setp.lt.u32 %p6, %r10, 512;
- @%p6 bra BB38_12;
+ @%p6 bra BB40_12;
setp.gt.u32 %p7, %r7, 255;
- @%p7 bra BB38_11;
+ @%p7 bra BB40_11;
ld.shared.f64 %fd33, [%r5+2048];
add.f64 %fd44, %fd44, %fd33;
st.shared.f64 [%r5], %fd44;
-BB38_11:
+BB40_11:
bar.sync 0;
-BB38_12:
+BB40_12:
setp.lt.u32 %p8, %r10, 256;
- @%p8 bra BB38_16;
+ @%p8 bra BB40_16;
setp.gt.u32 %p9, %r7, 127;
- @%p9 bra BB38_15;
+ @%p9 bra BB40_15;
ld.shared.f64 %fd34, [%r5+1024];
add.f64 %fd44, %fd44, %fd34;
st.shared.f64 [%r5], %fd44;
-BB38_15:
+BB40_15:
bar.sync 0;
-BB38_16:
+BB40_16:
setp.lt.u32 %p10, %r10, 128;
- @%p10 bra BB38_20;
+ @%p10 bra BB40_20;
setp.gt.u32 %p11, %r7, 63;
- @%p11 bra BB38_19;
+ @%p11 bra BB40_19;
ld.shared.f64 %fd35, [%r5+512];
add.f64 %fd44, %fd44, %fd35;
st.shared.f64 [%r5], %fd44;
-BB38_19:
+BB40_19:
bar.sync 0;
-BB38_20:
+BB40_20:
setp.gt.u32 %p12, %r7, 31;
- @%p12 bra BB38_33;
+ @%p12 bra BB40_33;
setp.lt.u32 %p13, %r10, 64;
- @%p13 bra BB38_23;
+ @%p13 bra BB40_23;
ld.volatile.shared.f64 %fd36, [%r5+256];
add.f64 %fd44, %fd44, %fd36;
st.volatile.shared.f64 [%r5], %fd44;
-BB38_23:
+BB40_23:
setp.lt.u32 %p14, %r10, 32;
- @%p14 bra BB38_25;
+ @%p14 bra BB40_25;
ld.volatile.shared.f64 %fd37, [%r5+128];
add.f64 %fd44, %fd44, %fd37;
st.volatile.shared.f64 [%r5], %fd44;
-BB38_25:
+BB40_25:
setp.lt.u32 %p15, %r10, 16;
- @%p15 bra BB38_27;
+ @%p15 bra BB40_27;
ld.volatile.shared.f64 %fd38, [%r5+64];
add.f64 %fd44, %fd44, %fd38;
st.volatile.shared.f64 [%r5], %fd44;
-BB38_27:
+BB40_27:
setp.lt.u32 %p16, %r10, 8;
- @%p16 bra BB38_29;
+ @%p16 bra BB40_29;
ld.volatile.shared.f64 %fd39, [%r5+32];
add.f64 %fd44, %fd44, %fd39;
st.volatile.shared.f64 [%r5], %fd44;
-BB38_29:
+BB40_29:
setp.lt.u32 %p17, %r10, 4;
- @%p17 bra BB38_31;
+ @%p17 bra BB40_31;
ld.volatile.shared.f64 %fd40, [%r5+16];
add.f64 %fd44, %fd44, %fd40;
st.volatile.shared.f64 [%r5], %fd44;
-BB38_31:
+BB40_31:
setp.lt.u32 %p18, %r10, 2;
- @%p18 bra BB38_33;
+ @%p18 bra BB40_33;
ld.volatile.shared.f64 %fd41, [%r5+8];
add.f64 %fd42, %fd44, %fd41;
st.volatile.shared.f64 [%r5], %fd42;
-BB38_33:
+BB40_33:
setp.ne.s32 %p19, %r7, 0;
- @%p19 bra BB38_35;
+ @%p19 bra BB40_35;
ld.shared.f64 %fd43, [my_sdata];
cvta.to.global.u64 %rd9, %rd2;
@@ -5116,7 +6849,7 @@ BB38_33:
add.s64 %rd11, %rd9, %rd10;
st.global.f64 [%rd11], %fd43;
-BB38_35:
+BB40_35:
ret;
}
@@ -5143,9 +6876,9 @@ BB38_35:
mad.lo.s32 %r35, %r9, %r10, %r7;
mov.f32 %f44, 0f00000000;
setp.ge.u32 %p1, %r35, %r6;
- @%p1 bra BB39_4;
+ @%p1 bra BB41_4;
-BB39_1:
+BB41_1:
cvta.to.global.u64 %rd3, %rd1;
mul.wide.u32 %rd4, %r35, 4;
add.s64 %rd5, %rd3, %rd4;
@@ -5153,135 +6886,135 @@ BB39_1:
add.f32 %f44, %f44, %f30;
add.s32 %r3, %r35, %r10;
setp.ge.u32 %p2, %r3, %r6;
- @%p2 bra BB39_3;
+ @%p2 bra BB41_3;
mul.wide.u32 %rd7, %r3, 4;
add.s64 %rd8, %rd3, %rd7;
ld.global.f32 %f31, [%rd8];
add.f32 %f44, %f44, %f31;
-BB39_3:
+BB41_3:
shl.b32 %r13, %r10, 1;
mov.u32 %r14, %nctaid.x;
mad.lo.s32 %r35, %r13, %r14, %r35;
setp.lt.u32 %p3, %r35, %r6;
- @%p3 bra BB39_1;
+ @%p3 bra BB41_1;
-BB39_4:
+BB41_4:
shl.b32 %r16, %r7, 2;
mov.u32 %r17, my_sdata;
add.s32 %r5, %r17, %r16;
st.shared.f32 [%r5], %f44;
bar.sync 0;
setp.lt.u32 %p4, %r10, 1024;
- @%p4 bra BB39_8;
+ @%p4 bra BB41_8;
setp.gt.u32 %p5, %r7, 511;
- @%p5 bra BB39_7;
+ @%p5 bra BB41_7;
ld.shared.f32 %f32, [%r5+2048];
add.f32 %f44, %f44, %f32;
st.shared.f32 [%r5], %f44;
-BB39_7:
+BB41_7:
bar.sync 0;
-BB39_8:
+BB41_8:
setp.lt.u32 %p6, %r10, 512;
- @%p6 bra BB39_12;
+ @%p6 bra BB41_12;
setp.gt.u32 %p7, %r7, 255;
- @%p7 bra BB39_11;
+ @%p7 bra BB41_11;
ld.shared.f32 %f33, [%r5+1024];
add.f32 %f44, %f44, %f33;
st.shared.f32 [%r5], %f44;
-BB39_11:
+BB41_11:
bar.sync 0;
-BB39_12:
+BB41_12:
setp.lt.u32 %p8, %r10, 256;
- @%p8 bra BB39_16;
+ @%p8 bra BB41_16;
setp.gt.u32 %p9, %r7, 127;
- @%p9 bra BB39_15;
+ @%p9 bra BB41_15;
ld.shared.f32 %f34, [%r5+512];
add.f32 %f44, %f44, %f34;
st.shared.f32 [%r5], %f44;
-BB39_15:
+BB41_15:
bar.sync 0;
-BB39_16:
+BB41_16:
setp.lt.u32 %p10, %r10, 128;
- @%p10 bra BB39_20;
+ @%p10 bra BB41_20;
setp.gt.u32 %p11, %r7, 63;
- @%p11 bra BB39_19;
+ @%p11 bra BB41_19;
ld.shared.f32 %f35, [%r5+256];
add.f32 %f44, %f44, %f35;
st.shared.f32 [%r5], %f44;
-BB39_19:
+BB41_19:
bar.sync 0;
-BB39_20:
+BB41_20:
setp.gt.u32 %p12, %r7, 31;
- @%p12 bra BB39_33;
+ @%p12 bra BB41_33;
setp.lt.u32 %p13, %r10, 64;
- @%p13 bra BB39_23;
+ @%p13 bra BB41_23;
ld.volatile.shared.f32 %f36, [%r5+128];
add.f32 %f44, %f44, %f36;
st.volatile.shared.f32 [%r5], %f44;
-BB39_23:
+BB41_23:
setp.lt.u32 %p14, %r10, 32;
- @%p14 bra BB39_25;
+ @%p14 bra BB41_25;
ld.volatile.shared.f32 %f37, [%r5+64];
add.f32 %f44, %f44, %f37;
st.volatile.shared.f32 [%r5], %f44;
-BB39_25:
+BB41_25:
setp.lt.u32 %p15, %r10, 16;
- @%p15 bra BB39_27;
+ @%p15 bra BB41_27;
ld.volatile.shared.f32 %f38, [%r5+32];
add.f32 %f44, %f44, %f38;
st.volatile.shared.f32 [%r5], %f44;
-BB39_27:
+BB41_27:
setp.lt.u32 %p16, %r10, 8;
- @%p16 bra BB39_29;
+ @%p16 bra BB41_29;
ld.volatile.shared.f32 %f39, [%r5+16];
add.f32 %f44, %f44, %f39;
st.volatile.shared.f32 [%r5], %f44;
-BB39_29:
+BB41_29:
setp.lt.u32 %p17, %r10, 4;
- @%p17 bra BB39_31;
+ @%p17 bra BB41_31;
ld.volatile.shared.f32 %f40, [%r5+8];
add.f32 %f44, %f44, %f40;
st.volatile.shared.f32 [%r5], %f44;
-BB39_31:
+BB41_31:
setp.lt.u32 %p18, %r10, 2;
- @%p18 bra BB39_33;
+ @%p18 bra BB41_33;
ld.volatile.shared.f32 %f41, [%r5+4];
add.f32 %f42, %f44, %f41;
st.volatile.shared.f32 [%r5], %f42;
-BB39_33:
+BB41_33:
setp.ne.s32 %p19, %r7, 0;
- @%p19 bra BB39_35;
+ @%p19 bra BB41_35;
ld.shared.f32 %f43, [my_sdata];
cvta.to.global.u64 %rd9, %rd2;
@@ -5289,7 +7022,7 @@ BB39_33:
add.s64 %rd11, %rd9, %rd10;
st.global.f32 [%rd11], %f43;
-BB39_35:
+BB41_35:
ret;
}
@@ -5313,16 +7046,16 @@ BB39_35:
ld.param.u32 %r4, [reduce_row_sum_d_param_3];
mov.u32 %r6, %ctaid.x;
setp.ge.u32 %p1, %r6, %r5;
- @%p1 bra BB40_35;
+ @%p1 bra BB42_35;
mov.u32 %r71, %tid.x;
mov.f64 %fd6, 0d0000000000000000;
setp.ge.u32 %p2, %r71, %r4;
- @%p2 bra BB40_4;
+ @%p2 bra BB42_4;
cvta.to.global.u64 %rd3, %rd1;
-BB40_3:
+BB42_3:
mad.lo.s32 %r8, %r6, %r4, %r71;
mul.wide.u32 %rd4, %r8, 8;
add.s64 %rd5, %rd3, %rd4;
@@ -5331,9 +7064,9 @@ BB40_3:
mov.u32 %r9, %ntid.x;
add.s32 %r71, %r9, %r71;
setp.lt.u32 %p3, %r71, %r4;
- @%p3 bra BB40_3;
+ @%p3 bra BB42_3;
-BB40_4:
+BB42_4:
mov.u32 %r10, %tid.x;
shl.b32 %r11, %r10, 3;
mov.u32 %r12, my_sdata;
@@ -5342,114 +7075,114 @@ BB40_4:
bar.sync 0;
mov.u32 %r14, %ntid.x;
setp.lt.u32 %p4, %r14, 1024;
- @%p4 bra BB40_8;
+ @%p4 bra BB42_8;
setp.gt.u32 %p5, %r10, 511;
- @%p5 bra BB40_7;
+ @%p5 bra BB42_7;
ld.shared.f64 %fd29, [%r13+4096];
add.f64 %fd6, %fd6, %fd29;
st.shared.f64 [%r13], %fd6;
-BB40_7:
+BB42_7:
bar.sync 0;
-BB40_8:
+BB42_8:
setp.lt.u32 %p6, %r14, 512;
- @%p6 bra BB40_12;
+ @%p6 bra BB42_12;
setp.gt.u32 %p7, %r10, 255;
- @%p7 bra BB40_11;
+ @%p7 bra BB42_11;
ld.shared.f64 %fd30, [%r13+2048];
add.f64 %fd6, %fd6, %fd30;
st.shared.f64 [%r13], %fd6;
-BB40_11:
+BB42_11:
bar.sync 0;
-BB40_12:
+BB42_12:
setp.lt.u32 %p8, %r14, 256;
- @%p8 bra BB40_16;
+ @%p8 bra BB42_16;
setp.gt.u32 %p9, %r10, 127;
- @%p9 bra BB40_15;
+ @%p9 bra BB42_15;
ld.shared.f64 %fd31, [%r13+1024];
add.f64 %fd6, %fd6, %fd31;
st.shared.f64 [%r13], %fd6;
-BB40_15:
+BB42_15:
bar.sync 0;
-BB40_16:
+BB42_16:
setp.lt.u32 %p10, %r14, 128;
- @%p10 bra BB40_20;
+ @%p10 bra BB42_20;
setp.gt.u32 %p11, %r10, 63;
- @%p11 bra BB40_19;
+ @%p11 bra BB42_19;
ld.shared.f64 %fd32, [%r13+512];
add.f64 %fd6, %fd6, %fd32;
st.shared.f64 [%r13], %fd6;
-BB40_19:
+BB42_19:
bar.sync 0;
-BB40_20:
+BB42_20:
setp.gt.u32 %p12, %r10, 31;
- @%p12 bra BB40_33;
+ @%p12 bra BB42_33;
setp.lt.u32 %p13, %r14, 64;
- @%p13 bra BB40_23;
+ @%p13 bra BB42_23;
ld.volatile.shared.f64 %fd33, [%r13+256];
add.f64 %fd6, %fd6, %fd33;
st.volatile.shared.f64 [%r13], %fd6;
-BB40_23:
+BB42_23:
setp.lt.u32 %p14, %r14, 32;
- @%p14 bra BB40_25;
+ @%p14 bra BB42_25;
ld.volatile.shared.f64 %fd34, [%r13+128];
add.f64 %fd6, %fd6, %fd34;
st.volatile.shared.f64 [%r13], %fd6;
-BB40_25:
+BB42_25:
setp.lt.u32 %p15, %r14, 16;
- @%p15 bra BB40_27;
+ @%p15 bra BB42_27;
ld.volatile.shared.f64 %fd35, [%r13+64];
add.f64 %fd6, %fd6, %fd35;
st.volatile.shared.f64 [%r13], %fd6;
-BB40_27:
+BB42_27:
setp.lt.u32 %p16, %r14, 8;
- @%p16 bra BB40_29;
+ @%p16 bra BB42_29;
ld.volatile.shared.f64 %fd36, [%r13+32];
add.f64 %fd6, %fd6, %fd36;
st.volatile.shared.f64 [%r13], %fd6;
-BB40_29:
+BB42_29:
setp.lt.u32 %p17, %r14, 4;
- @%p17 bra BB40_31;
+ @%p17 bra BB42_31;
ld.volatile.shared.f64 %fd37, [%r13+16];
add.f64 %fd6, %fd6, %fd37;
st.volatile.shared.f64 [%r13], %fd6;
-BB40_31:
+BB42_31:
setp.lt.u32 %p18, %r14, 2;
- @%p18 bra BB40_33;
+ @%p18 bra BB42_33;
ld.volatile.shared.f64 %fd38, [%r13+8];
add.f64 %fd39, %fd6, %fd38;
st.volatile.shared.f64 [%r13], %fd39;
-BB40_33:
+BB42_33:
setp.ne.s32 %p19, %r10, 0;
- @%p19 bra BB40_35;
+ @%p19 bra BB42_35;
ld.shared.f64 %fd40, [my_sdata];
cvta.to.global.u64 %rd6, %rd2;
@@ -5457,7 +7190,7 @@ BB40_33:
add.s64 %rd8, %rd6, %rd7;
st.global.f64 [%rd8], %fd40;
-BB40_35:
+BB42_35:
ret;
}
@@ -5481,16 +7214,16 @@ BB40_35:
ld.param.u32 %r4, [reduce_row_sum_f_param_3];
mov.u32 %r6, %ctaid.x;
setp.ge.u32 %p1, %r6, %r5;
- @%p1 bra BB41_35;
+ @%p1 bra BB43_35;
mov.u32 %r71, %tid.x;
mov.f32 %f6, 0f00000000;
setp.ge.u32 %p2, %r71, %r4;
- @%p2 bra BB41_4;
+ @%p2 bra BB43_4;
cvta.to.global.u64 %rd3, %rd1;
-BB41_3:
+BB43_3:
mad.lo.s32 %r8, %r6, %r4, %r71;
mul.wide.u32 %rd4, %r8, 4;
add.s64 %rd5, %rd3, %rd4;
@@ -5499,9 +7232,9 @@ BB41_3:
mov.u32 %r9, %ntid.x;
add.s32 %r71, %r9, %r71;
setp.lt.u32 %p3, %r71, %r4;
- @%p3 bra BB41_3;
+ @%p3 bra BB43_3;
-BB41_4:
+BB43_4:
mov.u32 %r10, %tid.x;
shl.b32 %r11, %r10, 2;
mov.u32 %r12, my_sdata;
@@ -5510,114 +7243,114 @@ BB41_4:
bar.sync 0;
mov.u32 %r14, %ntid.x;
setp.lt.u32 %p4, %r14, 1024;
- @%p4 bra BB41_8;
+ @%p4 bra BB43_8;
setp.gt.u32 %p5, %r10, 511;
- @%p5 bra BB41_7;
+ @%p5 bra BB43_7;
ld.shared.f32 %f29, [%r13+2048];
add.f32 %f6, %f6, %f29;
st.shared.f32 [%r13], %f6;
-BB41_7:
+BB43_7:
bar.sync 0;
-BB41_8:
+BB43_8:
setp.lt.u32 %p6, %r14, 512;
- @%p6 bra BB41_12;
+ @%p6 bra BB43_12;
setp.gt.u32 %p7, %r10, 255;
- @%p7 bra BB41_11;
+ @%p7 bra BB43_11;
ld.shared.f32 %f30, [%r13+1024];
add.f32 %f6, %f6, %f30;
st.shared.f32 [%r13], %f6;
-BB41_11:
+BB43_11:
bar.sync 0;
-BB41_12:
+BB43_12:
setp.lt.u32 %p8, %r14, 256;
- @%p8 bra BB41_16;
+ @%p8 bra BB43_16;
setp.gt.u32 %p9, %r10, 127;
- @%p9 bra BB41_15;
+ @%p9 bra BB43_15;
ld.shared.f32 %f31, [%r13+512];
add.f32 %f6, %f6, %f31;
st.shared.f32 [%r13], %f6;
-BB41_15:
+BB43_15:
bar.sync 0;
-BB41_16:
+BB43_16:
setp.lt.u32 %p10, %r14, 128;
- @%p10 bra BB41_20;
+ @%p10 bra BB43_20;
setp.gt.u32 %p11, %r10, 63;
- @%p11 bra BB41_19;
+ @%p11 bra BB43_19;
ld.shared.f32 %f32, [%r13+256];
add.f32 %f6, %f6, %f32;
st.shared.f32 [%r13], %f6;
-BB41_19:
+BB43_19:
bar.sync 0;
-BB41_20:
+BB43_20:
setp.gt.u32 %p12, %r10, 31;
- @%p12 bra BB41_33;
+ @%p12 bra BB43_33;
setp.lt.u32 %p13, %r14, 64;
- @%p13 bra BB41_23;
+ @%p13 bra BB43_23;
ld.volatile.shared.f32 %f33, [%r13+128];
add.f32 %f6, %f6, %f33;
st.volatile.shared.f32 [%r13], %f6;
-BB41_23:
+BB43_23:
setp.lt.u32 %p14, %r14, 32;
- @%p14 bra BB41_25;
+ @%p14 bra BB43_25;
ld.volatile.shared.f32 %f34, [%r13+64];
add.f32 %f6, %f6, %f34;
st.volatile.shared.f32 [%r13], %f6;
-BB41_25:
+BB43_25:
setp.lt.u32 %p15, %r14, 16;
- @%p15 bra BB41_27;
+ @%p15 bra BB43_27;
ld.volatile.shared.f32 %f35, [%r13+32];
add.f32 %f6, %f6, %f35;
st.volatile.shared.f32 [%r13], %f6;
-BB41_27:
+BB43_27:
setp.lt.u32 %p16, %r14, 8;
- @%p16 bra BB41_29;
+ @%p16 bra BB43_29;
ld.volatile.shared.f32 %f36, [%r13+16];
add.f32 %f6, %f6, %f36;
st.volatile.shared.f32 [%r13], %f6;
-BB41_29:
+BB43_29:
setp.lt.u32 %p17, %r14, 4;
- @%p17 bra BB41_31;
+ @%p17 bra BB43_31;
ld.volatile.shared.f32 %f37, [%r13+8];
add.f32 %f6, %f6, %f37;
st.volatile.shared.f32 [%r13], %f6;
-BB41_31:
+BB43_31:
setp.lt.u32 %p18, %r14, 2;
- @%p18 bra BB41_33;
+ @%p18 bra BB43_33;
ld.volatile.shared.f32 %f38, [%r13+4];
add.f32 %f39, %f6, %f38;
st.volatile.shared.f32 [%r13], %f39;
-BB41_33:
+BB43_33:
setp.ne.s32 %p19, %r10, 0;
- @%p19 bra BB41_35;
+ @%p19 bra BB43_35;
ld.shared.f32 %f40, [my_sdata];
cvta.to.global.u64 %rd6, %rd2;
@@ -5625,7 +7358,7 @@ BB41_33:
add.s64 %rd8, %rd6, %rd7;
st.global.f32 [%rd8], %f40;
-BB41_35:
+BB43_35:
ret;
}
@@ -5652,32 +7385,32 @@ BB41_35:
mov.u32 %r9, %tid.x;
mad.lo.s32 %r1, %r7, %r8, %r9;
setp.ge.u32 %p1, %r1, %r6;
- @%p1 bra BB42_5;
+ @%p1 bra BB44_5;
cvta.to.global.u64 %rd1, %rd2;
mul.lo.s32 %r2, %r6, %r5;
mov.f64 %fd8, 0d0000000000000000;
setp.ge.u32 %p2, %r1, %r2;
- @%p2 bra BB42_4;
+ @%p2 bra BB44_4;
mov.u32 %r10, %r1;
-BB42_3:
+BB44_3:
mul.wide.u32 %rd4, %r10, 8;
add.s64 %rd5, %rd1, %rd4;
ld.global.f64 %fd6, [%rd5];
add.f64 %fd8, %fd8, %fd6;
add.s32 %r10, %r10, %r6;
setp.lt.u32 %p3, %r10, %r2;
- @%p3 bra BB42_3;
+ @%p3 bra BB44_3;
-BB42_4:
+BB44_4:
cvta.to.global.u64 %rd6, %rd3;
mul.wide.u32 %rd7, %r1, 8;
add.s64 %rd8, %rd6, %rd7;
st.global.f64 [%rd8], %fd8;
-BB42_5:
+BB44_5:
ret;
}
@@ -5704,32 +7437,32 @@ BB42_5:
mov.u32 %r9, %tid.x;
mad.lo.s32 %r1, %r7, %r8, %r9;
setp.ge.u32 %p1, %r1, %r6;
- @%p1 bra BB43_5;
+ @%p1 bra BB45_5;
cvta.to.global.u64 %rd1, %rd2;
mul.lo.s32 %r2, %r6, %r5;
mov.f32 %f8, 0f00000000;
setp.ge.u32 %p2, %r1, %r2;
- @%p2 bra BB43_4;
+ @%p2 bra BB45_4;
mov.u32 %r10, %r1;
-BB43_3:
+BB45_3:
mul.wide.u32 %rd4, %r10, 4;
add.s64 %rd5, %rd1, %rd4;
ld.global.f32 %f6, [%rd5];
add.f32 %f8, %f8, %f6;
add.s32 %r10, %r10, %r6;
setp.lt.u32 %p3, %r10, %r2;
- @%p3 bra BB43_3;
+ @%p3 bra BB45_3;
-BB43_4:
+BB45_4:
cvta.to.global.u64 %rd6, %rd3;
mul.wide.u32 %rd7, %r1, 4;
add.s64 %rd8, %rd6, %rd7;
st.global.f32 [%rd8], %f8;
-BB43_5:
+BB45_5:
ret;
}
@@ -5756,9 +7489,9 @@ BB43_5:
mad.lo.s32 %r35, %r9, %r10, %r7;
mov.f64 %fd44, 0dFFEFFFFFFFFFFFFF;
setp.ge.u32 %p1, %r35, %r6;
- @%p1 bra BB44_4;
+ @%p1 bra BB46_4;
-BB44_1:
+BB46_1:
cvta.to.global.u64 %rd3, %rd1;
mul.wide.u32 %rd4, %r35, 8;
add.s64 %rd5, %rd3, %rd4;
@@ -5766,135 +7499,135 @@ BB44_1:
max.f64 %fd44, %fd44, %fd30;
add.s32 %r3, %r35, %r10;
setp.ge.u32 %p2, %r3, %r6;
- @%p2 bra BB44_3;
+ @%p2 bra BB46_3;
mul.wide.u32 %rd7, %r3, 8;
add.s64 %rd8, %rd3, %rd7;
ld.global.f64 %fd31, [%rd8];
max.f64 %fd44, %fd44, %fd31;
-BB44_3:
+BB46_3:
shl.b32 %r13, %r10, 1;
mov.u32 %r14, %nctaid.x;
mad.lo.s32 %r35, %r13, %r14, %r35;
setp.lt.u32 %p3, %r35, %r6;
- @%p3 bra BB44_1;
+ @%p3 bra BB46_1;
-BB44_4:
+BB46_4:
shl.b32 %r16, %r7, 3;
mov.u32 %r17, my_sdata;
add.s32 %r5, %r17, %r16;
st.shared.f64 [%r5], %fd44;
bar.sync 0;
setp.lt.u32 %p4, %r10, 1024;
- @%p4 bra BB44_8;
+ @%p4 bra BB46_8;
setp.gt.u32 %p5, %r7, 511;
- @%p5 bra BB44_7;
+ @%p5 bra BB46_7;
ld.shared.f64 %fd32, [%r5+4096];
max.f64 %fd44, %fd44, %fd32;
st.shared.f64 [%r5], %fd44;
-BB44_7:
+BB46_7:
bar.sync 0;
-BB44_8:
+BB46_8:
setp.lt.u32 %p6, %r10, 512;
- @%p6 bra BB44_12;
+ @%p6 bra BB46_12;
setp.gt.u32 %p7, %r7, 255;
- @%p7 bra BB44_11;
+ @%p7 bra BB46_11;
ld.shared.f64 %fd33, [%r5+2048];
max.f64 %fd44, %fd44, %fd33;
st.shared.f64 [%r5], %fd44;
-BB44_11:
+BB46_11:
bar.sync 0;
-BB44_12:
+BB46_12:
setp.lt.u32 %p8, %r10, 256;
- @%p8 bra BB44_16;
+ @%p8 bra BB46_16;
setp.gt.u32 %p9, %r7, 127;
- @%p9 bra BB44_15;
+ @%p9 bra BB46_15;
ld.shared.f64 %fd34, [%r5+1024];
max.f64 %fd44, %fd44, %fd34;
st.shared.f64 [%r5], %fd44;
-BB44_15:
+BB46_15:
bar.sync 0;
-BB44_16:
+BB46_16:
setp.lt.u32 %p10, %r10, 128;
- @%p10 bra BB44_20;
+ @%p10 bra BB46_20;
setp.gt.u32 %p11, %r7, 63;
- @%p11 bra BB44_19;
+ @%p11 bra BB46_19;
ld.shared.f64 %fd35, [%r5+512];
max.f64 %fd44, %fd44, %fd35;
st.shared.f64 [%r5], %fd44;
-BB44_19:
+BB46_19:
bar.sync 0;
-BB44_20:
+BB46_20:
setp.gt.u32 %p12, %r7, 31;
- @%p12 bra BB44_33;
+ @%p12 bra BB46_33;
setp.lt.u32 %p13, %r10, 64;
- @%p13 bra BB44_23;
+ @%p13 bra BB46_23;
ld.volatile.shared.f64 %fd36, [%r5+256];
max.f64 %fd44, %fd44, %fd36;
st.volatile.shared.f64 [%r5], %fd44;
-BB44_23:
+BB46_23:
setp.lt.u32 %p14, %r10, 32;
- @%p14 bra BB44_25;
+ @%p14 bra BB46_25;
ld.volatile.shared.f64 %fd37, [%r5+128];
max.f64 %fd44, %fd44, %fd37;
st.volatile.shared.f64 [%r5], %fd44;
-BB44_25:
+BB46_25:
setp.lt.u32 %p15, %r10, 16;
- @%p15 bra BB44_27;
+ @%p15 bra BB46_27;
ld.volatile.shared.f64 %fd38, [%r5+64];
max.f64 %fd44, %fd44, %fd38;
st.volatile.shared.f64 [%r5], %fd44;
-BB44_27:
+BB46_27:
setp.lt.u32 %p16, %r10, 8;
- @%p16 bra BB44_29;
+ @%p16 bra BB46_29;
ld.volatile.shared.f64 %fd39, [%r5+32];
max.f64 %fd44, %fd44, %fd39;
st.volatile.shared.f64 [%r5], %fd44;
-BB44_29:
+BB46_29:
setp.lt.u32 %p17, %r10, 4;
- @%p17 bra BB44_31;
+ @%p17 bra BB46_31;
ld.volatile.shared.f64 %fd40, [%r5+16];
max.f64 %fd44, %fd44, %fd40;
st.volatile.shared.f64 [%r5], %fd44;
-BB44_31:
+BB46_31:
setp.lt.u32 %p18, %r10, 2;
- @%p18 bra BB44_33;
+ @%p18 bra BB46_33;
ld.volatile.shared.f64 %fd41, [%r5+8];
max.f64 %fd42, %fd44, %fd41;
st.volatile.shared.f64 [%r5], %fd42;
-BB44_33:
+BB46_33:
setp.ne.s32 %p19, %r7, 0;
- @%p19 bra BB44_35;
+ @%p19 bra BB46_35;
ld.shared.f64 %fd43, [my_sdata];
cvta.to.global.u64 %rd9, %rd2;
@@ -5902,7 +7635,7 @@ BB44_33:
add.s64 %rd11, %rd9, %rd10;
st.global.f64 [%rd11], %fd43;
-BB44_35:
+BB46_35:
ret;
}
@@ -5929,9 +7662,9 @@ BB44_35:
mad.lo.s32 %r35, %r9, %r10, %r7;
mov.f32 %f44, 0fFF7FFFFF;
setp.ge.u32 %p1, %r35, %r6;
- @%p1 bra BB45_4;
+ @%p1 bra BB47_4;
-BB45_1:
+BB47_1:
cvta.to.global.u64 %rd3, %rd1;
mul.wide.u32 %rd4, %r35, 4;
add.s64 %rd5, %rd3, %rd4;
@@ -5939,135 +7672,135 @@ BB45_1:
max.f32 %f44, %f44, %f30;
add.s32 %r3, %r35, %r10;
setp.ge.u32 %p2, %r3, %r6;
- @%p2 bra BB45_3;
+ @%p2 bra BB47_3;
mul.wide.u32 %rd7, %r3, 4;
add.s64 %rd8, %rd3, %rd7;
ld.global.f32 %f31, [%rd8];
max.f32 %f44, %f44, %f31;
-BB45_3:
+BB47_3:
shl.b32 %r13, %r10, 1;
mov.u32 %r14, %nctaid.x;
mad.lo.s32 %r35, %r13, %r14, %r35;
setp.lt.u32 %p3, %r35, %r6;
- @%p3 bra BB45_1;
+ @%p3 bra BB47_1;
-BB45_4:
+BB47_4:
shl.b32 %r16, %r7, 2;
mov.u32 %r17, my_sdata;
add.s32 %r5, %r17, %r16;
st.shared.f32 [%r5], %f44;
bar.sync 0;
setp.lt.u32 %p4, %r10, 1024;
- @%p4 bra BB45_8;
+ @%p4 bra BB47_8;
setp.gt.u32 %p5, %r7, 511;
- @%p5 bra BB45_7;
+ @%p5 bra BB47_7;
ld.shared.f32 %f32, [%r5+2048];
max.f32 %f44, %f44, %f32;
st.shared.f32 [%r5], %f44;
-BB45_7:
+BB47_7:
bar.sync 0;
-BB45_8:
+BB47_8:
setp.lt.u32 %p6, %r10, 512;
- @%p6 bra BB45_12;
+ @%p6 bra BB47_12;
setp.gt.u32 %p7, %r7, 255;
- @%p7 bra BB45_11;
+ @%p7 bra BB47_11;
ld.shared.f32 %f33, [%r5+1024];
max.f32 %f44, %f44, %f33;
st.shared.f32 [%r5], %f44;
-BB45_11:
+BB47_11:
bar.sync 0;
-BB45_12:
+BB47_12:
setp.lt.u32 %p8, %r10, 256;
- @%p8 bra BB45_16;
+ @%p8 bra BB47_16;
setp.gt.u32 %p9, %r7, 127;
- @%p9 bra BB45_15;
+ @%p9 bra BB47_15;
ld.shared.f32 %f34, [%r5+512];
max.f32 %f44, %f44, %f34;
st.shared.f32 [%r5], %f44;
-BB45_15:
+BB47_15:
bar.sync 0;
-BB45_16:
+BB47_16:
setp.lt.u32 %p10, %r10, 128;
- @%p10 bra BB45_20;
+ @%p10 bra BB47_20;
setp.gt.u32 %p11, %r7, 63;
- @%p11 bra BB45_19;
+ @%p11 bra BB47_19;
ld.shared.f32 %f35, [%r5+256];
max.f32 %f44, %f44, %f35;
st.shared.f32 [%r5], %f44;
-BB45_19:
+BB47_19:
bar.sync 0;
-BB45_20:
+BB47_20:
setp.gt.u32 %p12, %r7, 31;
- @%p12 bra BB45_33;
+ @%p12 bra BB47_33;
setp.lt.u32 %p13, %r10, 64;
- @%p13 bra BB45_23;
+ @%p13 bra BB47_23;
ld.volatile.shared.f32 %f36, [%r5+128];
max.f32 %f44, %f44, %f36;
st.volatile.shared.f32 [%r5], %f44;
-BB45_23:
+BB47_23:
setp.lt.u32 %p14, %r10, 32;
- @%p14 bra BB45_25;
+ @%p14 bra BB47_25;
ld.volatile.shared.f32 %f37, [%r5+64];
max.f32 %f44, %f44, %f37;
st.volatile.shared.f32 [%r5], %f44;
-BB45_25:
+BB47_25:
setp.lt.u32 %p15, %r10, 16;
- @%p15 bra BB45_27;
+ @%p15 bra BB47_27;
ld.volatile.shared.f32 %f38, [%r5+32];
max.f32 %f44, %f44, %f38;
st.volatile.shared.f32 [%r5], %f44;
-BB45_27:
+BB47_27:
setp.lt.u32 %p16, %r10, 8;
- @%p16 bra BB45_29;
+ @%p16 bra BB47_29;
ld.volatile.shared.f32 %f39, [%r5+16];
max.f32 %f44, %f44, %f39;
st.volatile.shared.f32 [%r5], %f44;
-BB45_29:
+BB47_29:
setp.lt.u32 %p17, %r10, 4;
- @%p17 bra BB45_31;
+ @%p17 bra BB47_31;
ld.volatile.shared.f32 %f40, [%r5+8];
max.f32 %f44, %f44, %f40;
st.volatile.shared.f32 [%r5], %f44;
-BB45_31:
+BB47_31:
setp.lt.u32 %p18, %r10, 2;
- @%p18 bra BB45_33;
+ @%p18 bra BB47_33;
ld.volatile.shared.f32 %f41, [%r5+4];
max.f32 %f42, %f44, %f41;
st.volatile.shared.f32 [%r5], %f42;
-BB45_33:
+BB47_33:
setp.ne.s32 %p19, %r7, 0;
- @%p19 bra BB45_35;
+ @%p19 bra BB47_35;
ld.shared.f32 %f43, [my_sdata];
cvta.to.global.u64 %rd9, %rd2;
@@ -6075,7 +7808,7 @@ BB45_33:
add.s64 %rd11, %rd9, %rd10;
st.global.f32 [%rd11], %f43;
-BB45_35:
+BB47_35:
ret;
}
@@ -6099,16 +7832,16 @@ BB45_35:
ld.param.u32 %r4, [reduce_row_max_d_param_3];
mov.u32 %r6, %ctaid.x;
setp.ge.u32 %p1, %r6, %r5;
- @%p1 bra BB46_35;
+ @%p1 bra BB48_35;
mov.u32 %r71, %tid.x;
mov.f64 %fd6, 0dFFEFFFFFFFFFFFFF;
setp.ge.u32 %p2, %r71, %r4;
- @%p2 bra BB46_4;
+ @%p2 bra BB48_4;
cvta.to.global.u64 %rd3, %rd1;
-BB46_3:
+BB48_3:
mad.lo.s32 %r8, %r6, %r4, %r71;
mul.wide.u32 %rd4, %r8, 8;
add.s64 %rd5, %rd3, %rd4;
@@ -6117,9 +7850,9 @@ BB46_3:
mov.u32 %r9, %ntid.x;
add.s32 %r71, %r9, %r71;
setp.lt.u32 %p3, %r71, %r4;
- @%p3 bra BB46_3;
+ @%p3 bra BB48_3;
-BB46_4:
+BB48_4:
mov.u32 %r10, %tid.x;
shl.b32 %r11, %r10, 3;
mov.u32 %r12, my_sdata;
@@ -6128,114 +7861,114 @@ BB46_4:
bar.sync 0;
mov.u32 %r14, %ntid.x;
setp.lt.u32 %p4, %r14, 1024;
- @%p4 bra BB46_8;
+ @%p4 bra BB48_8;
setp.gt.u32 %p5, %r10, 511;
- @%p5 bra BB46_7;
+ @%p5 bra BB48_7;
ld.shared.f64 %fd29, [%r13+4096];
max.f64 %fd6, %fd6, %fd29;
st.shared.f64 [%r13], %fd6;
-BB46_7:
+BB48_7:
bar.sync 0;
-BB46_8:
+BB48_8:
setp.lt.u32 %p6, %r14, 512;
- @%p6 bra BB46_12;
+ @%p6 bra BB48_12;
setp.gt.u32 %p7, %r10, 255;
- @%p7 bra BB46_11;
+ @%p7 bra BB48_11;
ld.shared.f64 %fd30, [%r13+2048];
max.f64 %fd6, %fd6, %fd30;
st.shared.f64 [%r13], %fd6;
-BB46_11:
+BB48_11:
bar.sync 0;
-BB46_12:
+BB48_12:
setp.lt.u32 %p8, %r14, 256;
- @%p8 bra BB46_16;
+ @%p8 bra BB48_16;
setp.gt.u32 %p9, %r10, 127;
- @%p9 bra BB46_15;
+ @%p9 bra BB48_15;
ld.shared.f64 %fd31, [%r13+1024];
max.f64 %fd6, %fd6, %fd31;
st.shared.f64 [%r13], %fd6;
-BB46_15:
+BB48_15:
bar.sync 0;
-BB46_16:
+BB48_16:
setp.lt.u32 %p10, %r14, 128;
- @%p10 bra BB46_20;
+ @%p10 bra BB48_20;
setp.gt.u32 %p11, %r10, 63;
- @%p11 bra BB46_19;
+ @%p11 bra BB48_19;
ld.shared.f64 %fd32, [%r13+512];
max.f64 %fd6, %fd6, %fd32;
st.shared.f64 [%r13], %fd6;
-BB46_19:
+BB48_19:
bar.sync 0;
-BB46_20:
+BB48_20:
setp.gt.u32 %p12, %r10, 31;
- @%p12 bra BB46_33;
+ @%p12 bra BB48_33;
setp.lt.u32 %p13, %r14, 64;
- @%p13 bra BB46_23;
+ @%p13 bra BB48_23;
ld.volatile.shared.f64 %fd33, [%r13+256];
max.f64 %fd6, %fd6, %fd33;
st.volatile.shared.f64 [%r13], %fd6;
-BB46_23:
+BB48_23:
setp.lt.u32 %p14, %r14, 32;
- @%p14 bra BB46_25;
+ @%p14 bra BB48_25;
ld.volatile.shared.f64 %fd34, [%r13+128];
max.f64 %fd6, %fd6, %fd34;
st.volatile.shared.f64 [%r13], %fd6;
-BB46_25:
+BB48_25:
setp.lt.u32 %p15, %r14, 16;
- @%p15 bra BB46_27;
+ @%p15 bra BB48_27;
ld.volatile.shared.f64 %fd35, [%r13+64];
max.f64 %fd6, %fd6, %fd35;
st.volatile.shared.f64 [%r13], %fd6;
-BB46_27:
+BB48_27:
setp.lt.u32 %p16, %r14, 8;
- @%p16 bra BB46_29;
+ @%p16 bra BB48_29;
ld.volatile.shared.f64 %fd36, [%r13+32];
max.f64 %fd6, %fd6, %fd36;
st.volatile.shared.f64 [%r13], %fd6;
-BB46_29:
+BB48_29:
setp.lt.u32 %p17, %r14, 4;
- @%p17 bra BB46_31;
+ @%p17 bra BB48_31;
ld.volatile.shared.f64 %fd37, [%r13+16];
max.f64 %fd6, %fd6, %fd37;
st.volatile.shared.f64 [%r13], %fd6;
-BB46_31:
+BB48_31:
setp.lt.u32 %p18, %r14, 2;
- @%p18 bra BB46_33;
+ @%p18 bra BB48_33;
ld.volatile.shared.f64 %fd38, [%r13+8];
max.f64 %fd39, %fd6, %fd38;
st.volatile.shared.f64 [%r13], %fd39;
-BB46_33:
+BB48_33:
setp.ne.s32 %p19, %r10, 0;
- @%p19 bra BB46_35;
+ @%p19 bra BB48_35;
ld.shared.f64 %fd40, [my_sdata];
cvta.to.global.u64 %rd6, %rd2;
@@ -6243,7 +7976,7 @@ BB46_33:
add.s64 %rd8, %rd6, %rd7;
st.global.f64 [%rd8], %fd40;
-BB46_35:
+BB48_35:
ret;
}
@@ -6267,16 +8000,16 @@ BB46_35:
ld.param.u32 %r4, [reduce_row_max_f_param_3];
mov.u32 %r6, %ctaid.x;
setp.ge.u32 %p1, %r6, %r5;
- @%p1 bra BB47_35;
+ @%p1 bra BB49_35;
mov.u32 %r71, %tid.x;
mov.f32 %f6, 0fFF7FFFFF;
setp.ge.u32 %p2, %r71, %r4;
- @%p2 bra BB47_4;
+ @%p2 bra BB49_4;
cvta.to.global.u64 %rd3, %rd1;
-BB47_3:
+BB49_3:
mad.lo.s32 %r8, %r6, %r4, %r71;
mul.wide.u32 %rd4, %r8, 4;
add.s64 %rd5, %rd3, %rd4;
@@ -6285,9 +8018,9 @@ BB47_3:
mov.u32 %r9, %ntid.x;
add.s32 %r71, %r9, %r71;
setp.lt.u32 %p3, %r71, %r4;
- @%p3 bra BB47_3;
+ @%p3 bra BB49_3;
-BB47_4:
+BB49_4:
mov.u32 %r10, %tid.x;
shl.b32 %r11, %r10, 2;
mov.u32 %r12, my_sdata;
@@ -6296,114 +8029,114 @@ BB47_4:
bar.sync 0;
mov.u32 %r14, %ntid.x;
setp.lt.u32 %p4, %r14, 1024;
- @%p4 bra BB47_8;
+ @%p4 bra BB49_8;
setp.gt.u32 %p5, %r10, 511;
- @%p5 bra BB47_7;
+ @%p5 bra BB49_7;
ld.shared.f32 %f29, [%r13+2048];
max.f32 %f6, %f6, %f29;
st.shared.f32 [%r13], %f6;
-BB47_7:
+BB49_7:
bar.sync 0;
-BB47_8:
+BB49_8:
setp.lt.u32 %p6, %r14, 512;
- @%p6 bra BB47_12;
+ @%p6 bra BB49_12;
setp.gt.u32 %p7, %r10, 255;
- @%p7 bra BB47_11;
+ @%p7 bra BB49_11;
ld.shared.f32 %f30, [%r13+1024];
max.f32 %f6, %f6, %f30;
st.shared.f32 [%r13], %f6;
-BB47_11:
+BB49_11:
bar.sync 0;
-BB47_12:
+BB49_12:
setp.lt.u32 %p8, %r14, 256;
- @%p8 bra BB47_16;
+ @%p8 bra BB49_16;
setp.gt.u32 %p9, %r10, 127;
- @%p9 bra BB47_15;
+ @%p9 bra BB49_15;
ld.shared.f32 %f31, [%r13+512];
max.f32 %f6, %f6, %f31;
st.shared.f32 [%r13], %f6;
-BB47_15:
+BB49_15:
bar.sync 0;
-BB47_16:
+BB49_16:
setp.lt.u32 %p10, %r14, 128;
- @%p10 bra BB47_20;
+ @%p10 bra BB49_20;
setp.gt.u32 %p11, %r10, 63;
- @%p11 bra BB47_19;
+ @%p11 bra BB49_19;
ld.shared.f32 %f32, [%r13+256];
max.f32 %f6, %f6, %f32;
st.shared.f32 [%r13], %f6;
-BB47_19:
+BB49_19:
bar.sync 0;
-BB47_20:
+BB49_20:
setp.gt.u32 %p12, %r10, 31;
- @%p12 bra BB47_33;
+ @%p12 bra BB49_33;
setp.lt.u32 %p13, %r14, 64;
- @%p13 bra BB47_23;
+ @%p13 bra BB49_23;
ld.volatile.shared.f32 %f33, [%r13+128];
max.f32 %f6, %f6, %f33;
st.volatile.shared.f32 [%r13], %f6;
-BB47_23:
+BB49_23:
setp.lt.u32 %p14, %r14, 32;
- @%p14 bra BB47_25;
+ @%p14 bra BB49_25;
ld.volatile.shared.f32 %f34, [%r13+64];
max.f32 %f6, %f6, %f34;
st.volatile.shared.f32 [%r13], %f6;
-BB47_25:
+BB49_25:
setp.lt.u32 %p15, %r14, 16;
- @%p15 bra BB47_27;
+ @%p15 bra BB49_27;
ld.volatile.shared.f32 %f35, [%r13+32];
max.f32 %f6, %f6, %f35;
st.volatile.shared.f32 [%r13], %f6;
-BB47_27:
+BB49_27:
setp.lt.u32 %p16, %r14, 8;
- @%p16 bra BB47_29;
+ @%p16 bra BB49_29;
ld.volatile.shared.f32 %f36, [%r13+16];
max.f32 %f6, %f6, %f36;
st.volatile.shared.f32 [%r13], %f6;
-BB47_29:
+BB49_29:
setp.lt.u32 %p17, %r14, 4;
- @%p17 bra BB47_31;
+ @%p17 bra BB49_31;
ld.volatile.shared.f32 %f37, [%r13+8];
max.f32 %f6, %f6, %f37;
st.volatile.shared.f32 [%r13], %f6;
-BB47_31:
+BB49_31:
setp.lt.u32 %p18, %r14, 2;
- @%p18 bra BB47_33;
+ @%p18 bra BB49_33;
ld.volatile.shared.f32 %f38, [%r13+4];
max.f32 %f39, %f6, %f38;
st.volatile.shared.f32 [%r13], %f39;
-BB47_33:
+BB49_33:
setp.ne.s32 %p19, %r10, 0;
- @%p19 bra BB47_35;
+ @%p19 bra BB49_35;
ld.shared.f32 %f40, [my_sdata];
cvta.to.global.u64 %rd6, %rd2;
@@ -6411,7 +8144,7 @@ BB47_33:
add.s64 %rd8, %rd6, %rd7;
st.global.f32 [%rd8], %f40;
-BB47_35:
+BB49_35:
ret;
}
@@ -6438,32 +8171,32 @@ BB47_35:
mov.u32 %r9, %tid.x;
mad.lo.s32 %r1, %r7, %r8, %r9;
setp.ge.u32 %p1, %r1, %r6;
- @%p1 bra BB48_5;
+ @%p1 bra BB50_5;
cvta.to.global.u64 %rd1, %rd2;
mul.lo.s32 %r2, %r6, %r5;
mov.f64 %fd8, 0dFFEFFFFFFFFFFFFF;
setp.ge.u32 %p2, %r1, %r2;
- @%p2 bra BB48_4;
+ @%p2 bra BB50_4;
mov.u32 %r10, %r1;
-BB48_3:
+BB50_3:
mul.wide.u32 %rd4, %r10, 8;
add.s64 %rd5, %rd1, %rd4;
ld.global.f64 %fd6, [%rd5];
max.f64 %fd8, %fd8, %fd6;
add.s32 %r10, %r10, %r6;
setp.lt.u32 %p3, %r10, %r2;
- @%p3 bra BB48_3;
+ @%p3 bra BB50_3;
-BB48_4:
+BB50_4:
cvta.to.global.u64 %rd6, %rd3;
mul.wide.u32 %rd7, %r1, 8;
add.s64 %rd8, %rd6, %rd7;
st.global.f64 [%rd8], %fd8;
-BB48_5:
+BB50_5:
ret;
}
@@ -6490,32 +8223,32 @@ BB48_5:
mov.u32 %r9, %tid.x;
mad.lo.s32 %r1, %r7, %r8, %r9;
setp.ge.u32 %p1, %r1, %r6;
- @%p1 bra BB49_5;
+ @%p1 bra BB51_5;
cvta.to.global.u64 %rd1, %rd2;
mul.lo.s32 %r2, %r6, %r5;
mov.f32 %f8, 0fFF7FFFFF;
setp.ge.u32 %p2, %r1, %r2;
- @%p2 bra BB49_4;
+ @%p2 bra BB51_4;
mov.u32 %r10, %r1;
-BB49_3:
+BB51_3:
mul.wide.u32 %rd4, %r10, 4;
add.s64 %rd5, %rd1, %rd4;
ld.global.f32 %f6, [%rd5];
max.f32 %f8, %f8, %f6;
add.s32 %r10, %r10, %r6;
setp.lt.u32 %p3, %r10, %r2;
- @%p3 bra BB49_3;
+ @%p3 bra BB51_3;
-BB49_4:
+BB51_4:
cvta.to.global.u64 %rd6, %rd3;
mul.wide.u32 %rd7, %r1, 4;
add.s64 %rd8, %rd6, %rd7;
st.global.f32 [%rd8], %f8;
-BB49_5:
+BB51_5:
ret;
}
@@ -6542,9 +8275,9 @@ BB49_5:
mad.lo.s32 %r35, %r9, %r10, %r7;
mov.f64 %fd44, 0d7FEFFFFFFFFFFFFF;
setp.ge.u32 %p1, %r35, %r6;
- @%p1 bra BB50_4;
+ @%p1 bra BB52_4;
-BB50_1:
+BB52_1:
cvta.to.global.u64 %rd3, %rd1;
mul.wide.u32 %rd4, %r35, 8;
add.s64 %rd5, %rd3, %rd4;
@@ -6552,135 +8285,135 @@ BB50_1:
min.f64 %fd44, %fd44, %fd30;
add.s32 %r3, %r35, %r10;
setp.ge.u32 %p2, %r3, %r6;
- @%p2 bra BB50_3;
+ @%p2 bra BB52_3;
mul.wide.u32 %rd7, %r3, 8;
add.s64 %rd8, %rd3, %rd7;
ld.global.f64 %fd31, [%rd8];
min.f64 %fd44, %fd44, %fd31;
-BB50_3:
+BB52_3:
shl.b32 %r13, %r10, 1;
mov.u32 %r14, %nctaid.x;
mad.lo.s32 %r35, %r13, %r14, %r35;
setp.lt.u32 %p3, %r35, %r6;
- @%p3 bra BB50_1;
+ @%p3 bra BB52_1;
-BB50_4:
+BB52_4:
shl.b32 %r16, %r7, 3;
mov.u32 %r17, my_sdata;
add.s32 %r5, %r17, %r16;
st.shared.f64 [%r5], %fd44;
bar.sync 0;
setp.lt.u32 %p4, %r10, 1024;
- @%p4 bra BB50_8;
+ @%p4 bra BB52_8;
setp.gt.u32 %p5, %r7, 511;
- @%p5 bra BB50_7;
+ @%p5 bra BB52_7;
ld.shared.f64 %fd32, [%r5+4096];
min.f64 %fd44, %fd44, %fd32;
st.shared.f64 [%r5], %fd44;
-BB50_7:
+BB52_7:
bar.sync 0;
-BB50_8:
+BB52_8:
setp.lt.u32 %p6, %r10, 512;
- @%p6 bra BB50_12;
+ @%p6 bra BB52_12;
setp.gt.u32 %p7, %r7, 255;
- @%p7 bra BB50_11;
+ @%p7 bra BB52_11;
ld.shared.f64 %fd33, [%r5+2048];
min.f64 %fd44, %fd44, %fd33;
st.shared.f64 [%r5], %fd44;
-BB50_11:
+BB52_11:
bar.sync 0;
-BB50_12:
+BB52_12:
setp.lt.u32 %p8, %r10, 256;
- @%p8 bra BB50_16;
+ @%p8 bra BB52_16;
setp.gt.u32 %p9, %r7, 127;
- @%p9 bra BB50_15;
+ @%p9 bra BB52_15;
ld.shared.f64 %fd34, [%r5+1024];
min.f64 %fd44, %fd44, %fd34;
st.shared.f64 [%r5], %fd44;
-BB50_15:
+BB52_15:
bar.sync 0;
-BB50_16:
+BB52_16:
setp.lt.u32 %p10, %r10, 128;
- @%p10 bra BB50_20;
+ @%p10 bra BB52_20;
setp.gt.u32 %p11, %r7, 63;
- @%p11 bra BB50_19;
+ @%p11 bra BB52_19;
ld.shared.f64 %fd35, [%r5+512];
min.f64 %fd44, %fd44, %fd35;
st.shared.f64 [%r5], %fd44;
-BB50_19:
+BB52_19:
bar.sync 0;
-BB50_20:
+BB52_20:
setp.gt.u32 %p12, %r7, 31;
- @%p12 bra BB50_33;
+ @%p12 bra BB52_33;
setp.lt.u32 %p13, %r10, 64;
- @%p13 bra BB50_23;
+ @%p13 bra BB52_23;
ld.volatile.shared.f64 %fd36, [%r5+256];
min.f64 %fd44, %fd44, %fd36;
st.volatile.shared.f64 [%r5], %fd44;
-BB50_23:
+BB52_23:
setp.lt.u32 %p14, %r10, 32;
- @%p14 bra BB50_25;
+ @%p14 bra BB52_25;
ld.volatile.shared.f64 %fd37, [%r5+128];
min.f64 %fd44, %fd44, %fd37;
st.volatile.shared.f64 [%r5], %fd44;
-BB50_25:
+BB52_25:
setp.lt.u32 %p15, %r10, 16;
- @%p15 bra BB50_27;
+ @%p15 bra BB52_27;
ld.volatile.shared.f64 %fd38, [%r5+64];
min.f64 %fd44, %fd44, %fd38;
st.volatile.shared.f64 [%r5], %fd44;
-BB50_27:
+BB52_27:
setp.lt.u32 %p16, %r10, 8;
- @%p16 bra BB50_29;
+ @%p16 bra BB52_29;
ld.volatile.shared.f64 %fd39, [%r5+32];
min.f64 %fd44, %fd44, %fd39;
st.volatile.shared.f64 [%r5], %fd44;
-BB50_29:
+BB52_29:
setp.lt.u32 %p17, %r10, 4;
- @%p17 bra BB50_31;
+ @%p17 bra BB52_31;
ld.volatile.shared.f64 %fd40, [%r5+16];
min.f64 %fd44, %fd44, %fd40;
st.volatile.shared.f64 [%r5], %fd44;
-BB50_31:
+BB52_31:
setp.lt.u32 %p18, %r10, 2;
- @%p18 bra BB50_33;
+ @%p18 bra BB52_33;
ld.volatile.shared.f64 %fd41, [%r5+8];
min.f64 %fd42, %fd44, %fd41;
st.volatile.shared.f64 [%r5], %fd42;
-BB50_33:
+BB52_33:
setp.ne.s32 %p19, %r7, 0;
- @%p19 bra BB50_35;
+ @%p19 bra BB52_35;
ld.shared.f64 %fd43, [my_sdata];
cvta.to.global.u64 %rd9, %rd2;
@@ -6688,7 +8421,7 @@ BB50_33:
add.s64 %rd11, %rd9, %rd10;
st.global.f64 [%rd11], %fd43;
-BB50_35:
+BB52_35:
ret;
}
@@ -6715,9 +8448,9 @@ BB50_35:
mad.lo.s32 %r35, %r9, %r10, %r7;
mov.f32 %f44, 0f7F7FFFFF;
setp.ge.u32 %p1, %r35, %r6;
- @%p1 bra BB51_4;
+ @%p1 bra BB53_4;
-BB51_1:
+BB53_1:
cvta.to.global.u64 %rd3, %rd1;
mul.wide.u32 %rd4, %r35, 4;
add.s64 %rd5, %rd3, %rd4;
@@ -6725,135 +8458,135 @@ BB51_1:
min.f32 %f44, %f44, %f30;
add.s32 %r3, %r35, %r10;
setp.ge.u32 %p2, %r3, %r6;
- @%p2 bra BB51_3;
+ @%p2 bra BB53_3;
mul.wide.u32 %rd7, %r3, 4;
add.s64 %rd8, %rd3, %rd7;
ld.global.f32 %f31, [%rd8];
min.f32 %f44, %f44, %f31;
-BB51_3:
+BB53_3:
shl.b32 %r13, %r10, 1;
mov.u32 %r14, %nctaid.x;
mad.lo.s32 %r35, %r13, %r14, %r35;
setp.lt.u32 %p3, %r35, %r6;
- @%p3 bra BB51_1;
+ @%p3 bra BB53_1;
-BB51_4:
+BB53_4:
shl.b32 %r16, %r7, 2;
mov.u32 %r17, my_sdata;
add.s32 %r5, %r17, %r16;
st.shared.f32 [%r5], %f44;
bar.sync 0;
setp.lt.u32 %p4, %r10, 1024;
- @%p4 bra BB51_8;
+ @%p4 bra BB53_8;
setp.gt.u32 %p5, %r7, 511;
- @%p5 bra BB51_7;
+ @%p5 bra BB53_7;
ld.shared.f32 %f32, [%r5+2048];
min.f32 %f44, %f44, %f32;
st.shared.f32 [%r5], %f44;
-BB51_7:
+BB53_7:
bar.sync 0;
-BB51_8:
+BB53_8:
setp.lt.u32 %p6, %r10, 512;
- @%p6 bra BB51_12;
+ @%p6 bra BB53_12;
setp.gt.u32 %p7, %r7, 255;
- @%p7 bra BB51_11;
+ @%p7 bra BB53_11;
ld.shared.f32 %f33, [%r5+1024];
min.f32 %f44, %f44, %f33;
st.shared.f32 [%r5], %f44;
-BB51_11:
+BB53_11:
bar.sync 0;
-BB51_12:
+BB53_12:
setp.lt.u32 %p8, %r10, 256;
- @%p8 bra BB51_16;
+ @%p8 bra BB53_16;
setp.gt.u32 %p9, %r7, 127;
- @%p9 bra BB51_15;
+ @%p9 bra BB53_15;
ld.shared.f32 %f34, [%r5+512];
min.f32 %f44, %f44, %f34;
st.shared.f32 [%r5], %f44;
-BB51_15:
+BB53_15:
bar.sync 0;
-BB51_16:
+BB53_16:
setp.lt.u32 %p10, %r10, 128;
- @%p10 bra BB51_20;
+ @%p10 bra BB53_20;
setp.gt.u32 %p11, %r7, 63;
- @%p11 bra BB51_19;
+ @%p11 bra BB53_19;
ld.shared.f32 %f35, [%r5+256];
min.f32 %f44, %f44, %f35;
st.shared.f32 [%r5], %f44;
-BB51_19:
+BB53_19:
bar.sync 0;
-BB51_20:
+BB53_20:
setp.gt.u32 %p12, %r7, 31;
- @%p12 bra BB51_33;
+ @%p12 bra BB53_33;
setp.lt.u32 %p13, %r10, 64;
- @%p13 bra BB51_23;
+ @%p13 bra BB53_23;
ld.volatile.shared.f32 %f36, [%r5+128];
min.f32 %f44, %f44, %f36;
st.volatile.shared.f32 [%r5], %f44;
-BB51_23:
+BB53_23:
setp.lt.u32 %p14, %r10, 32;
- @%p14 bra BB51_25;
+ @%p14 bra BB53_25;
ld.volatile.shared.f32 %f37, [%r5+64];
min.f32 %f44, %f44, %f37;
st.volatile.shared.f32 [%r5], %f44;
-BB51_25:
+BB53_25:
setp.lt.u32 %p15, %r10, 16;
- @%p15 bra BB51_27;
+ @%p15 bra BB53_27;
ld.volatile.shared.f32 %f38, [%r5+32];
min.f32 %f44, %f44, %f38;
st.volatile.shared.f32 [%r5], %f44;
-BB51_27:
+BB53_27:
setp.lt.u32 %p16, %r10, 8;
- @%p16 bra BB51_29;
+ @%p16 bra BB53_29;
ld.volatile.shared.f32 %f39, [%r5+16];
min.f32 %f44, %f44, %f39;
st.volatile.shared.f32 [%r5], %f44;
-BB51_29:
+BB53_29:
setp.lt.u32 %p17, %r10, 4;
- @%p17 bra BB51_31;
+ @%p17 bra BB53_31;
ld.volatile.shared.f32 %f40, [%r5+8];
min.f32 %f44, %f44, %f40;
st.volatile.shared.f32 [%r5], %f44;
-BB51_31:
+BB53_31:
setp.lt.u32 %p18, %r10, 2;
- @%p18 bra BB51_33;
+ @%p18 bra BB53_33;
ld.volatile.shared.f32 %f41, [%r5+4];
min.f32 %f42, %f44, %f41;
st.volatile.shared.f32 [%r5], %f42;
-BB51_33:
+BB53_33:
setp.ne.s32 %p19, %r7, 0;
- @%p19 bra BB51_35;
+ @%p19 bra BB53_35;
ld.shared.f32 %f43, [my_sdata];
cvta.to.global.u64 %rd9, %rd2;
@@ -6861,7 +8594,7 @@ BB51_33:
add.s64 %rd11, %rd9, %rd10;
st.global.f32 [%rd11], %f43;
-BB51_35:
+BB53_35:
ret;
}
@@ -6885,16 +8618,16 @@ BB51_35:
ld.param.u32 %r4, [reduce_row_min_d_param_3];
mov.u32 %r6, %ctaid.x;
setp.ge.u32 %p1, %r6, %r5;
- @%p1 bra BB52_35;
+ @%p1 bra BB54_35;
mov.u32 %r71, %tid.x;
mov.f64 %fd6, 0d7FEFFFFFFFFFFFFF;
setp.ge.u32 %p2, %r71, %r4;
- @%p2 bra BB52_4;
+ @%p2 bra BB54_4;
cvta.to.global.u64 %rd3, %rd1;
-BB52_3:
+BB54_3:
mad.lo.s32 %r8, %r6, %r4, %r71;
mul.wide.u32 %rd4, %r8, 8;
add.s64 %rd5, %rd3, %rd4;
@@ -6903,9 +8636,9 @@ BB52_3:
mov.u32 %r9, %ntid.x;
add.s32 %r71, %r9, %r71;
setp.lt.u32 %p3, %r71, %r4;
- @%p3 bra BB52_3;
+ @%p3 bra BB54_3;
-BB52_4:
+BB54_4:
mov.u32 %r10, %tid.x;
shl.b32 %r11, %r10, 3;
mov.u32 %r12, my_sdata;
@@ -6914,114 +8647,114 @@ BB52_4:
bar.sync 0;
mov.u32 %r14, %ntid.x;
setp.lt.u32 %p4, %r14, 1024;
- @%p4 bra BB52_8;
+ @%p4 bra BB54_8;
setp.gt.u32 %p5, %r10, 511;
- @%p5 bra BB52_7;
+ @%p5 bra BB54_7;
ld.shared.f64 %fd29, [%r13+4096];
min.f64 %fd6, %fd6, %fd29;
st.shared.f64 [%r13], %fd6;
-BB52_7:
+BB54_7:
bar.sync 0;
-BB52_8:
+BB54_8:
setp.lt.u32 %p6, %r14, 512;
- @%p6 bra BB52_12;
+ @%p6 bra BB54_12;
setp.gt.u32 %p7, %r10, 255;
- @%p7 bra BB52_11;
+ @%p7 bra BB54_11;
ld.shared.f64 %fd30, [%r13+2048];
min.f64 %fd6, %fd6, %fd30;
st.shared.f64 [%r13], %fd6;
-BB52_11:
+BB54_11:
bar.sync 0;
-BB52_12:
+BB54_12:
setp.lt.u32 %p8, %r14, 256;
- @%p8 bra BB52_16;
+ @%p8 bra BB54_16;
setp.gt.u32 %p9, %r10, 127;
- @%p9 bra BB52_15;
+ @%p9 bra BB54_15;
ld.shared.f64 %fd31, [%r13+1024];
min.f64 %fd6, %fd6, %fd31;
st.shared.f64 [%r13], %fd6;
-BB52_15:
+BB54_15:
bar.sync 0;
-BB52_16:
+BB54_16:
setp.lt.u32 %p10, %r14, 128;
- @%p10 bra BB52_20;
+ @%p10 bra BB54_20;
setp.gt.u32 %p11, %r10, 63;
- @%p11 bra BB52_19;
+ @%p11 bra BB54_19;
ld.shared.f64 %fd32, [%r13+512];
min.f64 %fd6, %fd6, %fd32;
st.shared.f64 [%r13], %fd6;
-BB52_19:
+BB54_19:
bar.sync 0;
-BB52_20:
+BB54_20:
setp.gt.u32 %p12, %r10, 31;
- @%p12 bra BB52_33;
+ @%p12 bra BB54_33;
setp.lt.u32 %p13, %r14, 64;
- @%p13 bra BB52_23;
+ @%p13 bra BB54_23;
ld.volatile.shared.f64 %fd33, [%r13+256];
min.f64 %fd6, %fd6, %fd33;
st.volatile.shared.f64 [%r13], %fd6;
-BB52_23:
+BB54_23:
setp.lt.u32 %p14, %r14, 32;
- @%p14 bra BB52_25;
+ @%p14 bra BB54_25;
ld.volatile.shared.f64 %fd34, [%r13+128];
min.f64 %fd6, %fd6, %fd34;
st.volatile.shared.f64 [%r13], %fd6;
-BB52_25:
+BB54_25:
setp.lt.u32 %p15, %r14, 16;
- @%p15 bra BB52_27;
+ @%p15 bra BB54_27;
ld.volatile.shared.f64 %fd35, [%r13+64];
min.f64 %fd6, %fd6, %fd35;
st.volatile.shared.f64 [%r13], %fd6;
-BB52_27:
+BB54_27:
setp.lt.u32 %p16, %r14, 8;
- @%p16 bra BB52_29;
+ @%p16 bra BB54_29;
ld.volatile.shared.f64 %fd36, [%r13+32];
min.f64 %fd6, %fd6, %fd36;
st.volatile.shared.f64 [%r13], %fd6;
-BB52_29:
+BB54_29:
setp.lt.u32 %p17, %r14, 4;
- @%p17 bra BB52_31;
+ @%p17 bra BB54_31;
ld.volatile.shared.f64 %fd37, [%r13+16];
min.f64 %fd6, %fd6, %fd37;
st.volatile.shared.f64 [%r13], %fd6;
-BB52_31:
+BB54_31:
setp.lt.u32 %p18, %r14, 2;
- @%p18 bra BB52_33;
+ @%p18 bra BB54_33;
ld.volatile.shared.f64 %fd38, [%r13+8];
min.f64 %fd39, %fd6, %fd38;
st.volatile.shared.f64 [%r13], %fd39;
-BB52_33:
+BB54_33:
setp.ne.s32 %p19, %r10, 0;
- @%p19 bra BB52_35;
+ @%p19 bra BB54_35;
ld.shared.f64 %fd40, [my_sdata];
cvta.to.global.u64 %rd6, %rd2;
@@ -7029,7 +8762,7 @@ BB52_33:
add.s64 %rd8, %rd6, %rd7;
st.global.f64 [%rd8], %fd40;
-BB52_35:
+BB54_35:
ret;
}
@@ -7053,16 +8786,16 @@ BB52_35:
ld.param.u32 %r4, [reduce_row_min_f_param_3];
mov.u32 %r6, %ctaid.x;
setp.ge.u32 %p1, %r6, %r5;
- @%p1 bra BB53_35;
+ @%p1 bra BB55_35;
mov.u32 %r71, %tid.x;
mov.f32 %f6, 0f7F7FFFFF;
setp.ge.u32 %p2, %r71, %r4;
- @%p2 bra BB53_4;
+ @%p2 bra BB55_4;
cvta.to.global.u64 %rd3, %rd1;
-BB53_3:
+BB55_3:
mad.lo.s32 %r8, %r6, %r4, %r71;
mul.wide.u32 %rd4, %r8, 4;
add.s64 %rd5, %rd3, %rd4;
@@ -7071,9 +8804,9 @@ BB53_3:
mov.u32 %r9, %ntid.x;
add.s32 %r71, %r9, %r71;
setp.lt.u32 %p3, %r71, %r4;
- @%p3 bra BB53_3;
+ @%p3 bra BB55_3;
-BB53_4:
+BB55_4:
mov.u32 %r10, %tid.x;
shl.b32 %r11, %r10, 2;
mov.u32 %r12, my_sdata;
@@ -7082,114 +8815,114 @@ BB53_4:
bar.sync 0;
mov.u32 %r14, %ntid.x;
setp.lt.u32 %p4, %r14, 1024;
- @%p4 bra BB53_8;
+ @%p4 bra BB55_8;
setp.gt.u32 %p5, %r10, 511;
- @%p5 bra BB53_7;
+ @%p5 bra BB55_7;
ld.shared.f32 %f29, [%r13+2048];
min.f32 %f6, %f6, %f29;
st.shared.f32 [%r13], %f6;
-BB53_7:
+BB55_7:
bar.sync 0;
-BB53_8:
+BB55_8:
setp.lt.u32 %p6, %r14, 512;
- @%p6 bra BB53_12;
+ @%p6 bra BB55_12;
setp.gt.u32 %p7, %r10, 255;
- @%p7 bra BB53_11;
+ @%p7 bra BB55_11;
ld.shared.f32 %f30, [%r13+1024];
min.f32 %f6, %f6, %f30;
st.shared.f32 [%r13], %f6;
-BB53_11:
+BB55_11:
bar.sync 0;
-BB53_12:
+BB55_12:
setp.lt.u32 %p8, %r14, 256;
- @%p8 bra BB53_16;
+ @%p8 bra BB55_16;
setp.gt.u32 %p9, %r10, 127;
- @%p9 bra BB53_15;
+ @%p9 bra BB55_15;
ld.shared.f32 %f31, [%r13+512];
min.f32 %f6, %f6, %f31;
st.shared.f32 [%r13], %f6;
-BB53_15:
+BB55_15:
bar.sync 0;
-BB53_16:
+BB55_16:
setp.lt.u32 %p10, %r14, 128;
- @%p10 bra BB53_20;
+ @%p10 bra BB55_20;
setp.gt.u32 %p11, %r10, 63;
- @%p11 bra BB53_19;
+ @%p11 bra BB55_19;
ld.shared.f32 %f32, [%r13+256
<TRUNCATED>