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

.lr.ph:                                           ; preds = %11
  %38 = zext i32 %4 to i64
  %39 = shl nuw i64 %38, 32
  %40 = zext i32 %3 to i64
  %41 = or disjoint i64 %39, %40
  %42 = lshr i64 %41, 1
  %43 = and i64 %42, 1152921504606846975
  %44 = getelementptr [2 x i8], ptr addrspace(1) %0, i64 %43
  %45 = tail call range(i32 0, 896) i32 @llvm.nvvm.read.ptx.sreg.tid.x()
  %46 = and i32 %45, 31
  %47 = zext nneg i32 %46 to i64
  %48 = shl nuw nsw i32 %45, 2
  %49 = zext nneg i32 %48 to i64
  %invariant.gep = getelementptr [2 x i8], ptr addrspace(1) %44, i64 %49
  %.idx1 = mul nuw nsw i64 %49, 37888
  %50 = getelementptr i8, ptr addrspace(1) %1, i64 %.idx1
  %51 = mul nuw nsw i32 %45, 75776
  %52 = zext nneg i32 %51 to i64
  %invariant.gep5 = getelementptr [2 x i8], ptr addrspace(1) %1, i64 %52
  %53 = icmp eq i32 %46, 0
  %54 = lshr i32 %45, 5
  %55 = zext nneg i32 %54 to i64
  %56 = getelementptr [4 x i8], ptr addrspace(3) @__dynamic_shared_memory__, i64 %55
  %.scalar = icmp samesign ult i32 %46, 28
  %57 = insertelement <1 x i1> poison, i1 %.scalar, i64 0
  %58 = getelementptr [4 x i8], ptr addrspace(3) @__dynamic_shared_memory__, i64 %47
  %59 = icmp eq i32 %45, 0
  br label %60

60:                                               ; preds = %.lr.ph, %155
  %61 = phi i64 [ %34, %.lr.ph ], [ %156, %155 ]
  %.frozen = freeze i64 %61
  %62 = udiv i64 %.frozen, 18944
  %63 = mul i64 %62, 18944
  %.decomposed = sub i64 %.frozen, %63
  %.idx = mul nuw nsw i64 %62, 7168
  %gep = getelementptr i8, ptr addrspace(1) %invariant.gep, i64 %.idx
  %64 = load <4 x half>, ptr addrspace(1) %gep, align 2
  %65 = getelementptr [2 x i8], ptr addrspace(1) %50, i64 %.decomposed
  %66 = load <1 x half>, ptr addrspace(1) %65, align 2
  %gep6 = getelementptr [2 x i8], ptr addrspace(1) %invariant.gep5, i64 %.decomposed
  %67 = getelementptr i8, ptr addrspace(1) %gep6, i64 37888
  %68 = load <1 x half>, ptr addrspace(1) %67, align 2
  %69 = getelementptr i8, ptr addrspace(1) %gep6, i64 75776
  %70 = load <1 x half>, ptr addrspace(1) %69, align 2
  %71 = shufflevector <1 x half> %70, <1 x half> poison, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
  %72 = getelementptr i8, ptr addrspace(1) %gep6, i64 113664
  %73 = load <1 x half>, ptr addrspace(1) %72, align 2
  %74 = shufflevector <1 x half> %73, <1 x half> poison, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
  %75 = shufflevector <1 x half> %66, <1 x half> %68, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>
  %76 = shufflevector <4 x half> %75, <4 x half> %71, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>
  %77 = shufflevector <4 x half> %76, <4 x half> %74, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
  %78 = fpext <4 x half> %64 to <4 x float>
  %79 = fpext <4 x half> %77 to <4 x float>
  %80 = extractelement <4 x float> %78, i64 0
  %81 = extractelement <4 x float> %78, i64 1
  %82 = extractelement <4 x float> %78, i64 2
  %83 = extractelement <4 x float> %78, i64 3
  %84 = extractelement <4 x float> %79, i64 0
  %85 = extractelement <4 x float> %79, i64 1
  %86 = extractelement <4 x float> %79, i64 2
  %87 = extractelement <4 x float> %79, i64 3
  %88 = tail call float @llvm.fma.f32(float %83, float %87, float 0.000000e+00)
  %89 = tail call float @llvm.fma.f32(float %82, float %86, float %88)
  %90 = tail call float @llvm.fma.f32(float %81, float %85, float %89)
  %91 = tail call float @llvm.fma.f32(float %80, float %84, float %90)
  %92 = fadd float %91, 0.000000e+00
  %93 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %92, i32 1, i32 31)
  %94 = fadd float %93, %92
  %95 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %94, i32 2, i32 31)
  %96 = fadd float %95, %94
  %97 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %96, i32 4, i32 31)
  %98 = fadd float %97, %96
  %99 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %98, i32 8, i32 31)
  %100 = fadd float %99, %98
  %101 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %100, i32 16, i32 31)
  tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)
  br i1 %53, label %102, label %105

