define ptx_kernel void @main_graph_async_dispatch_33_softmax_Dx28xDxDxf32_generic(ptr addrspace(1) noalias noundef nonnull readonly align 16 captures(none) %0, ptr addrspace(1) noalias noundef nonnull readonly align 16 captures(none) %1, ptr addrspace(1) noalias nofree noundef nonnull writeonly align 16 captures(none) %2, i32 noundef %3, i32 noundef %4, i32 noundef %5, i32 noundef %6, i32 noundef %7, i32 noundef %8, i32 noundef %9, i32 noundef %10, i32 noundef %11, i32 noundef %12, i32 noundef %13, i32 noundef %14, i32 noundef %15, i32 noundef %16) local_unnamed_addr #9 { %18 = zext i32 %3 to i64 %19 = zext i32 %4 to i64 %20 = shl nuw i64 %19, 32 %21 = or disjoint i64 %20, %18 %22 = zext i32 %9 to i64 %23 = zext i32 %10 to i64 %24 = shl nuw i64 %23, 32 %25 = or disjoint i64 %24, %22 %26 = zext i32 %11 to i64 %27 = zext i32 %12 to i64 %28 = shl nuw i64 %27, 32 %29 = or disjoint i64 %28, %26 %30 = zext i32 %13 to i64 %31 = zext nneg i32 %14 to i64 %32 = shl nuw nsw i64 %31, 32 %33 = or disjoint i64 %32, %30 %34 = zext i32 %15 to i64 %35 = zext i32 %16 to i64 %36 = shl nuw i64 %35, 32 %37 = or disjoint i64 %36, %34 %38 = icmp ult i32 %10, 2097152 tail call void @llvm.assume(i1 %38) %39 = icmp ult i64 %29, 18014398509481983 tail call void @llvm.assume(i1 %39) %40 = icmp ult i32 %14, 2097152 tail call void @llvm.assume(i1 %40) %41 = icmp ult i64 %37, 18014398509481983 tail call void @llvm.assume(i1 %41) %42 = mul i64 %29, %25 call void @llvm.assume(i1 true) [ "align"(ptr addrspace(1) %0, i64 64) ] %43 = lshr i64 %21, 1 %44 = and i64 %43, 1152921504606846975 %45 = mul i64 %37, %25 %46 = mul i64 %45, 28 %47 = getelementptr [2 x i8], ptr addrspace(1) %1, i64 %44 call void @llvm.assume(i1 true) [ "align"(ptr addrspace(1) %47, i64 64) ] call void @llvm.assume(i1 true) [ "align"(ptr addrspace(1) %2, i64 64) ] %48 = mul nuw nsw i64 %25, 28 %49 = mul nuw nsw i64 %33, %48 %50 = tail call range(i32 0, 2147483647) i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() %51 = zext nneg i32 %50 to i64 %52 = tail call range(i32 1, -2147483648) i32 @llvm.nvvm.read.ptx.sreg.nctaid.x() %53 = zext nneg i32 %52 to i64 %54 = icmp samesign ugt i64 %49, %51 br i1 %54, label %.lr.ph23, label %._crit_edge24 .lr.ph23: ; preds = %17 %55 = tail call range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() %56 = icmp eq i64 %37, 0 %57 = tail call i64 @llvm.usub.sat.i64(i64 %37, i64 1) %58 = lshr i64 %57, 3 %59 = and i64 %58, 2251799813684224 %60 = shl nuw nsw i32 %55, 3 %61 = insertelement <8 x i32> poison, i32 %60, i64 0 %62 = shufflevector <8 x i32> %61, <8 x i32> poison, <8 x i32> zeroinitializer %63 = zext nneg i32 %60 to i64 %64 = insertelement <8 x i64> poison, i64 %37, i64 0 %65 = shufflevector <8 x i64> %64, <8 x i64> poison, <8 x i32> zeroinitializer %66 = lshr i32 %55, 5 %67 = zext nneg i32 %66 to i64 %68 = and i32 %55, 31 %69 = icmp eq i32 %68, 0 %70 = getelementptr [4 x i8], ptr addrspace(3) @__dynamic_shared_memory__, i64 %67 %71 = zext nneg i32 %68 to i64 %72 = getelementptr [4 x i8], ptr addrspace(3) @__dynamic_shared_memory__, i64 %71 %73 = shl nuw nsw i32 %66, 8 %74 = shl nuw nsw i32 %68, 3 %.scalar27 = or disjoint i32 %74, %73 %75 = insertelement <8 x i32> poison, i32 %.scalar27, i64 0 %76 = shufflevector <8 x i32> %75, <8 x i32> poison, <8 x i32> zeroinitializer %77 = or disjoint i32 %73, %74 %78 = zext nneg i32 %77 to i64 %79 = getelementptr [4 x i8], ptr addrspace(3) getelementptr (i8, ptr addrspace(3) @__dynamic_shared_memory__, i64 136), i64 %67 %80 = getelementptr [4 x i8], ptr addrspace(3) getelementptr (i8, ptr addrspace(3) @__dynamic_shared_memory__, i64 136), i64 %71 br label %81 81: ; preds = %.lr.ph23, %._crit_edge21 %82 = phi i64 [ %51, %.lr.ph23 ], [ %558, %._crit_edge21 ] %.frozen = freeze i64 %82 %83 = sdiv i64 %.frozen, %48 %84 = mul i64 %83, %48 %85 = icmp ne i64 %82, %84 %86 = icmp slt i64 %82, 0 %87 = and i1 %86, %85 %88 = sext i1 %87 to i64 %89 = add i64 %83, %88 %90 = mul i64 %83, %48 %.decomposed = sub i64 %.frozen, %90 %91 = icmp slt i64 %.decomposed, 0 %92 = select i1 %91, i64 %48, i64 0 %93 = add nsw i64 %92, %.decomposed %94 = sdiv i64 %93, %25 %95 = srem i64 %82, %25 %96 = icmp slt i64 %95, 0 %97 = select i1 %96, i64 %25, i64 0 %98 = add nsw i64 %97, %95 br i1 %56, label %._crit_edge, label %.lr.ph .lr.ph: ; preds = %81 %99 = tail call <8 x i64> @llvm.stepvector.v8i64() %100 = trunc <8 x i64> %99 to <8 x i32> %101 = add <8 x i32> %62, %100 %102 = zext <8 x i32> %101 to <8 x i64> %103 = mul i64 %89, %46 %104 = mul i64 %94, %45 %105 = mul i64 %98, %37 %106 = getelementptr [2 x i8], ptr addrspace(1) %47, i64 %103 %107 = getelementptr [2 x i8], ptr addrspace(1) %106, i64 %104 %108 = getelementptr [2 x i8], ptr addrspace(1) %107, i64 %105 %109 = mul i64 %89, %42 %110 = mul i64 %98, %29 %111 = getelementptr i8, ptr addrspace(1) %0, i64 %109 %112 = getelementptr i8, ptr addrspace(1) %111, i64 %110 br label %113 113: ; preds = %.lr.ph, %113 %114 = phi <1 x float> [ splat (float -qnan), %.lr.ph ], [ %140, %113 ] %115 = phi i64 [ 0, %.lr.ph ], [ %141, %113 ] %116 = shl nuw nsw i64 %115, 3 %117 = sub nsw i64 %37, %116 %118 = tail call i64 @llvm.smin.i64(i64 %117, i64 8192) %119 = insertelement <8 x i64> poison, i64 %118, i64 0 %120 = shufflevector <8 x i64> %119, <8 x i64> poison, <8 x i32> zeroinitializer %121 = icmp sgt <8 x i64> %120, %102 %122 = or disjoint i64 %116, %63 %123 = getelementptr [2 x i8], ptr addrspace(1) %108, i64 %122 %124 = tail call <8 x half> @llvm.masked.load.v8f16.p1(ptr addrspace(1) align 2 %123, <8 x i1> %121, <8 x half> poison) %125 = getelementptr i8, ptr addrspace(1) %112, i64 %122 %126 = tail call <8 x i8> @llvm.masked.load.v8i8.p1(ptr addrspace(1) align 1 %125, <8 x i1> %121, <8 x i8> poison) %.scalar29 = or disjoint i64 %116, %63 %127 = insertelement <8 x i64> poison, i64 %.scalar29, i64 0 %128 = shufflevector <8 x i64> %127, <8 x i64> poison, <8 x i32> zeroinitializer %129 = add <8 x i64> %99, %128 %130 = icmp ult <8 x i64> %129, %65 %131 = trunc <8 x i8> %126 to <8 x i1> %132 = select <8 x i1> %131, <8 x float> zeroinitializer, <8 x float> splat (float -inf) %133 = fpext <8 x half> %124 to <8 x float> %134 = fadd <8 x float> %132, %133 %135 = select <8 x i1> %130, <8 x float> %134, <8 x float> splat (float -qnan) %136 = tail call float @llvm.vector.reduce.fmax.v8f32(<8 x float> %135) %137 = tail call float @llvm.maxnum.f32(float %136, float -inf) %138 = extractelement <1 x float> %114, i64 0 %139 = tail call float @llvm.maximumnum.f32(float %137, float %138) %140 = insertelement <1 x float> poison, float %139, i64 0 %141 = add nuw nsw i64 %115, 1024 %142 = icmp samesign ult i64 %115, %59 br i1 %142, label %113, label %._crit_edge ._crit_edge: ; preds = %113, %81 %.lcssa = phi <1 x float> [ splat (float -qnan), %81 ], [ %140, %113 ] %143 = extractelement <1 x float> %.lcssa, i64 0 %144 = tail call float @llvm.maximumnum.f32(float %143, float -inf) %145 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %144, i32 1, i32 31) %146 = tail call float @llvm.maximumnum.f32(float %144, float %145) %147 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %146, i32 2, i32 31) %148 = tail call float @llvm.maximumnum.f32(float %146, float %147) %149 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %148, i32 4, i32 31) %150 = tail call float @llvm.maximumnum.f32(float %148, float %149) %151 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %150, i32 8, i32 31) %152 = tail call float @llvm.maximumnum.f32(float %150, float %151) %153 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %152, i32 16, i32 31) tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0) br i1 %69, label %154, label %157 154: ; preds = %._crit_edge %155 = tail call float @llvm.maximumnum.f32(float %152, float %153) %156 = insertelement <1 x float> poison, float %155, i64 0 store <1 x float> %156, ptr addrspace(3) %70, align 4 br label %157 157: ; preds = %154, %._crit_edge tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0) %158 = load float, ptr addrspace(3) %72, align 4 %159 = tail call float @llvm.maximumnum.f32(float %158, float -inf) %160 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %159, i32 1, i32 31) %161 = tail call float @llvm.maximumnum.f32(float %159, float %160) %162 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %161, i32 2, i32 31) %163 = tail call float @llvm.maximumnum.f32(float %161, float %162) %164 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %163, i32 4, i32 31) %165 = tail call float @llvm.maximumnum.f32(float %163, float %164) %166 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %165, i32 8, i32 31) %167 = tail call float @llvm.maximumnum.f32(float %165, float %166) %168 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %167, i32 16, i32 31) %169 = tail call float @llvm.maximumnum.f32(float %167, float %168) br i1 %56, label %._crit_edge17, label %.lr.ph16 .lr.ph16: ; preds = %157 %170 = tail call <8 x i64> @llvm.stepvector.v8i64() %171 = trunc <8 x i64> %170 to <8 x i32> %172 = add <8 x i32> %76, %171 %173 = zext <8 x i32> %172 to <8 x i64> %174 = mul i64 %89, %46 %175 = mul i64 %94, %45 %176 = mul i64 %98, %37 %177 = getelementptr [2 x i8], ptr addrspace(1) %47, i64 %174 %178 = getelementptr [2 x i8], ptr addrspace(1) %177, i64 %175 %179 = getelementptr [2 x i8], ptr addrspace(1) %178, i64 %176 %180 = mul i64 %89, %42 %181 = mul i64 %98, %29 %182 = getelementptr i8, ptr addrspace(1) %0, i64 %180 %183 = getelementptr i8, ptr addrspace(1) %182, i64 %181 %184 = insertelement <8 x float> poison, float %169, i64 0 %185 = shufflevector <8 x float> %184, <8 x float> poison, <8 x i32> zeroinitializer br label %186 186: ; preds = %.lr.ph16, %186 %187 = phi float [ 0.000000e+00, %.lr.ph16 ], [ %350, %186 ] %188 = phi i64 [ 0, %.lr.ph16 ], [ %351, %186 ] %189 = shl nuw nsw i64 %188, 3 %190 = sub nsw i64 %37, %189 %191 = tail call i64 @llvm.smin.i64(i64 %190, i64 8192) %192 = insertelement <8 x i64> poison, i64 %191, i64 0 %193 = shufflevector <8 x i64> %192, <8 x i64> poison, <8 x i32> zeroinitializer %194 = icmp sgt <8 x i64> %193, %173 %195 = or disjoint i64 %189, %78 %196 = getelementptr [2 x i8], ptr addrspace(1) %179, i64 %195 %197 = tail call <8 x half> @llvm.masked.load.v8f16.p1(ptr addrspace(1) align 2 %196, <8 x i1> %194, <8 x half> poison) %198 = getelementptr i8, ptr addrspace(1) %183, i64 %195 %199 = tail call <8 x i8> @llvm.masked.load.v8i8.p1(ptr addrspace(1) align 1 %198, <8 x i1> %194, <8 x i8> poison) %.scalar31 = or disjoint i64 %189, %63 %200 = insertelement <8 x i64> poison, i64 %.scalar31, i64 0 %201 = shufflevector <8 x i64> %200, <8 x i64> poison, <8 x i32> zeroinitializer %202 = add <8 x i64> %170, %201 %203 = icmp ult <8 x i64> %202, %65 %204 = trunc <8 x i8> %199 to <8 x i1> %205 = select <8 x i1> %204, <8 x float> zeroinitializer, <8 x float> splat (float -inf) %206 = fpext <8 x half> %197 to <8 x float> %207 = fadd <8 x float> %205, %206 %208 = fsub <8 x float> %207, %185 %.inv10 = fcmp olt <8 x float> %208, splat (float -8.780000e+01) %209 = select <8 x i1> %.inv10, <8 x float> splat (float -8.780000e+01), <8 x float> %208 %.inv11 = fcmp ogt <8 x float> %209, splat (float 8.880000e+01) %210 = select <8 x i1> %.inv11, <8 x float> splat (float 8.880000e+01), <8 x float> %209 %211 = extractelement <8 x float> %210, i64 0 %212 = tail call float @llvm.fma.f32(float %211, float f0x3FB8AA3B, float 5.000000e-01) %213 = extractelement <8 x float> %210, i64 1 %214 = tail call float @llvm.fma.f32(float %213, float f0x3FB8AA3B, float 5.000000e-01) %215 = extractelement <8 x float> %210, i64 2 %216 = tail call float @llvm.fma.f32(float %215, float f0x3FB8AA3B, float 5.000000e-01) %217 = extractelement <8 x float> %210, i64 3 %218 = tail call float @llvm.fma.f32(float %217, float f0x3FB8AA3B, float 5.000000e-01) %219 = extractelement <8 x float> %210, i64 4 %220 = tail call float @llvm.fma.f32(float %219, float f0x3FB8AA3B, float 5.000000e-01) %221 = extractelement <8 x float> %210, i64 5 %222 = tail call float @llvm.fma.f32(float %221, float f0x3FB8AA3B, float 5.000000e-01) %223 = extractelement <8 x float> %210, i64 6 %224 = tail call float @llvm.fma.f32(float %223, float f0x3FB8AA3B, float 5.000000e-01) %225 = extractelement <8 x float> %210, i64 7 %226 = tail call float @llvm.fma.f32(float %225, float f0x3FB8AA3B, float 5.000000e-01) %227 = tail call float @llvm.floor.f32(float %212) %228 = insertelement <8 x float> poison, float %227, i64 0 %229 = tail call float @llvm.floor.f32(float %214) %230 = insertelement <8 x float> %228, float %229, i64 1 %231 = tail call float @llvm.floor.f32(float %216) %232 = insertelement <8 x float> %230, float %231, i64 2 %233 = tail call float @llvm.floor.f32(float %218) %234 = insertelement <8 x float> %232, float %233, i64 3 %235 = tail call float @llvm.floor.f32(float %220) %236 = insertelement <8 x float> %234, float %235, i64 4 %237 = tail call float @llvm.floor.f32(float %222) %238 = insertelement <8 x float> %236, float %237, i64 5 %239 = tail call float @llvm.floor.f32(float %224) %240 = insertelement <8 x float> %238, float %239, i64 6 %241 = tail call float @llvm.floor.f32(float %226) %242 = insertelement <8 x float> %240, float %241, i64 7 %.inv12 = fcmp olt <8 x float> %242, splat (float -1.270000e+02) %243 = select <8 x i1> %.inv12, <8 x float> splat (float -1.270000e+02), <8 x float> %242 %.inv13 = fcmp ogt <8 x float> %243, splat (float 1.270000e+02) %244 = select <8 x i1> %.inv13, <8 x float> splat (float 1.270000e+02), <8 x float> %243 %245 = extractelement <8 x float> %244, i64 0 %246 = tail call float @llvm.fma.f32(float %245, float f0xBF318000, float %211) %247 = extractelement <8 x float> %244, i64 1 %248 = tail call float @llvm.fma.f32(float %247, float f0xBF318000, float %213) %249 = extractelement <8 x float> %244, i64 2 %250 = tail call float @llvm.fma.f32(float %249, float f0xBF318000, float %215) %251 = extractelement <8 x float> %244, i64 3 %252 = tail call float @llvm.fma.f32(float %251, float f0xBF318000, float %217) %253 = extractelement <8 x float> %244, i64 4 %254 = tail call float @llvm.fma.f32(float %253, float f0xBF318000, float %219) %255 = extractelement <8 x float> %244, i64 5 %256 = tail call float @llvm.fma.f32(float %255, float f0xBF318000, float %221) %257 = extractelement <8 x float> %244, i64 6 %258 = tail call float @llvm.fma.f32(float %257, float f0xBF318000, float %223) %259 = extractelement <8 x float> %244, i64 7 %260 = tail call float @llvm.fma.f32(float %259, float f0xBF318000, float %225) %261 = tail call float @llvm.fma.f32(float %245, float f0x395E8083, float %246) %262 = insertelement <8 x float> poison, float %261, i64 0 %263 = tail call float @llvm.fma.f32(float %247, float f0x395E8083, float %248) %264 = insertelement <8 x float> %262, float %263, i64 1 %265 = tail call float @llvm.fma.f32(float %249, float f0x395E8083, float %250) %266 = insertelement <8 x float> %264, float %265, i64 2 %267 = tail call float @llvm.fma.f32(float %251, float f0x395E8083, float %252) %268 = insertelement <8 x float> %266, float %267, i64 3 %269 = tail call float @llvm.fma.f32(float %253, float f0x395E8083, float %254) %270 = insertelement <8 x float> %268, float %269, i64 4 %271 = tail call float @llvm.fma.f32(float %255, float f0x395E8083, float %256) %272 = insertelement <8 x float> %270, float %271, i64 5 %273 = tail call float @llvm.fma.f32(float %257, float f0x395E8083, float %258) %274 = insertelement <8 x float> %272, float %273, i64 6 %275 = tail call float @llvm.fma.f32(float %259, float f0x395E8083, float %260) %276 = insertelement <8 x float> %274, float %275, i64 7 %277 = tail call float @llvm.fma.f32(float %261, float f0x39506967, float f0x3AB743CE) %278 = tail call float @llvm.fma.f32(float %263, float f0x39506967, float f0x3AB743CE) %279 = tail call float @llvm.fma.f32(float %265, float f0x39506967, float f0x3AB743CE) %280 = tail call float @llvm.fma.f32(float %267, float f0x39506967, float f0x3AB743CE) %281 = tail call float @llvm.fma.f32(float %269, float f0x39506967, float f0x3AB743CE) %282 = tail call float @llvm.fma.f32(float %271, float f0x39506967, float f0x3AB743CE) %283 = tail call float @llvm.fma.f32(float %273, float f0x39506967, float f0x3AB743CE) %284 = tail call float @llvm.fma.f32(float %275, float f0x39506967, float f0x3AB743CE) %285 = tail call float @llvm.fma.f32(float %277, float %261, float f0x3C088908) %286 = tail call float @llvm.fma.f32(float %278, float %263, float f0x3C088908) %287 = tail call float @llvm.fma.f32(float %279, float %265, float f0x3C088908) %288 = tail call float @llvm.fma.f32(float %280, float %267, float f0x3C088908) %289 = tail call float @llvm.fma.f32(float %281, float %269, float f0x3C088908) %290 = tail call float @llvm.fma.f32(float %282, float %271, float f0x3C088908) %291 = tail call float @llvm.fma.f32(float %283, float %273, float f0x3C088908) %292 = tail call float @llvm.fma.f32(float %284, float %275, float f0x3C088908) %293 = tail call float @llvm.fma.f32(float %285, float %261, float f0x3D2AA9C1) %294 = tail call float @llvm.fma.f32(float %286, float %263, float f0x3D2AA9C1) %295 = tail call float @llvm.fma.f32(float %287, float %265, float f0x3D2AA9C1) %296 = tail call float @llvm.fma.f32(float %288, float %267, float f0x3D2AA9C1) %297 = tail call float @llvm.fma.f32(float %289, float %269, float f0x3D2AA9C1) %298 = tail call float @llvm.fma.f32(float %290, float %271, float f0x3D2AA9C1) %299 = tail call float @llvm.fma.f32(float %291, float %273, float f0x3D2AA9C1) %300 = tail call float @llvm.fma.f32(float %292, float %275, float f0x3D2AA9C1) %301 = tail call float @llvm.fma.f32(float %293, float %261, float f0x3E2AAAAA) %302 = tail call float @llvm.fma.f32(float %294, float %263, float f0x3E2AAAAA) %303 = tail call float @llvm.fma.f32(float %295, float %265, float f0x3E2AAAAA) %304 = tail call float @llvm.fma.f32(float %296, float %267, float f0x3E2AAAAA) %305 = tail call float @llvm.fma.f32(float %297, float %269, float f0x3E2AAAAA) %306 = tail call float @llvm.fma.f32(float %298, float %271, float f0x3E2AAAAA) %307 = tail call float @llvm.fma.f32(float %299, float %273, float f0x3E2AAAAA) %308 = tail call float @llvm.fma.f32(float %300, float %275, float f0x3E2AAAAA) %309 = tail call float @llvm.fma.f32(float %301, float %261, float 5.000000e-01) %310 = tail call float @llvm.fma.f32(float %302, float %263, float 5.000000e-01) %311 = tail call float @llvm.fma.f32(float %303, float %265, float 5.000000e-01) %312 = tail call float @llvm.fma.f32(float %304, float %267, float 5.000000e-01) %313 = tail call float @llvm.fma.f32(float %305, float %269, float 5.000000e-01) %314 = tail call float @llvm.fma.f32(float %306, float %271, float 5.000000e-01) %315 = tail call float @llvm.fma.f32(float %307, float %273, float 5.000000e-01) %316 = tail call float @llvm.fma.f32(float %308, float %275, float 5.000000e-01) %317 = fmul <8 x float> %276, %276 %318 = extractelement <8 x float> %317, i64 0 %319 = tail call float @llvm.fma.f32(float %309, float %318, float %261) %320 = insertelement <8 x float> poison, float %319, i64 0 %321 = extractelement <8 x float> %317, i64 1 %322 = tail call float @llvm.fma.f32(float %310, float %321, float %263) %323 = insertelement <8 x float> %320, float %322, i64 1 %324 = extractelement <8 x float> %317, i64 2 %325 = tail call float @llvm.fma.f32(float %311, float %324, float %265) %326 = insertelement <8 x float> %323, float %325, i64 2 %327 = extractelement <8 x float> %317, i64 3 %328 = tail call float @llvm.fma.f32(float %312, float %327, float %267) %329 = insertelement <8 x float> %326, float %328, i64 3 %330 = extractelement <8 x float> %317, i64 4 %331 = tail call float @llvm.fma.f32(float %313, float %330, float %269) %332 = insertelement <8 x float> %329, float %331, i64 4 %333 = extractelement <8 x float> %317, i64 5 %334 = tail call float @llvm.fma.f32(float %314, float %333, float %271) %335 = insertelement <8 x float> %332, float %334, i64 5 %336 = extractelement <8 x float> %317, i64 6 %337 = tail call float @llvm.fma.f32(float %315, float %336, float %273) %338 = insertelement <8 x float> %335, float %337, i64 6 %339 = extractelement <8 x float> %317, i64 7 %340 = tail call float @llvm.fma.f32(float %316, float %339, float %275) %341 = insertelement <8 x float> %338, float %340, i64 7 %342 = fadd <8 x float> %341, splat (float 1.000000e+00) %343 = fptosi <8 x float> %244 to <8 x i32> %344 = shl <8 x i32> %343, splat (i32 23) %345 = add <8 x i32> %344, splat (i32 1065353216) %346 = bitcast <8 x i32> %345 to <8 x float> %347 = fmul <8 x float> %342, %346 %348 = select <8 x i1> %203, <8 x float> %347, <8 x float> zeroinitializer %349 = tail call float @llvm.vector.reduce.fadd.v8f32(float 0.000000e+00, <8 x float> %348) %350 = fadd float %187, %349 %351 = add nuw nsw i64 %188, 1024 %352 = icmp samesign ult i64 %188, %59 br i1 %352, label %186, label %._crit_edge17.loopexit ._crit_edge17.loopexit: ; preds = %186 %353 = fadd float %350, 0.000000e+00 br label %._crit_edge17 ._crit_edge17: ; preds = %._crit_edge17.loopexit, %157 %.lcssa14 = phi float [ 0.000000e+00, %157 ], [ %353, %._crit_edge17.loopexit ] %354 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %.lcssa14, i32 1, i32 31) %355 = fadd float %.lcssa14, %354 %356 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %355, i32 2, i32 31) %357 = fadd float %355, %356 %358 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %357, i32 4, i32 31) %359 = fadd float %357, %358 %360 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %359, i32 8, i32 31) %361 = fadd float %359, %360 %362 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %361, i32 16, i32 31) tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0) br i1 %69, label %363, label %366 363: ; preds = %._crit_edge17 %364 = fadd float %361, %362 %365 = insertelement <1 x float> poison, float %364, i64 0 store <1 x float> %365, ptr addrspace(3) %79, align 4 br label %366 366: ; preds = %363, %._crit_edge17 tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0) %367 = load float, ptr addrspace(3) %80, align 4 %368 = fadd float %367, 0.000000e+00 %369 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %368, i32 1, i32 31) %370 = fadd float %369, %368 %371 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %370, i32 2, i32 31) %372 = fadd float %371, %370 %373 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %372, i32 4, i32 31) %374 = fadd float %373, %372 %375 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %374, i32 8, i32 31) %376 = fadd float %375, %374 %377 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %376, i32 16, i32 31) br i1 %56, label %._crit_edge21, label %.lr.ph20 .lr.ph20: ; preds = %366 %378 = fadd float %377, %376 %379 = tail call <8 x i64> @llvm.stepvector.v8i64() %380 = trunc <8 x i64> %379 to <8 x i32> %381 = add <8 x i32> %76, %380 %382 = zext <8 x i32> %381 to <8 x i64> %383 = mul i64 %89, %46 %384 = mul i64 %94, %45 %385 = add i64 %383, %384 %386 = mul i64 %98, %37 %387 = add i64 %385, %386 %388 = mul i64 %89, %42 %389 = mul i64 %98, %29 %390 = getelementptr i8, ptr addrspace(1) %0, i64 %388 %391 = getelementptr i8, ptr addrspace(1) %390, i64 %389 %392 = insertelement <8 x float> poison, float %169, i64 0 %393 = shufflevector <8 x float> %392, <8 x float> poison, <8 x i32> zeroinitializer %394 = insertelement <8 x float> poison, float %378, i64 0 %395 = shufflevector <8 x float> %394, <8 x float> poison, <8 x i32> zeroinitializer br label %396 396: ; preds = %.lr.ph20, %396 %397 = phi i64 [ 0, %.lr.ph20 ], [ %556, %396 ] %398 = sub nuw nsw i64 %37, %397 %399 = tail call i64 @llvm.smin.i64(i64 %398, i64 8192) %400 = insertelement <8 x i64> poison, i64 %399, i64 0 %401 = shufflevector <8 x i64> %400, <8 x i64> poison, <8 x i32> zeroinitializer %402 = icmp samesign ugt <8 x i64> %401, %382 %403 = or disjoint i64 %397, %78 %404 = add i64 %387, %403 %405 = getelementptr [2 x i8], ptr addrspace(1) %47, i64 %404 %406 = tail call <8 x half> @llvm.masked.load.v8f16.p1(ptr addrspace(1) align 2 %405, <8 x i1> %402, <8 x half> poison) %407 = getelementptr i8, ptr addrspace(1) %391, i64 %403 %408 = tail call <8 x i8> @llvm.masked.load.v8i8.p1(ptr addrspace(1) align 1 %407, <8 x i1> %402, <8 x i8> poison) %409 = trunc <8 x i8> %408 to <8 x i1> %410 = select <8 x i1> %409, <8 x float> zeroinitializer, <8 x float> splat (float -inf) %411 = fpext <8 x half> %406 to <8 x float> %412 = fadd <8 x float> %410, %411 %413 = fsub <8 x float> %412, %393 %.inv = fcmp olt <8 x float> %413, splat (float -8.780000e+01) %414 = select <8 x i1> %.inv, <8 x float> splat (float -8.780000e+01), <8 x float> %413 %.inv6 = fcmp ogt <8 x float> %414, splat (float 8.880000e+01) %415 = select <8 x i1> %.inv6, <8 x float> splat (float 8.880000e+01), <8 x float> %414 %416 = extractelement <8 x float> %415, i64 0 %417 = tail call float @llvm.fma.f32(float %416, float f0x3FB8AA3B, float 5.000000e-01) %418 = extractelement <8 x float> %415, i64 1 %419 = tail call float @llvm.fma.f32(float %418, float f0x3FB8AA3B, float 5.000000e-01) %420 = extractelement <8 x float> %415, i64 2 %421 = tail call float @llvm.fma.f32(float %420, float f0x3FB8AA3B, float 5.000000e-01) %422 = extractelement <8 x float> %415, i64 3 %423 = tail call float @llvm.fma.f32(float %422, float f0x3FB8AA3B, float 5.000000e-01) %424 = extractelement <8 x float> %415, i64 4 %425 = tail call float @llvm.fma.f32(float %424, float f0x3FB8AA3B, float 5.000000e-01) %426 = extractelement <8 x float> %415, i64 5 %427 = tail call float @llvm.fma.f32(float %426, float f0x3FB8AA3B, float 5.000000e-01) %428 = extractelement <8 x float> %415, i64 6 %429 = tail call float @llvm.fma.f32(float %428, float f0x3FB8AA3B, float 5.000000e-01) %430 = extractelement <8 x float> %415, i64 7 %431 = tail call float @llvm.fma.f32(float %430, float f0x3FB8AA3B, float 5.000000e-01) %432 = tail call float @llvm.floor.f32(float %417) %433 = insertelement <8 x float> poison, float %432, i64 0 %434 = tail call float @llvm.floor.f32(float %419) %435 = insertelement <8 x float> %433, float %434, i64 1 %436 = tail call float @llvm.floor.f32(float %421) %437 = insertelement <8 x float> %435, float %436, i64 2 %438 = tail call float @llvm.floor.f32(float %423) %439 = insertelement <8 x float> %437, float %438, i64 3 %440 = tail call float @llvm.floor.f32(float %425) %441 = insertelement <8 x float> %439, float %440, i64 4 %442 = tail call float @llvm.floor.f32(float %427) %443 = insertelement <8 x float> %441, float %442, i64 5 %444 = tail call float @llvm.floor.f32(float %429) %445 = insertelement <8 x float> %443, float %444, i64 6 %446 = tail call float @llvm.floor.f32(float %431) %447 = insertelement <8 x float> %445, float %446, i64 7 %.inv7 = fcmp olt <8 x float> %447, splat (float -1.270000e+02) %448 = select <8 x i1> %.inv7, <8 x float> splat (float -1.270000e+02), <8 x float> %447 %.inv8 = fcmp ogt <8 x float> %448, splat (float 1.270000e+02) %449 = select <8 x i1> %.inv8, <8 x float> splat (float 1.270000e+02), <8 x float> %448 %450 = extractelement <8 x float> %449, i64 0 %451 = tail call float @llvm.fma.f32(float %450, float f0xBF318000, float %416) %452 = extractelement <8 x float> %449, i64 1 %453 = tail call float @llvm.fma.f32(float %452, float f0xBF318000, float %418) %454 = extractelement <8 x float> %449, i64 2 %455 = tail call float @llvm.fma.f32(float %454, float f0xBF318000, float %420) %456 = extractelement <8 x float> %449, i64 3 %457 = tail call float @llvm.fma.f32(float %456, float f0xBF318000, float %422) %458 = extractelement <8 x float> %449, i64 4 %459 = tail call float @llvm.fma.f32(float %458, float f0xBF318000, float %424) %460 = extractelement <8 x float> %449, i64 5 %461 = tail call float @llvm.fma.f32(float %460, float f0xBF318000, float %426) %462 = extractelement <8 x float> %449, i64 6 %463 = tail call float @llvm.fma.f32(float %462, float f0xBF318000, float %428) %464 = extractelement <8 x float> %449, i64 7 %465 = tail call float @llvm.fma.f32(float %464, float f0xBF318000, float %430) %466 = tail call float @llvm.fma.f32(float %450, float f0x395E8083, float %451) %467 = insertelement <8 x float> poison, float %466, i64 0 %468 = tail call float @llvm.fma.f32(float %452, float f0x395E8083, float %453) %469 = insertelement <8 x float> %467, float %468, i64 1 %470 = tail call float @llvm.fma.f32(float %454, float f0x395E8083, float %455) %471 = insertelement <8 x float> %469, float %470, i64 2 %472 = tail call float @llvm.fma.f32(float %456, float f0x395E8083, float %457) %473 = insertelement <8 x float> %471, float %472, i64 3 %474 = tail call float @llvm.fma.f32(float %458, float f0x395E8083, float %459) %475 = insertelement <8 x float> %473, float %474, i64 4 %476 = tail call float @llvm.fma.f32(float %460, float f0x395E8083, float %461) %477 = insertelement <8 x float> %475, float %476, i64 5 %478 = tail call float @llvm.fma.f32(float %462, float f0x395E8083, float %463) %479 = insertelement <8 x float> %477, float %478, i64 6 %480 = tail call float @llvm.fma.f32(float %464, float f0x395E8083, float %465) %481 = insertelement <8 x float> %479, float %480, i64 7 %482 = tail call float @llvm.fma.f32(float %466, float f0x39506967, float f0x3AB743CE) %483 = tail call float @llvm.fma.f32(float %468, float f0x39506967, float f0x3AB743CE) %484 = tail call float @llvm.fma.f32(float %470, float f0x39506967, float f0x3AB743CE) %485 = tail call float @llvm.fma.f32(float %472, float f0x39506967, float f0x3AB743CE) %486 = tail call float @llvm.fma.f32(float %474, float f0x39506967, float f0x3AB743CE) %487 = tail call float @llvm.fma.f32(float %476, float f0x39506967, float f0x3AB743CE) %488 = tail call float @llvm.fma.f32(float %478, float f0x39506967, float f0x3AB743CE) %489 = tail call float @llvm.fma.f32(float %480, float f0x39506967, float f0x3AB743CE) %490 = tail call float @llvm.fma.f32(float %482, float %466, float f0x3C088908) %491 = tail call float @llvm.fma.f32(float %483, float %468, float f0x3C088908) %492 = tail call float @llvm.fma.f32(float %484, float %470, float f0x3C088908) %493 = tail call float @llvm.fma.f32(float %485, float %472, float f0x3C088908) %494 = tail call float @llvm.fma.f32(float %486, float %474, float f0x3C088908) %495 = tail call float @llvm.fma.f32(float %487, float %476, float f0x3C088908) %496 = tail call float @llvm.fma.f32(float %488, float %478, float f0x3C088908) %497 = tail call float @llvm.fma.f32(float %489, float %480, float f0x3C088908) %498 = tail call float @llvm.fma.f32(float %490, float %466, float f0x3D2AA9C1) %499 = tail call float @llvm.fma.f32(float %491, float %468, float f0x3D2AA9C1) %500 = tail call float @llvm.fma.f32(float %492, float %470, float f0x3D2AA9C1) %501 = tail call float @llvm.fma.f32(float %493, float %472, float f0x3D2AA9C1) %502 = tail call float @llvm.fma.f32(float %494, float %474, float f0x3D2AA9C1) %503 = tail call float @llvm.fma.f32(float %495, float %476, float f0x3D2AA9C1) %504 = tail call float @llvm.fma.f32(float %496, float %478, float f0x3D2AA9C1) %505 = tail call float @llvm.fma.f32(float %497, float %480, float f0x3D2AA9C1) %506 = tail call float @llvm.fma.f32(float %498, float %466, float f0x3E2AAAAA) %507 = tail call float @llvm.fma.f32(float %499, float %468, float f0x3E2AAAAA) %508 = tail call float @llvm.fma.f32(float %500, float %470, float f0x3E2AAAAA) %509 = tail call float @llvm.fma.f32(float %501, float %472, float f0x3E2AAAAA) %510 = tail call float @llvm.fma.f32(float %502, float %474, float f0x3E2AAAAA) %511 = tail call float @llvm.fma.f32(float %503, float %476, float f0x3E2AAAAA) %512 = tail call float @llvm.fma.f32(float %504, float %478, float f0x3E2AAAAA) %513 = tail call float @llvm.fma.f32(float %505, float %480, float f0x3E2AAAAA) %514 = tail call float @llvm.fma.f32(float %506, float %466, float 5.000000e-01) %515 = tail call float @llvm.fma.f32(float %507, float %468, float 5.000000e-01) %516 = tail call float @llvm.fma.f32(float %508, float %470, float 5.000000e-01) %517 = tail call float @llvm.fma.f32(float %509, float %472, float 5.000000e-01) %518 = tail call float @llvm.fma.f32(float %510, float %474, float 5.000000e-01) %519 = tail call float @llvm.fma.f32(float %511, float %476, float 5.000000e-01) %520 = tail call float @llvm.fma.f32(float %512, float %478, float 5.000000e-01) %521 = tail call float @llvm.fma.f32(float %513, float %480, float 5.000000e-01) %522 = fmul <8 x float> %481, %481 %523 = extractelement <8 x float> %522, i64 0 %524 = tail call float @llvm.fma.f32(float %514, float %523, float %466) %525 = insertelement <8 x float> poison, float %524, i64 0 %526 = extractelement <8 x float> %522, i64 1 %527 = tail call float @llvm.fma.f32(float %515, float %526, float %468) %528 = insertelement <8 x float> %525, float %527, i64 1 %529 = extractelement <8 x float> %522, i64 2 %530 = tail call float @llvm.fma.f32(float %516, float %529, float %470) %531 = insertelement <8 x float> %528, float %530, i64 2 %532 = extractelement <8 x float> %522, i64 3 %533 = tail call float @llvm.fma.f32(float %517, float %532, float %472) %534 = insertelement <8 x float> %531, float %533, i64 3 %535 = extractelement <8 x float> %522, i64 4 %536 = tail call float @llvm.fma.f32(float %518, float %535, float %474) %537 = insertelement <8 x float> %534, float %536, i64 4 %538 = extractelement <8 x float> %522, i64 5 %539 = tail call float @llvm.fma.f32(float %519, float %538, float %476) %540 = insertelement <8 x float> %537, float %539, i64 5 %541 = extractelement <8 x float> %522, i64 6 %542 = tail call float @llvm.fma.f32(float %520, float %541, float %478) %543 = insertelement <8 x float> %540, float %542, i64 6 %544 = extractelement <8 x float> %522, i64 7 %545 = tail call float @llvm.fma.f32(float %521, float %544, float %480) %546 = insertelement <8 x float> %543, float %545, i64 7 %547 = fadd <8 x float> %546, splat (float 1.000000e+00) %548 = fptosi <8 x float> %449 to <8 x i32> %549 = shl <8 x i32> %548, splat (i32 23) %550 = add <8 x i32> %549, splat (i32 1065353216) %551 = bitcast <8 x i32> %550 to <8 x float> %552 = fmul <8 x float> %547, %551 %553 = fdiv <8 x float> %552, %395 %.inv9 = fcmp ord <8 x float> %553, zeroinitializer %554 = select <8 x i1> %.inv9, <8 x float> %553, <8 x float> zeroinitializer %555 = getelementptr [4 x i8], ptr addrspace(1) %2, i64 %404 tail call void @llvm.masked.store.v8f32.p1(<8 x float> %554, ptr addrspace(1) align 4 %555, <8 x i1> %402) %556 = add nuw nsw i64 %397, 8192 %557 = icmp samesign ult i64 %556, %37 br i1 %557, label %396, label %._crit_edge21 ._crit_edge21: ; preds = %396, %366 %558 = add i64 %82, %53 %559 = icmp slt i64 %558, %49 br i1 %559, label %81, label %._crit_edge24 ._crit_edge24: ; preds = %._crit_edge21, %17 ret void }