define ptx_kernel void @main_graph_async_dispatch_15_matmul_Dx3584x3584_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 noundef nonnull readonly align 16 captures(none) %2, ptr addrspace(1) noalias nofree noundef nonnull writeonly align 16 captures(none) %3, i32 noundef %4, i32 noundef %5, i32 noundef %6, i32 noundef %7, i32 noundef %8, i32 noundef %9) local_unnamed_addr #6 {
  %11 = zext i32 %4 to i64
  %12 = zext i32 %6 to i64
  %13 = zext i32 %7 to i64
  %14 = zext i32 %8 to i64
  %15 = zext i32 %9 to i64
  %16 = shl nuw i64 %15, 32
  %17 = or disjoint i64 %16, %14
  %18 = and i64 %11, 255
  %19 = icmp eq i64 %18, 0
  tail call void @llvm.assume(i1 %19)
  %20 = and i64 %12, 255
  %21 = icmp eq i64 %20, 0
  tail call void @llvm.assume(i1 %21)
  %22 = icmp sgt i64 %16, -1
  tail call void @llvm.assume(i1 %22)
  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) ]
  %23 = shl nuw i64 %13, 32
  %24 = and i64 %23, 2305843004918726656
  %25 = getelementptr i8, ptr addrspace(1) %3, i64 %24
  %26 = getelementptr i8, ptr addrspace(1) %25, i64 %12
  call void @llvm.assume(i1 true) [ "align"(ptr addrspace(1) %26, i64 64) ]
  %27 = mul nuw nsw i64 %17, 3584
  %28 = tail call range(i32 0, 2147483647) i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
  %29 = zext nneg i32 %28 to i64
  %30 = tail call range(i32 1, -2147483648) i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()
  %31 = zext nneg i32 %30 to i64
  %32 = icmp samesign ugt i64 %27, %29
  br i1 %32, label %.lr.ph, label %._crit_edge

.lr.ph:                                           ; preds = %10
  %33 = zext i32 %5 to i64
  %34 = shl nuw i64 %33, 32
  %35 = and i64 %34, 2305843004918726656
  %36 = getelementptr i8, ptr addrspace(1) %0, i64 %35
  %37 = getelementptr i8, ptr addrspace(1) %36, i64 %11
  %38 = tail call range(i32 0, 896) i32 @llvm.nvvm.read.ptx.sreg.tid.x()
  %39 = and i32 %38, 31
  %40 = zext nneg i32 %39 to i64
  %41 = shl nuw nsw i32 %38, 2
  %42 = zext nneg i32 %41 to i64
  %invariant.gep = getelementptr [2 x i8], ptr addrspace(1) %37, i64 %42
  %.idx = mul nuw nsw i64 %42, 7168
  %43 = getelementptr i8, ptr addrspace(1) %1, i64 %.idx
  %44 = mul nuw nsw i32 %38, 14336
  %45 = zext nneg i32 %44 to i64
  %invariant.gep2 = getelementptr [2 x i8], ptr addrspace(1) %1, i64 %45
  %46 = icmp eq i32 %39, 0
  %47 = lshr i32 %38, 5
  %48 = zext nneg i32 %47 to i64
  %49 = getelementptr [4 x i8], ptr addrspace(3) @__dynamic_shared_memory__, i64 %48
  %.scalar = icmp samesign ult i32 %39, 28
  %50 = insertelement <1 x i1> poison, i1 %.scalar, i64 0
  %51 = getelementptr [4 x i8], ptr addrspace(3) @__dynamic_shared_memory__, i64 %40
  %52 = icmp eq i32 %38, 0
  br label %53