102:                                              ; preds = %60
  %103 = fadd float %101, %100
  %104 = insertelement <1 x float> poison, float %103, i64 0
  store <1 x float> %104, ptr addrspace(3) %56, align 4
  br label %105

105:                                              ; preds = %102, %60
  tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)
  %106 = tail call <1 x float> @llvm.masked.load.v1f32.p3(ptr addrspace(3) align 4 %58, <1 x i1> %57, <1 x float> zeroinitializer)
  %107 = extractelement <1 x float> %106, i64 0
  %108 = fadd float %107, 0.000000e+00
  %109 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %108, i32 1, i32 31)
  %110 = fadd float %109, %108
  %111 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %110, i32 2, i32 31)
  %112 = fadd float %111, %110
  %113 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %112, i32 4, i32 31)
  %114 = fadd float %113, %112
  %115 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %114, i32 8, i32 31)
  %116 = fadd float %115, %114
  %117 = tail call float @llvm.nvvm.shfl.sync.bfly.f32(i32 -1, float %116, i32 16, i32 31)
  br i1 %59, label %118, label %155

118:                                              ; preds = %105
  %119 = fadd float %117, %116
  %120 = fptrunc float %119 to half
  %121 = fneg half %120
  %122 = fpext half %121 to float
  %.inv = fcmp olt float %122, -8.780000e+01
  %123 = select i1 %.inv, float -8.780000e+01, float %122
  %.inv2 = fcmp ogt float %123, 8.880000e+01
  %124 = select i1 %.inv2, float 8.880000e+01, float %123
  %125 = tail call float @llvm.fma.f32(float %124, float f0x3FB8AA3B, float 5.000000e-01)
  %126 = tail call float @llvm.floor.f32(float %125)
  %.inv3 = fcmp olt float %126, -1.270000e+02
  %127 = select i1 %.inv3, float -1.270000e+02, float %126
  %.inv4 = fcmp ogt float %127, 1.270000e+02
  %128 = select i1 %.inv4, float 1.270000e+02, float %127
  %129 = tail call float @llvm.fma.f32(float %128, float f0xBF318000, float %124)
  %130 = tail call float @llvm.fma.f32(float %128, float f0x395E8083, float %129)
  %131 = tail call float @llvm.fma.f32(float %130, float f0x39506967, float f0x3AB743CE)
  %132 = tail call float @llvm.fma.f32(float %131, float %130, float f0x3C088908)
  %133 = tail call float @llvm.fma.f32(float %132, float %130, float f0x3D2AA9C1)
  %134 = tail call float @llvm.fma.f32(float %133, float %130, float f0x3E2AAAAA)
  %135 = tail call float @llvm.fma.f32(float %134, float %130, float 5.000000e-01)
  %136 = fmul float %130, %130
  %137 = tail call float @llvm.fma.f32(float %135, float %136, float %130)
  %138 = mul nuw i64 %62, 18944
  %139 = add i64 %138, %.decomposed
  %140 = getelementptr [2 x i8], ptr addrspace(1) %27, i64 %139
  %141 = load <1 x half>, ptr addrspace(1) %140, align 2
  %142 = fadd float %137, 1.000000e+00
  %143 = fptosi float %128 to i32
  %144 = shl i32 %143, 23
  %145 = add i32 %144, 1065353216
  %146 = bitcast i32 %145 to float
  %147 = fmul float %142, %146
  %148 = fptrunc float %147 to half
  %149 = fadd half %148, 1.000000e+00
  %150 = fdiv half 1.000000e+00, %149
  %151 = fmul half %150, %120
  %152 = insertelement <1 x half> poison, half %151, i64 0
  %153 = fmul <1 x half> %152, %141
  %154 = getelementptr [2 x i8], ptr addrspace(1) %31, i64 %139
  store <1 x half> %153, ptr addrspace(1) %154, align 2
  br label %155

155:                                              ; preds = %118, %105
  %156 = add i64 %61, %36
  %157 = icmp slt i64 %156, %32
  br i1 %157, label %60, label %._crit_edge

._crit_edge:                                      ; preds = %155, %11
  ret void
}