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

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

50:                                               ; preds = %.lr.ph, %114
  %51 = phi i64 [ %24, %.lr.ph ], [ %115, %114 ]
  %.frozen = freeze i64 %51
  %52 = udiv i64 %.frozen, 18944
  %53 = mul i64 %52, 18944
  %.decomposed = sub i64 %.frozen, %53
  %.idx = mul nuw nsw i64 %52, 7168
  %gep = getelementptr i8, ptr addrspace(1) %invariant.gep, i64 %.idx
  %54 = load <4 x half>, ptr addrspace(1) %gep, align 2
  %55 = getelementptr [2 x i8], ptr addrspace(1) %40, i64 %.decomposed
  %56 = load <1 x half>, ptr addrspace(1) %55, align 2
  %gep4 = getelementptr [2 x i8], ptr addrspace(1) %invariant.gep3, i64 %.decomposed
  %57 = getelementptr i8, ptr addrspace(1) %gep4, i64 37888
  %58 = load <1 x half>, ptr addrspace(1) %57, align 2
  %59 = getelementptr i8, ptr addrspace(1) %gep4, i64 75776
  %60 = load <1 x half>, ptr addrspace(1) %59, align 2
  %61 = shufflevector <1 x half> %60, <1 x half> poison, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
  %62 = getelementptr i8, ptr addrspace(1) %gep4, i64 113664
  %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 = shufflevector <1 x half> %56, <1 x half> %58, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>
  %66 = shufflevector <4 x half> %65, <4 x half> %61, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>
  %67 = shufflevector <4 x half> %66, <4 x half> %64, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
  %68 = fpext <4 x half> %54 to <4 x float>
  %69 = fpext <4 x half> %67 to <4 x float>
  %70 = extractelement <4 x float> %68, i64 0
  %71 = extractelement <4 x float> %68, i64 1
  %72 = extractelement <4 x float> %68, i64 2
  %73 = extractelement <4 x float> %68, i64 3
  %74 = extractelement <4 x float> %69, i64 0
  %75 = extractelement <4 x float> %69, i64 1
  %76 = extractelement <4 x float> %69, i64 2
  %77 = extractelement <4 x float> %69, i64 3
  %78 = tail call float @llvm.fma.f32(float %73, float %77, float 0.000000e+00)
  %79 = tail call float @llvm.fma.f32(float %72, float %76, float %78)
  %80 = tail call float @llvm.fma.f32(float %71, float %75, float %79)
  %81 = tail call float @llvm.fma.f32(float %70, float %74, float %80)
  %82 = fadd float %81, 0.000000e+00
  %83 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %82, i32 1, i32 31)
  %84 = fadd float %83, %82
  %85 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %84, i32 2, i32 31)
  %86 = fadd float %85, %84
  %87 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %86, i32 4, i32 31)
  %88 = fadd float %87, %86
  %89 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %88, i32 8, i32 31)
  %90 = fadd float %89, %88
  %91 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %90, i32 16, i32 31)
  tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)
  br i1 %43, label %92, label %95

92:                                               ; preds = %50
  %93 = fadd float %91, %90
  %94 = insertelement <1 x float> poison, float %93, i64 0
  store <1 x float> %94, ptr addrspace(3) %46, align 4
  br label %95

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

108:                                              ; preds = %95
  %.idx2 = mul i64 %52, 37888
  %109 = getelementptr i8, ptr addrspace(1) %21, i64 %.idx2
  %110 = getelementptr [2 x i8], ptr addrspace(1) %109, i64 %.decomposed
  %111 = fadd float %107, %106
  %112 = fptrunc float %111 to half
  %113 = insertelement <1 x half> poison, half %112, i64 0
  store <1 x half> %113, ptr addrspace(1) %110, align 2
  br label %114

114:                                              ; preds = %108, %95
  %115 = add i64 %51, %26
  %116 = icmp slt i64 %115, %22
  br i1 %116, label %50, label %._crit_edge

._crit_edge:                                      ; preds = %114, %9
  ret void
}