define ptx_kernel void @main_graph_async_dispatch_743_matmul_Dx152064x3584_f16xf16xf32(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) local_unnamed_addr #6 {
  %8 = zext i32 %5 to i64
  %9 = zext i32 %6 to i64
  %10 = shl nuw i64 %9, 32
  %11 = or disjoint i64 %10, %8
  %12 = icmp sgt i64 %10, -1
  tail call void @llvm.assume(i1 %12)
  call void @llvm.assume(i1 true) [ "align"(ptr addrspace(1) %1, i64 64) ]
  call void @llvm.assume(i1 true) [ "align"(ptr addrspace(1) %2, i64 64) ]
  %13 = mul nuw nsw i64 %11, 152064
  %14 = tail call range(i32 0, 2147483647) i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
  %15 = zext nneg i32 %14 to i64
  %16 = tail call range(i32 1, -2147483648) i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()
  %17 = zext nneg i32 %16 to i64
  %18 = icmp samesign ugt i64 %13, %15
  br i1 %18, label %.lr.ph, label %._crit_edge

.lr.ph:                                           ; preds = %7
  %19 = zext i32 %4 to i64
  %20 = shl nuw i64 %19, 32
  %21 = zext i32 %3 to i64
  %22 = or disjoint i64 %20, %21
  %23 = lshr i64 %22, 1
  %24 = and i64 %23, 1152921504606846975
  %25 = getelementptr [2 x i8], ptr addrspace(1) %0, i64 %24
  %26 = tail call range(i32 0, 896) i32 @llvm.nvvm.read.ptx.sreg.tid.x()
  %27 = and i32 %26, 31
  %28 = zext nneg i32 %27 to i64
  %29 = shl nuw nsw i32 %26, 2
  %30 = zext nneg i32 %29 to i64
  %invariant.gep = getelementptr [2 x i8], ptr addrspace(1) %25, i64 %30
  %.idx1 = mul nuw nsw i64 %30, 304128
  %31 = getelementptr i8, ptr addrspace(1) %1, i64 %.idx1
  %32 = mul nuw nsw i32 %26, 608256
  %33 = zext nneg i32 %32 to i64
  %invariant.gep3 = getelementptr [2 x i8], ptr addrspace(1) %1, i64 %33
  %34 = icmp eq i32 %27, 0
  %35 = lshr i32 %26, 5
  %36 = zext nneg i32 %35 to i64
  %37 = getelementptr [4 x i8], ptr addrspace(3) @__dynamic_shared_memory__, i64 %36
  %.scalar = icmp samesign ult i32 %27, 28
  %38 = insertelement <1 x i1> poison, i1 %.scalar, i64 0
  %39 = getelementptr [4 x i8], ptr addrspace(3) @__dynamic_shared_memory__, i64 %28
  %40 = icmp eq i32 %26, 0
  br label %41

