; ModuleID = 'LLVMDialectModule'
source_filename = "LLVMDialectModule"

@assertFunc_1 = internal constant [25 x i8] c"_call_with_frames_removed"
@assertFile_1 = internal constant [38 x i8] c"<frozen importlib._bootstrap_external>"
@assertMessage_1 = internal constant [39 x i8] c"index out of bounds: 0 <= tmp13 < 50257"
@assertFunc_0 = internal constant [25 x i8] c"_call_with_frames_removed"
@assertFile_0 = internal constant [38 x i8] c"<frozen importlib._bootstrap_external>"
@assertMessage_0 = internal constant [38 x i8] c"index out of bounds: 0 <= tmp3 < 50257"
@global_smem = external local_unnamed_addr addrspace(3) global [0 x i8]
@.str = private unnamed_addr constant [11 x i8] c"__CUDA_FTZ\00", align 1

declare void @__assertfail(ptr, ptr, i32, ptr, i64) local_unnamed_addr

define void @triton__0d1d2d3d4d5de6de(ptr addrspace(1) %0, ptr addrspace(1) %1, ptr addrspace(1) %2, ptr addrspace(1) %3, ptr addrspace(1) %4, i32 %5, i32 %6) local_unnamed_addr !dbg !7 {
  %8 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !10
  %9 = lshr i32 %8, 5, !dbg !10
  %10 = and i32 %9, 7, !dbg !10
  %11 = and i32 %8, 15, !dbg !10
  %12 = shl i32 %8, 3, !dbg !11
  %13 = and i32 %12, 248, !dbg !11
  %14 = or i32 %13, 4, !dbg !11
  %urem = and i32 %8, 255, !dbg !11
  %15 = tail call i32 asm "mov.u32 $0, %ctaid.x;", "=r"() #6, !dbg !12
  %16 = shl i32 %15, 4, !dbg !13
  %17 = or i32 %16, %10, !dbg !14
  %18 = or i32 %17, 8, !dbg !14
  %19 = or i32 %16, %11, !dbg !14
  %20 = sext i32 %17 to i64, !dbg !15
  %21 = getelementptr i64, ptr addrspace(1) %0, i64 %20, !dbg !15
  %22 = sext i32 %18 to i64, !dbg !15
  %23 = getelementptr i64, ptr addrspace(1) %0, i64 %22, !dbg !15
  %24 = sext i32 %19 to i64, !dbg !15
  %25 = getelementptr i64, ptr addrspace(1) %0, i64 %24, !dbg !15
  %26 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$2 ld.global.L1::evict_last.b64 { $0 }, [ $1 + 0 ];", "=l,l,b"(ptr addrspace(1) %21, i1 true) #6, !dbg !16
  %27 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$2 ld.global.L1::evict_last.b64 { $0 }, [ $1 + 0 ];", "=l,l,b"(ptr addrspace(1) %21, i1 true) #6, !dbg !16
  %28 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$2 ld.global.L1::evict_last.b64 { $0 }, [ $1 + 0 ];", "=l,l,b"(ptr addrspace(1) %21, i1 true) #6, !dbg !16
  %29 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$2 ld.global.L1::evict_last.b64 { $0 }, [ $1 + 0 ];", "=l,l,b"(ptr addrspace(1) %21, i1 true) #6, !dbg !16
  %30 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$2 ld.global.L1::evict_last.b64 { $0 }, [ $1 + 0 ];", "=l,l,b"(ptr addrspace(1) %21, i1 true) #6, !dbg !16
  %31 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$2 ld.global.L1::evict_last.b64 { $0 }, [ $1 + 0 ];", "=l,l,b"(ptr addrspace(1) %21, i1 true) #6, !dbg !16
  %32 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$2 ld.global.L1::evict_last.b64 { $0 }, [ $1 + 0 ];", "=l,l,b"(ptr addrspace(1) %21, i1 true) #6, !dbg !16
  %33 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$2 ld.global.L1::evict_last.b64 { $0 }, [ $1 + 0 ];", "=l,l,b"(ptr addrspace(1) %21, i1 true) #6, !dbg !16
  %34 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$2 ld.global.L1::evict_last.b64 { $0 }, [ $1 + 0 ];", "=l,l,b"(ptr addrspace(1) %23, i1 true) #6, !dbg !16
  %35 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$2 ld.global.L1::evict_last.b64 { $0 }, [ $1 + 0 ];", "=l,l,b"(ptr addrspace(1) %23, i1 true) #6, !dbg !16
  %36 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$2 ld.global.L1::evict_last.b64 { $0 }, [ $1 + 0 ];", "=l,l,b"(ptr addrspace(1) %23, i1 true) #6, !dbg !16
  %37 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$2 ld.global.L1::evict_last.b64 { $0 }, [ $1 + 0 ];", "=l,l,b"(ptr addrspace(1) %23, i1 true) #6, !dbg !16
  %38 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$2 ld.global.L1::evict_last.b64 { $0 }, [ $1 + 0 ];", "=l,l,b"(ptr addrspace(1) %23, i1 true) #6, !dbg !16
  %39 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$2 ld.global.L1::evict_last.b64 { $0 }, [ $1 + 0 ];", "=l,l,b"(ptr addrspace(1) %23, i1 true) #6, !dbg !16
  %40 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$2 ld.global.L1::evict_last.b64 { $0 }, [ $1 + 0 ];", "=l,l,b"(ptr addrspace(1) %23, i1 true) #6, !dbg !16
  %41 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$2 ld.global.L1::evict_last.b64 { $0 }, [ $1 + 0 ];", "=l,l,b"(ptr addrspace(1) %23, i1 true) #6, !dbg !16
  %42 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$2 ld.global.L1::evict_last.b64 { $0 }, [ $1 + 0 ];", "=l,l,b"(ptr addrspace(1) %25, i1 true) #6, !dbg !16
  %43 = srem i32 %17, 512, !dbg !17
  %44 = srem i32 %18, 512, !dbg !17
  %45 = shl nsw i32 %43, 8, !dbg !18
  %46 = shl nsw i32 %44, 8, !dbg !18
  %47 = or i32 %45, %13, !dbg !19
  %48 = or i32 %45, %14, !dbg !19
  %49 = or i32 %46, %13, !dbg !19
  %50 = or i32 %46, %14, !dbg !19
  %51 = sext i32 %47 to i64, !dbg !20
  %52 = getelementptr float, ptr addrspace(1) %2, i64 %51, !dbg !20
  %53 = sext i32 %48 to i64, !dbg !20
  %54 = getelementptr float, ptr addrspace(1) %2, i64 %53, !dbg !20
  %55 = sext i32 %49 to i64, !dbg !20
  %56 = getelementptr float, ptr addrspace(1) %2, i64 %55, !dbg !20
  %57 = sext i32 %50 to i64, !dbg !20
  %58 = getelementptr float, ptr addrspace(1) %2, i64 %57, !dbg !20
  %59 = tail call { i32, i32, i32, i32 } asm sideeffect "mov.u32 $0, 0x0;\0A\09mov.u32 $1, 0x0;\0A\09mov.u32 $2, 0x0;\0A\09mov.u32 $3, 0x0;\0A\09@$5 ld.global.L1::evict_last.v4.b32 { $0, $1, $2, $3 }, [ $4 + 0 ];\0A\09@!$7 mov.u32 $0, $6;\0A\09@!$9 mov.u32 $1, $8;\0A\09@!$11 mov.u32 $2, $10;\0A\09@!$13 mov.u32 $3, $12;", "=r,=r,=r,=r,l,b,r,b,r,b,r,b,r,b"(ptr addrspace(1) %52, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true) #6, !dbg !21
  %60 = extractvalue { i32, i32, i32, i32 } %59, 0, !dbg !21
  %61 = extractvalue { i32, i32, i32, i32 } %59, 1, !dbg !21
  %62 = extractvalue { i32, i32, i32, i32 } %59, 2, !dbg !21
  %63 = extractvalue { i32, i32, i32, i32 } %59, 3, !dbg !21
  %64 = bitcast i32 %60 to float, !dbg !21
  %65 = bitcast i32 %61 to float, !dbg !21
  %66 = bitcast i32 %62 to float, !dbg !21
  %67 = bitcast i32 %63 to float, !dbg !21
  %68 = tail call { i32, i32, i32, i32 } asm sideeffect "mov.u32 $0, 0x0;\0A\09mov.u32 $1, 0x0;\0A\09mov.u32 $2, 0x0;\0A\09mov.u32 $3, 0x0;\0A\09@$5 ld.global.L1::evict_last.v4.b32 { $0, $1, $2, $3 }, [ $4 + 0 ];\0A\09@!$7 mov.u32 $0, $6;\0A\09@!$9 mov.u32 $1, $8;\0A\09@!$11 mov.u32 $2, $10;\0A\09@!$13 mov.u32 $3, $12;", "=r,=r,=r,=r,l,b,r,b,r,b,r,b,r,b"(ptr addrspace(1) %54, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true) #6, !dbg !21
  %69 = extractvalue { i32, i32, i32, i32 } %68, 0, !dbg !21
  %70 = extractvalue { i32, i32, i32, i32 } %68, 1, !dbg !21
  %71 = extractvalue { i32, i32, i32, i32 } %68, 2, !dbg !21
  %72 = extractvalue { i32, i32, i32, i32 } %68, 3, !dbg !21
  %73 = bitcast i32 %69 to float, !dbg !21
  %74 = bitcast i32 %70 to float, !dbg !21
  %75 = bitcast i32 %71 to float, !dbg !21
  %76 = bitcast i32 %72 to float, !dbg !21
  %77 = tail call { i32, i32, i32, i32 } asm sideeffect "mov.u32 $0, 0x0;\0A\09mov.u32 $1, 0x0;\0A\09mov.u32 $2, 0x0;\0A\09mov.u32 $3, 0x0;\0A\09@$5 ld.global.L1::evict_last.v4.b32 { $0, $1, $2, $3 }, [ $4 + 0 ];\0A\09@!$7 mov.u32 $0, $6;\0A\09@!$9 mov.u32 $1, $8;\0A\09@!$11 mov.u32 $2, $10;\0A\09@!$13 mov.u32 $3, $12;", "=r,=r,=r,=r,l,b,r,b,r,b,r,b,r,b"(ptr addrspace(1) %56, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true) #6, !dbg !21
  %78 = extractvalue { i32, i32, i32, i32 } %77, 0, !dbg !21
  %79 = extractvalue { i32, i32, i32, i32 } %77, 1, !dbg !21
  %80 = extractvalue { i32, i32, i32, i32 } %77, 2, !dbg !21
  %81 = extractvalue { i32, i32, i32, i32 } %77, 3, !dbg !21
  %82 = bitcast i32 %78 to float, !dbg !21
  %83 = bitcast i32 %79 to float, !dbg !21
  %84 = bitcast i32 %80 to float, !dbg !21
  %85 = bitcast i32 %81 to float, !dbg !21
  %86 = tail call { i32, i32, i32, i32 } asm sideeffect "mov.u32 $0, 0x0;\0A\09mov.u32 $1, 0x0;\0A\09mov.u32 $2, 0x0;\0A\09mov.u32 $3, 0x0;\0A\09@$5 ld.global.L1::evict_last.v4.b32 { $0, $1, $2, $3 }, [ $4 + 0 ];\0A\09@!$7 mov.u32 $0, $6;\0A\09@!$9 mov.u32 $1, $8;\0A\09@!$11 mov.u32 $2, $10;\0A\09@!$13 mov.u32 $3, $12;", "=r,=r,=r,=r,l,b,r,b,r,b,r,b,r,b"(ptr addrspace(1) %58, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true) #6, !dbg !21
  %87 = extractvalue { i32, i32, i32, i32 } %86, 0, !dbg !21
  %88 = extractvalue { i32, i32, i32, i32 } %86, 1, !dbg !21
  %89 = extractvalue { i32, i32, i32, i32 } %86, 2, !dbg !21
  %90 = extractvalue { i32, i32, i32, i32 } %86, 3, !dbg !21
  %91 = bitcast i32 %87 to float, !dbg !21
  %92 = bitcast i32 %88 to float, !dbg !21
  %93 = bitcast i32 %89 to float, !dbg !21
  %94 = bitcast i32 %90 to float, !dbg !21
  %95 = add i64 %42, 50257, !dbg !22
  %96 = icmp slt i64 %26, 0, !dbg !23
  %97 = icmp slt i64 %34, 0, !dbg !23
  %98 = icmp slt i64 %42, 0, !dbg !23
  %99 = select i1 %98, i64 %95, i64 %42, !dbg !24
  %100 = icmp ugt i64 %99, 50256, !dbg !25
  br i1 %100, label %101, label %102, !dbg !26

101:                                              ; preds = %7
  tail call void @__assertfail(ptr nonnull @assertMessage_0, ptr nonnull @assertFile_0, i32 883, ptr nonnull @assertFunc_0, i64 1), !dbg !26
  br label %102, !dbg !26

102:                                              ; preds = %101, %7
  %103 = shl i64 %26, 8, !dbg !27
  %104 = add i64 %103, 12865792, !dbg !27
  %105 = select i1 %96, i64 %104, i64 %103, !dbg !27
  %106 = shl i64 %34, 8, !dbg !27
  %107 = add i64 %106, 12865792, !dbg !27
  %108 = select i1 %97, i64 %107, i64 %106, !dbg !27
  %109 = zext nneg i32 %13 to i64
  %110 = zext nneg i32 %14 to i64
  %111 = or i64 %105, %109, !dbg !28
  %112 = or i64 %105, %110, !dbg !28
  %113 = or i64 %108, %109, !dbg !28
  %114 = or i64 %108, %110, !dbg !28
  %115 = getelementptr float, ptr addrspace(1) %1, i64 %111, !dbg !29
  %116 = getelementptr float, ptr addrspace(1) %1, i64 %112, !dbg !29
  %117 = getelementptr float, ptr addrspace(1) %1, i64 %113, !dbg !29
  %118 = getelementptr float, ptr addrspace(1) %1, i64 %114, !dbg !29
  %119 = tail call { i32, i32, i32, i32 } asm sideeffect "mov.u32 $0, 0x0;\0A\09mov.u32 $1, 0x0;\0A\09mov.u32 $2, 0x0;\0A\09mov.u32 $3, 0x0;\0A\09@$5 ld.global.L1::evict_last.v4.b32 { $0, $1, $2, $3 }, [ $4 + 0 ];\0A\09@!$7 mov.u32 $0, $6;\0A\09@!$9 mov.u32 $1, $8;\0A\09@!$11 mov.u32 $2, $10;\0A\09@!$13 mov.u32 $3, $12;", "=r,=r,=r,=r,l,b,r,b,r,b,r,b,r,b"(ptr addrspace(1) %115, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true) #6, !dbg !30
  %120 = extractvalue { i32, i32, i32, i32 } %119, 0, !dbg !30
  %121 = extractvalue { i32, i32, i32, i32 } %119, 1, !dbg !30
  %122 = extractvalue { i32, i32, i32, i32 } %119, 2, !dbg !30
  %123 = extractvalue { i32, i32, i32, i32 } %119, 3, !dbg !30
  %124 = bitcast i32 %120 to float, !dbg !30
  %125 = bitcast i32 %121 to float, !dbg !30
  %126 = bitcast i32 %122 to float, !dbg !30
  %127 = bitcast i32 %123 to float, !dbg !30
  %128 = tail call { i32, i32, i32, i32 } asm sideeffect "mov.u32 $0, 0x0;\0A\09mov.u32 $1, 0x0;\0A\09mov.u32 $2, 0x0;\0A\09mov.u32 $3, 0x0;\0A\09@$5 ld.global.L1::evict_last.v4.b32 { $0, $1, $2, $3 }, [ $4 + 0 ];\0A\09@!$7 mov.u32 $0, $6;\0A\09@!$9 mov.u32 $1, $8;\0A\09@!$11 mov.u32 $2, $10;\0A\09@!$13 mov.u32 $3, $12;", "=r,=r,=r,=r,l,b,r,b,r,b,r,b,r,b"(ptr addrspace(1) %116, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true) #6, !dbg !30
  %129 = extractvalue { i32, i32, i32, i32 } %128, 0, !dbg !30
  %130 = extractvalue { i32, i32, i32, i32 } %128, 1, !dbg !30
  %131 = extractvalue { i32, i32, i32, i32 } %128, 2, !dbg !30
  %132 = extractvalue { i32, i32, i32, i32 } %128, 3, !dbg !30
  %133 = bitcast i32 %129 to float, !dbg !30
  %134 = bitcast i32 %130 to float, !dbg !30
  %135 = bitcast i32 %131 to float, !dbg !30
  %136 = bitcast i32 %132 to float, !dbg !30
  %137 = tail call { i32, i32, i32, i32 } asm sideeffect "mov.u32 $0, 0x0;\0A\09mov.u32 $1, 0x0;\0A\09mov.u32 $2, 0x0;\0A\09mov.u32 $3, 0x0;\0A\09@$5 ld.global.L1::evict_last.v4.b32 { $0, $1, $2, $3 }, [ $4 + 0 ];\0A\09@!$7 mov.u32 $0, $6;\0A\09@!$9 mov.u32 $1, $8;\0A\09@!$11 mov.u32 $2, $10;\0A\09@!$13 mov.u32 $3, $12;", "=r,=r,=r,=r,l,b,r,b,r,b,r,b,r,b"(ptr addrspace(1) %117, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true) #6, !dbg !30
  %138 = extractvalue { i32, i32, i32, i32 } %137, 0, !dbg !30
  %139 = extractvalue { i32, i32, i32, i32 } %137, 1, !dbg !30
  %140 = extractvalue { i32, i32, i32, i32 } %137, 2, !dbg !30
  %141 = extractvalue { i32, i32, i32, i32 } %137, 3, !dbg !30
  %142 = bitcast i32 %138 to float, !dbg !30
  %143 = bitcast i32 %139 to float, !dbg !30
  %144 = bitcast i32 %140 to float, !dbg !30
  %145 = bitcast i32 %141 to float, !dbg !30
  %146 = tail call { i32, i32, i32, i32 } asm sideeffect "mov.u32 $0, 0x0;\0A\09mov.u32 $1, 0x0;\0A\09mov.u32 $2, 0x0;\0A\09mov.u32 $3, 0x0;\0A\09@$5 ld.global.L1::evict_last.v4.b32 { $0, $1, $2, $3 }, [ $4 + 0 ];\0A\09@!$7 mov.u32 $0, $6;\0A\09@!$9 mov.u32 $1, $8;\0A\09@!$11 mov.u32 $2, $10;\0A\09@!$13 mov.u32 $3, $12;", "=r,=r,=r,=r,l,b,r,b,r,b,r,b,r,b"(ptr addrspace(1) %118, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true) #6, !dbg !30
  %147 = extractvalue { i32, i32, i32, i32 } %146, 0, !dbg !30
  %148 = extractvalue { i32, i32, i32, i32 } %146, 1, !dbg !30
  %149 = extractvalue { i32, i32, i32, i32 } %146, 2, !dbg !30
  %150 = extractvalue { i32, i32, i32, i32 } %146, 3, !dbg !30
  %151 = bitcast i32 %147 to float, !dbg !30
  %152 = bitcast i32 %148 to float, !dbg !30
  %153 = bitcast i32 %149 to float, !dbg !30
  %154 = bitcast i32 %150 to float, !dbg !30
  %155 = fadd float %64, %124, !dbg !31
  %156 = fadd float %65, %125, !dbg !31
  %157 = fadd float %66, %126, !dbg !31
  %158 = fadd float %67, %127, !dbg !31
  %159 = fadd float %73, %133, !dbg !31
  %160 = fadd float %74, %134, !dbg !31
  %161 = fadd float %75, %135, !dbg !31
  %162 = fadd float %76, %136, !dbg !31
  %163 = fadd float %82, %142, !dbg !31
  %164 = fadd float %83, %143, !dbg !31
  %165 = fadd float %84, %144, !dbg !31
  %166 = fadd float %85, %145, !dbg !31
  %167 = fadd float %91, %151, !dbg !31
  %168 = fadd float %92, %152, !dbg !31
  %169 = fadd float %93, %153, !dbg !31
  %170 = fadd float %94, %154, !dbg !31
  %171 = fadd float %155, 0.000000e+00, !dbg !32
  %172 = fadd float %156, 0.000000e+00, !dbg !32
  %173 = fadd float %157, 0.000000e+00, !dbg !32
  %174 = fadd float %158, 0.000000e+00, !dbg !32
  %175 = fadd float %159, 0.000000e+00, !dbg !32
  %176 = fadd float %160, 0.000000e+00, !dbg !32
  %177 = fadd float %161, 0.000000e+00, !dbg !32
  %178 = fadd float %162, 0.000000e+00, !dbg !32
  %179 = fadd float %163, 0.000000e+00, !dbg !32
  %180 = fadd float %164, 0.000000e+00, !dbg !32
  %181 = fadd float %165, 0.000000e+00, !dbg !32
  %182 = fadd float %166, 0.000000e+00, !dbg !32
  %183 = fadd float %167, 0.000000e+00, !dbg !32
  %184 = fadd float %168, 0.000000e+00, !dbg !32
  %185 = fadd float %169, 0.000000e+00, !dbg !32
  %186 = fadd float %170, 0.000000e+00, !dbg !32
  %187 = fsub float %155, %171, !dbg !36
  %188 = fsub float %156, %172, !dbg !36
  %189 = fsub float %157, %173, !dbg !36
  %190 = fsub float %158, %174, !dbg !36
  %191 = fsub float %159, %175, !dbg !36
  %192 = fsub float %160, %176, !dbg !36
  %193 = fsub float %161, %177, !dbg !36
  %194 = fsub float %162, %178, !dbg !36
  %195 = fsub float %163, %179, !dbg !36
  %196 = fsub float %164, %180, !dbg !36
  %197 = fsub float %165, %181, !dbg !36
  %198 = fsub float %166, %182, !dbg !36
  %199 = fsub float %167, %183, !dbg !36
  %200 = fsub float %168, %184, !dbg !36
  %201 = fsub float %169, %185, !dbg !36
  %202 = fsub float %170, %186, !dbg !36
  %203 = fmul float %155, %187, !dbg !37
  %204 = fmul float %156, %188, !dbg !37
  %205 = fmul float %157, %189, !dbg !37
  %206 = fmul float %158, %190, !dbg !37
  %207 = fmul float %159, %191, !dbg !37
  %208 = fmul float %160, %192, !dbg !37
  %209 = fmul float %161, %193, !dbg !37
  %210 = fmul float %162, %194, !dbg !37
  %211 = fmul float %163, %195, !dbg !37
  %212 = fmul float %164, %196, !dbg !37
  %213 = fmul float %165, %197, !dbg !37
  %214 = fmul float %166, %198, !dbg !37
  %215 = fmul float %167, %199, !dbg !37
  %216 = fmul float %168, %200, !dbg !37
  %217 = fmul float %169, %201, !dbg !37
  %218 = fmul float %170, %202, !dbg !37
  %219 = fadd float %203, 0.000000e+00, !dbg !38
  %220 = fadd float %204, 0.000000e+00, !dbg !38
  %221 = fadd float %205, 0.000000e+00, !dbg !38
  %222 = fadd float %206, 0.000000e+00, !dbg !38
  %223 = fadd float %207, 0.000000e+00, !dbg !38
  %224 = fadd float %208, 0.000000e+00, !dbg !38
  %225 = fadd float %209, 0.000000e+00, !dbg !38
  %226 = fadd float %210, 0.000000e+00, !dbg !38
  %227 = fadd float %211, 0.000000e+00, !dbg !38
  %228 = fadd float %212, 0.000000e+00, !dbg !38
  %229 = fadd float %213, 0.000000e+00, !dbg !38
  %230 = fadd float %214, 0.000000e+00, !dbg !38
  %231 = fadd float %215, 0.000000e+00, !dbg !38
  %232 = fadd float %216, 0.000000e+00, !dbg !38
  %233 = fadd float %217, 0.000000e+00, !dbg !38
  %234 = fadd float %218, 0.000000e+00, !dbg !38
  %235 = fsub float %172, %171, !dbg !39
  %236 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float 1.000000e+00, float 2.000000e+00) #6, !dbg !43
  %237 = fmul float %236, %235, !dbg !44
  %238 = fadd float %171, %237, !dbg !45
  %239 = fadd float %219, %220, !dbg !46
  %240 = fmul float %235, %235, !dbg !47
  %241 = fmul float %236, %240, !dbg !48
  %242 = fadd float %241, %239, !dbg !49
  %243 = fsub float %173, %238, !dbg !39
  %244 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float 1.000000e+00, float 3.000000e+00) #6, !dbg !43
  %245 = fmul float %244, %243, !dbg !44
  %246 = fadd float %238, %245, !dbg !45
  %247 = fadd float %221, %242, !dbg !46
  %248 = fmul float %243, %243, !dbg !47
  %249 = fmul float %248, 2.000000e+00, !dbg !50
  %250 = fmul float %244, %249, !dbg !48
  %251 = fadd float %247, %250, !dbg !49
  %252 = fsub float %174, %246, !dbg !39
  %253 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float 1.000000e+00, float 4.000000e+00) #6, !dbg !43
  %254 = fmul float %253, %252, !dbg !44
  %255 = fadd float %246, %254, !dbg !45
  %256 = fadd float %222, %251, !dbg !46
  %257 = fmul float %252, %252, !dbg !47
  %258 = fmul float %257, 3.000000e+00, !dbg !50
  %259 = fmul float %253, %258, !dbg !48
  %260 = fadd float %256, %259, !dbg !49
  %261 = fsub float %175, %255, !dbg !39
  %262 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float 1.000000e+00, float 5.000000e+00) #6, !dbg !43
  %263 = fmul float %262, %261, !dbg !44
  %264 = fadd float %255, %263, !dbg !45
  %265 = fadd float %223, %260, !dbg !46
  %266 = fmul float %261, %261, !dbg !47
  %267 = fmul float %266, 4.000000e+00, !dbg !50
  %268 = fmul float %262, %267, !dbg !48
  %269 = fadd float %265, %268, !dbg !49
  %270 = fsub float %176, %264, !dbg !39
  %271 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float 1.000000e+00, float 6.000000e+00) #6, !dbg !43
  %272 = fmul float %271, %270, !dbg !44
  %273 = fadd float %264, %272, !dbg !45
  %274 = fadd float %224, %269, !dbg !46
  %275 = fmul float %270, %270, !dbg !47
  %276 = fmul float %275, 5.000000e+00, !dbg !50
  %277 = fmul float %271, %276, !dbg !48
  %278 = fadd float %274, %277, !dbg !49
  %279 = fsub float %177, %273, !dbg !39
  %280 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float 1.000000e+00, float 7.000000e+00) #6, !dbg !43
  %281 = fmul float %280, %279, !dbg !44
  %282 = fadd float %273, %281, !dbg !45
  %283 = fadd float %225, %278, !dbg !46
  %284 = fmul float %279, %279, !dbg !47
  %285 = fmul float %284, 6.000000e+00, !dbg !50
  %286 = fmul float %280, %285, !dbg !48
  %287 = fadd float %283, %286, !dbg !49
  %288 = fsub float %178, %282, !dbg !39
  %289 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float 1.000000e+00, float 8.000000e+00) #6, !dbg !43
  %290 = fmul float %289, %288, !dbg !44
  %291 = fadd float %282, %290, !dbg !45
  %292 = fadd float %226, %287, !dbg !46
  %293 = fmul float %288, %288, !dbg !47
  %294 = fmul float %293, 7.000000e+00, !dbg !50
  %295 = fmul float %289, %294, !dbg !48
  %296 = fadd float %292, %295, !dbg !49
  %297 = fsub float %180, %179, !dbg !39
  %298 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float 1.000000e+00, float 2.000000e+00) #6, !dbg !43
  %299 = fmul float %297, %298, !dbg !44
  %300 = fadd float %179, %299, !dbg !45
  %301 = fadd float %227, %228, !dbg !46
  %302 = fmul float %297, %297, !dbg !47
  %303 = fmul float %302, %298, !dbg !48
  %304 = fadd float %301, %303, !dbg !49
  %305 = fsub float %181, %300, !dbg !39
  %306 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float 1.000000e+00, float 3.000000e+00) #6, !dbg !43
  %307 = fmul float %306, %305, !dbg !44
  %308 = fadd float %300, %307, !dbg !45
  %309 = fadd float %229, %304, !dbg !46
  %310 = fmul float %305, %305, !dbg !47
  %311 = fmul float %310, 2.000000e+00, !dbg !50
  %312 = fmul float %306, %311, !dbg !48
  %313 = fadd float %309, %312, !dbg !49
  %314 = fsub float %182, %308, !dbg !39
  %315 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float 1.000000e+00, float 4.000000e+00) #6, !dbg !43
  %316 = fmul float %315, %314, !dbg !44
  %317 = fadd float %308, %316, !dbg !45
  %318 = fadd float %230, %313, !dbg !46
  %319 = fmul float %314, %314, !dbg !47
  %320 = fmul float %319, 3.000000e+00, !dbg !50
  %321 = fmul float %315, %320, !dbg !48
  %322 = fadd float %318, %321, !dbg !49
  %323 = fsub float %183, %317, !dbg !39
  %324 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float 1.000000e+00, float 5.000000e+00) #6, !dbg !43
  %325 = fmul float %324, %323, !dbg !44
  %326 = fadd float %317, %325, !dbg !45
  %327 = fadd float %231, %322, !dbg !46
  %328 = fmul float %323, %323, !dbg !47
  %329 = fmul float %328, 4.000000e+00, !dbg !50
  %330 = fmul float %324, %329, !dbg !48
  %331 = fadd float %327, %330, !dbg !49
  %332 = fsub float %184, %326, !dbg !39
  %333 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float 1.000000e+00, float 6.000000e+00) #6, !dbg !43
  %334 = fmul float %333, %332, !dbg !44
  %335 = fadd float %326, %334, !dbg !45
  %336 = fadd float %232, %331, !dbg !46
  %337 = fmul float %332, %332, !dbg !47
  %338 = fmul float %337, 5.000000e+00, !dbg !50
  %339 = fmul float %333, %338, !dbg !48
  %340 = fadd float %336, %339, !dbg !49
  %341 = fsub float %185, %335, !dbg !39
  %342 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float 1.000000e+00, float 7.000000e+00) #6, !dbg !43
  %343 = fmul float %342, %341, !dbg !44
  %344 = fadd float %335, %343, !dbg !45
  %345 = fadd float %233, %340, !dbg !46
  %346 = fmul float %341, %341, !dbg !47
  %347 = fmul float %346, 6.000000e+00, !dbg !50
  %348 = fmul float %342, %347, !dbg !48
  %349 = fadd float %345, %348, !dbg !49
  %350 = fsub float %186, %344, !dbg !39
  %351 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float 1.000000e+00, float 8.000000e+00) #6, !dbg !43
  %352 = fmul float %351, %350, !dbg !44
  %353 = fadd float %344, %352, !dbg !45
  %354 = fadd float %234, %349, !dbg !46
  %355 = fmul float %350, %350, !dbg !47
  %356 = fmul float %355, 7.000000e+00, !dbg !50
  %357 = fmul float %351, %356, !dbg !48
  %358 = fadd float %354, %357, !dbg !49
  %359 = bitcast float %291 to i32, !dbg !51
  %360 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %359, i32 16, i32 31), !dbg !51
  %361 = bitcast i32 %360 to float, !dbg !51
  %362 = bitcast float %296 to i32, !dbg !51
  %363 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %362, i32 16, i32 31), !dbg !51
  %364 = bitcast i32 %363 to float, !dbg !51
  %365 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 1090519040, i32 16, i32 31), !dbg !51
  %366 = bitcast i32 %365 to float, !dbg !51
  %367 = fsub float %361, %291, !dbg !39
  %368 = fadd float %366, 8.000000e+00, !dbg !53
  %369 = fcmp oeq float %368, 0.000000e+00, !dbg !54
  %370 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %366, float %368) #6, !dbg !43
  %371 = select i1 %369, float 0.000000e+00, float %370, !dbg !55
  %372 = fmul float %371, %367, !dbg !44
  %373 = fadd float %291, %372, !dbg !45
  %374 = fadd float %296, %364, !dbg !46
  %375 = fmul float %367, %367, !dbg !47
  %376 = fmul float %375, 8.000000e+00, !dbg !50
  %377 = fmul float %371, %376, !dbg !48
  %378 = fadd float %374, %377, !dbg !49
  %379 = bitcast float %373 to i32, !dbg !51
  %380 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %379, i32 8, i32 31), !dbg !51
  %381 = bitcast i32 %380 to float, !dbg !51
  %382 = bitcast float %378 to i32, !dbg !51
  %383 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %382, i32 8, i32 31), !dbg !51
  %384 = bitcast i32 %383 to float, !dbg !51
  %385 = bitcast float %368 to i32, !dbg !51
  %386 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %385, i32 8, i32 31), !dbg !51
  %387 = bitcast i32 %386 to float, !dbg !51
  %388 = fsub float %381, %373, !dbg !39
  %389 = fadd float %368, %387, !dbg !53
  %390 = fcmp oeq float %389, 0.000000e+00, !dbg !54
  %391 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %387, float %389) #6, !dbg !43
  %392 = select i1 %390, float 0.000000e+00, float %391, !dbg !55
  %393 = fmul float %392, %388, !dbg !44
  %394 = fadd float %373, %393, !dbg !45
  %395 = fadd float %378, %384, !dbg !46
  %396 = fmul float %388, %388, !dbg !47
  %397 = fmul float %368, %396, !dbg !50
  %398 = fmul float %392, %397, !dbg !48
  %399 = fadd float %395, %398, !dbg !49
  %400 = bitcast float %394 to i32, !dbg !51
  %401 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %400, i32 4, i32 31), !dbg !51
  %402 = bitcast i32 %401 to float, !dbg !51
  %403 = bitcast float %399 to i32, !dbg !51
  %404 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %403, i32 4, i32 31), !dbg !51
  %405 = bitcast i32 %404 to float, !dbg !51
  %406 = bitcast float %389 to i32, !dbg !51
  %407 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %406, i32 4, i32 31), !dbg !51
  %408 = bitcast i32 %407 to float, !dbg !51
  %409 = fsub float %402, %394, !dbg !39
  %410 = fadd float %389, %408, !dbg !53
  %411 = fcmp oeq float %410, 0.000000e+00, !dbg !54
  %412 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %408, float %410) #6, !dbg !43
  %413 = select i1 %411, float 0.000000e+00, float %412, !dbg !55
  %414 = fmul float %409, %413, !dbg !44
  %415 = fadd float %394, %414, !dbg !45
  %416 = fadd float %399, %405, !dbg !46
  %417 = fmul float %409, %409, !dbg !47
  %418 = fmul float %389, %417, !dbg !50
  %419 = fmul float %413, %418, !dbg !48
  %420 = fadd float %416, %419, !dbg !49
  %421 = bitcast float %415 to i32, !dbg !51
  %422 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %421, i32 2, i32 31), !dbg !51
  %423 = bitcast i32 %422 to float, !dbg !51
  %424 = bitcast float %420 to i32, !dbg !51
  %425 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %424, i32 2, i32 31), !dbg !51
  %426 = bitcast i32 %425 to float, !dbg !51
  %427 = bitcast float %410 to i32, !dbg !51
  %428 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %427, i32 2, i32 31), !dbg !51
  %429 = bitcast i32 %428 to float, !dbg !51
  %430 = fsub float %423, %415, !dbg !39
  %431 = fadd float %410, %429, !dbg !53
  %432 = fcmp oeq float %431, 0.000000e+00, !dbg !54
  %433 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %429, float %431) #6, !dbg !43
  %434 = select i1 %432, float 0.000000e+00, float %433, !dbg !55
  %435 = fmul float %430, %434, !dbg !44
  %436 = fadd float %415, %435, !dbg !45
  %437 = fadd float %420, %426, !dbg !46
  %438 = fmul float %430, %430, !dbg !47
  %439 = fmul float %410, %438, !dbg !50
  %440 = fmul float %434, %439, !dbg !48
  %441 = fadd float %437, %440, !dbg !49
  %442 = bitcast float %436 to i32, !dbg !51
  %443 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %442, i32 1, i32 31), !dbg !51
  %444 = bitcast float %441 to i32, !dbg !51
  %445 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %444, i32 1, i32 31), !dbg !51
  %446 = bitcast float %431 to i32, !dbg !51
  %447 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %446, i32 1, i32 31), !dbg !51
  %448 = bitcast i32 %447 to float, !dbg !51
  %449 = fadd float %431, %448, !dbg !53
  %450 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %448, float %449) #6, !dbg !43
  %451 = bitcast float %353 to i32, !dbg !51
  %452 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %451, i32 16, i32 31), !dbg !51
  %453 = bitcast i32 %452 to float, !dbg !51
  %454 = bitcast float %358 to i32, !dbg !51
  %455 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %454, i32 16, i32 31), !dbg !51
  %456 = bitcast i32 %455 to float, !dbg !51
  %457 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 1090519040, i32 16, i32 31), !dbg !51
  %458 = bitcast i32 %457 to float, !dbg !51
  %459 = fsub float %453, %353, !dbg !39
  %460 = fadd float %458, 8.000000e+00, !dbg !53
  %461 = fcmp oeq float %460, 0.000000e+00, !dbg !54
  %462 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %458, float %460) #6, !dbg !43
  %463 = select i1 %461, float 0.000000e+00, float %462, !dbg !55
  %464 = fmul float %459, %463, !dbg !44
  %465 = fadd float %353, %464, !dbg !45
  %466 = fadd float %358, %456, !dbg !46
  %467 = fmul float %459, %459, !dbg !47
  %468 = fmul float %467, 8.000000e+00, !dbg !50
  %469 = fmul float %468, %463, !dbg !48
  %470 = fadd float %466, %469, !dbg !49
  %471 = bitcast float %465 to i32, !dbg !51
  %472 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %471, i32 8, i32 31), !dbg !51
  %473 = bitcast i32 %472 to float, !dbg !51
  %474 = bitcast float %470 to i32, !dbg !51
  %475 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %474, i32 8, i32 31), !dbg !51
  %476 = bitcast i32 %475 to float, !dbg !51
  %477 = bitcast float %460 to i32, !dbg !51
  %478 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %477, i32 8, i32 31), !dbg !51
  %479 = bitcast i32 %478 to float, !dbg !51
  %480 = fsub float %473, %465, !dbg !39
  %481 = fadd float %460, %479, !dbg !53
  %482 = fcmp oeq float %481, 0.000000e+00, !dbg !54
  %483 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %479, float %481) #6, !dbg !43
  %484 = select i1 %482, float 0.000000e+00, float %483, !dbg !55
  %485 = fmul float %480, %484, !dbg !44
  %486 = fadd float %465, %485, !dbg !45
  %487 = fadd float %470, %476, !dbg !46
  %488 = fmul float %480, %480, !dbg !47
  %489 = fmul float %460, %488, !dbg !50
  %490 = fmul float %484, %489, !dbg !48
  %491 = fadd float %487, %490, !dbg !49
  %492 = bitcast float %486 to i32, !dbg !51
  %493 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %492, i32 4, i32 31), !dbg !51
  %494 = bitcast i32 %493 to float, !dbg !51
  %495 = bitcast float %491 to i32, !dbg !51
  %496 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %495, i32 4, i32 31), !dbg !51
  %497 = bitcast i32 %496 to float, !dbg !51
  %498 = bitcast float %481 to i32, !dbg !51
  %499 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %498, i32 4, i32 31), !dbg !51
  %500 = bitcast i32 %499 to float, !dbg !51
  %501 = fsub float %494, %486, !dbg !39
  %502 = fadd float %481, %500, !dbg !53
  %503 = fcmp oeq float %502, 0.000000e+00, !dbg !54
  %504 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %500, float %502) #6, !dbg !43
  %505 = select i1 %503, float 0.000000e+00, float %504, !dbg !55
  %506 = fmul float %501, %505, !dbg !44
  %507 = fadd float %486, %506, !dbg !45
  %508 = fadd float %491, %497, !dbg !46
  %509 = fmul float %501, %501, !dbg !47
  %510 = fmul float %481, %509, !dbg !50
  %511 = fmul float %505, %510, !dbg !48
  %512 = fadd float %508, %511, !dbg !49
  %513 = bitcast float %507 to i32, !dbg !51
  %514 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %513, i32 2, i32 31), !dbg !51
  %515 = bitcast i32 %514 to float, !dbg !51
  %516 = bitcast float %512 to i32, !dbg !51
  %517 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %516, i32 2, i32 31), !dbg !51
  %518 = bitcast i32 %517 to float, !dbg !51
  %519 = bitcast float %502 to i32, !dbg !51
  %520 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %519, i32 2, i32 31), !dbg !51
  %521 = bitcast i32 %520 to float, !dbg !51
  %522 = fsub float %515, %507, !dbg !39
  %523 = fadd float %502, %521, !dbg !53
  %524 = fcmp oeq float %523, 0.000000e+00, !dbg !54
  %525 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %521, float %523) #6, !dbg !43
  %526 = select i1 %524, float 0.000000e+00, float %525, !dbg !55
  %527 = fmul float %522, %526, !dbg !44
  %528 = fadd float %507, %527, !dbg !45
  %529 = fadd float %512, %518, !dbg !46
  %530 = fmul float %522, %522, !dbg !47
  %531 = fmul float %502, %530, !dbg !50
  %532 = fmul float %526, %531, !dbg !48
  %533 = fadd float %529, %532, !dbg !49
  %534 = bitcast float %528 to i32, !dbg !51
  %535 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %534, i32 1, i32 31), !dbg !51
  %536 = bitcast float %533 to i32, !dbg !51
  %537 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %536, i32 1, i32 31), !dbg !51
  %538 = bitcast float %523 to i32, !dbg !51
  %539 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %538, i32 1, i32 31), !dbg !51
  %540 = bitcast i32 %539 to float, !dbg !51
  %541 = fadd float %523, %540, !dbg !53
  %542 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %540, float %541) #6, !dbg !43
  %543 = tail call { i32, i32, i32, i32 } asm sideeffect "mov.u32 $0, 0x0;\0A\09mov.u32 $1, 0x0;\0A\09mov.u32 $2, 0x0;\0A\09mov.u32 $3, 0x0;\0A\09@$5 ld.global.L1::evict_last.v4.b32 { $0, $1, $2, $3 }, [ $4 + 0 ];\0A\09@!$7 mov.u32 $0, $6;\0A\09@!$9 mov.u32 $1, $8;\0A\09@!$11 mov.u32 $2, $10;\0A\09@!$13 mov.u32 $3, $12;", "=r,=r,=r,=r,l,b,r,b,r,b,r,b,r,b"(ptr addrspace(1) %52, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true) #6, !dbg !56
  %544 = tail call { i32, i32, i32, i32 } asm sideeffect "mov.u32 $0, 0x0;\0A\09mov.u32 $1, 0x0;\0A\09mov.u32 $2, 0x0;\0A\09mov.u32 $3, 0x0;\0A\09@$5 ld.global.L1::evict_last.v4.b32 { $0, $1, $2, $3 }, [ $4 + 0 ];\0A\09@!$7 mov.u32 $0, $6;\0A\09@!$9 mov.u32 $1, $8;\0A\09@!$11 mov.u32 $2, $10;\0A\09@!$13 mov.u32 $3, $12;", "=r,=r,=r,=r,l,b,r,b,r,b,r,b,r,b"(ptr addrspace(1) %54, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true) #6, !dbg !56
  %545 = tail call { i32, i32, i32, i32 } asm sideeffect "mov.u32 $0, 0x0;\0A\09mov.u32 $1, 0x0;\0A\09mov.u32 $2, 0x0;\0A\09mov.u32 $3, 0x0;\0A\09@$5 ld.global.L1::evict_last.v4.b32 { $0, $1, $2, $3 }, [ $4 + 0 ];\0A\09@!$7 mov.u32 $0, $6;\0A\09@!$9 mov.u32 $1, $8;\0A\09@!$11 mov.u32 $2, $10;\0A\09@!$13 mov.u32 $3, $12;", "=r,=r,=r,=r,l,b,r,b,r,b,r,b,r,b"(ptr addrspace(1) %56, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true) #6, !dbg !56
  %546 = tail call { i32, i32, i32, i32 } asm sideeffect "mov.u32 $0, 0x0;\0A\09mov.u32 $1, 0x0;\0A\09mov.u32 $2, 0x0;\0A\09mov.u32 $3, 0x0;\0A\09@$5 ld.global.L1::evict_last.v4.b32 { $0, $1, $2, $3 }, [ $4 + 0 ];\0A\09@!$7 mov.u32 $0, $6;\0A\09@!$9 mov.u32 $1, $8;\0A\09@!$11 mov.u32 $2, $10;\0A\09@!$13 mov.u32 $3, $12;", "=r,=r,=r,=r,l,b,r,b,r,b,r,b,r,b"(ptr addrspace(1) %58, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true) #6, !dbg !56
  %547 = zext nneg i32 %urem to i64, !dbg !57
  %548 = getelementptr float, ptr addrspace(1) %3, i64 %547, !dbg !57
  %549 = tail call i32 asm sideeffect "mov.u32 $0, 0x0;\0A\09@$2 ld.global.L1::evict_last.b32 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u32 $0, $3;", "=r,l,b,r,b"(ptr addrspace(1) %548, i1 true, i32 0, i1 true) #6, !dbg !58
  br i1 %100, label %550, label %551, !dbg !59

