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

.lr.ph:                                           ; preds = %6
  %18 = tail call range(i32 0, 896) i32 @llvm.nvvm.read.ptx.sreg.tid.x()
  %19 = and i32 %18, 31
  %20 = zext nneg i32 %19 to i64
  %21 = shl nuw nsw i32 %18, 2
  %22 = zext nneg i32 %21 to i64
  %invariant.gep = getelementptr [4 x i8], ptr addrspace(1) %1, i64 %22
  %.idx = mul nuw nsw i64 %22, 14336
  %23 = getelementptr i8, ptr addrspace(1) %2, i64 %.idx
  %24 = mul nuw nsw i32 %18, 14336
  %25 = zext nneg i32 %24 to i64
  %invariant.gep2 = getelementptr [4 x i8], ptr addrspace(1) %2, i64 %25
  %26 = icmp eq i32 %19, 0
  %27 = lshr i32 %18, 5
  %28 = zext nneg i32 %27 to i64
  %29 = getelementptr [4 x i8], ptr addrspace(3) @__dynamic_shared_memory__, i64 %28
  %.scalar = icmp samesign ult i32 %19, 28
  %30 = insertelement <1 x i1> poison, i1 %.scalar, i64 0
  %31 = getelementptr [4 x i8], ptr addrspace(3) @__dynamic_shared_memory__, i64 %20
  %32 = icmp eq i32 %18, 0
  br label %33

33:                                               ; preds = %.lr.ph, %90
  %34 = phi i64 [ %14, %.lr.ph ], [ %91, %90 ]
  %35 = urem i64 %34, 3584
  %36 = sub nuw i64 %34, %35
  %gep = getelementptr [4 x i8], ptr addrspace(1) %invariant.gep, i64 %36
  %37 = getelementptr [4 x i8], ptr addrspace(1) %23, i64 %35
  %gep3 = getelementptr [4 x i8], ptr addrspace(1) %invariant.gep2, i64 %35
  %38 = getelementptr i8, ptr addrspace(1) %gep3, i64 14336
  %39 = getelementptr i8, ptr addrspace(1) %gep3, i64 28672
  %40 = getelementptr i8, ptr addrspace(1) %gep3, i64 43008
  %41 = load float, ptr addrspace(1) %37, align 4
  %42 = load float, ptr addrspace(1) %38, align 4
  %43 = load float, ptr addrspace(1) %39, align 4
  %44 = load float, ptr addrspace(1) %40, align 4
  %45 = load float, ptr addrspace(1) %gep, align 4
  %46 = getelementptr inbounds nuw i8, ptr addrspace(1) %gep, i64 4
  %47 = load float, ptr addrspace(1) %46, align 4
  %48 = getelementptr inbounds nuw i8, ptr addrspace(1) %gep, i64 8
  %49 = load float, ptr addrspace(1) %48, align 4
  %50 = getelementptr inbounds nuw i8, ptr addrspace(1) %gep, i64 12
  %51 = load float, ptr addrspace(1) %50, align 4
  %52 = tail call float @llvm.fma.f32(float %51, float %44, float 0.000000e+00)
  %53 = tail call float @llvm.fma.f32(float %49, float %43, float %52)
  %54 = tail call float @llvm.fma.f32(float %47, float %42, float %53)
  %55 = tail call float @llvm.fma.f32(float %45, float %41, float %54)
  %56 = fadd float %55, 0.000000e+00
  %57 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %56, i32 1, i32 31)
  %58 = fadd float %57, %56
  %59 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %58, i32 2, i32 31)
  %60 = fadd float %59, %58
  %61 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %60, i32 4, i32 31)
  %62 = fadd float %61, %60
  %63 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %62, i32 8, i32 31)
  %64 = fadd float %63, %62
  %65 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %64, i32 16, i32 31)
  tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)
  br i1 %26, label %66, label %69

66:                                               ; preds = %33
  %67 = fadd float %65, %64
  %68 = insertelement <1 x float> poison, float %67, i64 0
  store <1 x float> %68, ptr addrspace(3) %29, align 4
  br label %69

69:                                               ; preds = %66, %33
  tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)
  %70 = tail call <1 x float> @llvm.masked.load.v1f32.p3(ptr addrspace(3) align 4 %31, <1 x i1> %30, <1 x float> zeroinitializer)
  %71 = extractelement <1 x float> %70, i64 0
  %72 = fadd float %71, 0.000000e+00
  %73 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %72, i32 1, i32 31)
  %74 = fadd float %73, %72
  %75 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %74, i32 2, i32 31)
  %76 = fadd float %75, %74
  %77 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %76, i32 4, i32 31)
  %78 = fadd float %77, %76
  %79 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %78, i32 8, i32 31)
  %80 = fadd float %79, %78
  %81 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %80, i32 16, i32 31)
  br i1 %32, label %82, label %90

82:                                               ; preds = %69
  %83 = getelementptr [4 x i8], ptr addrspace(1) %3, i64 %34
  %84 = fadd float %81, %80
  %85 = insertelement <1 x float> poison, float %84, i64 0
  %86 = getelementptr [2 x i8], ptr addrspace(1) %0, i64 %34
  %87 = load <1 x half>, ptr addrspace(1) %86, align 2
  %88 = fpext <1 x half> %87 to <1 x float>
  %89 = fadd <1 x float> %85, %88
  store <1 x float> %89, ptr addrspace(1) %83, align 4
  br label %90

90:                                               ; preds = %82, %69
  %91 = add i64 %34, %16
  %92 = icmp slt i64 %91, %12
  br i1 %92, label %33, label %._crit_edge

._crit_edge:                                      ; preds = %90, %6
  ret void
}