define ptx_kernel void @main_graph_async_dispatch_35_matmul_Dx3584x3584_f32(ptr addrspace(1) noalias noundef nonnull readonly align 16 %0, ptr addrspace(1) noalias noundef nonnull readonly align 16 %1, ptr addrspace(1) noalias noundef nonnull readonly align 16 %2, ptr addrspace(1) noalias noundef nonnull align 16 %3, i32 noundef %4, i32 noundef %5) #3 {
  %7 = call range(i32 0, 896) i32 @llvm.nvvm.read.ptx.sreg.tid.x()
  %8 = sext i32 %7 to i64
  %9 = zext i32 %4 to i64
  %10 = zext i32 %5 to i64
  %11 = shl i64 %10, 32
  %12 = or i64 %9, %11
  %13 = icmp ule i64 %12, 9223372036854775807
  call void @llvm.assume(i1 %13)
  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) ]
  %14 = mul nsw i64 %12, 3584
  %15 = call range(i32 0, 2147483647) i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
  %16 = sext i32 %15 to i64
  %17 = call range(i32 1, -2147483648) i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()
  %18 = sext i32 %17 to i64
  br label %19

19:                                               ; preds = %123, %6
  %20 = phi i64 [ %124, %123 ], [ %16, %6 ]
  %21 = icmp slt i64 %20, %14
  br i1 %21, label %22, label %125

22:                                               ; preds = %19
  %23 = udiv i64 %20, 3584
  %24 = urem i64 %20, 3584
  %25 = trunc i64 %8 to i32
  %26 = udiv i32 %25, 32
  %27 = zext i32 %26 to i64
  %28 = urem i32 %25, 32
  %29 = zext i32 %28 to i64
  %30 = mul nsw i32 %26, 128
  %31 = mul nsw i32 %28, 4
  %32 = add nsw i32 %30, %31
  %33 = zext i32 %32 to i64
  %34 = mul i64 %23, 3584
  %35 = add i64 %34, %33
  %36 = getelementptr float, ptr addrspace(1) %1, i64 %35
  %37 = load <4 x float>, ptr addrspace(1) %36, align 4
  %38 = mul i64 %33, 3584
  %39 = add i64 %38, %24
  %40 = getelementptr float, ptr addrspace(1) %2, i64 %39
  %41 = load <1 x float>, ptr addrspace(1) %40, align 4
  %42 = add nsw i32 %32, 1
  %43 = zext i32 %42 to i64
  %44 = mul i64 %43, 3584
  %45 = add i64 %44, %24
  %46 = getelementptr float, ptr addrspace(1) %2, i64 %45
  %47 = load <1 x float>, ptr addrspace(1) %46, align 4
  %48 = add nsw i32 %32, 2
  %49 = zext i32 %48 to i64
  %50 = mul i64 %49, 3584
  %51 = add i64 %50, %24
  %52 = getelementptr float, ptr addrspace(1) %2, i64 %51
  %53 = load <1 x float>, ptr addrspace(1) %52, align 4
  %54 = add nsw i32 %32, 3
  %55 = zext i32 %54 to i64
  %56 = mul i64 %55, 3584
  %57 = add i64 %56, %24
  %58 = getelementptr float, ptr addrspace(1) %2, i64 %57
  %59 = load <1 x float>, ptr addrspace(1) %58, align 4
  %60 = extractelement <1 x float> %41, i64 0
  %61 = extractelement <1 x float> %47, i64 0
  %62 = extractelement <1 x float> %53, i64 0
  %63 = extractelement <1 x float> %59, i64 0
  %64 = extractelement <4 x float> %37, i64 0
  %65 = extractelement <4 x float> %37, i64 1
  %66 = extractelement <4 x float> %37, i64 2
  %67 = extractelement <4 x float> %37, i64 3
  %68 = call float @__nv_fmaf(float %67, float %63, float 0.000000e+00)
  %69 = call float @__nv_fmaf(float %66, float %62, float %68)
  %70 = call float @__nv_fmaf(float %65, float %61, float %69)
  %71 = call float @__nv_fmaf(float %64, float %60, float %70)
  %72 = fadd float %71, 0.000000e+00
  %73 = call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %72, i32 1, i32 31)
  %74 = fadd float %72, %73
  %75 = call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %74, i32 2, i32 31)
  %76 = fadd float %74, %75
  %77 = call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %76, i32 4, i32 31)
  %78 = fadd float %76, %77
  %79 = call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %78, i32 8, i32 31)
  %80 = fadd float %78, %79
  %81 = call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %80, i32 16, i32 31)
  %82 = fadd float %80, %81
  call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)
  %83 = icmp eq i32 %28, 0
  br i1 %83, label %84, label %88

84:                                               ; preds = %22
  %85 = insertelement <1 x float> poison, float %82, i32 0
  %86 = add i64 0, %27
  %87 = getelementptr float, ptr addrspace(3) @__dynamic_shared_memory__, i64 %86
  store <1 x float> %85, ptr addrspace(3) %87, align 4
  br label %88

88:                                               ; preds = %84, %22
  call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)
  %89 = icmp eq i32 %28, 27
  %90 = icmp slt i32 %28, 27
  %91 = select i1 %90, i64 1, i64 0
  %92 = select i1 %89, i64 1, i64 %91
  %93 = insertelement <1 x i64> poison, i64 %92, i32 0
  %94 = icmp sgt <1 x i64> %93, zeroinitializer
  %95 = add i64 0, %29
  %96 = getelementptr float, ptr addrspace(3) @__dynamic_shared_memory__, i64 %95
  %97 = call <1 x float> @llvm.masked.load.v1f32.p3(ptr addrspace(3) align 4 %96, <1 x i1> %94, <1 x float> zeroinitializer)
  %98 = extractelement <1 x float> %97, i64 0
  %99 = fadd float %98, 0.000000e+00
  %100 = call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %99, i32 1, i32 31)
  %101 = fadd float %99, %100
  %102 = call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %101, i32 2, i32 31)
  %103 = fadd float %101, %102
  %104 = call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %103, i32 4, i32 31)
  %105 = fadd float %103, %104
  %106 = call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %105, i32 8, i32 31)
  %107 = fadd float %105, %106
  %108 = call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %107, i32 16, i32 31)
  %109 = fadd float %107, %108
  %110 = fadd float %109, 0.000000e+00
  %111 = insertelement <1 x float> poison, float %110, i32 0
  %112 = mul i64 %23, 3584
  %113 = add i64 %112, %24
  %114 = getelementptr half, ptr addrspace(1) %0, i64 %113
  %115 = load <1 x half>, ptr addrspace(1) %114, align 2
  %116 = fpext <1 x half> %115 to <1 x float>
  %117 = fadd <1 x float> %116, %111
  %118 = icmp eq i32 %25, 0
  br i1 %118, label %119, label %123

119:                                              ; preds = %88
  %120 = mul i64 %23, 3584
  %121 = add i64 %120, %24
  %122 = getelementptr float, ptr addrspace(1) %3, i64 %121
  store <1 x float> %117, ptr addrspace(1) %122, align 4
  br label %123

123:                                              ; preds = %119, %88
  %124 = add i64 %20, %18
  br label %19

125:                                              ; preds = %19
  ret void
}