550:                                              ; preds = %102
  tail call void @__assertfail(ptr nonnull @assertMessage_1, ptr nonnull @assertFile_1, i32 883, ptr nonnull @assertFunc_1, i64 1), !dbg !59
  br label %551, !dbg !59

551:                                              ; preds = %550, %102
  %552 = bitcast i32 %537 to float, !dbg !51
  %553 = fadd float %533, %552, !dbg !46
  %554 = bitcast i32 %535 to float, !dbg !51
  %555 = fsub float %554, %528, !dbg !39
  %556 = fmul float %555, %555, !dbg !47
  %557 = fmul float %523, %556, !dbg !50
  %558 = fcmp oeq float %541, 0.000000e+00, !dbg !54
  %559 = select i1 %558, float 0.000000e+00, float %542, !dbg !55
  %560 = fmul float %559, %557, !dbg !48
  %561 = fadd float %553, %560, !dbg !49
  %562 = bitcast i32 %445 to float, !dbg !51
  %563 = fadd float %441, %562, !dbg !46
  %564 = bitcast i32 %443 to float, !dbg !51
  %565 = fsub float %564, %436, !dbg !39
  %566 = fmul float %565, %565, !dbg !47
  %567 = fmul float %431, %566, !dbg !50
  %568 = fcmp oeq float %449, 0.000000e+00, !dbg !54
  %569 = select i1 %568, float 0.000000e+00, float %450, !dbg !55
  %570 = fmul float %569, %567, !dbg !48
  %571 = fadd float %563, %570, !dbg !49
  %572 = tail call { i32, i32, i32, i32 } asm sideeffect "mov.u32 $0, 0x0;\0A\09mov.u32 $1, 0x0;\0A\09mov.u32 $2, 0x0;\0A\09mov.u32 $3, 0x0;\0A\09@$5 ld.global.L1::evict_first.v4.b32 { $0, $1, $2, $3 }, [ $4 + 0 ];\0A\09@!$7 mov.u32 $0, $6;\0A\09@!$9 mov.u32 $1, $8;\0A\09@!$11 mov.u32 $2, $10;\0A\09@!$13 mov.u32 $3, $12;", "=r,=r,=r,=r,l,b,r,b,r,b,r,b,r,b"(ptr addrspace(1) %115, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true) #6, !dbg !60
  %573 = tail call { i32, i32, i32, i32 } asm sideeffect "mov.u32 $0, 0x0;\0A\09mov.u32 $1, 0x0;\0A\09mov.u32 $2, 0x0;\0A\09mov.u32 $3, 0x0;\0A\09@$5 ld.global.L1::evict_first.v4.b32 { $0, $1, $2, $3 }, [ $4 + 0 ];\0A\09@!$7 mov.u32 $0, $6;\0A\09@!$9 mov.u32 $1, $8;\0A\09@!$11 mov.u32 $2, $10;\0A\09@!$13 mov.u32 $3, $12;", "=r,=r,=r,=r,l,b,r,b,r,b,r,b,r,b"(ptr addrspace(1) %116, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true) #6, !dbg !60
  %574 = tail call { i32, i32, i32, i32 } asm sideeffect "mov.u32 $0, 0x0;\0A\09mov.u32 $1, 0x0;\0A\09mov.u32 $2, 0x0;\0A\09mov.u32 $3, 0x0;\0A\09@$5 ld.global.L1::evict_first.v4.b32 { $0, $1, $2, $3 }, [ $4 + 0 ];\0A\09@!$7 mov.u32 $0, $6;\0A\09@!$9 mov.u32 $1, $8;\0A\09@!$11 mov.u32 $2, $10;\0A\09@!$13 mov.u32 $3, $12;", "=r,=r,=r,=r,l,b,r,b,r,b,r,b,r,b"(ptr addrspace(1) %117, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true) #6, !dbg !60
  %575 = tail call { i32, i32, i32, i32 } asm sideeffect "mov.u32 $0, 0x0;\0A\09mov.u32 $1, 0x0;\0A\09mov.u32 $2, 0x0;\0A\09mov.u32 $3, 0x0;\0A\09@$5 ld.global.L1::evict_first.v4.b32 { $0, $1, $2, $3 }, [ $4 + 0 ];\0A\09@!$7 mov.u32 $0, $6;\0A\09@!$9 mov.u32 $1, $8;\0A\09@!$11 mov.u32 $2, $10;\0A\09@!$13 mov.u32 $3, $12;", "=r,=r,=r,=r,l,b,r,b,r,b,r,b,r,b"(ptr addrspace(1) %118, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true, i32 0, i1 true) #6, !dbg !60
  %576 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %571, float 2.560000e+02) #6, !dbg !61
  %577 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %571, float 2.560000e+02) #6, !dbg !61
  %578 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %571, float 2.560000e+02) #6, !dbg !61
  %579 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %571, float 2.560000e+02) #6, !dbg !61
  %580 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %571, float 2.560000e+02) #6, !dbg !61
  %581 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %571, float 2.560000e+02) #6, !dbg !61
  %582 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %571, float 2.560000e+02) #6, !dbg !61
  %583 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %571, float 2.560000e+02) #6, !dbg !61
  %584 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %561, float 2.560000e+02) #6, !dbg !61
  %585 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %561, float 2.560000e+02) #6, !dbg !61
  %586 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %561, float 2.560000e+02) #6, !dbg !61
  %587 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %561, float 2.560000e+02) #6, !dbg !61
  %588 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %561, float 2.560000e+02) #6, !dbg !61
  %589 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %561, float 2.560000e+02) #6, !dbg !61
  %590 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %561, float 2.560000e+02) #6, !dbg !61
  %591 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %561, float 2.560000e+02) #6, !dbg !61
  %592 = fadd float %576, 0x3EE4F8B580000000, !dbg !62
  %593 = fadd float %584, 0x3EE4F8B580000000, !dbg !62
  %594 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #6, !dbg !63
  %.not.i = icmp eq i32 %594, 0, !dbg !63
  br i1 %.not.i, label %597, label %595, !dbg !63

