118 ld.param.u64 %rd4, [_Z11mm2_kernel2PfS_S_Pd_param_0];
119 ld.param.u64 %rd5, [_Z11mm2_kernel2PfS_S_Pd_param_1];
120 ld.param.u64 %rd6, [_Z11mm2_kernel2PfS_S_Pd_param_2];
121 mov.u32 %r6, %ntid.x;
122 mov.u32 %r7, %ctaid.x;
123 mov.u32 %r8, %tid.x;
124 mad.lo.s32 %r1, %r6, %r7, %r8;
125 mov.u32 %r9, %ntid.y;
126 mov.u32 %r10, %ctaid.y;
127 mov.u32 %r11, %tid.y;
128 mad.lo.s32 %r2, %r9, %r10, %r11;
129 setp.lt.s32%p1, %r2, 256;
130 setp.lt.s32%p2, %r1, 256;
131 and.pred %p3, %p1, %p2;
132 @!%p3 bra BB1_3;
133 bra.uni BB1_1;
136 cvta.to.global.u64 %rd1, %rd5;
137 cvta.to.global.u64 %rd2, %rd4;
138 cvta.to.global.u64 %rd7, %rd6;
139 shl.b32 %r3, %r2, 8;
140 add.s32 %r13, %r3, %r1;
141 mul.wide.s32 %rd8, %r13, 4;
142 add.s64 %rd3, %rd7, %rd8;
143 ld.global.f32 %f27, [%rd3];
144 mov.u32 %r17, 0;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
147 add.s32 %r14, %r17, %r3;
148 mul.wide.s32 %rd9, %r14, 4;
149 add.s64 %rd10, %rd2, %rd9;
150 shl.b32 %r15, %r17, 8;
151 add.s32 %r16, %r15, %r1;
152 mul.wide.s32 %rd11, %r16, 4;
153 add.s64 %rd12, %rd1, %rd11;
154 ld.global.f32 %f4, [%rd12];
155 ld.global.f32 %f5, [%rd10];
156 fma.rn.f32 %f6, %f5, %f4, %f27;
157 st.global.f32 [%rd3], %f6;
158 ld.global.f32 %f7, [%rd12+1024];
159 ld.global.f32 %f8, [%rd10+4];
160 fma.rn.f32 %f9, %f8, %f7, %f6;
161 st.global.f32 [%rd3], %f9;
162 ld.global.f32 %f10, [%rd12+2048];
163 ld.global.f32 %f11, [%rd10+8];
164 fma.rn.f32 %f12, %f11, %f10, %f9;
165 st.global.f32 [%rd3], %f12;
166 ld.global.f32 %f13, [%rd12+3072];
167 ld.global.f32 %f14, [%rd10+12];
168 fma.rn.f32 %f15, %f14, %f13, %f12;
169 st.global.f32 [%rd3], %f15;
170 ld.global.f32 %f16, [%rd12+4096];
171 ld.global.f32 %f17, [%rd10+16];
172 fma.rn.f32 %f18, %f17, %f16, %f15;
173 st.global.f32 [%rd3], %f18;
174 ld.global.f32 %f19, [%rd12+5120];
175 ld.global.f32 %f20, [%rd10+20];
176 fma.rn.f32 %f21, %f20, %f19, %f18;
177 st.global.f32 [%rd3], %f21;
178 ld.global.f32 %f22, [%rd12+6144];
179 ld.global.f32 %f23, [%rd10+24];
180 fma.rn.f32 %f24, %f23, %f22, %f21;
181 st.global.f32 [%rd3], %f24;
182 ld.global.f32 %f25, [%rd12+7168];
183 ld.global.f32 %f26, [%rd10+28];
184 fma.rn.f32 %f27, %f26, %f25, %f24;
185 st.global.f32 [%rd3], %f27;
186 add.s32 %r17, %r17, 8;
187 setp.ne.s32%p4, %r17, 256;
188 @%p4 bra BB1_2;
191 ret;
