|
|
|
source_filename = "LLVMDialectModule" |
|
|
|
@.str = private unnamed_addr constant [11 x i8] c"__CUDA_FTZ\00", align 1 |
|
|
|
define void @triton__0d1de(ptr addrspace(1) %0, i32 %1) local_unnamed_addr !dbg !7 { |
|
%3 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !10 |
|
%4 = shl i32 %3, 3, !dbg !10 |
|
%5 = and i32 %4, 1016, !dbg !10 |
|
%6 = tail call i32 asm "mov.u32 $0, %ctaid.x;", "=r"() #4, !dbg !11 |
|
%7 = shl i32 %6, 10, !dbg !12 |
|
%8 = or i32 %7, %5, !dbg !13 |
|
%9 = sext i32 %8 to i64, !dbg !14 |
|
%10 = getelementptr i16, ptr addrspace(1) %0, i64 %9, !dbg !14 |
|
%11 = 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.v4.b32 { $0, $1, $2, $3 }, [ $4 + 0 ];", "=r,=r,=r,=r,l,b"(ptr addrspace(1) %10, i1 true) #4, !dbg !15 |
|
%12 = extractvalue { i32, i32, i32, i32 } %11, 0, !dbg !15 |
|
%13 = extractvalue { i32, i32, i32, i32 } %11, 1, !dbg !15 |
|
%14 = extractvalue { i32, i32, i32, i32 } %11, 2, !dbg !15 |
|
%15 = extractvalue { i32, i32, i32, i32 } %11, 3, !dbg !15 |
|
%16 = trunc i32 %12 to i16, !dbg !15 |
|
%extelt.offset = lshr i32 %12, 16, !dbg !15 |
|
%17 = trunc i32 %extelt.offset to i16, !dbg !15 |
|
%18 = trunc i32 %13 to i16, !dbg !15 |
|
%extelt.offset1 = lshr i32 %13, 16, !dbg !15 |
|
%19 = trunc i32 %extelt.offset1 to i16, !dbg !15 |
|
%20 = trunc i32 %14 to i16, !dbg !15 |
|
%extelt.offset2 = lshr i32 %14, 16, !dbg !15 |
|
%21 = trunc i32 %extelt.offset2 to i16, !dbg !15 |
|
%22 = trunc i32 %15 to i16, !dbg !15 |
|
%extelt.offset3 = lshr i32 %15, 16, !dbg !15 |
|
%23 = trunc i32 %extelt.offset3 to i16, !dbg !15 |
|
%24 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %16) #4, !dbg !16 |
|
%25 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %17) #4, !dbg !16 |
|
%26 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %18) #4, !dbg !16 |
|
%27 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %19) #4, !dbg !16 |
|
%28 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %20) #4, !dbg !16 |
|
%29 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %21) #4, !dbg !16 |
|
%30 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %22) #4, !dbg !16 |
|
%31 = tail call float asm "cvt.f32.bf16 $0, $1;", "=r,h"(i16 %23) #4, !dbg !16 |
|
%32 = fmul float %24, 0x3FE6A09E60000000, !dbg !17 |
|
%33 = fmul float %25, 0x3FE6A09E60000000, !dbg !17 |
|
%34 = fmul float %26, 0x3FE6A09E60000000, !dbg !17 |
|
%35 = fmul float %27, 0x3FE6A09E60000000, !dbg !17 |
|
%36 = fmul float %28, 0x3FE6A09E60000000, !dbg !17 |
|
%37 = fmul float %29, 0x3FE6A09E60000000, !dbg !17 |
|
%38 = fmul float %30, 0x3FE6A09E60000000, !dbg !17 |
|
%39 = fmul float %31, 0x3FE6A09E60000000, !dbg !17 |
|
%40 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not.i = icmp eq i32 %40, 0, !dbg !18 |
|
%41 = tail call float @llvm.nvvm.fabs.ftz.f(float %32) #4, !dbg !18 |
|
%42 = tail call float @llvm.nvvm.fabs.f(float %32) #4, !dbg !18 |
|
%.0.i = select i1 %.not.i, float %42, float %41, !dbg !18 |
|
%43 = fcmp oge float %.0.i, 0x3FF00C1FC0000000, !dbg !18 |
|
br i1 %43, label %__nv_fabsf.exit1.i, label %45, !dbg !18 |
|
|
|
__nv_fabsf.exit1.i: |
|
%44 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not1.i = icmp eq i32 %44, 0, !dbg !18 |
|
%.01.i = select i1 %.not1.i, float %42, float %41, !dbg !18 |
|
br label %__internal_fmad.exit.i, !dbg !18 |
|
|
|
45: |
|
%46 = fmul float %32, %32, !dbg !18 |
|
br label %__internal_fmad.exit.i, !dbg !18 |
|
|
|
__internal_fmad.exit.i: |
|
%47 = phi float [ 0x3FE41B0840000000, %__nv_fabsf.exit1.i ], [ 0x3FC06EBA60000000, %45 ], !dbg !18 |
|
%48 = phi float [ 0x3FED526FC0000000, %__nv_fabsf.exit1.i ], [ 0xBFD8127580000000, %45 ], !dbg !18 |
|
%49 = phi float [ 0x3FC39F20C0000000, %__nv_fabsf.exit1.i ], [ 0x3FBCE315E0000000, %45 ], !dbg !18 |
|
%50 = phi float [ 0xBFA1902C40000000, %__nv_fabsf.exit1.i ], [ 0xBF9B837CE0000000, %45 ], !dbg !18 |
|
%51 = phi float [ 0x3F75908160000000, %__nv_fabsf.exit1.i ], [ 0x3F755ABD40000000, %45 ], !dbg !18 |
|
%52 = phi float [ 0xBF3EAC1720000000, %__nv_fabsf.exit1.i ], [ 0xBF4AE9A400000000, %45 ], !dbg !18 |
|
%53 = phi float [ 0x3EF1394780000000, %__nv_fabsf.exit1.i ], [ 0x3F163D2D40000000, %45 ], !dbg !18 |
|
%54 = phi float [ %.01.i, %__nv_fabsf.exit1.i ], [ %46, %45 ], !dbg !18 |
|
%55 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not2.i = icmp eq i32 %55, 0, !dbg !18 |
|
%56 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %53, float %54, float %52) #4, !dbg !18 |
|
%57 = tail call float @llvm.nvvm.fma.rn.f(float %53, float %54, float %52) #4, !dbg !18 |
|
%.02.i = select i1 %.not2.i, float %57, float %56, !dbg !18 |
|
%58 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not3.i = icmp eq i32 %58, 0, !dbg !18 |
|
%59 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.02.i, float %54, float %51) #4, !dbg !18 |
|
%60 = tail call float @llvm.nvvm.fma.rn.f(float %.02.i, float %54, float %51) #4, !dbg !18 |
|
%.03.i = select i1 %.not3.i, float %60, float %59, !dbg !18 |
|
%61 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not4.i = icmp eq i32 %61, 0, !dbg !18 |
|
%62 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.03.i, float %54, float %50) #4, !dbg !18 |
|
%63 = tail call float @llvm.nvvm.fma.rn.f(float %.03.i, float %54, float %50) #4, !dbg !18 |
|
%.04.i = select i1 %.not4.i, float %63, float %62, !dbg !18 |
|
%64 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not5.i = icmp eq i32 %64, 0, !dbg !18 |
|
%65 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.04.i, float %54, float %49) #4, !dbg !18 |
|
%66 = tail call float @llvm.nvvm.fma.rn.f(float %.04.i, float %54, float %49) #4, !dbg !18 |
|
%.05.i = select i1 %.not5.i, float %66, float %65, !dbg !18 |
|
%67 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not6.i = icmp eq i32 %67, 0, !dbg !18 |
|
%68 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.05.i, float %54, float %48) #4, !dbg !18 |
|
%69 = tail call float @llvm.nvvm.fma.rn.f(float %.05.i, float %54, float %48) #4, !dbg !18 |
|
%.06.i = select i1 %.not6.i, float %69, float %68, !dbg !18 |
|
%70 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not7.i = icmp eq i32 %70, 0, !dbg !18 |
|
%71 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.06.i, float %54, float %47) #4, !dbg !18 |
|
%72 = tail call float @llvm.nvvm.fma.rn.f(float %.06.i, float %54, float %47) #4, !dbg !18 |
|
%.07.i = select i1 %.not7.i, float %72, float %71, !dbg !18 |
|
%73 = fneg float %54, !dbg !18 |
|
%74 = select i1 %43, float %73, float %32, !dbg !18 |
|
%75 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not8.i = icmp eq i32 %75, 0, !dbg !18 |
|
%76 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.07.i, float %74, float %74) #4, !dbg !18 |
|
%77 = tail call float @llvm.nvvm.fma.rn.f(float %.07.i, float %74, float %74) #4, !dbg !18 |
|
%.08.i = select i1 %.not8.i, float %77, float %76, !dbg !18 |
|
br i1 %43, label %78, label %__nv_erff.exit, !dbg !18 |
|
|
|
78: |
|
%79 = tail call float @llvm.nvvm.ex2.approx.ftz.f(float %.08.i) #4, !dbg !18 |
|
%80 = fsub float 1.000000e+00, %79, !dbg !18 |
|
%81 = bitcast float %80 to i32, !dbg !18 |
|
%82 = bitcast float %32 to i32, !dbg !18 |
|
%83 = and i32 %82, -2147483648, !dbg !18 |
|
%84 = or i32 %83, %81, !dbg !18 |
|
%85 = bitcast i32 %84 to float, !dbg !18 |
|
br label %__nv_erff.exit, !dbg !18 |
|
|
|
__nv_erff.exit: |
|
%r.0.i = phi float [ %85, %78 ], [ %.08.i, %__internal_fmad.exit.i ], !dbg !18 |
|
%86 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not.i4 = icmp eq i32 %86, 0, !dbg !18 |
|
%87 = tail call float @llvm.nvvm.fabs.ftz.f(float %33) #4, !dbg !18 |
|
%88 = tail call float @llvm.nvvm.fabs.f(float %33) #4, !dbg !18 |
|
%.0.i5 = select i1 %.not.i4, float %88, float %87, !dbg !18 |
|
%89 = fcmp oge float %.0.i5, 0x3FF00C1FC0000000, !dbg !18 |
|
br i1 %89, label %__nv_fabsf.exit1.i22, label %91, !dbg !18 |
|
|
|
__nv_fabsf.exit1.i22: |
|
%90 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not1.i23 = icmp eq i32 %90, 0, !dbg !18 |
|
%.01.i24 = select i1 %.not1.i23, float %88, float %87, !dbg !18 |
|
br label %__internal_fmad.exit.i6, !dbg !18 |
|
|
|
91: |
|
%92 = fmul float %33, %33, !dbg !18 |
|
br label %__internal_fmad.exit.i6, !dbg !18 |
|
|
|
__internal_fmad.exit.i6: |
|
%93 = phi float [ 0x3FE41B0840000000, %__nv_fabsf.exit1.i22 ], [ 0x3FC06EBA60000000, %91 ], !dbg !18 |
|
%94 = phi float [ 0x3FED526FC0000000, %__nv_fabsf.exit1.i22 ], [ 0xBFD8127580000000, %91 ], !dbg !18 |
|
%95 = phi float [ 0x3FC39F20C0000000, %__nv_fabsf.exit1.i22 ], [ 0x3FBCE315E0000000, %91 ], !dbg !18 |
|
%96 = phi float [ 0xBFA1902C40000000, %__nv_fabsf.exit1.i22 ], [ 0xBF9B837CE0000000, %91 ], !dbg !18 |
|
%97 = phi float [ 0x3F75908160000000, %__nv_fabsf.exit1.i22 ], [ 0x3F755ABD40000000, %91 ], !dbg !18 |
|
%98 = phi float [ 0xBF3EAC1720000000, %__nv_fabsf.exit1.i22 ], [ 0xBF4AE9A400000000, %91 ], !dbg !18 |
|
%99 = phi float [ 0x3EF1394780000000, %__nv_fabsf.exit1.i22 ], [ 0x3F163D2D40000000, %91 ], !dbg !18 |
|
%100 = phi float [ %.01.i24, %__nv_fabsf.exit1.i22 ], [ %92, %91 ], !dbg !18 |
|
%101 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not2.i7 = icmp eq i32 %101, 0, !dbg !18 |
|
%102 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %99, float %100, float %98) #4, !dbg !18 |
|
%103 = tail call float @llvm.nvvm.fma.rn.f(float %99, float %100, float %98) #4, !dbg !18 |
|
%.02.i8 = select i1 %.not2.i7, float %103, float %102, !dbg !18 |
|
%104 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not3.i9 = icmp eq i32 %104, 0, !dbg !18 |
|
%105 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.02.i8, float %100, float %97) #4, !dbg !18 |
|
%106 = tail call float @llvm.nvvm.fma.rn.f(float %.02.i8, float %100, float %97) #4, !dbg !18 |
|
%.03.i10 = select i1 %.not3.i9, float %106, float %105, !dbg !18 |
|
%107 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not4.i11 = icmp eq i32 %107, 0, !dbg !18 |
|
%108 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.03.i10, float %100, float %96) #4, !dbg !18 |
|
%109 = tail call float @llvm.nvvm.fma.rn.f(float %.03.i10, float %100, float %96) #4, !dbg !18 |
|
%.04.i12 = select i1 %.not4.i11, float %109, float %108, !dbg !18 |
|
%110 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not5.i13 = icmp eq i32 %110, 0, !dbg !18 |
|
%111 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.04.i12, float %100, float %95) #4, !dbg !18 |
|
%112 = tail call float @llvm.nvvm.fma.rn.f(float %.04.i12, float %100, float %95) #4, !dbg !18 |
|
%.05.i14 = select i1 %.not5.i13, float %112, float %111, !dbg !18 |
|
%113 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not6.i15 = icmp eq i32 %113, 0, !dbg !18 |
|
%114 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.05.i14, float %100, float %94) #4, !dbg !18 |
|
%115 = tail call float @llvm.nvvm.fma.rn.f(float %.05.i14, float %100, float %94) #4, !dbg !18 |
|
%.06.i16 = select i1 %.not6.i15, float %115, float %114, !dbg !18 |
|
%116 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not7.i17 = icmp eq i32 %116, 0, !dbg !18 |
|
%117 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.06.i16, float %100, float %93) #4, !dbg !18 |
|
%118 = tail call float @llvm.nvvm.fma.rn.f(float %.06.i16, float %100, float %93) #4, !dbg !18 |
|
%.07.i18 = select i1 %.not7.i17, float %118, float %117, !dbg !18 |
|
%119 = fneg float %100, !dbg !18 |
|
%120 = select i1 %89, float %119, float %33, !dbg !18 |
|
%121 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not8.i19 = icmp eq i32 %121, 0, !dbg !18 |
|
%122 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.07.i18, float %120, float %120) #4, !dbg !18 |
|
%123 = tail call float @llvm.nvvm.fma.rn.f(float %.07.i18, float %120, float %120) #4, !dbg !18 |
|
%.08.i20 = select i1 %.not8.i19, float %123, float %122, !dbg !18 |
|
br i1 %89, label %124, label %__nv_erff.exit25, !dbg !18 |
|
|
|
124: |
|
%125 = tail call float @llvm.nvvm.ex2.approx.ftz.f(float %.08.i20) #4, !dbg !18 |
|
%126 = fsub float 1.000000e+00, %125, !dbg !18 |
|
%127 = bitcast float %126 to i32, !dbg !18 |
|
%128 = bitcast float %33 to i32, !dbg !18 |
|
%129 = and i32 %128, -2147483648, !dbg !18 |
|
%130 = or i32 %129, %127, !dbg !18 |
|
%131 = bitcast i32 %130 to float, !dbg !18 |
|
br label %__nv_erff.exit25, !dbg !18 |
|
|
|
__nv_erff.exit25: |
|
%r.0.i21 = phi float [ %131, %124 ], [ %.08.i20, %__internal_fmad.exit.i6 ], !dbg !18 |
|
%132 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not.i26 = icmp eq i32 %132, 0, !dbg !18 |
|
%133 = tail call float @llvm.nvvm.fabs.ftz.f(float %34) #4, !dbg !18 |
|
%134 = tail call float @llvm.nvvm.fabs.f(float %34) #4, !dbg !18 |
|
%.0.i27 = select i1 %.not.i26, float %134, float %133, !dbg !18 |
|
%135 = fcmp oge float %.0.i27, 0x3FF00C1FC0000000, !dbg !18 |
|
br i1 %135, label %__nv_fabsf.exit1.i44, label %137, !dbg !18 |
|
|
|
__nv_fabsf.exit1.i44: |
|
%136 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not1.i45 = icmp eq i32 %136, 0, !dbg !18 |
|
%.01.i46 = select i1 %.not1.i45, float %134, float %133, !dbg !18 |
|
br label %__internal_fmad.exit.i28, !dbg !18 |
|
|
|
137: |
|
%138 = fmul float %34, %34, !dbg !18 |
|
br label %__internal_fmad.exit.i28, !dbg !18 |
|
|
|
__internal_fmad.exit.i28: |
|
%139 = phi float [ 0x3FE41B0840000000, %__nv_fabsf.exit1.i44 ], [ 0x3FC06EBA60000000, %137 ], !dbg !18 |
|
%140 = phi float [ 0x3FED526FC0000000, %__nv_fabsf.exit1.i44 ], [ 0xBFD8127580000000, %137 ], !dbg !18 |
|
%141 = phi float [ 0x3FC39F20C0000000, %__nv_fabsf.exit1.i44 ], [ 0x3FBCE315E0000000, %137 ], !dbg !18 |
|
%142 = phi float [ 0xBFA1902C40000000, %__nv_fabsf.exit1.i44 ], [ 0xBF9B837CE0000000, %137 ], !dbg !18 |
|
%143 = phi float [ 0x3F75908160000000, %__nv_fabsf.exit1.i44 ], [ 0x3F755ABD40000000, %137 ], !dbg !18 |
|
%144 = phi float [ 0xBF3EAC1720000000, %__nv_fabsf.exit1.i44 ], [ 0xBF4AE9A400000000, %137 ], !dbg !18 |
|
%145 = phi float [ 0x3EF1394780000000, %__nv_fabsf.exit1.i44 ], [ 0x3F163D2D40000000, %137 ], !dbg !18 |
|
%146 = phi float [ %.01.i46, %__nv_fabsf.exit1.i44 ], [ %138, %137 ], !dbg !18 |
|
%147 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not2.i29 = icmp eq i32 %147, 0, !dbg !18 |
|
%148 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %145, float %146, float %144) #4, !dbg !18 |
|
%149 = tail call float @llvm.nvvm.fma.rn.f(float %145, float %146, float %144) #4, !dbg !18 |
|
%.02.i30 = select i1 %.not2.i29, float %149, float %148, !dbg !18 |
|
%150 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not3.i31 = icmp eq i32 %150, 0, !dbg !18 |
|
%151 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.02.i30, float %146, float %143) #4, !dbg !18 |
|
%152 = tail call float @llvm.nvvm.fma.rn.f(float %.02.i30, float %146, float %143) #4, !dbg !18 |
|
%.03.i32 = select i1 %.not3.i31, float %152, float %151, !dbg !18 |
|
%153 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not4.i33 = icmp eq i32 %153, 0, !dbg !18 |
|
%154 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.03.i32, float %146, float %142) #4, !dbg !18 |
|
%155 = tail call float @llvm.nvvm.fma.rn.f(float %.03.i32, float %146, float %142) #4, !dbg !18 |
|
%.04.i34 = select i1 %.not4.i33, float %155, float %154, !dbg !18 |
|
%156 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not5.i35 = icmp eq i32 %156, 0, !dbg !18 |
|
%157 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.04.i34, float %146, float %141) #4, !dbg !18 |
|
%158 = tail call float @llvm.nvvm.fma.rn.f(float %.04.i34, float %146, float %141) #4, !dbg !18 |
|
%.05.i36 = select i1 %.not5.i35, float %158, float %157, !dbg !18 |
|
%159 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not6.i37 = icmp eq i32 %159, 0, !dbg !18 |
|
%160 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.05.i36, float %146, float %140) #4, !dbg !18 |
|
%161 = tail call float @llvm.nvvm.fma.rn.f(float %.05.i36, float %146, float %140) #4, !dbg !18 |
|
%.06.i38 = select i1 %.not6.i37, float %161, float %160, !dbg !18 |
|
%162 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not7.i39 = icmp eq i32 %162, 0, !dbg !18 |
|
%163 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.06.i38, float %146, float %139) #4, !dbg !18 |
|
%164 = tail call float @llvm.nvvm.fma.rn.f(float %.06.i38, float %146, float %139) #4, !dbg !18 |
|
%.07.i40 = select i1 %.not7.i39, float %164, float %163, !dbg !18 |
|
%165 = fneg float %146, !dbg !18 |
|
%166 = select i1 %135, float %165, float %34, !dbg !18 |
|
%167 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not8.i41 = icmp eq i32 %167, 0, !dbg !18 |
|
%168 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.07.i40, float %166, float %166) #4, !dbg !18 |
|
%169 = tail call float @llvm.nvvm.fma.rn.f(float %.07.i40, float %166, float %166) #4, !dbg !18 |
|
%.08.i42 = select i1 %.not8.i41, float %169, float %168, !dbg !18 |
|
br i1 %135, label %170, label %__nv_erff.exit47, !dbg !18 |
|
|
|
170: |
|
%171 = tail call float @llvm.nvvm.ex2.approx.ftz.f(float %.08.i42) #4, !dbg !18 |
|
%172 = fsub float 1.000000e+00, %171, !dbg !18 |
|
%173 = bitcast float %172 to i32, !dbg !18 |
|
%174 = bitcast float %34 to i32, !dbg !18 |
|
%175 = and i32 %174, -2147483648, !dbg !18 |
|
%176 = or i32 %175, %173, !dbg !18 |
|
%177 = bitcast i32 %176 to float, !dbg !18 |
|
br label %__nv_erff.exit47, !dbg !18 |
|
|
|
__nv_erff.exit47: |
|
%r.0.i43 = phi float [ %177, %170 ], [ %.08.i42, %__internal_fmad.exit.i28 ], !dbg !18 |
|
%178 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not.i48 = icmp eq i32 %178, 0, !dbg !18 |
|
%179 = tail call float @llvm.nvvm.fabs.ftz.f(float %35) #4, !dbg !18 |
|
%180 = tail call float @llvm.nvvm.fabs.f(float %35) #4, !dbg !18 |
|
%.0.i49 = select i1 %.not.i48, float %180, float %179, !dbg !18 |
|
%181 = fcmp oge float %.0.i49, 0x3FF00C1FC0000000, !dbg !18 |
|
br i1 %181, label %__nv_fabsf.exit1.i66, label %183, !dbg !18 |
|
|
|
__nv_fabsf.exit1.i66: |
|
%182 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not1.i67 = icmp eq i32 %182, 0, !dbg !18 |
|
%.01.i68 = select i1 %.not1.i67, float %180, float %179, !dbg !18 |
|
br label %__internal_fmad.exit.i50, !dbg !18 |
|
|
|
183: |
|
%184 = fmul float %35, %35, !dbg !18 |
|
br label %__internal_fmad.exit.i50, !dbg !18 |
|
|
|
__internal_fmad.exit.i50: |
|
%185 = phi float [ 0x3FE41B0840000000, %__nv_fabsf.exit1.i66 ], [ 0x3FC06EBA60000000, %183 ], !dbg !18 |
|
%186 = phi float [ 0x3FED526FC0000000, %__nv_fabsf.exit1.i66 ], [ 0xBFD8127580000000, %183 ], !dbg !18 |
|
%187 = phi float [ 0x3FC39F20C0000000, %__nv_fabsf.exit1.i66 ], [ 0x3FBCE315E0000000, %183 ], !dbg !18 |
|
%188 = phi float [ 0xBFA1902C40000000, %__nv_fabsf.exit1.i66 ], [ 0xBF9B837CE0000000, %183 ], !dbg !18 |
|
%189 = phi float [ 0x3F75908160000000, %__nv_fabsf.exit1.i66 ], [ 0x3F755ABD40000000, %183 ], !dbg !18 |
|
%190 = phi float [ 0xBF3EAC1720000000, %__nv_fabsf.exit1.i66 ], [ 0xBF4AE9A400000000, %183 ], !dbg !18 |
|
%191 = phi float [ 0x3EF1394780000000, %__nv_fabsf.exit1.i66 ], [ 0x3F163D2D40000000, %183 ], !dbg !18 |
|
%192 = phi float [ %.01.i68, %__nv_fabsf.exit1.i66 ], [ %184, %183 ], !dbg !18 |
|
%193 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not2.i51 = icmp eq i32 %193, 0, !dbg !18 |
|
%194 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %191, float %192, float %190) #4, !dbg !18 |
|
%195 = tail call float @llvm.nvvm.fma.rn.f(float %191, float %192, float %190) #4, !dbg !18 |
|
%.02.i52 = select i1 %.not2.i51, float %195, float %194, !dbg !18 |
|
%196 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not3.i53 = icmp eq i32 %196, 0, !dbg !18 |
|
%197 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.02.i52, float %192, float %189) #4, !dbg !18 |
|
%198 = tail call float @llvm.nvvm.fma.rn.f(float %.02.i52, float %192, float %189) #4, !dbg !18 |
|
%.03.i54 = select i1 %.not3.i53, float %198, float %197, !dbg !18 |
|
%199 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not4.i55 = icmp eq i32 %199, 0, !dbg !18 |
|
%200 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.03.i54, float %192, float %188) #4, !dbg !18 |
|
%201 = tail call float @llvm.nvvm.fma.rn.f(float %.03.i54, float %192, float %188) #4, !dbg !18 |
|
%.04.i56 = select i1 %.not4.i55, float %201, float %200, !dbg !18 |
|
%202 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not5.i57 = icmp eq i32 %202, 0, !dbg !18 |
|
%203 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.04.i56, float %192, float %187) #4, !dbg !18 |
|
%204 = tail call float @llvm.nvvm.fma.rn.f(float %.04.i56, float %192, float %187) #4, !dbg !18 |
|
%.05.i58 = select i1 %.not5.i57, float %204, float %203, !dbg !18 |
|
%205 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not6.i59 = icmp eq i32 %205, 0, !dbg !18 |
|
%206 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.05.i58, float %192, float %186) #4, !dbg !18 |
|
%207 = tail call float @llvm.nvvm.fma.rn.f(float %.05.i58, float %192, float %186) #4, !dbg !18 |
|
%.06.i60 = select i1 %.not6.i59, float %207, float %206, !dbg !18 |
|
%208 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not7.i61 = icmp eq i32 %208, 0, !dbg !18 |
|
%209 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.06.i60, float %192, float %185) #4, !dbg !18 |
|
%210 = tail call float @llvm.nvvm.fma.rn.f(float %.06.i60, float %192, float %185) #4, !dbg !18 |
|
%.07.i62 = select i1 %.not7.i61, float %210, float %209, !dbg !18 |
|
%211 = fneg float %192, !dbg !18 |
|
%212 = select i1 %181, float %211, float %35, !dbg !18 |
|
%213 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not8.i63 = icmp eq i32 %213, 0, !dbg !18 |
|
%214 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.07.i62, float %212, float %212) #4, !dbg !18 |
|
%215 = tail call float @llvm.nvvm.fma.rn.f(float %.07.i62, float %212, float %212) #4, !dbg !18 |
|
%.08.i64 = select i1 %.not8.i63, float %215, float %214, !dbg !18 |
|
br i1 %181, label %216, label %__nv_erff.exit69, !dbg !18 |
|
|
|
216: |
|
%217 = tail call float @llvm.nvvm.ex2.approx.ftz.f(float %.08.i64) #4, !dbg !18 |
|
%218 = fsub float 1.000000e+00, %217, !dbg !18 |
|
%219 = bitcast float %218 to i32, !dbg !18 |
|
%220 = bitcast float %35 to i32, !dbg !18 |
|
%221 = and i32 %220, -2147483648, !dbg !18 |
|
%222 = or i32 %221, %219, !dbg !18 |
|
%223 = bitcast i32 %222 to float, !dbg !18 |
|
br label %__nv_erff.exit69, !dbg !18 |
|
|
|
__nv_erff.exit69: |
|
%r.0.i65 = phi float [ %223, %216 ], [ %.08.i64, %__internal_fmad.exit.i50 ], !dbg !18 |
|
%224 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not.i70 = icmp eq i32 %224, 0, !dbg !18 |
|
%225 = tail call float @llvm.nvvm.fabs.ftz.f(float %36) #4, !dbg !18 |
|
%226 = tail call float @llvm.nvvm.fabs.f(float %36) #4, !dbg !18 |
|
%.0.i71 = select i1 %.not.i70, float %226, float %225, !dbg !18 |
|
%227 = fcmp oge float %.0.i71, 0x3FF00C1FC0000000, !dbg !18 |
|
br i1 %227, label %__nv_fabsf.exit1.i88, label %229, !dbg !18 |
|
|
|
__nv_fabsf.exit1.i88: |
|
%228 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not1.i89 = icmp eq i32 %228, 0, !dbg !18 |
|
%.01.i90 = select i1 %.not1.i89, float %226, float %225, !dbg !18 |
|
br label %__internal_fmad.exit.i72, !dbg !18 |
|
|
|
229: |
|
%230 = fmul float %36, %36, !dbg !18 |
|
br label %__internal_fmad.exit.i72, !dbg !18 |
|
|
|
__internal_fmad.exit.i72: |
|
%231 = phi float [ 0x3FE41B0840000000, %__nv_fabsf.exit1.i88 ], [ 0x3FC06EBA60000000, %229 ], !dbg !18 |
|
%232 = phi float [ 0x3FED526FC0000000, %__nv_fabsf.exit1.i88 ], [ 0xBFD8127580000000, %229 ], !dbg !18 |
|
%233 = phi float [ 0x3FC39F20C0000000, %__nv_fabsf.exit1.i88 ], [ 0x3FBCE315E0000000, %229 ], !dbg !18 |
|
%234 = phi float [ 0xBFA1902C40000000, %__nv_fabsf.exit1.i88 ], [ 0xBF9B837CE0000000, %229 ], !dbg !18 |
|
%235 = phi float [ 0x3F75908160000000, %__nv_fabsf.exit1.i88 ], [ 0x3F755ABD40000000, %229 ], !dbg !18 |
|
%236 = phi float [ 0xBF3EAC1720000000, %__nv_fabsf.exit1.i88 ], [ 0xBF4AE9A400000000, %229 ], !dbg !18 |
|
%237 = phi float [ 0x3EF1394780000000, %__nv_fabsf.exit1.i88 ], [ 0x3F163D2D40000000, %229 ], !dbg !18 |
|
%238 = phi float [ %.01.i90, %__nv_fabsf.exit1.i88 ], [ %230, %229 ], !dbg !18 |
|
%239 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not2.i73 = icmp eq i32 %239, 0, !dbg !18 |
|
%240 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %237, float %238, float %236) #4, !dbg !18 |
|
%241 = tail call float @llvm.nvvm.fma.rn.f(float %237, float %238, float %236) #4, !dbg !18 |
|
%.02.i74 = select i1 %.not2.i73, float %241, float %240, !dbg !18 |
|
%242 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not3.i75 = icmp eq i32 %242, 0, !dbg !18 |
|
%243 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.02.i74, float %238, float %235) #4, !dbg !18 |
|
%244 = tail call float @llvm.nvvm.fma.rn.f(float %.02.i74, float %238, float %235) #4, !dbg !18 |
|
%.03.i76 = select i1 %.not3.i75, float %244, float %243, !dbg !18 |
|
%245 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not4.i77 = icmp eq i32 %245, 0, !dbg !18 |
|
%246 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.03.i76, float %238, float %234) #4, !dbg !18 |
|
%247 = tail call float @llvm.nvvm.fma.rn.f(float %.03.i76, float %238, float %234) #4, !dbg !18 |
|
%.04.i78 = select i1 %.not4.i77, float %247, float %246, !dbg !18 |
|
%248 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not5.i79 = icmp eq i32 %248, 0, !dbg !18 |
|
%249 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.04.i78, float %238, float %233) #4, !dbg !18 |
|
%250 = tail call float @llvm.nvvm.fma.rn.f(float %.04.i78, float %238, float %233) #4, !dbg !18 |
|
%.05.i80 = select i1 %.not5.i79, float %250, float %249, !dbg !18 |
|
%251 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not6.i81 = icmp eq i32 %251, 0, !dbg !18 |
|
%252 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.05.i80, float %238, float %232) #4, !dbg !18 |
|
%253 = tail call float @llvm.nvvm.fma.rn.f(float %.05.i80, float %238, float %232) #4, !dbg !18 |
|
%.06.i82 = select i1 %.not6.i81, float %253, float %252, !dbg !18 |
|
%254 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not7.i83 = icmp eq i32 %254, 0, !dbg !18 |
|
%255 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.06.i82, float %238, float %231) #4, !dbg !18 |
|
%256 = tail call float @llvm.nvvm.fma.rn.f(float %.06.i82, float %238, float %231) #4, !dbg !18 |
|
%.07.i84 = select i1 %.not7.i83, float %256, float %255, !dbg !18 |
|
%257 = fneg float %238, !dbg !18 |
|
%258 = select i1 %227, float %257, float %36, !dbg !18 |
|
%259 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not8.i85 = icmp eq i32 %259, 0, !dbg !18 |
|
%260 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.07.i84, float %258, float %258) #4, !dbg !18 |
|
%261 = tail call float @llvm.nvvm.fma.rn.f(float %.07.i84, float %258, float %258) #4, !dbg !18 |
|
%.08.i86 = select i1 %.not8.i85, float %261, float %260, !dbg !18 |
|
br i1 %227, label %262, label %__nv_erff.exit91, !dbg !18 |
|
|
|
262: |
|
%263 = tail call float @llvm.nvvm.ex2.approx.ftz.f(float %.08.i86) #4, !dbg !18 |
|
%264 = fsub float 1.000000e+00, %263, !dbg !18 |
|
%265 = bitcast float %264 to i32, !dbg !18 |
|
%266 = bitcast float %36 to i32, !dbg !18 |
|
%267 = and i32 %266, -2147483648, !dbg !18 |
|
%268 = or i32 %267, %265, !dbg !18 |
|
%269 = bitcast i32 %268 to float, !dbg !18 |
|
br label %__nv_erff.exit91, !dbg !18 |
|
|
|
__nv_erff.exit91: |
|
%r.0.i87 = phi float [ %269, %262 ], [ %.08.i86, %__internal_fmad.exit.i72 ], !dbg !18 |
|
%270 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not.i92 = icmp eq i32 %270, 0, !dbg !18 |
|
%271 = tail call float @llvm.nvvm.fabs.ftz.f(float %37) #4, !dbg !18 |
|
%272 = tail call float @llvm.nvvm.fabs.f(float %37) #4, !dbg !18 |
|
%.0.i93 = select i1 %.not.i92, float %272, float %271, !dbg !18 |
|
%273 = fcmp oge float %.0.i93, 0x3FF00C1FC0000000, !dbg !18 |
|
br i1 %273, label %__nv_fabsf.exit1.i110, label %275, !dbg !18 |
|
|
|
__nv_fabsf.exit1.i110: |
|
%274 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not1.i111 = icmp eq i32 %274, 0, !dbg !18 |
|
%.01.i112 = select i1 %.not1.i111, float %272, float %271, !dbg !18 |
|
br label %__internal_fmad.exit.i94, !dbg !18 |
|
|
|
275: |
|
%276 = fmul float %37, %37, !dbg !18 |
|
br label %__internal_fmad.exit.i94, !dbg !18 |
|
|
|
__internal_fmad.exit.i94: |
|
%277 = phi float [ 0x3FE41B0840000000, %__nv_fabsf.exit1.i110 ], [ 0x3FC06EBA60000000, %275 ], !dbg !18 |
|
%278 = phi float [ 0x3FED526FC0000000, %__nv_fabsf.exit1.i110 ], [ 0xBFD8127580000000, %275 ], !dbg !18 |
|
%279 = phi float [ 0x3FC39F20C0000000, %__nv_fabsf.exit1.i110 ], [ 0x3FBCE315E0000000, %275 ], !dbg !18 |
|
%280 = phi float [ 0xBFA1902C40000000, %__nv_fabsf.exit1.i110 ], [ 0xBF9B837CE0000000, %275 ], !dbg !18 |
|
%281 = phi float [ 0x3F75908160000000, %__nv_fabsf.exit1.i110 ], [ 0x3F755ABD40000000, %275 ], !dbg !18 |
|
%282 = phi float [ 0xBF3EAC1720000000, %__nv_fabsf.exit1.i110 ], [ 0xBF4AE9A400000000, %275 ], !dbg !18 |
|
%283 = phi float [ 0x3EF1394780000000, %__nv_fabsf.exit1.i110 ], [ 0x3F163D2D40000000, %275 ], !dbg !18 |
|
%284 = phi float [ %.01.i112, %__nv_fabsf.exit1.i110 ], [ %276, %275 ], !dbg !18 |
|
%285 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not2.i95 = icmp eq i32 %285, 0, !dbg !18 |
|
%286 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %283, float %284, float %282) #4, !dbg !18 |
|
%287 = tail call float @llvm.nvvm.fma.rn.f(float %283, float %284, float %282) #4, !dbg !18 |
|
%.02.i96 = select i1 %.not2.i95, float %287, float %286, !dbg !18 |
|
%288 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not3.i97 = icmp eq i32 %288, 0, !dbg !18 |
|
%289 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.02.i96, float %284, float %281) #4, !dbg !18 |
|
%290 = tail call float @llvm.nvvm.fma.rn.f(float %.02.i96, float %284, float %281) #4, !dbg !18 |
|
%.03.i98 = select i1 %.not3.i97, float %290, float %289, !dbg !18 |
|
%291 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not4.i99 = icmp eq i32 %291, 0, !dbg !18 |
|
%292 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.03.i98, float %284, float %280) #4, !dbg !18 |
|
%293 = tail call float @llvm.nvvm.fma.rn.f(float %.03.i98, float %284, float %280) #4, !dbg !18 |
|
%.04.i100 = select i1 %.not4.i99, float %293, float %292, !dbg !18 |
|
%294 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not5.i101 = icmp eq i32 %294, 0, !dbg !18 |
|
%295 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.04.i100, float %284, float %279) #4, !dbg !18 |
|
%296 = tail call float @llvm.nvvm.fma.rn.f(float %.04.i100, float %284, float %279) #4, !dbg !18 |
|
%.05.i102 = select i1 %.not5.i101, float %296, float %295, !dbg !18 |
|
%297 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not6.i103 = icmp eq i32 %297, 0, !dbg !18 |
|
%298 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.05.i102, float %284, float %278) #4, !dbg !18 |
|
%299 = tail call float @llvm.nvvm.fma.rn.f(float %.05.i102, float %284, float %278) #4, !dbg !18 |
|
%.06.i104 = select i1 %.not6.i103, float %299, float %298, !dbg !18 |
|
%300 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not7.i105 = icmp eq i32 %300, 0, !dbg !18 |
|
%301 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.06.i104, float %284, float %277) #4, !dbg !18 |
|
%302 = tail call float @llvm.nvvm.fma.rn.f(float %.06.i104, float %284, float %277) #4, !dbg !18 |
|
%.07.i106 = select i1 %.not7.i105, float %302, float %301, !dbg !18 |
|
%303 = fneg float %284, !dbg !18 |
|
%304 = select i1 %273, float %303, float %37, !dbg !18 |
|
%305 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not8.i107 = icmp eq i32 %305, 0, !dbg !18 |
|
%306 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.07.i106, float %304, float %304) #4, !dbg !18 |
|
%307 = tail call float @llvm.nvvm.fma.rn.f(float %.07.i106, float %304, float %304) #4, !dbg !18 |
|
%.08.i108 = select i1 %.not8.i107, float %307, float %306, !dbg !18 |
|
br i1 %273, label %308, label %__nv_erff.exit113, !dbg !18 |
|
|
|
308: |
|
%309 = tail call float @llvm.nvvm.ex2.approx.ftz.f(float %.08.i108) #4, !dbg !18 |
|
%310 = fsub float 1.000000e+00, %309, !dbg !18 |
|
%311 = bitcast float %310 to i32, !dbg !18 |
|
%312 = bitcast float %37 to i32, !dbg !18 |
|
%313 = and i32 %312, -2147483648, !dbg !18 |
|
%314 = or i32 %313, %311, !dbg !18 |
|
%315 = bitcast i32 %314 to float, !dbg !18 |
|
br label %__nv_erff.exit113, !dbg !18 |
|
|
|
__nv_erff.exit113: |
|
%r.0.i109 = phi float [ %315, %308 ], [ %.08.i108, %__internal_fmad.exit.i94 ], !dbg !18 |
|
%316 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not.i114 = icmp eq i32 %316, 0, !dbg !18 |
|
%317 = tail call float @llvm.nvvm.fabs.ftz.f(float %38) #4, !dbg !18 |
|
%318 = tail call float @llvm.nvvm.fabs.f(float %38) #4, !dbg !18 |
|
%.0.i115 = select i1 %.not.i114, float %318, float %317, !dbg !18 |
|
%319 = fcmp oge float %.0.i115, 0x3FF00C1FC0000000, !dbg !18 |
|
br i1 %319, label %__nv_fabsf.exit1.i132, label %321, !dbg !18 |
|
|
|
__nv_fabsf.exit1.i132: |
|
%320 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not1.i133 = icmp eq i32 %320, 0, !dbg !18 |
|
%.01.i134 = select i1 %.not1.i133, float %318, float %317, !dbg !18 |
|
br label %__internal_fmad.exit.i116, !dbg !18 |
|
|
|
321: |
|
%322 = fmul float %38, %38, !dbg !18 |
|
br label %__internal_fmad.exit.i116, !dbg !18 |
|
|
|
__internal_fmad.exit.i116: |
|
%323 = phi float [ 0x3FE41B0840000000, %__nv_fabsf.exit1.i132 ], [ 0x3FC06EBA60000000, %321 ], !dbg !18 |
|
%324 = phi float [ 0x3FED526FC0000000, %__nv_fabsf.exit1.i132 ], [ 0xBFD8127580000000, %321 ], !dbg !18 |
|
%325 = phi float [ 0x3FC39F20C0000000, %__nv_fabsf.exit1.i132 ], [ 0x3FBCE315E0000000, %321 ], !dbg !18 |
|
%326 = phi float [ 0xBFA1902C40000000, %__nv_fabsf.exit1.i132 ], [ 0xBF9B837CE0000000, %321 ], !dbg !18 |
|
%327 = phi float [ 0x3F75908160000000, %__nv_fabsf.exit1.i132 ], [ 0x3F755ABD40000000, %321 ], !dbg !18 |
|
%328 = phi float [ 0xBF3EAC1720000000, %__nv_fabsf.exit1.i132 ], [ 0xBF4AE9A400000000, %321 ], !dbg !18 |
|
%329 = phi float [ 0x3EF1394780000000, %__nv_fabsf.exit1.i132 ], [ 0x3F163D2D40000000, %321 ], !dbg !18 |
|
%330 = phi float [ %.01.i134, %__nv_fabsf.exit1.i132 ], [ %322, %321 ], !dbg !18 |
|
%331 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not2.i117 = icmp eq i32 %331, 0, !dbg !18 |
|
%332 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %329, float %330, float %328) #4, !dbg !18 |
|
%333 = tail call float @llvm.nvvm.fma.rn.f(float %329, float %330, float %328) #4, !dbg !18 |
|
%.02.i118 = select i1 %.not2.i117, float %333, float %332, !dbg !18 |
|
%334 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not3.i119 = icmp eq i32 %334, 0, !dbg !18 |
|
%335 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.02.i118, float %330, float %327) #4, !dbg !18 |
|
%336 = tail call float @llvm.nvvm.fma.rn.f(float %.02.i118, float %330, float %327) #4, !dbg !18 |
|
%.03.i120 = select i1 %.not3.i119, float %336, float %335, !dbg !18 |
|
%337 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not4.i121 = icmp eq i32 %337, 0, !dbg !18 |
|
%338 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.03.i120, float %330, float %326) #4, !dbg !18 |
|
%339 = tail call float @llvm.nvvm.fma.rn.f(float %.03.i120, float %330, float %326) #4, !dbg !18 |
|
%.04.i122 = select i1 %.not4.i121, float %339, float %338, !dbg !18 |
|
%340 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not5.i123 = icmp eq i32 %340, 0, !dbg !18 |
|
%341 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.04.i122, float %330, float %325) #4, !dbg !18 |
|
%342 = tail call float @llvm.nvvm.fma.rn.f(float %.04.i122, float %330, float %325) #4, !dbg !18 |
|
%.05.i124 = select i1 %.not5.i123, float %342, float %341, !dbg !18 |
|
%343 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not6.i125 = icmp eq i32 %343, 0, !dbg !18 |
|
%344 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.05.i124, float %330, float %324) #4, !dbg !18 |
|
%345 = tail call float @llvm.nvvm.fma.rn.f(float %.05.i124, float %330, float %324) #4, !dbg !18 |
|
%.06.i126 = select i1 %.not6.i125, float %345, float %344, !dbg !18 |
|
%346 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not7.i127 = icmp eq i32 %346, 0, !dbg !18 |
|
%347 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.06.i126, float %330, float %323) #4, !dbg !18 |
|
%348 = tail call float @llvm.nvvm.fma.rn.f(float %.06.i126, float %330, float %323) #4, !dbg !18 |
|
%.07.i128 = select i1 %.not7.i127, float %348, float %347, !dbg !18 |
|
%349 = fneg float %330, !dbg !18 |
|
%350 = select i1 %319, float %349, float %38, !dbg !18 |
|
%351 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not8.i129 = icmp eq i32 %351, 0, !dbg !18 |
|
%352 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.07.i128, float %350, float %350) #4, !dbg !18 |
|
%353 = tail call float @llvm.nvvm.fma.rn.f(float %.07.i128, float %350, float %350) #4, !dbg !18 |
|
%.08.i130 = select i1 %.not8.i129, float %353, float %352, !dbg !18 |
|
br i1 %319, label %354, label %__nv_erff.exit135, !dbg !18 |
|
|
|
354: |
|
%355 = tail call float @llvm.nvvm.ex2.approx.ftz.f(float %.08.i130) #4, !dbg !18 |
|
%356 = fsub float 1.000000e+00, %355, !dbg !18 |
|
%357 = bitcast float %356 to i32, !dbg !18 |
|
%358 = bitcast float %38 to i32, !dbg !18 |
|
%359 = and i32 %358, -2147483648, !dbg !18 |
|
%360 = or i32 %359, %357, !dbg !18 |
|
%361 = bitcast i32 %360 to float, !dbg !18 |
|
br label %__nv_erff.exit135, !dbg !18 |
|
|
|
__nv_erff.exit135: |
|
%r.0.i131 = phi float [ %361, %354 ], [ %.08.i130, %__internal_fmad.exit.i116 ], !dbg !18 |
|
%362 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not.i136 = icmp eq i32 %362, 0, !dbg !18 |
|
%363 = tail call float @llvm.nvvm.fabs.ftz.f(float %39) #4, !dbg !18 |
|
%364 = tail call float @llvm.nvvm.fabs.f(float %39) #4, !dbg !18 |
|
%.0.i137 = select i1 %.not.i136, float %364, float %363, !dbg !18 |
|
%365 = fcmp oge float %.0.i137, 0x3FF00C1FC0000000, !dbg !18 |
|
br i1 %365, label %__nv_fabsf.exit1.i154, label %367, !dbg !18 |
|
|
|
__nv_fabsf.exit1.i154: |
|
%366 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not1.i155 = icmp eq i32 %366, 0, !dbg !18 |
|
%.01.i156 = select i1 %.not1.i155, float %364, float %363, !dbg !18 |
|
br label %__internal_fmad.exit.i138, !dbg !18 |
|
|
|
367: |
|
%368 = fmul float %39, %39, !dbg !18 |
|
br label %__internal_fmad.exit.i138, !dbg !18 |
|
|
|
__internal_fmad.exit.i138: |
|
%369 = phi float [ 0x3FE41B0840000000, %__nv_fabsf.exit1.i154 ], [ 0x3FC06EBA60000000, %367 ], !dbg !18 |
|
%370 = phi float [ 0x3FED526FC0000000, %__nv_fabsf.exit1.i154 ], [ 0xBFD8127580000000, %367 ], !dbg !18 |
|
%371 = phi float [ 0x3FC39F20C0000000, %__nv_fabsf.exit1.i154 ], [ 0x3FBCE315E0000000, %367 ], !dbg !18 |
|
%372 = phi float [ 0xBFA1902C40000000, %__nv_fabsf.exit1.i154 ], [ 0xBF9B837CE0000000, %367 ], !dbg !18 |
|
%373 = phi float [ 0x3F75908160000000, %__nv_fabsf.exit1.i154 ], [ 0x3F755ABD40000000, %367 ], !dbg !18 |
|
%374 = phi float [ 0xBF3EAC1720000000, %__nv_fabsf.exit1.i154 ], [ 0xBF4AE9A400000000, %367 ], !dbg !18 |
|
%375 = phi float [ 0x3EF1394780000000, %__nv_fabsf.exit1.i154 ], [ 0x3F163D2D40000000, %367 ], !dbg !18 |
|
%376 = phi float [ %.01.i156, %__nv_fabsf.exit1.i154 ], [ %368, %367 ], !dbg !18 |
|
%377 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not2.i139 = icmp eq i32 %377, 0, !dbg !18 |
|
%378 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %375, float %376, float %374) #4, !dbg !18 |
|
%379 = tail call float @llvm.nvvm.fma.rn.f(float %375, float %376, float %374) #4, !dbg !18 |
|
%.02.i140 = select i1 %.not2.i139, float %379, float %378, !dbg !18 |
|
%380 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not3.i141 = icmp eq i32 %380, 0, !dbg !18 |
|
%381 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.02.i140, float %376, float %373) #4, !dbg !18 |
|
%382 = tail call float @llvm.nvvm.fma.rn.f(float %.02.i140, float %376, float %373) #4, !dbg !18 |
|
%.03.i142 = select i1 %.not3.i141, float %382, float %381, !dbg !18 |
|
%383 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not4.i143 = icmp eq i32 %383, 0, !dbg !18 |
|
%384 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.03.i142, float %376, float %372) #4, !dbg !18 |
|
%385 = tail call float @llvm.nvvm.fma.rn.f(float %.03.i142, float %376, float %372) #4, !dbg !18 |
|
%.04.i144 = select i1 %.not4.i143, float %385, float %384, !dbg !18 |
|
%386 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not5.i145 = icmp eq i32 %386, 0, !dbg !18 |
|
%387 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.04.i144, float %376, float %371) #4, !dbg !18 |
|
%388 = tail call float @llvm.nvvm.fma.rn.f(float %.04.i144, float %376, float %371) #4, !dbg !18 |
|
%.05.i146 = select i1 %.not5.i145, float %388, float %387, !dbg !18 |
|
%389 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not6.i147 = icmp eq i32 %389, 0, !dbg !18 |
|
%390 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.05.i146, float %376, float %370) #4, !dbg !18 |
|
%391 = tail call float @llvm.nvvm.fma.rn.f(float %.05.i146, float %376, float %370) #4, !dbg !18 |
|
%.06.i148 = select i1 %.not6.i147, float %391, float %390, !dbg !18 |
|
%392 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not7.i149 = icmp eq i32 %392, 0, !dbg !18 |
|
%393 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.06.i148, float %376, float %369) #4, !dbg !18 |
|
%394 = tail call float @llvm.nvvm.fma.rn.f(float %.06.i148, float %376, float %369) #4, !dbg !18 |
|
%.07.i150 = select i1 %.not7.i149, float %394, float %393, !dbg !18 |
|
%395 = fneg float %376, !dbg !18 |
|
%396 = select i1 %365, float %395, float %39, !dbg !18 |
|
%397 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4, !dbg !18 |
|
%.not8.i151 = icmp eq i32 %397, 0, !dbg !18 |
|
%398 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.07.i150, float %396, float %396) #4, !dbg !18 |
|
%399 = tail call float @llvm.nvvm.fma.rn.f(float %.07.i150, float %396, float %396) #4, !dbg !18 |
|
%.08.i152 = select i1 %.not8.i151, float %399, float %398, !dbg !18 |
|
br i1 %365, label %400, label %__nv_erff.exit157, !dbg !18 |
|
|
|
400: |
|
%401 = tail call float @llvm.nvvm.ex2.approx.ftz.f(float %.08.i152) #4, !dbg !18 |
|
%402 = fsub float 1.000000e+00, %401, !dbg !18 |
|
%403 = bitcast float %402 to i32, !dbg !18 |
|
%404 = bitcast float %39 to i32, !dbg !18 |
|
%405 = and i32 %404, -2147483648, !dbg !18 |
|
%406 = or i32 %405, %403, !dbg !18 |
|
%407 = bitcast i32 %406 to float, !dbg !18 |
|
br label %__nv_erff.exit157, !dbg !18 |
|
|
|
__nv_erff.exit157: |
|
%r.0.i153 = phi float [ %407, %400 ], [ %.08.i152, %__internal_fmad.exit.i138 ], !dbg !18 |
|
%408 = fmul float %31, 5.000000e-01, !dbg !19 |
|
%409 = fmul float %30, 5.000000e-01, !dbg !19 |
|
%410 = fmul float %29, 5.000000e-01, !dbg !19 |
|
%411 = fmul float %28, 5.000000e-01, !dbg !19 |
|
%412 = fmul float %27, 5.000000e-01, !dbg !19 |
|
%413 = fmul float %26, 5.000000e-01, !dbg !19 |
|
%414 = fmul float %25, 5.000000e-01, !dbg !19 |
|
%415 = fmul float %24, 5.000000e-01, !dbg !19 |
|
%416 = fadd float %r.0.i, 1.000000e+00, !dbg !20 |
|
%417 = fadd float %r.0.i21, 1.000000e+00, !dbg !20 |
|
%418 = fadd float %r.0.i43, 1.000000e+00, !dbg !20 |
|
%419 = fadd float %r.0.i65, 1.000000e+00, !dbg !20 |
|
%420 = fadd float %r.0.i87, 1.000000e+00, !dbg !20 |
|
%421 = fadd float %r.0.i109, 1.000000e+00, !dbg !20 |
|
%422 = fadd float %r.0.i131, 1.000000e+00, !dbg !20 |
|
%423 = fadd float %r.0.i153, 1.000000e+00, !dbg !20 |
|
%424 = fmul float %415, %416, !dbg !21 |
|
%425 = fmul float %414, %417, !dbg !21 |
|
%426 = fmul float %413, %418, !dbg !21 |
|
%427 = fmul float %412, %419, !dbg !21 |
|
%428 = fmul float %411, %420, !dbg !21 |
|
%429 = fmul float %410, %421, !dbg !21 |
|
%430 = fmul float %409, %422, !dbg !21 |
|
%431 = fmul float %408, %423, !dbg !21 |
|
%432 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %424) #4, !dbg !22 |
|
%433 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %425) #4, !dbg !22 |
|
%434 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %426) #4, !dbg !22 |
|
%435 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %427) #4, !dbg !22 |
|
%436 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %428) #4, !dbg !22 |
|
%437 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %429) #4, !dbg !22 |
|
%438 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %430) #4, !dbg !22 |
|
%439 = tail call i16 asm "cvt.rn.bf16.f32 $0, $1;", "=h,r"(float %431) #4, !dbg !22 |
|
%440 = insertelement <2 x i16> undef, i16 %432, i64 0, !dbg !22 |
|
%441 = insertelement <2 x i16> %440, i16 %433, i64 1, !dbg !22 |
|
%442 = bitcast <2 x i16> %441 to i32, !dbg !22 |
|
%443 = insertelement <2 x i16> undef, i16 %434, i64 0, !dbg !22 |
|
%444 = insertelement <2 x i16> %443, i16 %435, i64 1, !dbg !22 |
|
%445 = bitcast <2 x i16> %444 to i32, !dbg !22 |
|
%446 = insertelement <2 x i16> undef, i16 %436, i64 0, !dbg !22 |
|
%447 = insertelement <2 x i16> %446, i16 %437, i64 1, !dbg !22 |
|
%448 = bitcast <2 x i16> %447 to i32, !dbg !22 |
|
%449 = insertelement <2 x i16> undef, i16 %438, i64 0, !dbg !22 |
|
%450 = insertelement <2 x i16> %449, i16 %439, i64 1, !dbg !22 |
|
%451 = bitcast <2 x i16> %450 to i32, !dbg !22 |
|
tail call void asm sideeffect "@$5 st.global.v4.b32 [ $4 + 0 ], { $0, $1, $2, $3 };", "r,r,r,r,l,b"(i32 %442, i32 %445, i32 %448, i32 %451, ptr addrspace(1) %10, i1 true) #4, !dbg !22 |
|
ret void, !dbg !23 |
|
} |
|
|
|
|
|
declare noundef i32 @llvm.nvvm.read.ptx.sreg.tid.x() #0 |
|
|
|
|
|
define float @__nv_erff(float %a) local_unnamed_addr #1 { |
|
__nv_fabsf.exit: |
|
%0 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4 |
|
%.not = icmp eq i32 %0, 0 |
|
%1 = tail call float @llvm.nvvm.fabs.ftz.f(float %a) #4 |
|
%2 = tail call float @llvm.nvvm.fabs.f(float %a) #4 |
|
%.0 = select i1 %.not, float %2, float %1 |
|
%3 = fcmp oge float %.0, 0x3FF00C1FC0000000 |
|
br i1 %3, label %__nv_fabsf.exit1, label %5 |
|
|
|
__nv_fabsf.exit1: |
|
%4 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4 |
|
%.not1 = icmp eq i32 %4, 0 |
|
%.01 = select i1 %.not1, float %2, float %1 |
|
br label %__internal_fmad.exit |
|
|
|
5: |
|
%6 = fmul float %a, %a |
|
br label %__internal_fmad.exit |
|
|
|
__internal_fmad.exit: |
|
%7 = phi float [ 0x3FE41B0840000000, %__nv_fabsf.exit1 ], [ 0x3FC06EBA60000000, %5 ] |
|
%8 = phi float [ 0x3FED526FC0000000, %__nv_fabsf.exit1 ], [ 0xBFD8127580000000, %5 ] |
|
%9 = phi float [ 0x3FC39F20C0000000, %__nv_fabsf.exit1 ], [ 0x3FBCE315E0000000, %5 ] |
|
%10 = phi float [ 0xBFA1902C40000000, %__nv_fabsf.exit1 ], [ 0xBF9B837CE0000000, %5 ] |
|
%11 = phi float [ 0x3F75908160000000, %__nv_fabsf.exit1 ], [ 0x3F755ABD40000000, %5 ] |
|
%12 = phi float [ 0xBF3EAC1720000000, %__nv_fabsf.exit1 ], [ 0xBF4AE9A400000000, %5 ] |
|
%13 = phi float [ 0x3EF1394780000000, %__nv_fabsf.exit1 ], [ 0x3F163D2D40000000, %5 ] |
|
%14 = phi float [ %.01, %__nv_fabsf.exit1 ], [ %6, %5 ] |
|
%15 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4 |
|
%.not2 = icmp eq i32 %15, 0 |
|
%16 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %13, float %14, float %12) #4 |
|
%17 = tail call float @llvm.nvvm.fma.rn.f(float %13, float %14, float %12) #4 |
|
%.02 = select i1 %.not2, float %17, float %16 |
|
%18 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4 |
|
%.not3 = icmp eq i32 %18, 0 |
|
%19 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.02, float %14, float %11) #4 |
|
%20 = tail call float @llvm.nvvm.fma.rn.f(float %.02, float %14, float %11) #4 |
|
%.03 = select i1 %.not3, float %20, float %19 |
|
%21 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4 |
|
%.not4 = icmp eq i32 %21, 0 |
|
%22 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.03, float %14, float %10) #4 |
|
%23 = tail call float @llvm.nvvm.fma.rn.f(float %.03, float %14, float %10) #4 |
|
%.04 = select i1 %.not4, float %23, float %22 |
|
%24 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4 |
|
%.not5 = icmp eq i32 %24, 0 |
|
%25 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.04, float %14, float %9) #4 |
|
%26 = tail call float @llvm.nvvm.fma.rn.f(float %.04, float %14, float %9) #4 |
|
%.05 = select i1 %.not5, float %26, float %25 |
|
%27 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4 |
|
%.not6 = icmp eq i32 %27, 0 |
|
%28 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.05, float %14, float %8) #4 |
|
%29 = tail call float @llvm.nvvm.fma.rn.f(float %.05, float %14, float %8) #4 |
|
%.06 = select i1 %.not6, float %29, float %28 |
|
%30 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4 |
|
%.not7 = icmp eq i32 %30, 0 |
|
%31 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.06, float %14, float %7) #4 |
|
%32 = tail call float @llvm.nvvm.fma.rn.f(float %.06, float %14, float %7) #4 |
|
%.07 = select i1 %.not7, float %32, float %31 |
|
%33 = fneg float %14 |
|
%34 = select i1 %3, float %33, float %a |
|
%35 = tail call i32 @__nvvm_reflect(ptr nonnull @.str) #4 |
|
%.not8 = icmp eq i32 %35, 0 |
|
%36 = tail call float @llvm.nvvm.fma.rn.ftz.f(float %.07, float %34, float %34) #4 |
|
%37 = tail call float @llvm.nvvm.fma.rn.f(float %.07, float %34, float %34) #4 |
|
%.08 = select i1 %.not8, float %37, float %36 |
|
br i1 %3, label %38, label %46 |
|
|
|
38: |
|
%39 = tail call float @llvm.nvvm.ex2.approx.ftz.f(float %.08) #4 |
|
%40 = fsub float 1.000000e+00, %39 |
|
%41 = bitcast float %40 to i32 |
|
%42 = bitcast float %a to i32 |
|
%43 = and i32 %42, -2147483648 |
|
%44 = or i32 %43, %41 |
|
%45 = bitcast i32 %44 to float |
|
br label %46 |
|
|
|
46: |
|
%r.0 = phi float [ %45, %38 ], [ %.08, %__internal_fmad.exit ] |
|
ret float %r.0 |
|
} |
|
|
|
declare i32 @__nvvm_reflect(ptr) local_unnamed_addr #2 |
|
|
|
|
|
declare float @llvm.nvvm.fabs.ftz.f(float) #0 |
|
|
|
|
|
declare float @llvm.nvvm.fabs.f(float) #0 |
|
|
|
|
|
declare float @llvm.nvvm.fma.rn.ftz.f(float, float, float) #0 |
|
|
|
|
|
declare float @llvm.nvvm.fma.rn.f(float, float, float) #0 |
|
|
|
|
|
declare float @llvm.nvvm.ex2.approx.ftz.f(float) #3 |
|
|
|
attributes #0 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) } |
|
attributes #1 = { 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 #2 = { "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 #3 = { mustprogress nocallback nofree nosync nounwind willreturn memory(none) } |
|
attributes #4 = { 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: "ckphrtdpgsxl7sfarkkzylhv4st3uhmzvg3u6z5excfp6ydybq74.py", directory: "/tmp/torchinductor_root/kp") |
|
!4 = !{ptr @triton__0d1de, !"kernel", i32 1} |
|
!5 = !{ptr @triton__0d1de, !"maxntidx", i32 128} |
|
!6 = !{!"clang version 3.8.0 (tags/RELEASE_380/final)"} |
|
!7 = distinct !DISubprogram(name: "triton__0d1de", linkageName: "triton__0d1de", 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: 21, column: 36, scope: !7) |
|
!11 = !DILocation(line: 20, column: 28, scope: !7) |
|
!12 = !DILocation(line: 20, column: 33, scope: !7) |
|
!13 = !DILocation(line: 21, column: 23, scope: !7) |
|
!14 = !DILocation(line: 24, column: 34, scope: !7) |
|
!15 = !DILocation(line: 24, column: 39, scope: !7) |
|
!16 = !DILocation(line: 24, column: 48, scope: !7) |
|
!17 = !DILocation(line: 29, column: 18, scope: !7) |
|
!18 = !DILocation(line: 30, column: 23, scope: !7) |
|
!19 = !DILocation(line: 27, column: 18, scope: !7) |
|
!20 = !DILocation(line: 32, column: 18, scope: !7) |
|
!21 = !DILocation(line: 33, column: 18, scope: !7) |
|
!22 = !DILocation(line: 35, column: 40, scope: !7) |
|
!23 = !DILocation(line: 35, column: 4, scope: !7) |
|
|