595:                                              ; preds = %551
  %596 = tail call float @llvm.nvvm.rsqrt.approx.ftz.f(float %592), !dbg !63
  br label %__nv_rsqrtf.exit, !dbg !63

597:                                              ; preds = %551
  %598 = tail call float @llvm.nvvm.rsqrt.approx.f(float %592), !dbg !63
  br label %__nv_rsqrtf.exit, !dbg !63

__nv_rsqrtf.exit:                                 ; preds = %595, %597
  %.0.i = phi float [ %596, %595 ], [ %598, %597 ], !dbg !63
  %599 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #6, !dbg !63
  %600 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #6, !dbg !63
  %601 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #6, !dbg !63
  %602 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #6, !dbg !63
  %603 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #6, !dbg !63
  %604 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #6, !dbg !63
  %605 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #6, !dbg !63
  %606 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #6, !dbg !63
  %.not.i22 = icmp eq i32 %606, 0, !dbg !63
  br i1 %.not.i22, label %609, label %607, !dbg !63

607:                                              ; preds = %__nv_rsqrtf.exit
  %608 = tail call float @llvm.nvvm.rsqrt.approx.ftz.f(float %593), !dbg !63
  br label %__nv_rsqrtf.exit24, !dbg !63

609:                                              ; preds = %__nv_rsqrtf.exit
  %610 = tail call float @llvm.nvvm.rsqrt.approx.f(float %593), !dbg !63
  br label %__nv_rsqrtf.exit24, !dbg !63

