|
|
|
source_filename = "LLVMDialectModule" |
|
|
|
@global_smem = external addrspace(3) global [0 x i8] |
|
|
|
define void @triton__0d1d2d3d4d5d6d7de8(ptr addrspace(1) %0, ptr addrspace(1) %1, ptr addrspace(1) %2, ptr addrspace(1) %3, ptr addrspace(1) %4, ptr addrspace(1) %5, ptr addrspace(1) %6, i64 %7, i64 %8) local_unnamed_addr !dbg !5 { |
|
%10 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !8 |
|
%11 = and i32 %10, 31, !dbg !8 |
|
%12 = lshr i32 %10, 6, !dbg !8 |
|
%13 = and i32 %12, 3, !dbg !8 |
|
%14 = or i32 %13, 4, !dbg !8 |
|
%15 = or i32 %13, 8, !dbg !8 |
|
%16 = or i32 %13, 12, !dbg !8 |
|
%17 = or i32 %13, 16, !dbg !8 |
|
%18 = or i32 %13, 20, !dbg !8 |
|
%19 = or i32 %13, 24, !dbg !8 |
|
%20 = or i32 %13, 28, !dbg !8 |
|
%21 = or i32 %13, 32, !dbg !8 |
|
%22 = or i32 %13, 36, !dbg !8 |
|
%23 = or i32 %13, 40, !dbg !8 |
|
%24 = or i32 %13, 44, !dbg !8 |
|
%25 = or i32 %13, 48, !dbg !8 |
|
%26 = or i32 %13, 52, !dbg !8 |
|
%27 = or i32 %13, 56, !dbg !8 |
|
%28 = or i32 %13, 60, !dbg !8 |
|
%urem = and i32 %10, 63, !dbg !9 |
|
%29 = tail call i32 asm "mov.u32 $0, %ctaid.x;", "=r"() #3, !dbg !10 |
|
%30 = sext i32 %29 to i64, !dbg !11 |
|
%31 = shl nsw i64 %30, 6, !dbg !12 |
|
%32 = zext nneg i32 %13 to i64 |
|
%33 = zext nneg i32 %14 to i64 |
|
%34 = zext nneg i32 %15 to i64 |
|
%35 = zext nneg i32 %16 to i64 |
|
%36 = zext nneg i32 %17 to i64 |
|
%37 = zext nneg i32 %18 to i64 |
|
%38 = zext nneg i32 %19 to i64 |
|
%39 = zext nneg i32 %20 to i64 |
|
%40 = zext nneg i32 %21 to i64 |
|
%41 = zext nneg i32 %22 to i64 |
|
%42 = zext nneg i32 %23 to i64 |
|
%43 = zext nneg i32 %24 to i64 |
|
%44 = zext nneg i32 %25 to i64 |
|
%45 = zext nneg i32 %26 to i64 |
|
%46 = zext nneg i32 %27 to i64 |
|
%47 = zext nneg i32 %28 to i64 |
|
%48 = or i64 %31, %32, !dbg !13 |
|
%49 = or i64 %31, %33, !dbg !13 |
|
%50 = or i64 %31, %34, !dbg !13 |
|
%51 = or i64 %31, %35, !dbg !13 |
|
%52 = or i64 %31, %36, !dbg !13 |
|
%53 = or i64 %31, %37, !dbg !13 |
|
%54 = or i64 %31, %38, !dbg !13 |
|
%55 = or i64 %31, %39, !dbg !13 |
|
%56 = or i64 %31, %40, !dbg !13 |
|
%57 = or i64 %31, %41, !dbg !13 |
|
%58 = or i64 %31, %42, !dbg !13 |
|
%59 = or i64 %31, %43, !dbg !13 |
|
%60 = or i64 %31, %44, !dbg !13 |
|
%61 = or i64 %31, %45, !dbg !13 |
|
%62 = or i64 %31, %46, !dbg !13 |
|
%63 = or i64 %31, %47, !dbg !13 |
|
%64 = getelementptr i64, ptr addrspace(1) %1, i64 %48, !dbg !14 |
|
%65 = getelementptr i64, ptr addrspace(1) %1, i64 %49, !dbg !14 |
|
%66 = getelementptr i64, ptr addrspace(1) %1, i64 %50, !dbg !14 |
|
%67 = getelementptr i64, ptr addrspace(1) %1, i64 %51, !dbg !14 |
|
%68 = getelementptr i64, ptr addrspace(1) %1, i64 %52, !dbg !14 |
|
%69 = getelementptr i64, ptr addrspace(1) %1, i64 %53, !dbg !14 |
|
%70 = getelementptr i64, ptr addrspace(1) %1, i64 %54, !dbg !14 |
|
%71 = getelementptr i64, ptr addrspace(1) %1, i64 %55, !dbg !14 |
|
%72 = getelementptr i64, ptr addrspace(1) %1, i64 %56, !dbg !14 |
|
%73 = getelementptr i64, ptr addrspace(1) %1, i64 %57, !dbg !14 |
|
%74 = getelementptr i64, ptr addrspace(1) %1, i64 %58, !dbg !14 |
|
%75 = getelementptr i64, ptr addrspace(1) %1, i64 %59, !dbg !14 |
|
%76 = getelementptr i64, ptr addrspace(1) %1, i64 %60, !dbg !14 |
|
%77 = getelementptr i64, ptr addrspace(1) %1, i64 %61, !dbg !14 |
|
%78 = getelementptr i64, ptr addrspace(1) %1, i64 %62, !dbg !14 |
|
%79 = getelementptr i64, ptr addrspace(1) %1, i64 %63, !dbg !14 |
|
%80 = 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) %64, i1 true) #3, !dbg !15 |
|
%81 = 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) %65, i1 true) #3, !dbg !15 |
|
%82 = 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) %66, i1 true) #3, !dbg !15 |
|
%83 = 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) %67, i1 true) #3, !dbg !15 |
|
%84 = 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) %68, i1 true) #3, !dbg !15 |
|
%85 = 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) %69, i1 true) #3, !dbg !15 |
|
%86 = 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) %70, i1 true) #3, !dbg !15 |
|
%87 = 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) %71, i1 true) #3, !dbg !15 |
|
%88 = 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) %72, i1 true) #3, !dbg !15 |
|
%89 = 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) %73, i1 true) #3, !dbg !15 |
|
%90 = 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) %74, i1 true) #3, !dbg !15 |
|
%91 = 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) %75, i1 true) #3, !dbg !15 |
|
%92 = 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) %76, i1 true) #3, !dbg !15 |
|
%93 = 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) %77, i1 true) #3, !dbg !15 |
|
%94 = 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) %78, i1 true) #3, !dbg !15 |
|
%95 = 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) %79, i1 true) #3, !dbg !15 |
|
%96 = tail call i32 asm sideeffect "mov.u32 $0, 0x0;\0A\09@$2 ld.global.b32 { $0 }, [ $1 + 0 ];", "=r,l,b"(ptr addrspace(1) %2, i1 true) #3, !dbg !16 |
|
%97 = bitcast i32 %96 to float, !dbg !16 |
|
%98 = tail call i32 asm sideeffect "mov.u32 $0, 0x0;\0A\09@$2 ld.global.b32 { $0 }, [ $1 + 0 ];", "=r,l,b"(ptr addrspace(1) %3, i1 true) #3, !dbg !17 |
|
%99 = bitcast i32 %98 to float, !dbg !17 |
|
%100 = mul nsw i64 %48, 50257, !dbg !18 |
|
%101 = mul nsw i64 %49, 50257, !dbg !18 |
|
%102 = mul nsw i64 %50, 50257, !dbg !18 |
|
%103 = mul nsw i64 %51, 50257, !dbg !18 |
|
%104 = mul nsw i64 %52, 50257, !dbg !18 |
|
%105 = mul nsw i64 %53, 50257, !dbg !18 |
|
%106 = mul nsw i64 %54, 50257, !dbg !18 |
|
%107 = mul nsw i64 %55, 50257, !dbg !18 |
|
%108 = mul nsw i64 %56, 50257, !dbg !18 |
|
%109 = mul nsw i64 %57, 50257, !dbg !18 |
|
%110 = mul nsw i64 %58, 50257, !dbg !18 |
|
%111 = mul nsw i64 %59, 50257, !dbg !18 |
|
%112 = mul nsw i64 %60, 50257, !dbg !18 |
|
%113 = mul nsw i64 %61, 50257, !dbg !18 |
|
%114 = mul nsw i64 %62, 50257, !dbg !18 |
|
%115 = mul nsw i64 %63, 50257, !dbg !18 |
|
%116 = insertelement <16 x i64> poison, i64 %80, i64 0, !dbg !19 |
|
%117 = insertelement <16 x i64> %116, i64 %81, i64 1, !dbg !19 |
|
%118 = insertelement <16 x i64> %117, i64 %82, i64 2, !dbg !19 |
|
%119 = insertelement <16 x i64> %118, i64 %83, i64 3, !dbg !19 |
|
%120 = insertelement <16 x i64> %119, i64 %84, i64 4, !dbg !19 |
|
%121 = insertelement <16 x i64> %120, i64 %85, i64 5, !dbg !19 |
|
%122 = insertelement <16 x i64> %121, i64 %86, i64 6, !dbg !19 |
|
%123 = insertelement <16 x i64> %122, i64 %87, i64 7, !dbg !19 |
|
%124 = insertelement <16 x i64> %123, i64 %88, i64 8, !dbg !19 |
|
%125 = insertelement <16 x i64> %124, i64 %89, i64 9, !dbg !19 |
|
%126 = insertelement <16 x i64> %125, i64 %90, i64 10, !dbg !19 |
|
%127 = insertelement <16 x i64> %126, i64 %91, i64 11, !dbg !19 |
|
%128 = insertelement <16 x i64> %127, i64 %92, i64 12, !dbg !19 |
|
%129 = insertelement <16 x i64> %128, i64 %93, i64 13, !dbg !19 |
|
%130 = insertelement <16 x i64> %129, i64 %94, i64 14, !dbg !19 |
|
%131 = insertelement <16 x i64> %130, i64 %95, i64 15, !dbg !19 |
|
%132 = icmp eq <16 x i64> %131, <i64 -1, i64 -1, i64 -1, i64 -1, i64 -1, i64 -1, i64 -1, i64 -1, i64 -1, i64 -1, i64 -1, i64 -1, i64 -1, i64 -1, i64 -1, i64 -1>, !dbg !19 |
|
%133 = tail call float asm "div.full.f32 $0, $1, $2;", "=r,r,r"(float %97, float %99) #3, !dbg !20 |
|
%134 = insertelement <16 x float> poison, float %133, i64 0, !dbg !21 |
|
%135 = shufflevector <16 x float> %134, <16 x float> poison, <16 x i32> zeroinitializer, !dbg !21 |
|
%136 = select <16 x i1> %132, <16 x float> zeroinitializer, <16 x float> %135, !dbg !21 |
|
%137 = getelementptr float, ptr addrspace(1) %0, i64 %100 |
|
%138 = getelementptr float, ptr addrspace(1) %0, i64 %101 |
|
%139 = getelementptr float, ptr addrspace(1) %0, i64 %102 |
|
%140 = getelementptr float, ptr addrspace(1) %0, i64 %103 |
|
%141 = getelementptr float, ptr addrspace(1) %0, i64 %104 |
|
%142 = getelementptr float, ptr addrspace(1) %0, i64 %105 |
|
%143 = getelementptr float, ptr addrspace(1) %0, i64 %106 |
|
%144 = getelementptr float, ptr addrspace(1) %0, i64 %107 |
|
%145 = getelementptr float, ptr addrspace(1) %0, i64 %108 |
|
%146 = getelementptr float, ptr addrspace(1) %0, i64 %109 |
|
%147 = getelementptr float, ptr addrspace(1) %0, i64 %110 |
|
%148 = getelementptr float, ptr addrspace(1) %0, i64 %111 |
|
%149 = getelementptr float, ptr addrspace(1) %0, i64 %112 |
|
%150 = getelementptr float, ptr addrspace(1) %0, i64 %113 |
|
%151 = getelementptr float, ptr addrspace(1) %0, i64 %114 |
|
%152 = getelementptr float, ptr addrspace(1) %0, i64 %115 |
|
br label %153, !dbg !22 |
|
|
|
153: |
|
%154 = phi i32 [ 0, %9 ], [ %213, %153 ] |
|
%155 = phi <16 x float> [ zeroinitializer, %9 ], [ %212, %153 ] |
|
%156 = or i32 %154, %urem, !dbg !23 |
|
%157 = zext nneg i32 %156 to i64, !dbg !23 |
|
%158 = icmp ult i32 %156, 50257, !dbg !24 |
|
%159 = getelementptr float, ptr addrspace(1) %137, i64 %157, !dbg !25 |
|
%160 = getelementptr float, ptr addrspace(1) %138, i64 %157, !dbg !25 |
|
%161 = getelementptr float, ptr addrspace(1) %139, i64 %157, !dbg !25 |
|
%162 = getelementptr float, ptr addrspace(1) %140, i64 %157, !dbg !25 |
|
%163 = getelementptr float, ptr addrspace(1) %141, i64 %157, !dbg !25 |
|
%164 = getelementptr float, ptr addrspace(1) %142, i64 %157, !dbg !25 |
|
%165 = getelementptr float, ptr addrspace(1) %143, i64 %157, !dbg !25 |
|
%166 = getelementptr float, ptr addrspace(1) %144, i64 %157, !dbg !25 |
|
%167 = getelementptr float, ptr addrspace(1) %145, i64 %157, !dbg !25 |
|
%168 = getelementptr float, ptr addrspace(1) %146, i64 %157, !dbg !25 |
|
%169 = getelementptr float, ptr addrspace(1) %147, i64 %157, !dbg !25 |
|
%170 = getelementptr float, ptr addrspace(1) %148, i64 %157, !dbg !25 |
|
%171 = getelementptr float, ptr addrspace(1) %149, i64 %157, !dbg !25 |
|
%172 = getelementptr float, ptr addrspace(1) %150, i64 %157, !dbg !25 |
|
%173 = getelementptr float, ptr addrspace(1) %151, i64 %157, !dbg !25 |
|
%174 = getelementptr float, ptr addrspace(1) %152, i64 %157, !dbg !25 |
|
%175 = 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) %159, i1 %158, i32 0, i1 %158) #3, !dbg !26 |
|
%176 = 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) %160, i1 %158, i32 0, i1 %158) #3, !dbg !26 |
|
%177 = 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) %161, i1 %158, i32 0, i1 %158) #3, !dbg !26 |
|
%178 = 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) %162, i1 %158, i32 0, i1 %158) #3, !dbg !26 |
|
%179 = 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) %163, i1 %158, i32 0, i1 %158) #3, !dbg !26 |
|
%180 = 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) %164, i1 %158, i32 0, i1 %158) #3, !dbg !26 |
|
%181 = 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) %165, i1 %158, i32 0, i1 %158) #3, !dbg !26 |
|
%182 = 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) %166, i1 %158, i32 0, i1 %158) #3, !dbg !26 |
|
%183 = 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) %167, i1 %158, i32 0, i1 %158) #3, !dbg !26 |
|
%184 = 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) %168, i1 %158, i32 0, i1 %158) #3, !dbg !26 |
|
%185 = 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) %169, i1 %158, i32 0, i1 %158) #3, !dbg !26 |
|
%186 = 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) %170, i1 %158, i32 0, i1 %158) #3, !dbg !26 |
|
%187 = 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) %171, i1 %158, i32 0, i1 %158) #3, !dbg !26 |
|
%188 = 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) %172, i1 %158, i32 0, i1 %158) #3, !dbg !26 |
|
%189 = 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) %173, i1 %158, i32 0, i1 %158) #3, !dbg !26 |
|
%190 = 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) %174, i1 %158, i32 0, i1 %158) #3, !dbg !26 |
|
%191 = insertelement <16 x i32> poison, i32 %175, i64 0, !dbg !26 |
|
%192 = insertelement <16 x i32> %191, i32 %176, i64 1, !dbg !26 |
|
%193 = insertelement <16 x i32> %192, i32 %177, i64 2, !dbg !26 |
|
%194 = insertelement <16 x i32> %193, i32 %178, i64 3, !dbg !26 |
|
%195 = insertelement <16 x i32> %194, i32 %179, i64 4, !dbg !26 |
|
%196 = insertelement <16 x i32> %195, i32 %180, i64 5, !dbg !26 |
|
%197 = insertelement <16 x i32> %196, i32 %181, i64 6, !dbg !26 |
|
%198 = insertelement <16 x i32> %197, i32 %182, i64 7, !dbg !26 |
|
%199 = insertelement <16 x i32> %198, i32 %183, i64 8, !dbg !26 |
|
%200 = insertelement <16 x i32> %199, i32 %184, i64 9, !dbg !26 |
|
%201 = insertelement <16 x i32> %200, i32 %185, i64 10, !dbg !26 |
|
%202 = insertelement <16 x i32> %201, i32 %186, i64 11, !dbg !26 |
|
%203 = insertelement <16 x i32> %202, i32 %187, i64 12, !dbg !26 |
|
%204 = insertelement <16 x i32> %203, i32 %188, i64 13, !dbg !26 |
|
%205 = insertelement <16 x i32> %204, i32 %189, i64 14, !dbg !26 |
|
%206 = insertelement <16 x i32> %205, i32 %190, i64 15, !dbg !26 |
|
%207 = bitcast <16 x i32> %206 to <16 x float>, !dbg !26 |
|
%208 = fmul <16 x float> %136, %207, !dbg !27 |
|
%209 = insertelement <16 x i1> poison, i1 %158, i64 0, !dbg !28 |
|
%210 = shufflevector <16 x i1> %209, <16 x i1> poison, <16 x i32> zeroinitializer, !dbg !28 |
|
%211 = select <16 x i1> %210, <16 x float> %208, <16 x float> <float -0.000000e+00, float -0.000000e+00, float -0.000000e+00, float -0.000000e+00, float -0.000000e+00, float -0.000000e+00, float -0.000000e+00, float -0.000000e+00, float -0.000000e+00, float -0.000000e+00, float -0.000000e+00, float -0.000000e+00, float -0.000000e+00, float -0.000000e+00, float -0.000000e+00, float -0.000000e+00>, !dbg !28 |
|
%212 = fadd <16 x float> %155, %211, !dbg !28 |
|
%213 = add nuw nsw i32 %154, 64, !dbg !22 |
|
%214 = icmp ult i32 %154, 50193, !dbg !22 |
|
br i1 %214, label %153, label %215, !dbg !22 |
|
|
|
215: |
|
%216 = lshr i32 %10, 5, !dbg !8 |
|
%217 = and i32 %216, 1, !dbg !9 |
|
%218 = extractelement <16 x float> %212, i64 0, !dbg !29 |
|
%219 = bitcast float %218 to i32, !dbg !29 |
|
%220 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %219, i32 16, i32 31), !dbg !29 |
|
%221 = bitcast i32 %220 to float, !dbg !29 |
|
%222 = fadd float %218, %221, !dbg !33 |
|
%223 = bitcast float %222 to i32, !dbg !29 |
|
%224 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %223, i32 8, i32 31), !dbg !29 |
|
%225 = bitcast i32 %224 to float, !dbg !29 |
|
%226 = fadd float %222, %225, !dbg !33 |
|
%227 = bitcast float %226 to i32, !dbg !29 |
|
%228 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %227, i32 4, i32 31), !dbg !29 |
|
%229 = bitcast i32 %228 to float, !dbg !29 |
|
%230 = fadd float %226, %229, !dbg !33 |
|
%231 = bitcast float %230 to i32, !dbg !29 |
|
%232 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %231, i32 2, i32 31), !dbg !29 |
|
%233 = bitcast i32 %232 to float, !dbg !29 |
|
%234 = fadd float %230, %233, !dbg !33 |
|
%235 = bitcast float %234 to i32, !dbg !29 |
|
%236 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %235, i32 1, i32 31), !dbg !29 |
|
%237 = bitcast i32 %236 to float, !dbg !29 |
|
%238 = fadd float %234, %237, !dbg !33 |
|
%239 = extractelement <16 x float> %212, i64 1, !dbg !29 |
|
%240 = bitcast float %239 to i32, !dbg !29 |
|
%241 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %240, i32 16, i32 31), !dbg !29 |
|
%242 = bitcast i32 %241 to float, !dbg !29 |
|
%243 = fadd float %239, %242, !dbg !33 |
|
%244 = bitcast float %243 to i32, !dbg !29 |
|
%245 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %244, i32 8, i32 31), !dbg !29 |
|
%246 = bitcast i32 %245 to float, !dbg !29 |
|
%247 = fadd float %243, %246, !dbg !33 |
|
%248 = bitcast float %247 to i32, !dbg !29 |
|
%249 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %248, i32 4, i32 31), !dbg !29 |
|
%250 = bitcast i32 %249 to float, !dbg !29 |
|
%251 = fadd float %247, %250, !dbg !33 |
|
%252 = bitcast float %251 to i32, !dbg !29 |
|
%253 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %252, i32 2, i32 31), !dbg !29 |
|
%254 = bitcast i32 %253 to float, !dbg !29 |
|
%255 = fadd float %251, %254, !dbg !33 |
|
%256 = bitcast float %255 to i32, !dbg !29 |
|
%257 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %256, i32 1, i32 31), !dbg !29 |
|
%258 = bitcast i32 %257 to float, !dbg !29 |
|
%259 = fadd float %255, %258, !dbg !33 |
|
%260 = extractelement <16 x float> %212, i64 2, !dbg !29 |
|
%261 = bitcast float %260 to i32, !dbg !29 |
|
%262 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %261, i32 16, i32 31), !dbg !29 |
|
%263 = bitcast i32 %262 to float, !dbg !29 |
|
%264 = fadd float %260, %263, !dbg !33 |
|
%265 = bitcast float %264 to i32, !dbg !29 |
|
%266 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %265, i32 8, i32 31), !dbg !29 |
|
%267 = bitcast i32 %266 to float, !dbg !29 |
|
%268 = fadd float %264, %267, !dbg !33 |
|
%269 = bitcast float %268 to i32, !dbg !29 |
|
%270 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %269, i32 4, i32 31), !dbg !29 |
|
%271 = bitcast i32 %270 to float, !dbg !29 |
|
%272 = fadd float %268, %271, !dbg !33 |
|
%273 = bitcast float %272 to i32, !dbg !29 |
|
%274 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %273, i32 2, i32 31), !dbg !29 |
|
%275 = bitcast i32 %274 to float, !dbg !29 |
|
%276 = fadd float %272, %275, !dbg !33 |
|
%277 = bitcast float %276 to i32, !dbg !29 |
|
%278 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %277, i32 1, i32 31), !dbg !29 |
|
%279 = bitcast i32 %278 to float, !dbg !29 |
|
%280 = fadd float %276, %279, !dbg !33 |
|
%281 = extractelement <16 x float> %212, i64 3, !dbg !29 |
|
%282 = bitcast float %281 to i32, !dbg !29 |
|
%283 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %282, i32 16, i32 31), !dbg !29 |
|
%284 = bitcast i32 %283 to float, !dbg !29 |
|
%285 = fadd float %281, %284, !dbg !33 |
|
%286 = bitcast float %285 to i32, !dbg !29 |
|
%287 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %286, i32 8, i32 31), !dbg !29 |
|
%288 = bitcast i32 %287 to float, !dbg !29 |
|
%289 = fadd float %285, %288, !dbg !33 |
|
%290 = bitcast float %289 to i32, !dbg !29 |
|
%291 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %290, i32 4, i32 31), !dbg !29 |
|
%292 = bitcast i32 %291 to float, !dbg !29 |
|
%293 = fadd float %289, %292, !dbg !33 |
|
%294 = bitcast float %293 to i32, !dbg !29 |
|
%295 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %294, i32 2, i32 31), !dbg !29 |
|
%296 = bitcast i32 %295 to float, !dbg !29 |
|
%297 = fadd float %293, %296, !dbg !33 |
|
%298 = bitcast float %297 to i32, !dbg !29 |
|
%299 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %298, i32 1, i32 31), !dbg !29 |
|
%300 = bitcast i32 %299 to float, !dbg !29 |
|
%301 = fadd float %297, %300, !dbg !33 |
|
%302 = extractelement <16 x float> %212, i64 4, !dbg !29 |
|
%303 = bitcast float %302 to i32, !dbg !29 |
|
%304 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %303, i32 16, i32 31), !dbg !29 |
|
%305 = bitcast i32 %304 to float, !dbg !29 |
|
%306 = fadd float %302, %305, !dbg !33 |
|
%307 = bitcast float %306 to i32, !dbg !29 |
|
%308 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %307, i32 8, i32 31), !dbg !29 |
|
%309 = bitcast i32 %308 to float, !dbg !29 |
|
%310 = fadd float %306, %309, !dbg !33 |
|
%311 = bitcast float %310 to i32, !dbg !29 |
|
%312 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %311, i32 4, i32 31), !dbg !29 |
|
%313 = bitcast i32 %312 to float, !dbg !29 |
|
%314 = fadd float %310, %313, !dbg !33 |
|
%315 = bitcast float %314 to i32, !dbg !29 |
|
%316 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %315, i32 2, i32 31), !dbg !29 |
|
%317 = bitcast i32 %316 to float, !dbg !29 |
|
%318 = fadd float %314, %317, !dbg !33 |
|
%319 = bitcast float %318 to i32, !dbg !29 |
|
%320 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %319, i32 1, i32 31), !dbg !29 |
|
%321 = bitcast i32 %320 to float, !dbg !29 |
|
%322 = fadd float %318, %321, !dbg !33 |
|
%323 = extractelement <16 x float> %212, i64 5, !dbg !29 |
|
%324 = bitcast float %323 to i32, !dbg !29 |
|
%325 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %324, i32 16, i32 31), !dbg !29 |
|
%326 = bitcast i32 %325 to float, !dbg !29 |
|
%327 = fadd float %323, %326, !dbg !33 |
|
%328 = bitcast float %327 to i32, !dbg !29 |
|
%329 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %328, i32 8, i32 31), !dbg !29 |
|
%330 = bitcast i32 %329 to float, !dbg !29 |
|
%331 = fadd float %327, %330, !dbg !33 |
|
%332 = bitcast float %331 to i32, !dbg !29 |
|
%333 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %332, i32 4, i32 31), !dbg !29 |
|
%334 = bitcast i32 %333 to float, !dbg !29 |
|
%335 = fadd float %331, %334, !dbg !33 |
|
%336 = bitcast float %335 to i32, !dbg !29 |
|
%337 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %336, i32 2, i32 31), !dbg !29 |
|
%338 = bitcast i32 %337 to float, !dbg !29 |
|
%339 = fadd float %335, %338, !dbg !33 |
|
%340 = bitcast float %339 to i32, !dbg !29 |
|
%341 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %340, i32 1, i32 31), !dbg !29 |
|
%342 = bitcast i32 %341 to float, !dbg !29 |
|
%343 = fadd float %339, %342, !dbg !33 |
|
%344 = extractelement <16 x float> %212, i64 6, !dbg !29 |
|
%345 = bitcast float %344 to i32, !dbg !29 |
|
%346 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %345, i32 16, i32 31), !dbg !29 |
|
%347 = bitcast i32 %346 to float, !dbg !29 |
|
%348 = fadd float %344, %347, !dbg !33 |
|
%349 = bitcast float %348 to i32, !dbg !29 |
|
%350 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %349, i32 8, i32 31), !dbg !29 |
|
%351 = bitcast i32 %350 to float, !dbg !29 |
|
%352 = fadd float %348, %351, !dbg !33 |
|
%353 = bitcast float %352 to i32, !dbg !29 |
|
%354 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %353, i32 4, i32 31), !dbg !29 |
|
%355 = bitcast i32 %354 to float, !dbg !29 |
|
%356 = fadd float %352, %355, !dbg !33 |
|
%357 = bitcast float %356 to i32, !dbg !29 |
|
%358 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %357, i32 2, i32 31), !dbg !29 |
|
%359 = bitcast i32 %358 to float, !dbg !29 |
|
%360 = fadd float %356, %359, !dbg !33 |
|
%361 = bitcast float %360 to i32, !dbg !29 |
|
%362 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %361, i32 1, i32 31), !dbg !29 |
|
%363 = bitcast i32 %362 to float, !dbg !29 |
|
%364 = fadd float %360, %363, !dbg !33 |
|
%365 = extractelement <16 x float> %212, i64 7, !dbg !29 |
|
%366 = bitcast float %365 to i32, !dbg !29 |
|
%367 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %366, i32 16, i32 31), !dbg !29 |
|
%368 = bitcast i32 %367 to float, !dbg !29 |
|
%369 = fadd float %365, %368, !dbg !33 |
|
%370 = bitcast float %369 to i32, !dbg !29 |
|
%371 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %370, i32 8, i32 31), !dbg !29 |
|
%372 = bitcast i32 %371 to float, !dbg !29 |
|
%373 = fadd float %369, %372, !dbg !33 |
|
%374 = bitcast float %373 to i32, !dbg !29 |
|
%375 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %374, i32 4, i32 31), !dbg !29 |
|
%376 = bitcast i32 %375 to float, !dbg !29 |
|
%377 = fadd float %373, %376, !dbg !33 |
|
%378 = bitcast float %377 to i32, !dbg !29 |
|
%379 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %378, i32 2, i32 31), !dbg !29 |
|
%380 = bitcast i32 %379 to float, !dbg !29 |
|
%381 = fadd float %377, %380, !dbg !33 |
|
%382 = bitcast float %381 to i32, !dbg !29 |
|
%383 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %382, i32 1, i32 31), !dbg !29 |
|
%384 = bitcast i32 %383 to float, !dbg !29 |
|
%385 = fadd float %381, %384, !dbg !33 |
|
%386 = extractelement <16 x float> %212, i64 8, !dbg !29 |
|
%387 = bitcast float %386 to i32, !dbg !29 |
|
%388 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %387, i32 16, i32 31), !dbg !29 |
|
%389 = bitcast i32 %388 to float, !dbg !29 |
|
%390 = fadd float %386, %389, !dbg !33 |
|
%391 = bitcast float %390 to i32, !dbg !29 |
|
%392 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %391, i32 8, i32 31), !dbg !29 |
|
%393 = bitcast i32 %392 to float, !dbg !29 |
|
%394 = fadd float %390, %393, !dbg !33 |
|
%395 = bitcast float %394 to i32, !dbg !29 |
|
%396 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %395, i32 4, i32 31), !dbg !29 |
|
%397 = bitcast i32 %396 to float, !dbg !29 |
|
%398 = fadd float %394, %397, !dbg !33 |
|
%399 = bitcast float %398 to i32, !dbg !29 |
|
%400 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %399, i32 2, i32 31), !dbg !29 |
|
%401 = bitcast i32 %400 to float, !dbg !29 |
|
%402 = fadd float %398, %401, !dbg !33 |
|
%403 = bitcast float %402 to i32, !dbg !29 |
|
%404 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %403, i32 1, i32 31), !dbg !29 |
|
%405 = bitcast i32 %404 to float, !dbg !29 |
|
%406 = fadd float %402, %405, !dbg !33 |
|
%407 = extractelement <16 x float> %212, i64 9, !dbg !29 |
|
%408 = bitcast float %407 to i32, !dbg !29 |
|
%409 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %408, i32 16, i32 31), !dbg !29 |
|
%410 = bitcast i32 %409 to float, !dbg !29 |
|
%411 = fadd float %407, %410, !dbg !33 |
|
%412 = bitcast float %411 to i32, !dbg !29 |
|
%413 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %412, i32 8, i32 31), !dbg !29 |
|
%414 = bitcast i32 %413 to float, !dbg !29 |
|
%415 = fadd float %411, %414, !dbg !33 |
|
%416 = bitcast float %415 to i32, !dbg !29 |
|
%417 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %416, i32 4, i32 31), !dbg !29 |
|
%418 = bitcast i32 %417 to float, !dbg !29 |
|
%419 = fadd float %415, %418, !dbg !33 |
|
%420 = bitcast float %419 to i32, !dbg !29 |
|
%421 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %420, i32 2, i32 31), !dbg !29 |
|
%422 = bitcast i32 %421 to float, !dbg !29 |
|
%423 = fadd float %419, %422, !dbg !33 |
|
%424 = bitcast float %423 to i32, !dbg !29 |
|
%425 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %424, i32 1, i32 31), !dbg !29 |
|
%426 = bitcast i32 %425 to float, !dbg !29 |
|
%427 = fadd float %423, %426, !dbg !33 |
|
%428 = extractelement <16 x float> %212, i64 10, !dbg !29 |
|
%429 = bitcast float %428 to i32, !dbg !29 |
|
%430 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %429, i32 16, i32 31), !dbg !29 |
|
%431 = bitcast i32 %430 to float, !dbg !29 |
|
%432 = fadd float %428, %431, !dbg !33 |
|
%433 = bitcast float %432 to i32, !dbg !29 |
|
%434 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %433, i32 8, i32 31), !dbg !29 |
|
%435 = bitcast i32 %434 to float, !dbg !29 |
|
%436 = fadd float %432, %435, !dbg !33 |
|
%437 = bitcast float %436 to i32, !dbg !29 |
|
%438 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %437, i32 4, i32 31), !dbg !29 |
|
%439 = bitcast i32 %438 to float, !dbg !29 |
|
%440 = fadd float %436, %439, !dbg !33 |
|
%441 = bitcast float %440 to i32, !dbg !29 |
|
%442 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %441, i32 2, i32 31), !dbg !29 |
|
%443 = bitcast i32 %442 to float, !dbg !29 |
|
%444 = fadd float %440, %443, !dbg !33 |
|
%445 = bitcast float %444 to i32, !dbg !29 |
|
%446 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %445, i32 1, i32 31), !dbg !29 |
|
%447 = bitcast i32 %446 to float, !dbg !29 |
|
%448 = fadd float %444, %447, !dbg !33 |
|
%449 = extractelement <16 x float> %212, i64 11, !dbg !29 |
|
%450 = bitcast float %449 to i32, !dbg !29 |
|
%451 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %450, i32 16, i32 31), !dbg !29 |
|
%452 = bitcast i32 %451 to float, !dbg !29 |
|
%453 = fadd float %449, %452, !dbg !33 |
|
%454 = bitcast float %453 to i32, !dbg !29 |
|
%455 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %454, i32 8, i32 31), !dbg !29 |
|
%456 = bitcast i32 %455 to float, !dbg !29 |
|
%457 = fadd float %453, %456, !dbg !33 |
|
%458 = bitcast float %457 to i32, !dbg !29 |
|
%459 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %458, i32 4, i32 31), !dbg !29 |
|
%460 = bitcast i32 %459 to float, !dbg !29 |
|
%461 = fadd float %457, %460, !dbg !33 |
|
%462 = bitcast float %461 to i32, !dbg !29 |
|
%463 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %462, i32 2, i32 31), !dbg !29 |
|
%464 = bitcast i32 %463 to float, !dbg !29 |
|
%465 = fadd float %461, %464, !dbg !33 |
|
%466 = bitcast float %465 to i32, !dbg !29 |
|
%467 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %466, i32 1, i32 31), !dbg !29 |
|
%468 = bitcast i32 %467 to float, !dbg !29 |
|
%469 = fadd float %465, %468, !dbg !33 |
|
%470 = extractelement <16 x float> %212, i64 12, !dbg !29 |
|
%471 = bitcast float %470 to i32, !dbg !29 |
|
%472 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %471, i32 16, i32 31), !dbg !29 |
|
%473 = bitcast i32 %472 to float, !dbg !29 |
|
%474 = fadd float %470, %473, !dbg !33 |
|
%475 = bitcast float %474 to i32, !dbg !29 |
|
%476 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %475, i32 8, i32 31), !dbg !29 |
|
%477 = bitcast i32 %476 to float, !dbg !29 |
|
%478 = fadd float %474, %477, !dbg !33 |
|
%479 = bitcast float %478 to i32, !dbg !29 |
|
%480 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %479, i32 4, i32 31), !dbg !29 |
|
%481 = bitcast i32 %480 to float, !dbg !29 |
|
%482 = fadd float %478, %481, !dbg !33 |
|
%483 = bitcast float %482 to i32, !dbg !29 |
|
%484 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %483, i32 2, i32 31), !dbg !29 |
|
%485 = bitcast i32 %484 to float, !dbg !29 |
|
%486 = fadd float %482, %485, !dbg !33 |
|
%487 = bitcast float %486 to i32, !dbg !29 |
|
%488 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %487, i32 1, i32 31), !dbg !29 |
|
%489 = bitcast i32 %488 to float, !dbg !29 |
|
%490 = fadd float %486, %489, !dbg !33 |
|
%491 = extractelement <16 x float> %212, i64 13, !dbg !29 |
|
%492 = bitcast float %491 to i32, !dbg !29 |
|
%493 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %492, i32 16, i32 31), !dbg !29 |
|
%494 = bitcast i32 %493 to float, !dbg !29 |
|
%495 = fadd float %491, %494, !dbg !33 |
|
%496 = bitcast float %495 to i32, !dbg !29 |
|
%497 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %496, i32 8, i32 31), !dbg !29 |
|
%498 = bitcast i32 %497 to float, !dbg !29 |
|
%499 = fadd float %495, %498, !dbg !33 |
|
%500 = bitcast float %499 to i32, !dbg !29 |
|
%501 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %500, i32 4, i32 31), !dbg !29 |
|
%502 = bitcast i32 %501 to float, !dbg !29 |
|
%503 = fadd float %499, %502, !dbg !33 |
|
%504 = bitcast float %503 to i32, !dbg !29 |
|
%505 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %504, i32 2, i32 31), !dbg !29 |
|
%506 = bitcast i32 %505 to float, !dbg !29 |
|
%507 = fadd float %503, %506, !dbg !33 |
|
%508 = bitcast float %507 to i32, !dbg !29 |
|
%509 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %508, i32 1, i32 31), !dbg !29 |
|
%510 = bitcast i32 %509 to float, !dbg !29 |
|
%511 = fadd float %507, %510, !dbg !33 |
|
%512 = extractelement <16 x float> %212, i64 14, !dbg !29 |
|
%513 = bitcast float %512 to i32, !dbg !29 |
|
%514 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %513, i32 16, i32 31), !dbg !29 |
|
%515 = bitcast i32 %514 to float, !dbg !29 |
|
%516 = fadd float %512, %515, !dbg !33 |
|
%517 = bitcast float %516 to i32, !dbg !29 |
|
%518 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %517, i32 8, i32 31), !dbg !29 |
|
%519 = bitcast i32 %518 to float, !dbg !29 |
|
%520 = fadd float %516, %519, !dbg !33 |
|
%521 = bitcast float %520 to i32, !dbg !29 |
|
%522 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %521, i32 4, i32 31), !dbg !29 |
|
%523 = bitcast i32 %522 to float, !dbg !29 |
|
%524 = fadd float %520, %523, !dbg !33 |
|
%525 = bitcast float %524 to i32, !dbg !29 |
|
%526 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %525, i32 2, i32 31), !dbg !29 |
|
%527 = bitcast i32 %526 to float, !dbg !29 |
|
%528 = fadd float %524, %527, !dbg !33 |
|
%529 = bitcast float %528 to i32, !dbg !29 |
|
%530 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %529, i32 1, i32 31), !dbg !29 |
|
%531 = bitcast i32 %530 to float, !dbg !29 |
|
%532 = fadd float %528, %531, !dbg !33 |
|
%533 = extractelement <16 x float> %212, i64 15, !dbg !29 |
|
%534 = bitcast float %533 to i32, !dbg !29 |
|
%535 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %534, i32 16, i32 31), !dbg !29 |
|
%536 = bitcast i32 %535 to float, !dbg !29 |
|
%537 = fadd float %533, %536, !dbg !33 |
|
%538 = bitcast float %537 to i32, !dbg !29 |
|
%539 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %538, i32 8, i32 31), !dbg !29 |
|
%540 = bitcast i32 %539 to float, !dbg !29 |
|
%541 = fadd float %537, %540, !dbg !33 |
|
%542 = bitcast float %541 to i32, !dbg !29 |
|
%543 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %542, i32 4, i32 31), !dbg !29 |
|
%544 = bitcast i32 %543 to float, !dbg !29 |
|
%545 = fadd float %541, %544, !dbg !33 |
|
%546 = bitcast float %545 to i32, !dbg !29 |
|
%547 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %546, i32 2, i32 31), !dbg !29 |
|
%548 = bitcast i32 %547 to float, !dbg !29 |
|
%549 = fadd float %545, %548, !dbg !33 |
|
%550 = bitcast float %549 to i32, !dbg !29 |
|
%551 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %550, i32 1, i32 31), !dbg !29 |
|
%552 = bitcast i32 %551 to float, !dbg !29 |
|
%553 = fadd float %549, %552, !dbg !33 |
|
%554 = icmp eq i32 %11, 0, !dbg !29 |
|
%555 = shl nuw nsw i32 %13, 1, !dbg !29 |
|
%556 = or i32 %555, %217, !dbg !29 |
|
%557 = zext nneg i32 %556 to i64, !dbg !29 |
|
%558 = getelementptr float, ptr addrspace(3) @global_smem, i64 %557, !dbg !29 |
|
tail call void asm sideeffect "@$2 st.shared.b32 [ $0 + 0 ], $1;", "r,r,b"(ptr addrspace(3) %558, float %238, i1 %554) #3, !dbg !29 |
|
%559 = shl nuw nsw i32 %14, 1, !dbg !29 |
|
%560 = or i32 %559, %217, !dbg !29 |
|
%561 = zext nneg i32 %560 to i64, !dbg !29 |
|
%562 = getelementptr float, ptr addrspace(3) @global_smem, i64 %561, !dbg !29 |
|
tail call void asm sideeffect "@$2 st.shared.b32 [ $0 + 0 ], $1;", "r,r,b"(ptr addrspace(3) %562, float %259, i1 %554) #3, !dbg !29 |
|
%563 = shl nuw nsw i32 %15, 1, !dbg !29 |
|
%564 = or i32 %563, %217, !dbg !29 |
|
%565 = zext nneg i32 %564 to i64, !dbg !29 |
|
%566 = getelementptr float, ptr addrspace(3) @global_smem, i64 %565, !dbg !29 |
|
tail call void asm sideeffect "@$2 st.shared.b32 [ $0 + 0 ], $1;", "r,r,b"(ptr addrspace(3) %566, float %280, i1 %554) #3, !dbg !29 |
|
%567 = shl nuw nsw i32 %16, 1, !dbg !29 |
|
%568 = or i32 %567, %217, !dbg !29 |
|
%569 = zext nneg i32 %568 to i64, !dbg !29 |
|
%570 = getelementptr float, ptr addrspace(3) @global_smem, i64 %569, !dbg !29 |
|
tail call void asm sideeffect "@$2 st.shared.b32 [ $0 + 0 ], $1;", "r,r,b"(ptr addrspace(3) %570, float %301, i1 %554) #3, !dbg !29 |
|
%571 = shl nuw nsw i32 %17, 1, !dbg !29 |
|
%572 = or i32 %571, %217, !dbg !29 |
|
%573 = zext nneg i32 %572 to i64, !dbg !29 |
|
%574 = getelementptr float, ptr addrspace(3) @global_smem, i64 %573, !dbg !29 |
|
tail call void asm sideeffect "@$2 st.shared.b32 [ $0 + 0 ], $1;", "r,r,b"(ptr addrspace(3) %574, float %322, i1 %554) #3, !dbg !29 |
|
%575 = shl nuw nsw i32 %18, 1, !dbg !29 |
|
%576 = or i32 %575, %217, !dbg !29 |
|
%577 = zext nneg i32 %576 to i64, !dbg !29 |
|
%578 = getelementptr float, ptr addrspace(3) @global_smem, i64 %577, !dbg !29 |
|
tail call void asm sideeffect "@$2 st.shared.b32 [ $0 + 0 ], $1;", "r,r,b"(ptr addrspace(3) %578, float %343, i1 %554) #3, !dbg !29 |
|
%579 = shl nuw nsw i32 %19, 1, !dbg !29 |
|
%580 = or i32 %579, %217, !dbg !29 |
|
%581 = zext nneg i32 %580 to i64, !dbg !29 |
|
%582 = getelementptr float, ptr addrspace(3) @global_smem, i64 %581, !dbg !29 |
|
tail call void asm sideeffect "@$2 st.shared.b32 [ $0 + 0 ], $1;", "r,r,b"(ptr addrspace(3) %582, float %364, i1 %554) #3, !dbg !29 |
|
%583 = shl nuw nsw i32 %20, 1, !dbg !29 |
|
%584 = or i32 %583, %217, !dbg !29 |
|
%585 = zext nneg i32 %584 to i64, !dbg !29 |
|
%586 = getelementptr float, ptr addrspace(3) @global_smem, i64 %585, !dbg !29 |
|
tail call void asm sideeffect "@$2 st.shared.b32 [ $0 + 0 ], $1;", "r,r,b"(ptr addrspace(3) %586, float %385, i1 %554) #3, !dbg !29 |
|
%587 = shl nuw nsw i32 %21, 1, !dbg !29 |
|
%588 = or i32 %587, %217, !dbg !29 |
|
%589 = zext nneg i32 %588 to i64, !dbg !29 |
|
%590 = getelementptr float, ptr addrspace(3) @global_smem, i64 %589, !dbg !29 |
|
tail call void asm sideeffect "@$2 st.shared.b32 [ $0 + 0 ], $1;", "r,r,b"(ptr addrspace(3) %590, float %406, i1 %554) #3, !dbg !29 |
|
%591 = shl nuw nsw i32 %22, 1, !dbg !29 |
|
%592 = or i32 %591, %217, !dbg !29 |
|
%593 = zext nneg i32 %592 to i64, !dbg !29 |
|
%594 = getelementptr float, ptr addrspace(3) @global_smem, i64 %593, !dbg !29 |
|
tail call void asm sideeffect "@$2 st.shared.b32 [ $0 + 0 ], $1;", "r,r,b"(ptr addrspace(3) %594, float %427, i1 %554) #3, !dbg !29 |
|
%595 = shl nuw nsw i32 %23, 1, !dbg !29 |
|
%596 = or i32 %595, %217, !dbg !29 |
|
%597 = zext nneg i32 %596 to i64, !dbg !29 |
|
%598 = getelementptr float, ptr addrspace(3) @global_smem, i64 %597, !dbg !29 |
|
tail call void asm sideeffect "@$2 st.shared.b32 [ $0 + 0 ], $1;", "r,r,b"(ptr addrspace(3) %598, float %448, i1 %554) #3, !dbg !29 |
|
%599 = shl nuw nsw i32 %24, 1, !dbg !29 |
|
%600 = or i32 %599, %217, !dbg !29 |
|
%601 = zext nneg i32 %600 to i64, !dbg !29 |
|
%602 = getelementptr float, ptr addrspace(3) @global_smem, i64 %601, !dbg !29 |
|
tail call void asm sideeffect "@$2 st.shared.b32 [ $0 + 0 ], $1;", "r,r,b"(ptr addrspace(3) %602, float %469, i1 %554) #3, !dbg !29 |
|
%603 = shl nuw nsw i32 %25, 1, !dbg !29 |
|
%604 = or i32 %603, %217, !dbg !29 |
|
%605 = zext nneg i32 %604 to i64, !dbg !29 |
|
%606 = getelementptr float, ptr addrspace(3) @global_smem, i64 %605, !dbg !29 |
|
tail call void asm sideeffect "@$2 st.shared.b32 [ $0 + 0 ], $1;", "r,r,b"(ptr addrspace(3) %606, float %490, i1 %554) #3, !dbg !29 |
|
%607 = shl nuw nsw i32 %26, 1, !dbg !29 |
|
%608 = or i32 %607, %217, !dbg !29 |
|
%609 = zext nneg i32 %608 to i64, !dbg !29 |
|
%610 = getelementptr float, ptr addrspace(3) @global_smem, i64 %609, !dbg !29 |
|
tail call void asm sideeffect "@$2 st.shared.b32 [ $0 + 0 ], $1;", "r,r,b"(ptr addrspace(3) %610, float %511, i1 %554) #3, !dbg !29 |
|
%611 = shl nuw nsw i32 %27, 1, !dbg !29 |
|
%612 = or i32 %611, %217, !dbg !29 |
|
%613 = zext nneg i32 %612 to i64, !dbg !29 |
|
%614 = getelementptr float, ptr addrspace(3) @global_smem, i64 %613, !dbg !29 |
|
tail call void asm sideeffect "@$2 st.shared.b32 [ $0 + 0 ], $1;", "r,r,b"(ptr addrspace(3) %614, float %532, i1 %554) #3, !dbg !29 |
|
%615 = shl nuw nsw i32 %28, 1, !dbg !29 |
|
%616 = or i32 %615, %217, !dbg !29 |
|
%617 = zext nneg i32 %616 to i64, !dbg !29 |
|
%618 = getelementptr float, ptr addrspace(3) @global_smem, i64 %617, !dbg !29 |
|
tail call void asm sideeffect "@$2 st.shared.b32 [ $0 + 0 ], $1;", "r,r,b"(ptr addrspace(3) %618, float %553, i1 %554) #3, !dbg !29 |
|
tail call void @llvm.nvvm.barrier0(), !dbg !29 |
|
%619 = icmp slt i32 %10, 128, !dbg !29 |
|
%620 = sext i32 %10 to i64, !dbg !29 |
|
%621 = getelementptr float, ptr addrspace(3) @global_smem, i64 %620, !dbg !29 |
|
%622 = tail call float asm sideeffect "@$2 ld.shared.b32 $0, [ $1 + 0 ];", "=r,r,b"(ptr addrspace(3) %621, i1 %619) #3, !dbg !29 |
|
%623 = bitcast float %622 to i32, !dbg !29 |
|
%624 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %623, i32 1, i32 31), !dbg !29 |
|
%625 = bitcast i32 %624 to float, !dbg !29 |
|
%626 = fadd float %622, %625, !dbg !33 |
|
%627 = and i32 %10, 1, !dbg !29 |
|
%628 = icmp eq i32 %627, 0, !dbg !29 |
|
%629 = and i1 %619, %628, !dbg !29 |
|
tail call void asm sideeffect "@$2 st.shared.b32 [ $0 + 0 ], $1;", "r,r,b"(ptr addrspace(3) %621, float %626, i1 %629) #3, !dbg !29 |
|
tail call void @llvm.nvvm.barrier0(), !dbg !29 |
|
%630 = zext nneg i32 %555 to i64, !dbg !29 |
|
%631 = getelementptr float, ptr addrspace(3) @global_smem, i64 %630, !dbg !29 |
|
%632 = load float, ptr addrspace(3) %631, align 4, !dbg !29 |
|
%633 = zext nneg i32 %559 to i64, !dbg !29 |
|
%634 = getelementptr float, ptr addrspace(3) @global_smem, i64 %633, !dbg !29 |
|
%635 = load float, ptr addrspace(3) %634, align 4, !dbg !29 |
|
%636 = zext nneg i32 %563 to i64, !dbg !29 |
|
%637 = getelementptr float, ptr addrspace(3) @global_smem, i64 %636, !dbg !29 |
|
%638 = load float, ptr addrspace(3) %637, align 4, !dbg !29 |
|
%639 = zext nneg i32 %567 to i64, !dbg !29 |
|
%640 = getelementptr float, ptr addrspace(3) @global_smem, i64 %639, !dbg !29 |
|
%641 = load float, ptr addrspace(3) %640, align 4, !dbg !29 |
|
%642 = zext nneg i32 %571 to i64, !dbg !29 |
|
%643 = getelementptr float, ptr addrspace(3) @global_smem, i64 %642, !dbg !29 |
|
%644 = load float, ptr addrspace(3) %643, align 4, !dbg !29 |
|
%645 = zext nneg i32 %575 to i64, !dbg !29 |
|
%646 = getelementptr float, ptr addrspace(3) @global_smem, i64 %645, !dbg !29 |
|
%647 = load float, ptr addrspace(3) %646, align 4, !dbg !29 |
|
%648 = zext nneg i32 %579 to i64, !dbg !29 |
|
%649 = getelementptr float, ptr addrspace(3) @global_smem, i64 %648, !dbg !29 |
|
%650 = load float, ptr addrspace(3) %649, align 4, !dbg !29 |
|
%651 = zext nneg i32 %583 to i64, !dbg !29 |
|
%652 = getelementptr float, ptr addrspace(3) @global_smem, i64 %651, !dbg !29 |
|
%653 = load float, ptr addrspace(3) %652, align 4, !dbg !29 |
|
%654 = zext nneg i32 %587 to i64, !dbg !29 |
|
%655 = getelementptr float, ptr addrspace(3) @global_smem, i64 %654, !dbg !29 |
|
%656 = load float, ptr addrspace(3) %655, align 4, !dbg !29 |
|
%657 = zext nneg i32 %591 to i64, !dbg !29 |
|
%658 = getelementptr float, ptr addrspace(3) @global_smem, i64 %657, !dbg !29 |
|
%659 = load float, ptr addrspace(3) %658, align 4, !dbg !29 |
|
%660 = zext nneg i32 %595 to i64, !dbg !29 |
|
%661 = getelementptr float, ptr addrspace(3) @global_smem, i64 %660, !dbg !29 |
|
%662 = load float, ptr addrspace(3) %661, align 4, !dbg !29 |
|
%663 = zext nneg i32 %599 to i64, !dbg !29 |
|
%664 = getelementptr float, ptr addrspace(3) @global_smem, i64 %663, !dbg !29 |
|
%665 = load float, ptr addrspace(3) %664, align 4, !dbg !29 |
|
%666 = zext nneg i32 %603 to i64, !dbg !29 |
|
%667 = getelementptr float, ptr addrspace(3) @global_smem, i64 %666, !dbg !29 |
|
%668 = load float, ptr addrspace(3) %667, align 4, !dbg !29 |
|
%669 = zext nneg i32 %607 to i64, !dbg !29 |
|
%670 = getelementptr float, ptr addrspace(3) @global_smem, i64 %669, !dbg !29 |
|
%671 = load float, ptr addrspace(3) %670, align 4, !dbg !29 |
|
%672 = zext nneg i32 %611 to i64, !dbg !29 |
|
%673 = getelementptr float, ptr addrspace(3) @global_smem, i64 %672, !dbg !29 |
|
%674 = load float, ptr addrspace(3) %673, align 4, !dbg !29 |
|
%675 = zext nneg i32 %615 to i64, !dbg !29 |
|
%676 = getelementptr float, ptr addrspace(3) @global_smem, i64 %675, !dbg !29 |
|
%677 = load float, ptr addrspace(3) %676, align 4, !dbg !29 |
|
%678 = extractelement <16 x float> %136, i64 0, !dbg !37 |
|
%679 = extractelement <16 x float> %136, i64 1, !dbg !37 |
|
%680 = extractelement <16 x float> %136, i64 2, !dbg !37 |
|
%681 = extractelement <16 x float> %136, i64 3, !dbg !37 |
|
%682 = extractelement <16 x float> %136, i64 4, !dbg !37 |
|
%683 = extractelement <16 x float> %136, i64 5, !dbg !37 |
|
%684 = extractelement <16 x float> %136, i64 6, !dbg !37 |
|
%685 = extractelement <16 x float> %136, i64 7, !dbg !37 |
|
%686 = extractelement <16 x float> %136, i64 8, !dbg !37 |
|
%687 = extractelement <16 x float> %136, i64 9, !dbg !37 |
|
%688 = extractelement <16 x float> %136, i64 10, !dbg !37 |
|
%689 = extractelement <16 x float> %136, i64 11, !dbg !37 |
|
%690 = extractelement <16 x float> %136, i64 12, !dbg !37 |
|
%691 = extractelement <16 x float> %136, i64 13, !dbg !37 |
|
%692 = extractelement <16 x float> %136, i64 14, !dbg !37 |
|
%693 = extractelement <16 x float> %136, i64 15, !dbg !37 |
|
br label %694, !dbg !38 |
|
|
|
694: |
|
%695 = phi i32 [ 0, %215 ], [ %987, %694 ] |
|
%696 = or i32 %695, %urem, !dbg !39 |
|
%697 = zext nneg i32 %696 to i64, !dbg !39 |
|
%698 = icmp ult i32 %696, 50257, !dbg !40 |
|
%699 = add nsw i64 %100, %697, !dbg !41 |
|
%700 = add nsw i64 %101, %697, !dbg !41 |
|
%701 = add nsw i64 %102, %697, !dbg !41 |
|
%702 = add nsw i64 %103, %697, !dbg !41 |
|
%703 = add nsw i64 %104, %697, !dbg !41 |
|
%704 = add nsw i64 %105, %697, !dbg !41 |
|
%705 = add nsw i64 %106, %697, !dbg !41 |
|
%706 = add nsw i64 %107, %697, !dbg !41 |
|
%707 = add nsw i64 %108, %697, !dbg !41 |
|
%708 = add nsw i64 %109, %697, !dbg !41 |
|
%709 = add nsw i64 %110, %697, !dbg !41 |
|
%710 = add nsw i64 %111, %697, !dbg !41 |
|
%711 = add nsw i64 %112, %697, !dbg !41 |
|
%712 = add nsw i64 %113, %697, !dbg !41 |
|
%713 = add nsw i64 %114, %697, !dbg !41 |
|
%714 = add nsw i64 %115, %697, !dbg !41 |
|
%715 = getelementptr i16, ptr addrspace(1) %4, i64 %699, !dbg !42 |
|
%716 = getelementptr i16, ptr addrspace(1) %4, i64 %700, !dbg !42 |
|
%717 = getelementptr i16, ptr addrspace(1) %4, i64 %701, !dbg !42 |
|
%718 = getelementptr i16, ptr addrspace(1) %4, i64 %702, !dbg !42 |
|
%719 = getelementptr i16, ptr addrspace(1) %4, i64 %703, !dbg !42 |
|
%720 = getelementptr i16, ptr addrspace(1) %4, i64 %704, !dbg !42 |
|
%721 = getelementptr i16, ptr addrspace(1) %4, i64 %705, !dbg !42 |
|
%722 = getelementptr i16, ptr addrspace(1) %4, i64 %706, !dbg !42 |
|
%723 = getelementptr i16, ptr addrspace(1) %4, i64 %707, !dbg !42 |
|
%724 = getelementptr i16, ptr addrspace(1) %4, i64 %708, !dbg !42 |
|
%725 = getelementptr i16, ptr addrspace(1) %4, i64 %709, !dbg !42 |
|
%726 = getelementptr i16, ptr addrspace(1) %4, i64 %710, !dbg !42 |
|
%727 = getelementptr i16, ptr addrspace(1) %4, i64 %711, !dbg !42 |
|
%728 = getelementptr i16, ptr addrspace(1) %4, i64 %712, !dbg !42 |
|
%729 = getelementptr i16, ptr addrspace(1) %4, i64 %713, !dbg !42 |
|
%730 = getelementptr i16, ptr addrspace(1) %4, i64 %714, !dbg !42 |
|
%731 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %715, i1 %698, i16 0, i1 %698) #3, !dbg !43 |
|
%732 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %716, i1 %698, i16 0, i1 %698) #3, !dbg !43 |
|
%733 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %717, i1 %698, i16 0, i1 %698) #3, !dbg !43 |
|
%734 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %718, i1 %698, i16 0, i1 %698) #3, !dbg !43 |
|
%735 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %719, i1 %698, i16 0, i1 %698) #3, !dbg !43 |
|
%736 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %720, i1 %698, i16 0, i1 %698) #3, !dbg !43 |
|
%737 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %721, i1 %698, i16 0, i1 %698) #3, !dbg !43 |
|
%738 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %722, i1 %698, i16 0, i1 %698) #3, !dbg !43 |
|
%739 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %723, i1 %698, i16 0, i1 %698) #3, !dbg !43 |
|
%740 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %724, i1 %698, i16 0, i1 %698) #3, !dbg !43 |
|
%741 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %725, i1 %698, i16 0, i1 %698) #3, !dbg !43 |
|
%742 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %726, i1 %698, i16 0, i1 %698) #3, !dbg !43 |
|
%743 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %727, i1 %698, i16 0, i1 %698) #3, !dbg !43 |
|
%744 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %728, i1 %698, i16 0, i1 %698) #3, !dbg !43 |
|
%745 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %729, i1 %698, i16 0, i1 %698) #3, !dbg !43 |
|
%746 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %730, i1 %698, i16 0, i1 %698) #3, !dbg !43 |
|
%747 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %731) #3, !dbg !44 |
|
%748 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %732) #3, !dbg !44 |
|
%749 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %733) #3, !dbg !44 |
|
%750 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %734) #3, !dbg !44 |
|
%751 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %735) #3, !dbg !44 |
|
%752 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %736) #3, !dbg !44 |
|
%753 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %737) #3, !dbg !44 |
|
%754 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %738) #3, !dbg !44 |
|
%755 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %739) #3, !dbg !44 |
|
%756 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %740) #3, !dbg !44 |
|
%757 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %741) #3, !dbg !44 |
|
%758 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %742) #3, !dbg !44 |
|
%759 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %743) #3, !dbg !44 |
|
%760 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %744) #3, !dbg !44 |
|
%761 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %745) #3, !dbg !44 |
|
%762 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %746) #3, !dbg !44 |
|
%763 = getelementptr float, ptr addrspace(1) %0, i64 %699, !dbg !45 |
|
%764 = getelementptr float, ptr addrspace(1) %0, i64 %700, !dbg !45 |
|
%765 = getelementptr float, ptr addrspace(1) %0, i64 %701, !dbg !45 |
|
%766 = getelementptr float, ptr addrspace(1) %0, i64 %702, !dbg !45 |
|
%767 = getelementptr float, ptr addrspace(1) %0, i64 %703, !dbg !45 |
|
%768 = getelementptr float, ptr addrspace(1) %0, i64 %704, !dbg !45 |
|
%769 = getelementptr float, ptr addrspace(1) %0, i64 %705, !dbg !45 |
|
%770 = getelementptr float, ptr addrspace(1) %0, i64 %706, !dbg !45 |
|
%771 = getelementptr float, ptr addrspace(1) %0, i64 %707, !dbg !45 |
|
%772 = getelementptr float, ptr addrspace(1) %0, i64 %708, !dbg !45 |
|
%773 = getelementptr float, ptr addrspace(1) %0, i64 %709, !dbg !45 |
|
%774 = getelementptr float, ptr addrspace(1) %0, i64 %710, !dbg !45 |
|
%775 = getelementptr float, ptr addrspace(1) %0, i64 %711, !dbg !45 |
|
%776 = getelementptr float, ptr addrspace(1) %0, i64 %712, !dbg !45 |
|
%777 = getelementptr float, ptr addrspace(1) %0, i64 %713, !dbg !45 |
|
%778 = getelementptr float, ptr addrspace(1) %0, i64 %714, !dbg !45 |
|
%779 = tail call i32 asm sideeffect "mov.u32 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b32 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u32 $0, $3;", "=r,l,b,r,b"(ptr addrspace(1) %763, i1 %698, i32 0, i1 %698) #3, !dbg !46 |
|
%780 = bitcast i32 %779 to float, !dbg !46 |
|
%781 = tail call i32 asm sideeffect "mov.u32 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b32 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u32 $0, $3;", "=r,l,b,r,b"(ptr addrspace(1) %764, i1 %698, i32 0, i1 %698) #3, !dbg !46 |
|
%782 = bitcast i32 %781 to float, !dbg !46 |
|
%783 = tail call i32 asm sideeffect "mov.u32 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b32 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u32 $0, $3;", "=r,l,b,r,b"(ptr addrspace(1) %765, i1 %698, i32 0, i1 %698) #3, !dbg !46 |
|
%784 = bitcast i32 %783 to float, !dbg !46 |
|
%785 = tail call i32 asm sideeffect "mov.u32 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b32 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u32 $0, $3;", "=r,l,b,r,b"(ptr addrspace(1) %766, i1 %698, i32 0, i1 %698) #3, !dbg !46 |
|
%786 = bitcast i32 %785 to float, !dbg !46 |
|
%787 = tail call i32 asm sideeffect "mov.u32 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b32 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u32 $0, $3;", "=r,l,b,r,b"(ptr addrspace(1) %767, i1 %698, i32 0, i1 %698) #3, !dbg !46 |
|
%788 = bitcast i32 %787 to float, !dbg !46 |
|
%789 = tail call i32 asm sideeffect "mov.u32 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b32 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u32 $0, $3;", "=r,l,b,r,b"(ptr addrspace(1) %768, i1 %698, i32 0, i1 %698) #3, !dbg !46 |
|
%790 = bitcast i32 %789 to float, !dbg !46 |
|
%791 = tail call i32 asm sideeffect "mov.u32 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b32 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u32 $0, $3;", "=r,l,b,r,b"(ptr addrspace(1) %769, i1 %698, i32 0, i1 %698) #3, !dbg !46 |
|
%792 = bitcast i32 %791 to float, !dbg !46 |
|
%793 = tail call i32 asm sideeffect "mov.u32 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b32 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u32 $0, $3;", "=r,l,b,r,b"(ptr addrspace(1) %770, i1 %698, i32 0, i1 %698) #3, !dbg !46 |
|
%794 = bitcast i32 %793 to float, !dbg !46 |
|
%795 = tail call i32 asm sideeffect "mov.u32 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b32 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u32 $0, $3;", "=r,l,b,r,b"(ptr addrspace(1) %771, i1 %698, i32 0, i1 %698) #3, !dbg !46 |
|
%796 = bitcast i32 %795 to float, !dbg !46 |
|
%797 = tail call i32 asm sideeffect "mov.u32 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b32 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u32 $0, $3;", "=r,l,b,r,b"(ptr addrspace(1) %772, i1 %698, i32 0, i1 %698) #3, !dbg !46 |
|
%798 = bitcast i32 %797 to float, !dbg !46 |
|
%799 = tail call i32 asm sideeffect "mov.u32 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b32 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u32 $0, $3;", "=r,l,b,r,b"(ptr addrspace(1) %773, i1 %698, i32 0, i1 %698) #3, !dbg !46 |
|
%800 = bitcast i32 %799 to float, !dbg !46 |
|
%801 = tail call i32 asm sideeffect "mov.u32 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b32 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u32 $0, $3;", "=r,l,b,r,b"(ptr addrspace(1) %774, i1 %698, i32 0, i1 %698) #3, !dbg !46 |
|
%802 = bitcast i32 %801 to float, !dbg !46 |
|
%803 = tail call i32 asm sideeffect "mov.u32 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b32 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u32 $0, $3;", "=r,l,b,r,b"(ptr addrspace(1) %775, i1 %698, i32 0, i1 %698) #3, !dbg !46 |
|
%804 = bitcast i32 %803 to float, !dbg !46 |
|
%805 = tail call i32 asm sideeffect "mov.u32 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b32 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u32 $0, $3;", "=r,l,b,r,b"(ptr addrspace(1) %776, i1 %698, i32 0, i1 %698) #3, !dbg !46 |
|
%806 = bitcast i32 %805 to float, !dbg !46 |
|
%807 = tail call i32 asm sideeffect "mov.u32 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b32 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u32 $0, $3;", "=r,l,b,r,b"(ptr addrspace(1) %777, i1 %698, i32 0, i1 %698) #3, !dbg !46 |
|
%808 = bitcast i32 %807 to float, !dbg !46 |
|
%809 = tail call i32 asm sideeffect "mov.u32 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b32 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u32 $0, $3;", "=r,l,b,r,b"(ptr addrspace(1) %778, i1 %698, i32 0, i1 %698) #3, !dbg !46 |
|
%810 = bitcast i32 %809 to float, !dbg !46 |
|
%811 = getelementptr i16, ptr addrspace(1) %5, i64 %699, !dbg !47 |
|
%812 = getelementptr i16, ptr addrspace(1) %5, i64 %700, !dbg !47 |
|
%813 = getelementptr i16, ptr addrspace(1) %5, i64 %701, !dbg !47 |
|
%814 = getelementptr i16, ptr addrspace(1) %5, i64 %702, !dbg !47 |
|
%815 = getelementptr i16, ptr addrspace(1) %5, i64 %703, !dbg !47 |
|
%816 = getelementptr i16, ptr addrspace(1) %5, i64 %704, !dbg !47 |
|
%817 = getelementptr i16, ptr addrspace(1) %5, i64 %705, !dbg !47 |
|
%818 = getelementptr i16, ptr addrspace(1) %5, i64 %706, !dbg !47 |
|
%819 = getelementptr i16, ptr addrspace(1) %5, i64 %707, !dbg !47 |
|
%820 = getelementptr i16, ptr addrspace(1) %5, i64 %708, !dbg !47 |
|
%821 = getelementptr i16, ptr addrspace(1) %5, i64 %709, !dbg !47 |
|
%822 = getelementptr i16, ptr addrspace(1) %5, i64 %710, !dbg !47 |
|
%823 = getelementptr i16, ptr addrspace(1) %5, i64 %711, !dbg !47 |
|
%824 = getelementptr i16, ptr addrspace(1) %5, i64 %712, !dbg !47 |
|
%825 = getelementptr i16, ptr addrspace(1) %5, i64 %713, !dbg !47 |
|
%826 = getelementptr i16, ptr addrspace(1) %5, i64 %714, !dbg !47 |
|
%827 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %811, i1 %698, i16 0, i1 %698) #3, !dbg !48 |
|
%828 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %812, i1 %698, i16 0, i1 %698) #3, !dbg !48 |
|
%829 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %813, i1 %698, i16 0, i1 %698) #3, !dbg !48 |
|
%830 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %814, i1 %698, i16 0, i1 %698) #3, !dbg !48 |
|
%831 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %815, i1 %698, i16 0, i1 %698) #3, !dbg !48 |
|
%832 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %816, i1 %698, i16 0, i1 %698) #3, !dbg !48 |
|
%833 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %817, i1 %698, i16 0, i1 %698) #3, !dbg !48 |
|
%834 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %818, i1 %698, i16 0, i1 %698) #3, !dbg !48 |
|
%835 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %819, i1 %698, i16 0, i1 %698) #3, !dbg !48 |
|
%836 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %820, i1 %698, i16 0, i1 %698) #3, !dbg !48 |
|
%837 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %821, i1 %698, i16 0, i1 %698) #3, !dbg !48 |
|
%838 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %822, i1 %698, i16 0, i1 %698) #3, !dbg !48 |
|
%839 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %823, i1 %698, i16 0, i1 %698) #3, !dbg !48 |
|
%840 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %824, i1 %698, i16 0, i1 %698) #3, !dbg !48 |
|
%841 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %825, i1 %698, i16 0, i1 %698) #3, !dbg !48 |
|
%842 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$2 ld.global.L1::evict_first.b16 { $0 }, [ $1 + 0 ];\0A\09@!$4 mov.u16 $0, $3;", "=c,l,b,c,b"(ptr addrspace(1) %826, i1 %698, i16 0, i1 %698) #3, !dbg !48 |
|
%843 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %827) #3, !dbg !49 |
|
%844 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %828) #3, !dbg !49 |
|
%845 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %829) #3, !dbg !49 |
|
%846 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %830) #3, !dbg !49 |
|
%847 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %831) #3, !dbg !49 |
|
%848 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %832) #3, !dbg !49 |
|
%849 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %833) #3, !dbg !49 |
|
%850 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %834) #3, !dbg !49 |
|
%851 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %835) #3, !dbg !49 |
|
%852 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %836) #3, !dbg !49 |
|
%853 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %837) #3, !dbg !49 |
|
%854 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %838) #3, !dbg !49 |
|
%855 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %839) #3, !dbg !49 |
|
%856 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %840) #3, !dbg !49 |
|
%857 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %841) #3, !dbg !49 |
|
%858 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %842) #3, !dbg !49 |
|
%859 = fmul float %678, %780, !dbg !37 |
|
%860 = fmul float %679, %782, !dbg !37 |
|
%861 = fmul float %680, %784, !dbg !37 |
|
%862 = fmul float %681, %786, !dbg !37 |
|
%863 = fmul float %682, %788, !dbg !37 |
|
%864 = fmul float %683, %790, !dbg !37 |
|
%865 = fmul float %684, %792, !dbg !37 |
|
%866 = fmul float %685, %794, !dbg !37 |
|
%867 = fmul float %686, %796, !dbg !37 |
|
%868 = fmul float %687, %798, !dbg !37 |
|
%869 = fmul float %688, %800, !dbg !37 |
|
%870 = fmul float %689, %802, !dbg !37 |
|
%871 = fmul float %690, %804, !dbg !37 |
|
%872 = fmul float %691, %806, !dbg !37 |
|
%873 = fmul float %692, %808, !dbg !37 |
|
%874 = fmul float %693, %810, !dbg !37 |
|
%875 = fmul float %843, 0x3FF7154760000000, !dbg !50 |
|
%876 = tail call float asm "ex2.approx.f32 $0, $1;", "=f,f"(float %875) #3, !dbg !50 |
|
%877 = fmul float %844, 0x3FF7154760000000, !dbg !50 |
|
%878 = tail call float asm "ex2.approx.f32 $0, $1;", "=f,f"(float %877) #3, !dbg !50 |
|
%879 = fmul float %845, 0x3FF7154760000000, !dbg !50 |
|
%880 = tail call float asm "ex2.approx.f32 $0, $1;", "=f,f"(float %879) #3, !dbg !50 |
|
%881 = fmul float %846, 0x3FF7154760000000, !dbg !50 |
|
%882 = tail call float asm "ex2.approx.f32 $0, $1;", "=f,f"(float %881) #3, !dbg !50 |
|
%883 = fmul float %847, 0x3FF7154760000000, !dbg !50 |
|
%884 = tail call float asm "ex2.approx.f32 $0, $1;", "=f,f"(float %883) #3, !dbg !50 |
|
%885 = fmul float %848, 0x3FF7154760000000, !dbg !50 |
|
%886 = tail call float asm "ex2.approx.f32 $0, $1;", "=f,f"(float %885) #3, !dbg !50 |
|
%887 = fmul float %849, 0x3FF7154760000000, !dbg !50 |
|
%888 = tail call float asm "ex2.approx.f32 $0, $1;", "=f,f"(float %887) #3, !dbg !50 |
|
%889 = fmul float %850, 0x3FF7154760000000, !dbg !50 |
|
%890 = tail call float asm "ex2.approx.f32 $0, $1;", "=f,f"(float %889) #3, !dbg !50 |
|
%891 = fmul float %851, 0x3FF7154760000000, !dbg !50 |
|
%892 = tail call float asm "ex2.approx.f32 $0, $1;", "=f,f"(float %891) #3, !dbg !50 |
|
%893 = fmul float %852, 0x3FF7154760000000, !dbg !50 |
|
%894 = tail call float asm "ex2.approx.f32 $0, $1;", "=f,f"(float %893) #3, !dbg !50 |
|
%895 = fmul float %853, 0x3FF7154760000000, !dbg !50 |
|
%896 = tail call float asm "ex2.approx.f32 $0, $1;", "=f,f"(float %895) #3, !dbg !50 |
|
%897 = fmul float %854, 0x3FF7154760000000, !dbg !50 |
|
%898 = tail call float asm "ex2.approx.f32 $0, $1;", "=f,f"(float %897) #3, !dbg !50 |
|
%899 = fmul float %855, 0x3FF7154760000000, !dbg !50 |
|
%900 = tail call float asm "ex2.approx.f32 $0, $1;", "=f,f"(float %899) #3, !dbg !50 |
|
%901 = fmul float %856, 0x3FF7154760000000, !dbg !50 |
|
%902 = tail call float asm "ex2.approx.f32 $0, $1;", "=f,f"(float %901) #3, !dbg !50 |
|
%903 = fmul float %857, 0x3FF7154760000000, !dbg !50 |
|
%904 = tail call float asm "ex2.approx.f32 $0, $1;", "=f,f"(float %903) #3, !dbg !50 |
|
%905 = fmul float %858, 0x3FF7154760000000, !dbg !50 |
|
%906 = tail call float asm "ex2.approx.f32 $0, $1;", "=f,f"(float %905) #3, !dbg !50 |
|
%907 = fmul float %632, %876, !dbg !51 |
|
%908 = fmul float %635, %878, !dbg !51 |
|
%909 = fmul float %638, %880, !dbg !51 |
|
%910 = fmul float %641, %882, !dbg !51 |
|
%911 = fmul float %644, %884, !dbg !51 |
|
%912 = fmul float %647, %886, !dbg !51 |
|
%913 = fmul float %650, %888, !dbg !51 |
|
%914 = fmul float %653, %890, !dbg !51 |
|
%915 = fmul float %656, %892, !dbg !51 |
|
%916 = fmul float %659, %894, !dbg !51 |
|
%917 = fmul float %662, %896, !dbg !51 |
|
%918 = fmul float %665, %898, !dbg !51 |
|
%919 = fmul float %668, %900, !dbg !51 |
|
%920 = fmul float %671, %902, !dbg !51 |
|
%921 = fmul float %674, %904, !dbg !51 |
|
%922 = fmul float %677, %906, !dbg !51 |
|
%923 = fsub float %859, %907, !dbg !52 |
|
%924 = fsub float %860, %908, !dbg !52 |
|
%925 = fsub float %861, %909, !dbg !52 |
|
%926 = fsub float %862, %910, !dbg !52 |
|
%927 = fsub float %863, %911, !dbg !52 |
|
%928 = fsub float %864, %912, !dbg !52 |
|
%929 = fsub float %865, %913, !dbg !52 |
|
%930 = fsub float %866, %914, !dbg !52 |
|
%931 = fsub float %867, %915, !dbg !52 |
|
%932 = fsub float %868, %916, !dbg !52 |
|
%933 = fsub float %869, %917, !dbg !52 |
|
%934 = fsub float %870, %918, !dbg !52 |
|
%935 = fsub float %871, %919, !dbg !52 |
|
%936 = fsub float %872, %920, !dbg !52 |
|
%937 = fsub float %873, %921, !dbg !52 |
|
%938 = fsub float %874, %922, !dbg !52 |
|
%939 = fadd float %747, %923, !dbg !53 |
|
%940 = fadd float %748, %924, !dbg !53 |
|
%941 = fadd float %749, %925, !dbg !53 |
|
%942 = fadd float %750, %926, !dbg !53 |
|
%943 = fadd float %751, %927, !dbg !53 |
|
%944 = fadd float %752, %928, !dbg !53 |
|
%945 = fadd float %753, %929, !dbg !53 |
|
%946 = fadd float %754, %930, !dbg !53 |
|
%947 = fadd float %755, %931, !dbg !53 |
|
%948 = fadd float %756, %932, !dbg !53 |
|
%949 = fadd float %757, %933, !dbg !53 |
|
%950 = fadd float %758, %934, !dbg !53 |
|
%951 = fadd float %759, %935, !dbg !53 |
|
%952 = fadd float %760, %936, !dbg !53 |
|
%953 = fadd float %761, %937, !dbg !53 |
|
%954 = fadd float %762, %938, !dbg !53 |
|
%955 = getelementptr i16, ptr addrspace(1) %6, i64 %699, !dbg !54 |
|
%956 = getelementptr i16, ptr addrspace(1) %6, i64 %700, !dbg !54 |
|
%957 = getelementptr i16, ptr addrspace(1) %6, i64 %701, !dbg !54 |
|
%958 = getelementptr i16, ptr addrspace(1) %6, i64 %702, !dbg !54 |
|
%959 = getelementptr i16, ptr addrspace(1) %6, i64 %703, !dbg !54 |
|
%960 = getelementptr i16, ptr addrspace(1) %6, i64 %704, !dbg !54 |
|
%961 = getelementptr i16, ptr addrspace(1) %6, i64 %705, !dbg !54 |
|
%962 = getelementptr i16, ptr addrspace(1) %6, i64 %706, !dbg !54 |
|
%963 = getelementptr i16, ptr addrspace(1) %6, i64 %707, !dbg !54 |
|
%964 = getelementptr i16, ptr addrspace(1) %6, i64 %708, !dbg !54 |
|
%965 = getelementptr i16, ptr addrspace(1) %6, i64 %709, !dbg !54 |
|
%966 = getelementptr i16, ptr addrspace(1) %6, i64 %710, !dbg !54 |
|
%967 = getelementptr i16, ptr addrspace(1) %6, i64 %711, !dbg !54 |
|
%968 = getelementptr i16, ptr addrspace(1) %6, i64 %712, !dbg !54 |
|
%969 = getelementptr i16, ptr addrspace(1) %6, i64 %713, !dbg !54 |
|
%970 = getelementptr i16, ptr addrspace(1) %6, i64 %714, !dbg !54 |
|
%971 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %939) #3, !dbg !55 |
|
%972 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %940) #3, !dbg !55 |
|
%973 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %941) #3, !dbg !55 |
|
%974 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %942) #3, !dbg !55 |
|
%975 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %943) #3, !dbg !55 |
|
%976 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %944) #3, !dbg !55 |
|
%977 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %945) #3, !dbg !55 |
|
%978 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %946) #3, !dbg !55 |
|
%979 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %947) #3, !dbg !55 |
|
%980 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %948) #3, !dbg !55 |
|
%981 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %949) #3, !dbg !55 |
|
%982 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %950) #3, !dbg !55 |
|
%983 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %951) #3, !dbg !55 |
|
%984 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %952) #3, !dbg !55 |
|
%985 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %953) #3, !dbg !55 |
|
%986 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %954) #3, !dbg !55 |
|
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %971, ptr addrspace(1) %955, i1 %698) #3, !dbg !55 |
|
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %972, ptr addrspace(1) %956, i1 %698) #3, !dbg !55 |
|
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %973, ptr addrspace(1) %957, i1 %698) #3, !dbg !55 |
|
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %974, ptr addrspace(1) %958, i1 %698) #3, !dbg !55 |
|
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %975, ptr addrspace(1) %959, i1 %698) #3, !dbg !55 |
|
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %976, ptr addrspace(1) %960, i1 %698) #3, !dbg !55 |
|
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %977, ptr addrspace(1) %961, i1 %698) #3, !dbg !55 |
|
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %978, ptr addrspace(1) %962, i1 %698) #3, !dbg !55 |
|
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %979, ptr addrspace(1) %963, i1 %698) #3, !dbg !55 |
|
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %980, ptr addrspace(1) %964, i1 %698) #3, !dbg !55 |
|
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %981, ptr addrspace(1) %965, i1 %698) #3, !dbg !55 |
|
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %982, ptr addrspace(1) %966, i1 %698) #3, !dbg !55 |
|
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %983, ptr addrspace(1) %967, i1 %698) #3, !dbg !55 |
|
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %984, ptr addrspace(1) %968, i1 %698) #3, !dbg !55 |
|
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %985, ptr addrspace(1) %969, i1 %698) #3, !dbg !55 |
|
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %986, ptr addrspace(1) %970, i1 %698) #3, !dbg !55 |
|
%987 = add nuw nsw i32 %695, 64, !dbg !38 |
|
%988 = icmp ult i32 %695, 50193, !dbg !38 |
|
br i1 %988, label %694, label %989, !dbg !38 |
|
|
|
989: |
|
ret void, !dbg !56 |
|
} |
|
|
|
|
|
declare noundef i32 @llvm.nvvm.read.ptx.sreg.tid.x() #0 |
|
|
|
|
|
declare i32 @llvm.nvvm.shfl.sync.bfly.i32(i32, i32, i32, i32) #1 |
|
|
|
|
|
declare void @llvm.nvvm.barrier0() #2 |
|
|
|
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 = { nounwind } |
|
|
|
!llvm.module.flags = !{!0} |
|
!llvm.dbg.cu = !{!1} |
|
!nvvm.annotations = !{!3, !4, !4, !3} |
|
|
|
!0 = !{i32 2, !"Debug Info Version", i32 3} |
|
!1 = distinct !DICompileUnit(language: DW_LANG_C, file: !2, producer: "triton", isOptimized: true, runtimeVersion: 0, emissionKind: FullDebug) |
|
!2 = !DIFile(filename: "ckzgl7thb4xdfkfnd2tidks6mt5f3hauwfyjflbtzyepo5oxkvhk.py", directory: "/tmp/torchinductor_root/kz") |
|
!3 = !{ptr @triton__0d1d2d3d4d5d6d7de8, !"kernel", i32 1} |
|
!4 = !{ptr @triton__0d1d2d3d4d5d6d7de8, !"maxntidx", i32 256} |
|
!5 = distinct !DISubprogram(name: "triton__0d1d2d3d4d5d6d7de8", linkageName: "triton__0d1d2d3d4d5d6d7de8", scope: !2, file: !2, line: 18, type: !6, scopeLine: 18, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !1) |
|
!6 = !DISubroutineType(cc: DW_CC_normal, types: !7) |
|
!7 = !{} |
|
!8 = !DILocation(line: 22, column: 44, scope: !5) |
|
!9 = !DILocation(line: 24, column: 33, scope: !5) |
|
!10 = !DILocation(line: 21, column: 28, scope: !5) |
|
!11 = !DILocation(line: 21, column: 34, scope: !5) |
|
!12 = !DILocation(line: 21, column: 46, scope: !5) |
|
!13 = !DILocation(line: 22, column: 23, scope: !5) |
|
!14 = !DILocation(line: 26, column: 30, scope: !5) |
|
!15 = !DILocation(line: 26, column: 35, scope: !5) |
|
!16 = !DILocation(line: 27, column: 19, scope: !5) |
|
!17 = !DILocation(line: 29, column: 19, scope: !5) |
|
!18 = !DILocation(line: 36, column: 46, scope: !5) |
|
!19 = !DILocation(line: 38, column: 23, scope: !5) |
|
!20 = !DILocation(line: 39, column: 22, scope: !5) |
|
!21 = !DILocation(line: 41, column: 37, scope: !5) |
|
!22 = !DILocation(line: 32, column: 36, scope: !5) |
|
!23 = !DILocation(line: 33, column: 27, scope: !5) |
|
!24 = !DILocation(line: 34, column: 25, scope: !5) |
|
!25 = !DILocation(line: 36, column: 34, scope: !5) |
|
!26 = !DILocation(line: 36, column: 52, scope: !5) |
|
!27 = !DILocation(line: 42, column: 23, scope: !5) |
|
!28 = !DILocation(line: 45, column: 40, scope: !5) |
|
!29 = !DILocation(line: 243, column: 36, scope: !30, inlinedAt: !32) |
|
!30 = distinct !DILexicalBlockFile(scope: !5, file: !31, discriminator: 0) |
|
!31 = !DIFile(filename: "standard.py", directory: "/usr/local/lib/python3.10/dist-packages/triton/language") |
|
!32 = !DILocation(line: 46, column: 27, scope: !30) |
|
!33 = !DILocation(line: 233, column: 15, scope: !34, inlinedAt: !35) |
|
!34 = distinct !DILexicalBlockFile(scope: !30, file: !31, discriminator: 0) |
|
!35 = !DILocation(line: 243, column: 36, scope: !34, inlinedAt: !36) |
|
!36 = !DILocation(line: 46, column: 27, scope: !34) |
|
!37 = !DILocation(line: 63, column: 24, scope: !5) |
|
!38 = !DILocation(line: 51, column: 36, scope: !5) |
|
!39 = !DILocation(line: 52, column: 27, scope: !5) |
|
!40 = !DILocation(line: 53, column: 25, scope: !5) |
|
!41 = !DILocation(line: 55, column: 41, scope: !5) |
|
!42 = !DILocation(line: 55, column: 35, scope: !5) |
|
!43 = !DILocation(line: 55, column: 53, scope: !5) |
|
!44 = !DILocation(line: 55, column: 105, scope: !5) |
|
!45 = !DILocation(line: 56, column: 35, scope: !5) |
|
!46 = !DILocation(line: 56, column: 53, scope: !5) |
|
!47 = !DILocation(line: 57, column: 35, scope: !5) |
|
!48 = !DILocation(line: 57, column: 53, scope: !5) |
|
!49 = !DILocation(line: 57, column: 105, scope: !5) |
|
!50 = !DILocation(line: 65, column: 23, scope: !5) |
|
!51 = !DILocation(line: 66, column: 24, scope: !5) |
|
!52 = !DILocation(line: 67, column: 24, scope: !5) |
|
!53 = !DILocation(line: 69, column: 24, scope: !5) |
|
!54 = !DILocation(line: 70, column: 29, scope: !5) |
|
!55 = !DILocation(line: 70, column: 54, scope: !5) |
|
!56 = !DILocation(line: 51, column: 4, scope: !5) |
|
|