41:                                               ; preds = %.lr.ph, %105
  %42 = phi i64 [ %15, %.lr.ph ], [ %106, %105 ]
  %.frozen = freeze i64 %42
  %43 = udiv i64 %.frozen, 152064
  %44 = mul i64 %43, 152064
  %.decomposed = sub i64 %.frozen, %44
  %.idx = mul nuw nsw i64 %43, 7168
  %gep = getelementptr i8, ptr addrspace(1) %invariant.gep, i64 %.idx
  %45 = load <4 x half>, ptr addrspace(1) %gep, align 2
  %46 = getelementptr [2 x i8], ptr addrspace(1) %31, i64 %.decomposed
  %47 = load <1 x half>, ptr addrspace(1) %46, align 2
  %gep4 = getelementptr [2 x i8], ptr addrspace(1) %invariant.gep3, i64 %.decomposed
  %48 = getelementptr i8, ptr addrspace(1) %gep4, i64 304128
  %49 = load <1 x half>, ptr addrspace(1) %48, align 2
  %50 = getelementptr i8, ptr addrspace(1) %gep4, i64 608256
  %51 = load <1 x half>, ptr addrspace(1) %50, align 2
  %52 = shufflevector <1 x half> %51, <1 x half> poison, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
  %53 = getelementptr i8, ptr addrspace(1) %gep4, i64 912384
  %54 = load <1 x half>, ptr addrspace(1) %53, align 2
  %55 = shufflevector <1 x half> %54, <1 x half> poison, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
  %56 = shufflevector <1 x half> %47, <1 x half> %49, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>
  %57 = shufflevector <4 x half> %56, <4 x half> %52, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>
  %58 = shufflevector <4 x half> %57, <4 x half> %55, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
  %59 = fpext <4 x half> %45 to <4 x float>
  %60 = fpext <4 x half> %58 to <4 x float>
  %61 = extractelement <4 x float> %59, i64 0
  %62 = extractelement <4 x float> %59, i64 1
  %63 = extractelement <4 x float> %59, i64 2
  %64 = extractelement <4 x float> %59, i64 3
  %65 = extractelement <4 x float> %60, i64 0
  %66 = extractelement <4 x float> %60, i64 1
  %67 = extractelement <4 x float> %60, i64 2
  %68 = extractelement <4 x float> %60, i64 3
  %69 = tail call float @llvm.fma.f32(float %64, float %68, float 0.000000e+00)
  %70 = tail call float @llvm.fma.f32(float %63, float %67, float %69)
  %71 = tail call float @llvm.fma.f32(float %62, float %66, float %70)
  %72 = tail call float @llvm.fma.f32(float %61, float %65, float %71)
  %73 = fadd float %72, 0.000000e+00
  %74 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %73, i32 1, i32 31)
  %75 = fadd float %74, %73
  %76 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %75, i32 2, i32 31)
  %77 = fadd float %76, %75
  %78 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %77, i32 4, i32 31)
  %79 = fadd float %78, %77
  %80 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %79, i32 8, i32 31)
  %81 = fadd float %80, %79
  %82 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %81, i32 16, i32 31)
  tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)
  br i1 %34, label %83, label %86

83:                                               ; preds = %41
  %84 = fadd float %82, %81
  %85 = insertelement <1 x float> poison, float %84, i64 0
  store <1 x float> %85, ptr addrspace(3) %37, align 4
  br label %86

86:                                               ; preds = %83, %41
  tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)
  %87 = tail call <1 x float> @llvm.masked.load.v1f32.p3(ptr addrspace(3) align 4 %39, <1 x i1> %38, <1 x float> zeroinitializer)
  %88 = extractelement <1 x float> %87, i64 0
  %89 = fadd float %88, 0.000000e+00
  %90 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %89, i32 1, i32 31)
  %91 = fadd float %90, %89
  %92 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %91, i32 2, i32 31)
  %93 = fadd float %92, %91
  %94 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %93, i32 4, i32 31)
  %95 = fadd float %94, %93
  %96 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %95, i32 8, i32 31)
  %97 = fadd float %96, %95
  %98 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %97, i32 16, i32 31)
  br i1 %40, label %99, label %105

99:                                               ; preds = %86
  %.idx2 = mul i64 %43, 304128
  %100 = getelementptr i8, ptr addrspace(1) %2, i64 %.idx2
  %101 = getelementptr [2 x i8], ptr addrspace(1) %100, i64 %.decomposed
  %102 = fadd float %98, %97
  %103 = fptrunc float %102 to half
  %104 = insertelement <1 x half> poison, half %103, i64 0
  store <1 x half> %104, ptr addrspace(1) %101, align 2
  br label %105

105:                                              ; preds = %99, %86
  %106 = add i64 %42, %17
  %107 = icmp slt i64 %106, %13
  br i1 %107, label %41, label %._crit_edge

._crit_edge:                                      ; preds = %105, %7
  ret void
}