__nv_rsqrtf.exit24:                               ; preds = %607, %609
  %.0.i23 = phi float [ %608, %607 ], [ %610, %609 ], !dbg !63
  %611 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #6, !dbg !63
  %612 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #6, !dbg !63
  %613 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #6, !dbg !63
  %614 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #6, !dbg !63
  %615 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #6, !dbg !63
  %616 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #6, !dbg !63
  %617 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #6, !dbg !63
  %618 = extractvalue { i32, i32, i32, i32 } %575, 3, !dbg !60
  %619 = bitcast i32 %618 to float, !dbg !60
  %620 = extractvalue { i32, i32, i32, i32 } %546, 3, !dbg !56
  %621 = bitcast i32 %620 to float, !dbg !56
  %622 = fadd float %621, %619, !dbg !64
  %623 = fmul float %555, %559, !dbg !44
  %624 = fadd float %528, %623, !dbg !45
  %625 = fsub float %622, %624, !dbg !65
  %626 = extractvalue { i32, i32, i32, i32 } %575, 2, !dbg !60
  %627 = bitcast i32 %626 to float, !dbg !60
  %628 = extractvalue { i32, i32, i32, i32 } %546, 2, !dbg !56
  %629 = bitcast i32 %628 to float, !dbg !56
  %630 = fadd float %629, %627, !dbg !64
  %631 = fsub float %630, %624, !dbg !65
  %632 = extractvalue { i32, i32, i32, i32 } %575, 1, !dbg !60
  %633 = bitcast i32 %632 to float, !dbg !60
  %634 = extractvalue { i32, i32, i32, i32 } %546, 1, !dbg !56
  %635 = bitcast i32 %634 to float, !dbg !56
  %636 = fadd float %635, %633, !dbg !64
  %637 = fsub float %636, %624, !dbg !65
  %638 = extractvalue { i32, i32, i32, i32 } %575, 0, !dbg !60
  %639 = bitcast i32 %638 to float, !dbg !60
  %640 = extractvalue { i32, i32, i32, i32 } %546, 0, !dbg !56
  %641 = bitcast i32 %640 to float, !dbg !56
  %642 = fadd float %641, %639, !dbg !64
  %643 = fsub float %642, %624, !dbg !65
  %644 = extractvalue { i32, i32, i32, i32 } %574, 3, !dbg !60
  %645 = bitcast i32 %644 to float, !dbg !60
  %646 = extractvalue { i32, i32, i32, i32 } %545, 3, !dbg !56
  %647 = bitcast i32 %646 to float, !dbg !56
  %648 = fadd float %647, %645, !dbg !64
  %649 = fsub float %648, %624, !dbg !65
  %650 = extractvalue { i32, i32, i32, i32 } %574, 2, !dbg !60
  %651 = bitcast i32 %650 to float, !dbg !60
  %652 = extractvalue { i32, i32, i32, i32 } %545, 2, !dbg !56
  %653 = bitcast i32 %652 to float, !dbg !56
  %654 = fadd float %653, %651, !dbg !64
  %655 = fsub float %654, %624, !dbg !65
  %656 = extractvalue { i32, i32, i32, i32 } %574, 1, !dbg !60
  %657 = bitcast i32 %656 to float, !dbg !60
  %658 = extractvalue { i32, i32, i32, i32 } %545, 1, !dbg !56
  %659 = bitcast i32 %658 to float, !dbg !56
  %660 = fadd float %659, %657, !dbg !64
  %661 = fsub float %660, %624, !dbg !65
  %662 = extractvalue { i32, i32, i32, i32 } %574, 0, !dbg !60
  %663 = bitcast i32 %662 to float, !dbg !60
  %664 = extractvalue { i32, i32, i32, i32 } %545, 0, !dbg !56
  %665 = bitcast i32 %664 to float, !dbg !56
  %666 = fadd float %665, %663, !dbg !64
  %667 = fsub float %666, %624, !dbg !65
  %668 = extractvalue { i32, i32, i32, i32 } %573, 3, !dbg !60
  %669 = bitcast i32 %668 to float, !dbg !60
  %670 = extractvalue { i32, i32, i32, i32 } %544, 3, !dbg !56
  %671 = bitcast i32 %670 to float, !dbg !56
  %672 = fadd float %671, %669, !dbg !64
  %673 = fmul float %565, %569, !dbg !44
  %674 = fadd float %436, %673, !dbg !45
  %675 = fsub float %672, %674, !dbg !65
  %676 = extractvalue { i32, i32, i32, i32 } %573, 2, !dbg !60
  %677 = bitcast i32 %676 to float, !dbg !60
  %678 = extractvalue { i32, i32, i32, i32 } %544, 2, !dbg !56
  %679 = bitcast i32 %678 to float, !dbg !56
  %680 = fadd float %679, %677, !dbg !64
  %681 = fsub float %680, %674, !dbg !65
  %682 = extractvalue { i32, i32, i32, i32 } %573, 1, !dbg !60
  %683 = bitcast i32 %682 to float, !dbg !60
  %684 = extractvalue { i32, i32, i32, i32 } %544, 1, !dbg !56
  %685 = bitcast i32 %684 to float, !dbg !56
  %686 = fadd float %685, %683, !dbg !64
  %687 = fsub float %686, %674, !dbg !65
  %688 = extractvalue { i32, i32, i32, i32 } %573, 0, !dbg !60
  %689 = bitcast i32 %688 to float, !dbg !60
  %690 = extractvalue { i32, i32, i32, i32 } %544, 0, !dbg !56
  %691 = bitcast i32 %690 to float, !dbg !56
  %692 = fadd float %691, %689, !dbg !64
  %693 = fsub float %692, %674, !dbg !65
  %694 = extractvalue { i32, i32, i32, i32 } %572, 3, !dbg !60
  %695 = bitcast i32 %694 to float, !dbg !60
  %696 = extractvalue { i32, i32, i32, i32 } %543, 3, !dbg !56
  %697 = bitcast i32 %696 to float, !dbg !56
  %698 = fadd float %697, %695, !dbg !64
  %699 = fsub float %698, %674, !dbg !65
  %700 = extractvalue { i32, i32, i32, i32 } %572, 2, !dbg !60
  %701 = bitcast i32 %700 to float, !dbg !60
  %702 = extractvalue { i32, i32, i32, i32 } %543, 2, !dbg !56
  %703 = bitcast i32 %702 to float, !dbg !56
  %704 = fadd float %703, %701, !dbg !64
  %705 = fsub float %704, %674, !dbg !65
  %706 = extractvalue { i32, i32, i32, i32 } %572, 1, !dbg !60
  %707 = bitcast i32 %706 to float, !dbg !60
  %708 = extractvalue { i32, i32, i32, i32 } %543, 1, !dbg !56
  %709 = bitcast i32 %708 to float, !dbg !56
  %710 = fadd float %709, %707, !dbg !64
  %711 = fsub float %710, %674, !dbg !65
  %712 = extractvalue { i32, i32, i32, i32 } %572, 0, !dbg !60
  %713 = bitcast i32 %712 to float, !dbg !60
  %714 = extractvalue { i32, i32, i32, i32 } %543, 0, !dbg !56
  %715 = bitcast i32 %714 to float, !dbg !56
  %716 = fadd float %715, %713, !dbg !64
  %717 = fsub float %716, %674, !dbg !65
  %718 = fmul float %717, %.0.i, !dbg !66
  %719 = fmul float %711, %.0.i, !dbg !66
  %720 = fmul float %705, %.0.i, !dbg !66
  %721 = fmul float %699, %.0.i, !dbg !66
  %722 = fmul float %693, %.0.i, !dbg !66
  %723 = fmul float %687, %.0.i, !dbg !66
  %724 = fmul float %681, %.0.i, !dbg !66
  %725 = fmul float %675, %.0.i, !dbg !66
  %726 = fmul float %667, %.0.i23, !dbg !66
  %727 = fmul float %661, %.0.i23, !dbg !66
  %728 = fmul float %655, %.0.i23, !dbg !66
  %729 = fmul float %649, %.0.i23, !dbg !66
  %730 = fmul float %643, %.0.i23, !dbg !66
  %731 = fmul float %637, %.0.i23, !dbg !66
  %732 = fmul float %631, %.0.i23, !dbg !66
  %733 = fmul float %625, %.0.i23, !dbg !66
  %734 = getelementptr float, ptr addrspace(3) @global_smem, i64 %547, !dbg !67
  store i32 %549, ptr addrspace(3) %734, align 4, !dbg !67
  tail call void @llvm.nvvm.barrier0(), !dbg !67
  %735 = getelementptr float, ptr addrspace(3) @global_smem, i64 %109, !dbg !67
  %736 = load float, ptr addrspace(3) %735, align 32, !dbg !67
  %737 = getelementptr inbounds <8 x float>, ptr addrspace(3) %735, i64 0, i64 1, !dbg !67
  %738 = load float, ptr addrspace(3) %737, align 4, !dbg !67
  %739 = getelementptr inbounds <8 x float>, ptr addrspace(3) %735, i64 0, i64 2, !dbg !67
  %740 = load float, ptr addrspace(3) %739, align 8, !dbg !67
  %741 = getelementptr inbounds <8 x float>, ptr addrspace(3) %735, i64 0, i64 3, !dbg !67
  %742 = load float, ptr addrspace(3) %741, align 4, !dbg !67
  %743 = getelementptr inbounds <8 x float>, ptr addrspace(3) %735, i64 0, i64 4, !dbg !67
  %744 = load float, ptr addrspace(3) %743, align 16, !dbg !67
  %745 = getelementptr inbounds <8 x float>, ptr addrspace(3) %735, i64 0, i64 5, !dbg !67
  %746 = load float, ptr addrspace(3) %745, align 4, !dbg !67
  %747 = getelementptr inbounds <8 x float>, ptr addrspace(3) %735, i64 0, i64 6, !dbg !67
  %748 = load float, ptr addrspace(3) %747, align 8, !dbg !67
  %749 = getelementptr inbounds <8 x float>, ptr addrspace(3) %735, i64 0, i64 7, !dbg !67
  %750 = load float, ptr addrspace(3) %749, align 4, !dbg !67
  %751 = fmul float %718, %736, !dbg !67
  %752 = fmul float %719, %738, !dbg !67
  %753 = fmul float %720, %740, !dbg !67
  %754 = fmul float %721, %742, !dbg !67
  %755 = fmul float %722, %744, !dbg !67
  %756 = fmul float %723, %746, !dbg !67
  %757 = fmul float %724, %748, !dbg !67
  %758 = fmul float %725, %750, !dbg !67
  %759 = fmul float %726, %736, !dbg !67
  %760 = fmul float %727, %738, !dbg !67
  %761 = fmul float %728, %740, !dbg !67
  %762 = fmul float %729, %742, !dbg !67
  %763 = fmul float %730, %744, !dbg !67
  %764 = fmul float %731, %746, !dbg !67
  %765 = fmul float %732, %748, !dbg !67
  %766 = fmul float %733, %750, !dbg !67
  %767 = shl i32 %17, 8, !dbg !68
  %768 = shl i32 %18, 8, !dbg !68
  %769 = or i32 %767, %13, !dbg !69
  %770 = or i32 %768, %13, !dbg !69
  %771 = sext i32 %769 to i64, !dbg !70
  %772 = getelementptr i16, ptr addrspace(1) %4, i64 %771, !dbg !70
  %773 = sext i32 %770 to i64, !dbg !70
  %774 = getelementptr i16, ptr addrspace(1) %4, i64 %773, !dbg !70
  %775 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %751) #6, !dbg !71
  %776 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %752) #6, !dbg !71
  %777 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %753) #6, !dbg !71
  %778 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %754) #6, !dbg !71
  %779 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %755) #6, !dbg !71
  %780 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %756) #6, !dbg !71
  %781 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %757) #6, !dbg !71
  %782 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %758) #6, !dbg !71
  %783 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %759) #6, !dbg !71
  %784 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %760) #6, !dbg !71
  %785 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %761) #6, !dbg !71
  %786 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %762) #6, !dbg !71
  %787 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %763) #6, !dbg !71
  %788 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %764) #6, !dbg !71
  %789 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %765) #6, !dbg !71
  %790 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %766) #6, !dbg !71
  %791 = insertelement <2 x i16> undef, i16 %775, i64 0, !dbg !71
  %792 = insertelement <2 x i16> %791, i16 %776, i64 1, !dbg !71
  %793 = bitcast <2 x i16> %792 to i32, !dbg !71
  %794 = insertelement <2 x i16> undef, i16 %777, i64 0, !dbg !71
  %795 = insertelement <2 x i16> %794, i16 %778, i64 1, !dbg !71
  %796 = bitcast <2 x i16> %795 to i32, !dbg !71
  %797 = insertelement <2 x i16> undef, i16 %779, i64 0, !dbg !71
  %798 = insertelement <2 x i16> %797, i16 %780, i64 1, !dbg !71
  %799 = bitcast <2 x i16> %798 to i32, !dbg !71
  %800 = insertelement <2 x i16> undef, i16 %781, i64 0, !dbg !71
  %801 = insertelement <2 x i16> %800, i16 %782, i64 1, !dbg !71
  %802 = bitcast <2 x i16> %801 to i32, !dbg !71
  tail call void asm sideeffect "@$5 st.global.v4.b32 [ $4 + 0 ], { $0, $1, $2, $3 };", "r,r,r,r,l,b"(i32 %793, i32 %796, i32 %799, i32 %802, ptr addrspace(1) %772, i1 true) #6, !dbg !71
  %803 = insertelement <2 x i16> undef, i16 %783, i64 0, !dbg !71
  %804 = insertelement <2 x i16> %803, i16 %784, i64 1, !dbg !71
  %805 = bitcast <2 x i16> %804 to i32, !dbg !71
  %806 = insertelement <2 x i16> undef, i16 %785, i64 0, !dbg !71
  %807 = insertelement <2 x i16> %806, i16 %786, i64 1, !dbg !71
  %808 = bitcast <2 x i16> %807 to i32, !dbg !71
  %809 = insertelement <2 x i16> undef, i16 %787, i64 0, !dbg !71
  %810 = insertelement <2 x i16> %809, i16 %788, i64 1, !dbg !71
  %811 = bitcast <2 x i16> %810 to i32, !dbg !71
  %812 = insertelement <2 x i16> undef, i16 %789, i64 0, !dbg !71
  %813 = insertelement <2 x i16> %812, i16 %790, i64 1, !dbg !71
  %814 = bitcast <2 x i16> %813 to i32, !dbg !71
  tail call void asm sideeffect "@$5 st.global.v4.b32 [ $4 + 0 ], { $0, $1, $2, $3 };", "r,r,r,r,l,b"(i32 %805, i32 %808, i32 %811, i32 %814, ptr addrspace(1) %774, i1 true) #6, !dbg !71
  ret void, !dbg !72
}

; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
declare noundef i32 @llvm.nvvm.read.ptx.sreg.tid.x() #0

; Function Attrs: convergent nocallback nounwind memory(inaccessiblemem: readwrite)
declare i32 @llvm.nvvm.shfl.sync.bfly.i32(i32, i32, i32, i32) #1

; Function Attrs: convergent nocallback nounwind
declare void @llvm.nvvm.barrier0() #2

; Function Attrs: alwaysinline nounwind
define float @__nv_rsqrtf(float %x) local_unnamed_addr #3 {
  %1 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #6
  %.not = icmp eq i32 %1, 0
  br i1 %.not, label %4, label %2

2:                                                ; preds = %0
  %3 = tail call float @llvm.nvvm.rsqrt.approx.ftz.f(float %x)
  br label %6

4:                                                ; preds = %0
  %5 = tail call float @llvm.nvvm.rsqrt.approx.f(float %x)
  br label %6

6:                                                ; preds = %4, %2
  %.0 = phi float [ %3, %2 ], [ %5, %4 ]
  ret float %.0
}

declare i32 @__nvvm_reflect(ptr) local_unnamed_addr #4

; Function Attrs: mustprogress nocallback nofree nosync nounwind willreturn memory(none)
declare float @llvm.nvvm.rsqrt.approx.ftz.f(float) #5

