define ptx_kernel void @main_graph_async_dispatch_15_matmul_Dx3584x3584_f16xf16xf32(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, i32 noundef %6, i32 noundef %7, i32 noundef %8, i32 noundef %9) #3 {
  %11 = call range(i32 0, 896) i32 @llvm.nvvm.read.ptx.sreg.tid.x()
  %12 = sext i32 %11 to i64
  %13 = zext i32 %4 to i64
  %14 = zext i32 %5 to i64
  %15 = shl i64 %14, 32
  %16 = or i64 %13, %15
  %17 = zext i32 %6 to i64
  %18 = zext i32 %7 to i64
  %19 = shl i64 %18, 32
  %20 = or i64 %17, %19
  %21 = zext i32 %8 to i64
  %22 = zext i32 %9 to i64
  %23 = shl i64 %22, 32
  %24 = or i64 %21, %23
  %25 = urem i64 %16, 256
  %26 = icmp eq i64 %25, 0
  call void @llvm.assume(i1 %26)
  %27 = urem i64 %20, 256
  %28 = icmp eq i64 %27, 0
  call void @llvm.assume(i1 %28)
  %29 = icmp ule i64 %24, 9223372036854775807
  call void @llvm.assume(i1 %29)
  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) ]
  %30 = mul i64 %16, 8
  %31 = udiv i64 %30, 16
  %32 = getelementptr half, ptr addrspace(1) %0, i64 %31
  call void @llvm.assume(i1 true) [ "align"(ptr addrspace(1) %32, i64 64) ]
  %33 = mul i64 %20, 8
  %34 = udiv i64 %33, 16
  %35 = getelementptr half, ptr addrspace(1) %3, i64 %34
  call void @llvm.assume(i1 true) [ "align"(ptr addrspace(1) %35, i64 64) ]
  %36 = mul nsw i64 %24, 3584
  %37 = call range(i32 0, 2147483647) i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
  %38 = sext i32 %37 to i64
  %39 = call range(i32 1, -2147483648) i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()
  %40 = sext i32 %39 to i64
  br label %41

41:                                               ; preds = %155, %10
  %42 = phi i64 [ %156, %155 ], [ %38, %10 ]
  %43 = icmp slt i64 %42, %36
  br i1 %43, label %44, label %157

44:                                               ; preds = %41
  %45 = udiv i64 %42, 3584
  %46 = urem i64 %42, 3584
  %47 = trunc i64 %12 to i32
  %48 = udiv i32 %47, 32
  %49 = zext i32 %48 to i64
  %50 = urem i32 %47, 32
  %51 = zext i32 %50 to i64
  %52 = mul nsw i32 %48, 128
  %53 = mul nsw i32 %50, 4
  %54 = add nsw i32 %52, %53
  %55 = zext i32 %54 to i64
  %56 = getelementptr half, ptr addrspace(1) %0, i64 %31
  %57 = mul i64 %45, 3584
  %58 = add i64 %57, %55
  %59 = getelementptr half, ptr addrspace(1) %56, i64 %58
  %60 = load <4 x half>, ptr addrspace(1) %59, align 2
  %61 = mul i64 %55, 3584
  %62 = add i64 %61, %46
  %63 = getelementptr half, ptr addrspace(1) %1, i64 %62
  %64 = load <1 x half>, ptr addrspace(1) %63, align 2
  %65 = add nsw i32 %54, 1
  %66 = zext i32 %65 to i64
  %67 = mul i64 %66, 3584
  %68 = add i64 %67, %46
  %69 = getelementptr half, ptr addrspace(1) %1, i64 %68
  %70 = load <1 x half>, ptr addrspace(1) %69, align 2
  %71 = add nsw i32 %54, 2
  %72 = zext i32 %71 to i64
  %73 = mul i64 %72, 3584
  %74 = add i64 %73, %46
  %75 = getelementptr half, ptr addrspace(1) %1, i64 %74
  %76 = load <1 x half>, ptr addrspace(1) %75, align 2
  %77 = add nsw i32 %54, 3
  %78 = zext i32 %77 to i64
  %79 = mul i64 %78, 3584
  %80 = add i64 %79, %46
  %81 = getelementptr half, ptr addrspace(1) %1, i64 %80
  %82 = load <1 x half>, ptr addrspace(1) %81, align 2
  %83 = extractelement <1 x half> %64, i64 0
  %84 = extractelement <1 x half> %70, i64 0
  %85 = extractelement <1 x half> %76, i64 0
  %86 = extractelement <1 x half> %82, i64 0
  %87 = insertelement <4 x half> poison, half %83, i64 0
  %88 = insertelement <4 x half> %87, half %84, i64 1
  %89 = insertelement <4 x half> %88, half %85, i64 2
  %90 = insertelement <4 x half> %89, half %86, i64 3
  %91 = fpext <4 x half> %60 to <4 x float>
  %92 = fpext <4 x half> %90 to <4 x float>
  %93 = extractelement <4 x float> %91, i64 0
  %94 = extractelement <4 x float> %91, i64 1
  %95 = extractelement <4 x float> %91, i64 2
  %96 = extractelement <4 x float> %91, i64 3
  %97 = extractelement <4 x float> %92, i64 0
  %98 = extractelement <4 x float> %92, i64 1
  %99 = extractelement <4 x float> %92, i64 2
  %100 = extractelement <4 x float> %92, i64 3
  %101 = call float @__nv_fmaf(float %96, float %100, float 0.000000e+00)
  %102 = call float @__nv_fmaf(float %95, float %99, float %101)
  %103 = call float @__nv_fmaf(float %94, float %98, float %102)
  %104 = call float @__nv_fmaf(float %93, float %97, float %103)
  %105 = fadd float %104, 0.000000e+00
  %106 = call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %105, i32 1, i32 31)
  %107 = fadd float %105, %106
  %108 = call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %107, i32 2, i32 31)
  %109 = fadd float %107, %108
  %110 = call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %109, i32 4, i32 31)
  %111 = fadd float %109, %110
  %112 = call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %111, i32 8, i32 31)
  %113 = fadd float %111, %112
  %114 = call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %113, i32 16, i32 31)
  %115 = fadd float %113, %114
  call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)
  %116 = icmp eq i32 %50, 0
  br i1 %116, label %117, label %121