53:                                               ; preds = %.lr.ph, %119
  %54 = phi i64 [ %29, %.lr.ph ], [ %120, %119 ]
  %55 = urem i64 %54, 3584
  %56 = sub nuw i64 %54, %55
  %gep = getelementptr [2 x i8], ptr addrspace(1) %invariant.gep, i64 %56
  %57 = load <4 x half>, ptr addrspace(1) %gep, align 2
  %58 = getelementptr [2 x i8], ptr addrspace(1) %43, i64 %55
  %59 = load <1 x half>, ptr addrspace(1) %58, align 2
  %gep3 = getelementptr [2 x i8], ptr addrspace(1) %invariant.gep2, i64 %55
  %60 = getelementptr i8, ptr addrspace(1) %gep3, i64 7168
  %61 = load <1 x half>, ptr addrspace(1) %60, align 2
  %62 = getelementptr i8, ptr addrspace(1) %gep3, i64 14336
  %63 = load <1 x half>, ptr addrspace(1) %62, align 2
  %64 = shufflevector <1 x half> %63, <1 x half> poison, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
  %65 = getelementptr i8, ptr addrspace(1) %gep3, i64 21504
  %66 = load <1 x half>, ptr addrspace(1) %65, align 2
  %67 = shufflevector <1 x half> %66, <1 x half> poison, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
  %68 = shufflevector <1 x half> %59, <1 x half> %61, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>
  %69 = shufflevector <4 x half> %68, <4 x half> %64, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>
  %70 = shufflevector <4 x half> %69, <4 x half> %67, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
  %71 = fpext <4 x half> %57 to <4 x float>
  %72 = fpext <4 x half> %70 to <4 x float>
  %73 = extractelement <4 x float> %71, i64 0
  %74 = extractelement <4 x float> %71, i64 1
  %75 = extractelement <4 x float> %71, i64 2
  %76 = extractelement <4 x float> %71, i64 3
  %77 = extractelement <4 x float> %72, i64 0
  %78 = extractelement <4 x float> %72, i64 1
  %79 = extractelement <4 x float> %72, i64 2
  %80 = extractelement <4 x float> %72, i64 3
  %81 = tail call float @llvm.fma.f32(float %76, float %80, float 0.000000e+00)
  %82 = tail call float @llvm.fma.f32(float %75, float %79, float %81)
  %83 = tail call float @llvm.fma.f32(float %74, float %78, float %82)
  %84 = tail call float @llvm.fma.f32(float %73, float %77, float %83)
  %85 = fadd float %84, 0.000000e+00
  %86 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %85, i32 1, i32 31)
  %87 = fadd float %86, %85
  %88 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %87, i32 2, i32 31)
  %89 = fadd float %88, %87
  %90 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %89, i32 4, i32 31)
  %91 = fadd float %90, %89
  %92 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %91, i32 8, i32 31)
  %93 = fadd float %92, %91
  %94 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %93, i32 16, i32 31)
  tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)
  br i1 %46, label %95, label %98

95:                                               ; preds = %53
  %96 = fadd float %94, %93
  %97 = insertelement <1 x float> poison, float %96, i64 0
  store <1 x float> %97, ptr addrspace(3) %49, align 4
  br label %98

98:                                               ; preds = %95, %53
  tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)
  %99 = tail call <1 x float> @llvm.masked.load.v1f32.p3(ptr addrspace(3) align 4 %51, <1 x i1> %50, <1 x float> zeroinitializer)
  %100 = extractelement <1 x float> %99, i64 0
  %101 = fadd float %100, 0.000000e+00
  %102 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %101, i32 1, i32 31)
  %103 = fadd float %102, %101
  %104 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %103, i32 2, i32 31)
  %105 = fadd float %104, %103
  %106 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %105, i32 4, i32 31)
  %107 = fadd float %106, %105
  %108 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %107, i32 8, i32 31)
  %109 = fadd float %108, %107
  %110 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %109, i32 16, i32 31)
  br i1 %52, label %111, label %119

111:                                              ; preds = %98
  %112 = getelementptr [2 x i8], ptr addrspace(1) %26, i64 %54
  %113 = fadd float %110, %109
  %114 = fptrunc float %113 to half
  %115 = insertelement <1 x half> poison, half %114, i64 0
  %116 = getelementptr [2 x i8], ptr addrspace(1) %2, i64 %55
  %117 = load <1 x half>, ptr addrspace(1) %116, align 2
  %118 = fadd <1 x half> %115, %117
  store <1 x half> %118, ptr addrspace(1) %112, align 2
  br label %119

119:                                              ; preds = %111, %98
  %120 = add i64 %54, %31
  %121 = icmp slt i64 %120, %27
  br i1 %121, label %53, label %._crit_edge

._crit_edge:                                      ; preds = %119, %10
  ret void
}