; Function Attrs: mustprogress nocallback nofree nosync nounwind willreturn memory(none)
declare float @llvm.nvvm.rsqrt.approx.f(float) #5

attributes #0 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) }
attributes #1 = { convergent nocallback nounwind memory(inaccessiblemem: readwrite) }
attributes #2 = { convergent nocallback nounwind }
attributes #3 = { alwaysinline nounwind "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
attributes #4 = { "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
attributes #5 = { mustprogress nocallback nofree nosync nounwind willreturn memory(none) }
attributes #6 = { nounwind }

!llvm.module.flags = !{!0, !1}
!llvm.dbg.cu = !{!2}
!nvvm.annotations = !{!4, !5, !5, !4}
!llvm.ident = !{!6}

!0 = !{i32 2, !"Debug Info Version", i32 3}
!1 = !{i32 4, !"nvvm-reflect-ftz", i32 1}
!2 = distinct !DICompileUnit(language: DW_LANG_C, file: !3, producer: "triton", isOptimized: true, runtimeVersion: 0, emissionKind: FullDebug)
!3 = !DIFile(filename: "cgx5lxpuexpindj4dsmjz5x42uhyy7iskevq7ovzpwagb3t5powj.py", directory: "/tmp/torchinductor_root/gx")
!4 = !{ptr @triton__0d1d2d3d4d5de6de, !"kernel", i32 1}
!5 = !{ptr @triton__0d1d2d3d4d5de6de, !"maxntidx", i32 256}
!6 = !{!"clang version 3.8.0 (tags/RELEASE_380/final)"}
!7 = distinct !DISubprogram(name: "triton__0d1d2d3d4d5de6de", linkageName: "triton__0d1d2d3d4d5de6de", scope: !3, file: !3, line: 18, type: !8, scopeLine: 18, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !2)
!8 = !DISubroutineType(cc: DW_CC_normal, types: !9)
!9 = !{}
!10 = !DILocation(line: 22, column: 44, scope: !7)
!11 = !DILocation(line: 24, column: 33, scope: !7)
!12 = !DILocation(line: 21, column: 28, scope: !7)
!13 = !DILocation(line: 21, column: 33, scope: !7)
!14 = !DILocation(line: 22, column: 23, scope: !7)
!15 = !DILocation(line: 26, column: 30, scope: !7)
!16 = !DILocation(line: 26, column: 35, scope: !7)
!17 = !DILocation(line: 27, column: 18, scope: !7)
!18 = !DILocation(line: 35, column: 44, scope: !7)
!19 = !DILocation(line: 35, column: 40, scope: !7)
!20 = !DILocation(line: 35, column: 34, scope: !7)
!21 = !DILocation(line: 35, column: 50, scope: !7)
!22 = !DILocation(line: 36, column: 22, scope: !7)
!23 = !DILocation(line: 37, column: 22, scope: !7)
!24 = !DILocation(line: 38, column: 36, scope: !7)
!25 = !DILocation(line: 39, column: 40, scope: !7)
!26 = !DILocation(line: 39, column: 55, scope: !7)
!27 = !DILocation(line: 40, column: 44, scope: !7)
!28 = !DILocation(line: 40, column: 40, scope: !7)
!29 = !DILocation(line: 40, column: 34, scope: !7)
!30 = !DILocation(line: 40, column: 52, scope: !7)
!31 = !DILocation(line: 41, column: 22, scope: !7)
!32 = !DILocation(line: 98, column: 22, scope: !33, inlinedAt: !35)
!33 = distinct !DILexicalBlockFile(scope: !7, file: !34, discriminator: 0)
!34 = !DIFile(filename: "triton_helpers.py", directory: "/usr/local/lib/python3.10/dist-packages/torch/_inductor")
!35 = !DILocation(line: 44, column: 38, scope: !33)
!36 = !DILocation(line: 101, column: 30, scope: !33, inlinedAt: !35)
!37 = !DILocation(line: 101, column: 22, scope: !33, inlinedAt: !35)
!38 = !DILocation(line: 101, column: 13, scope: !33, inlinedAt: !35)
!39 = !DILocation(line: 108, column: 21, scope: !40, inlinedAt: !41)
!40 = distinct !DILexicalBlockFile(scope: !33, file: !34, discriminator: 0)
!41 = !DILocation(line: 120, column: 46, scope: !40, inlinedAt: !42)
!42 = !DILocation(line: 50, column: 41, scope: !40)
!43 = !DILocation(line: 110, column: 60, scope: !40, inlinedAt: !41)
!44 = !DILocation(line: 112, column: 25, scope: !40, inlinedAt: !41)
!45 = !DILocation(line: 112, column: 17, scope: !40, inlinedAt: !41)
!46 = !DILocation(line: 113, column: 15, scope: !40, inlinedAt: !41)
!47 = !DILocation(line: 113, column: 30, scope: !40, inlinedAt: !41)
!48 = !DILocation(line: 113, column: 49, scope: !40, inlinedAt: !41)
!49 = !DILocation(line: 113, column: 22, scope: !40, inlinedAt: !41)
!50 = !DILocation(line: 113, column: 38, scope: !40, inlinedAt: !41)
!51 = !DILocation(line: 120, column: 46, scope: !33, inlinedAt: !52)
!52 = !DILocation(line: 50, column: 41, scope: !33)
!53 = !DILocation(line: 109, column: 28, scope: !40, inlinedAt: !41)
!54 = !DILocation(line: 110, column: 39, scope: !40, inlinedAt: !41)
!55 = !DILocation(line: 110, column: 49, scope: !40, inlinedAt: !41)
!56 = !DILocation(line: 59, column: 51, scope: !7)
!57 = !DILocation(line: 60, column: 35, scope: !7)
!58 = !DILocation(line: 60, column: 40, scope: !7)
!59 = !DILocation(line: 64, column: 57, scope: !7)
!60 = !DILocation(line: 65, column: 54, scope: !7)
!61 = !DILocation(line: 69, column: 23, scope: !7)
!62 = !DILocation(line: 71, column: 24, scope: !7)
!63 = !DILocation(line: 72, column: 30, scope: !7)
!64 = !DILocation(line: 66, column: 24, scope: !7)
!65 = !DILocation(line: 67, column: 24, scope: !7)
!66 = !DILocation(line: 73, column: 24, scope: !7)
!67 = !DILocation(line: 74, column: 24, scope: !7)
!68 = !DILocation(line: 76, column: 39, scope: !7)
!69 = !DILocation(line: 76, column: 35, scope: !7)
!70 = !DILocation(line: 76, column: 29, scope: !7)
!71 = !DILocation(line: 76, column: 52, scope: !7)
!72 = !DILocation(line: 55, column: 4, scope: !7)