117:                                              ; preds = %44
  %118 = insertelement <1 x float> poison, float %115, i32 0
  %119 = add i64 0, %49
  %120 = getelementptr float, ptr addrspace(3) @__dynamic_shared_memory__, i64 %119
  store <1 x float> %118, ptr addrspace(3) %120, align 4
  br label %121

121:                                              ; preds = %117, %44
  call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)
  %122 = icmp eq i32 %50, 27
  %123 = icmp slt i32 %50, 27
  %124 = select i1 %123, i64 1, i64 0
  %125 = select i1 %122, i64 1, i64 %124
  %126 = insertelement <1 x i64> poison, i64 %125, i32 0
  %127 = icmp sgt <1 x i64> %126, zeroinitializer
  %128 = add i64 0, %51
  %129 = getelementptr float, ptr addrspace(3) @__dynamic_shared_memory__, i64 %128
  %130 = call <1 x float> @llvm.masked.load.v1f32.p3(ptr addrspace(3) align 4 %129, <1 x i1> %127, <1 x float> zeroinitializer)
  %131 = extractelement <1 x float> %130, i64 0
  %132 = fadd float %131, 0.000000e+00
  %133 = call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %132, i32 1, i32 31)
  %134 = fadd float %132, %133
  %135 = call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %134, i32 2, i32 31)
  %136 = fadd float %134, %135
  %137 = call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %136, i32 4, i32 31)
  %138 = fadd float %136, %137
  %139 = call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %138, i32 8, i32 31)
  %140 = fadd float %138, %139
  %141 = call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %140, i32 16, i32 31)
  %142 = fadd float %140, %141
  %143 = fadd float %142, 0.000000e+00
  %144 = getelementptr half, ptr addrspace(1) %2, i64 %46
  %145 = load <1 x half>, ptr addrspace(1) %144, align 2
  %146 = fptrunc float %143 to half
  %147 = insertelement <1 x half> poison, half %146, i32 0
  %148 = fadd <1 x half> %145, %147
  %149 = icmp eq i32 %47, 0
  br i1 %149, label %150, label %155

150:                                              ; preds = %121
  %151 = getelementptr half, ptr addrspace(1) %3, i64 %34
  %152 = mul i64 %45, 3584
  %153 = add i64 %152, %46
  %154 = getelementptr half, ptr addrspace(1) %151, i64 %153
  store <1 x half> %148, ptr addrspace(1) %154, align 2
  br label %155

155:                                              ; preds = %150, %121
  %156 = add i64 %42, %40
  br label %41

157:                                              ; preds = %41
  ret void
}