|
|
|
|
|
|
|
|
|
.version 8.2 |
|
.target sm_89 |
|
.address_size 64 |
|
|
|
|
|
.global .align 1 .b8 _$_str[11] = {95, 95, 67, 85, 68, 65, 95, 70, 84, 90, 0}; |
|
|
|
.visible .entry triton__0d1d2de( |
|
.param .u64 triton__0d1d2de_param_0, |
|
.param .u64 triton__0d1d2de_param_1, |
|
.param .u32 triton__0d1d2de_param_2 |
|
) |
|
.maxntid 128, 1, 1 |
|
{ |
|
.reg .pred %p<27>; |
|
.reg .b16 %rs<17>; |
|
.reg .b32 %r<67>; |
|
.reg .f32 %f<431>; |
|
.reg .b64 %rd<8>; |
|
.loc 1 18 0 |
|
$L__func_begin0: |
|
.loc 1 18 0 |
|
|
|
ld.param.u64 %rd4, [triton__0d1d2de_param_0]; |
|
$L__tmp0: |
|
.loc 1 21 36 |
|
mov.u32 %r14, %tid.x; |
|
shl.b32 %r15, %r14, 3; |
|
and.b32 %r16, %r15, 1016; |
|
.loc 1 20 28 |
|
mov.u32 %r1, %ctaid.x; |
|
.loc 1 20 33 |
|
shl.b32 %r17, %r1, 10; |
|
.loc 1 21 23 |
|
or.b32 %r18, %r17, %r16; |
|
.loc 1 24 30 |
|
mul.wide.s32 %rd5, %r18, 2; |
|
add.s64 %rd3, %rd4, %rd5; |
|
mov.pred %p1, -1; |
|
.loc 1 24 35 |
|
mov.u32 %r2, 0x0; |
|
mov.u32 %r3, 0x0; |
|
mov.u32 %r4, 0x0; |
|
mov.u32 %r5, 0x0; |
|
@%p1 ld.global.v4.b32 { %r2, %r3, %r4, %r5 }, [ %rd3 + 0 ]; |
|
cvt.u16.u32 %rs1, %r2; |
|
{ .reg .b16 tmp; mov.b32 {tmp, %rs2}, %r2; } |
|
cvt.u16.u32 %rs3, %r3; |
|
.loc 1 24 44 |
|
cvt.f32.bf16 %r6, %rs1; |
|
mov.b32 %f1, %r6; |
|
cvt.f32.bf16 %r7, %rs2; |
|
mov.b32 %f2, %r7; |
|
.loc 1 29 18 |
|
mul.f32 %f9, %f1, 0f3F3504F3; |
|
.loc 1 30 23 |
|
abs.ftz.f32 %f17, %f9; |
|
setp.ge.f32 %p2, %f17, 0f3F8060FE; |
|
mov.f32 %f365, 0f3789CA3C; |
|
mov.f32 %f364, 0fB9F560B9; |
|
mov.f32 %f363, 0f3BAC840B; |
|
mov.f32 %f362, 0fBD0C8162; |
|
mov.f32 %f361, 0f3E1CF906; |
|
mov.f32 %f360, 0f3F6A937E; |
|
mov.f32 %f359, 0f3F20D842; |
|
mov.f32 %f366, %f17; |
|
@%p2 bra $L__BB0_2; |
|
.loc 1 0 23 |
|
mov.f32 %f365, 0f38B1E96A; |
|
mov.f32 %f364, 0fBA574D20; |
|
mov.f32 %f363, 0f3BAAD5EA; |
|
mov.f32 %f362, 0fBCDC1BE7; |
|
mov.f32 %f361, 0f3DE718AF; |
|
mov.f32 %f360, 0fBEC093AC; |
|
mov.f32 %f359, 0f3E0375D3; |
|
.loc 1 30 23 |
|
mul.f32 %f366, %f9, %f9; |
|
$L__BB0_2: |
|
.loc 1 0 0 |
|
cvt.f32.bf16 %r8, %rs3; |
|
mul.f32 %f10, %f2, 0f3F3504F3; |
|
.loc 1 30 23 |
|
setp.ltu.f32 %p3, %f17, 0f3F8060FE; |
|
fma.rn.ftz.f32 %f135, %f365, %f366, %f364; |
|
fma.rn.ftz.f32 %f136, %f135, %f366, %f363; |
|
fma.rn.ftz.f32 %f137, %f136, %f366, %f362; |
|
fma.rn.ftz.f32 %f138, %f137, %f366, %f361; |
|
fma.rn.ftz.f32 %f139, %f138, %f366, %f360; |
|
fma.rn.ftz.f32 %f140, %f139, %f366, %f359; |
|
neg.f32 %f141, %f366; |
|
selp.f32 %f142, %f141, %f9, %p2; |
|
fma.rn.ftz.f32 %f367, %f140, %f142, %f142; |
|
mov.f32 %f358, 0f3F800000; |
|
@%p3 bra $L__BB0_4; |
|
ex2.approx.ftz.f32 %f143, %f367; |
|
sub.f32 %f145, %f358, %f143; |
|
mov.b32 %r19, %f145; |
|
mov.b32 %r20, %f9; |
|
and.b32 %r21, %r20, -2147483648; |
|
or.b32 %r22, %r21, %r19; |
|
mov.b32 %f367, %r22; |
|
$L__BB0_4: |
|
.loc 1 0 0 |
|
{ .reg .b16 tmp; mov.b32 {tmp, %rs4}, %r3; } |
|
mov.b32 %f3, %r8; |
|
.loc 1 30 23 |
|
abs.ftz.f32 %f30, %f10; |
|
setp.ge.f32 %p5, %f30, 0f3F8060FE; |
|
mov.f32 %f374, 0f3789CA3C; |
|
mov.f32 %f373, 0fB9F560B9; |
|
mov.f32 %f372, 0f3BAC840B; |
|
mov.f32 %f371, 0fBD0C8162; |
|
mov.f32 %f370, 0f3E1CF906; |
|
mov.f32 %f369, 0f3F6A937E; |
|
mov.f32 %f368, 0f3F20D842; |
|
mov.f32 %f375, %f30; |
|
@%p5 bra $L__BB0_6; |
|
mul.f32 %f375, %f10, %f10; |
|
mov.f32 %f374, 0f38B1E96A; |
|
mov.f32 %f373, 0fBA574D20; |
|
mov.f32 %f372, 0f3BAAD5EA; |
|
mov.f32 %f371, 0fBCDC1BE7; |
|
mov.f32 %f370, 0f3DE718AF; |
|
mov.f32 %f369, 0fBEC093AC; |
|
mov.f32 %f368, 0f3E0375D3; |
|
$L__BB0_6: |
|
.loc 1 0 0 |
|
cvt.f32.bf16 %r9, %rs4; |
|
mul.f32 %f11, %f3, 0f3F3504F3; |
|
.loc 1 30 23 |
|
setp.ltu.f32 %p6, %f30, 0f3F8060FE; |
|
fma.rn.ftz.f32 %f160, %f374, %f375, %f373; |
|
fma.rn.ftz.f32 %f161, %f160, %f375, %f372; |
|
fma.rn.ftz.f32 %f162, %f161, %f375, %f371; |
|
fma.rn.ftz.f32 %f163, %f162, %f375, %f370; |
|
fma.rn.ftz.f32 %f164, %f163, %f375, %f369; |
|
fma.rn.ftz.f32 %f165, %f164, %f375, %f368; |
|
neg.f32 %f166, %f375; |
|
selp.f32 %f167, %f166, %f10, %p5; |
|
fma.rn.ftz.f32 %f376, %f165, %f167, %f167; |
|
@%p6 bra $L__BB0_8; |
|
ex2.approx.ftz.f32 %f168, %f376; |
|
sub.f32 %f170, %f358, %f168; |
|
mov.b32 %r23, %f170; |
|
mov.b32 %r24, %f10; |
|
and.b32 %r25, %r24, -2147483648; |
|
or.b32 %r26, %r25, %r23; |
|
mov.b32 %f376, %r26; |
|
$L__BB0_8: |
|
.loc 1 0 0 |
|
cvt.u16.u32 %rs5, %r4; |
|
mov.b32 %f4, %r9; |
|
.loc 1 30 23 |
|
abs.ftz.f32 %f43, %f11; |
|
setp.ge.f32 %p8, %f43, 0f3F8060FE; |
|
mov.f32 %f383, 0f3789CA3C; |
|
mov.f32 %f382, 0fB9F560B9; |
|
mov.f32 %f381, 0f3BAC840B; |
|
mov.f32 %f380, 0fBD0C8162; |
|
mov.f32 %f379, 0f3E1CF906; |
|
mov.f32 %f378, 0f3F6A937E; |
|
mov.f32 %f377, 0f3F20D842; |
|
mov.f32 %f384, %f43; |
|
@%p8 bra $L__BB0_10; |
|
mul.f32 %f384, %f11, %f11; |
|
mov.f32 %f383, 0f38B1E96A; |
|
mov.f32 %f382, 0fBA574D20; |
|
mov.f32 %f381, 0f3BAAD5EA; |
|
mov.f32 %f380, 0fBCDC1BE7; |
|
mov.f32 %f379, 0f3DE718AF; |
|
mov.f32 %f378, 0fBEC093AC; |
|
mov.f32 %f377, 0f3E0375D3; |
|
$L__BB0_10: |
|
.loc 1 0 0 |
|
cvt.f32.bf16 %r10, %rs5; |
|
mul.f32 %f12, %f4, 0f3F3504F3; |
|
.loc 1 30 23 |
|
setp.ltu.f32 %p9, %f43, 0f3F8060FE; |
|
fma.rn.ftz.f32 %f185, %f383, %f384, %f382; |
|
fma.rn.ftz.f32 %f186, %f185, %f384, %f381; |
|
fma.rn.ftz.f32 %f187, %f186, %f384, %f380; |
|
fma.rn.ftz.f32 %f188, %f187, %f384, %f379; |
|
fma.rn.ftz.f32 %f189, %f188, %f384, %f378; |
|
fma.rn.ftz.f32 %f190, %f189, %f384, %f377; |
|
neg.f32 %f191, %f384; |
|
selp.f32 %f192, %f191, %f11, %p8; |
|
fma.rn.ftz.f32 %f385, %f190, %f192, %f192; |
|
@%p9 bra $L__BB0_12; |
|
ex2.approx.ftz.f32 %f193, %f385; |
|
sub.f32 %f195, %f358, %f193; |
|
mov.b32 %r27, %f195; |
|
mov.b32 %r28, %f11; |
|
and.b32 %r29, %r28, -2147483648; |
|
or.b32 %r30, %r29, %r27; |
|
mov.b32 %f385, %r30; |
|
$L__BB0_12: |
|
.loc 1 0 0 |
|
{ .reg .b16 tmp; mov.b32 {tmp, %rs6}, %r4; } |
|
mov.b32 %f5, %r10; |
|
.loc 1 30 23 |
|
abs.ftz.f32 %f56, %f12; |
|
setp.ge.f32 %p11, %f56, 0f3F8060FE; |
|
mov.f32 %f392, 0f3789CA3C; |
|
mov.f32 %f391, 0fB9F560B9; |
|
mov.f32 %f390, 0f3BAC840B; |
|
mov.f32 %f389, 0fBD0C8162; |
|
mov.f32 %f388, 0f3E1CF906; |
|
mov.f32 %f387, 0f3F6A937E; |
|
mov.f32 %f386, 0f3F20D842; |
|
mov.f32 %f393, %f56; |
|
@%p11 bra $L__BB0_14; |
|
mul.f32 %f393, %f12, %f12; |
|
mov.f32 %f392, 0f38B1E96A; |
|
mov.f32 %f391, 0fBA574D20; |
|
mov.f32 %f390, 0f3BAAD5EA; |
|
mov.f32 %f389, 0fBCDC1BE7; |
|
mov.f32 %f388, 0f3DE718AF; |
|
mov.f32 %f387, 0fBEC093AC; |
|
mov.f32 %f386, 0f3E0375D3; |
|
$L__BB0_14: |
|
.loc 1 0 0 |
|
cvt.f32.bf16 %r11, %rs6; |
|
mul.f32 %f13, %f5, 0f3F3504F3; |
|
.loc 1 30 23 |
|
setp.ltu.f32 %p12, %f56, 0f3F8060FE; |
|
fma.rn.ftz.f32 %f210, %f392, %f393, %f391; |
|
fma.rn.ftz.f32 %f211, %f210, %f393, %f390; |
|
fma.rn.ftz.f32 %f212, %f211, %f393, %f389; |
|
fma.rn.ftz.f32 %f213, %f212, %f393, %f388; |
|
fma.rn.ftz.f32 %f214, %f213, %f393, %f387; |
|
fma.rn.ftz.f32 %f215, %f214, %f393, %f386; |
|
neg.f32 %f216, %f393; |
|
selp.f32 %f217, %f216, %f12, %p11; |
|
fma.rn.ftz.f32 %f394, %f215, %f217, %f217; |
|
@%p12 bra $L__BB0_16; |
|
ex2.approx.ftz.f32 %f218, %f394; |
|
sub.f32 %f220, %f358, %f218; |
|
mov.b32 %r31, %f220; |
|
mov.b32 %r32, %f12; |
|
and.b32 %r33, %r32, -2147483648; |
|
or.b32 %r34, %r33, %r31; |
|
mov.b32 %f394, %r34; |
|
$L__BB0_16: |
|
.loc 1 0 0 |
|
cvt.u16.u32 %rs7, %r5; |
|
mov.b32 %f6, %r11; |
|
.loc 1 30 23 |
|
abs.ftz.f32 %f69, %f13; |
|
setp.ge.f32 %p14, %f69, 0f3F8060FE; |
|
mov.f32 %f401, 0f3789CA3C; |
|
mov.f32 %f400, 0fB9F560B9; |
|
mov.f32 %f399, 0f3BAC840B; |
|
mov.f32 %f398, 0fBD0C8162; |
|
mov.f32 %f397, 0f3E1CF906; |
|
mov.f32 %f396, 0f3F6A937E; |
|
mov.f32 %f395, 0f3F20D842; |
|
mov.f32 %f402, %f69; |
|
@%p14 bra $L__BB0_18; |
|
mul.f32 %f402, %f13, %f13; |
|
mov.f32 %f401, 0f38B1E96A; |
|
mov.f32 %f400, 0fBA574D20; |
|
mov.f32 %f399, 0f3BAAD5EA; |
|
mov.f32 %f398, 0fBCDC1BE7; |
|
mov.f32 %f397, 0f3DE718AF; |
|
mov.f32 %f396, 0fBEC093AC; |
|
mov.f32 %f395, 0f3E0375D3; |
|
$L__BB0_18: |
|
.loc 1 0 0 |
|
cvt.f32.bf16 %r12, %rs7; |
|
mul.f32 %f14, %f6, 0f3F3504F3; |
|
.loc 1 30 23 |
|
setp.ltu.f32 %p15, %f69, 0f3F8060FE; |
|
fma.rn.ftz.f32 %f235, %f401, %f402, %f400; |
|
fma.rn.ftz.f32 %f236, %f235, %f402, %f399; |
|
fma.rn.ftz.f32 %f237, %f236, %f402, %f398; |
|
fma.rn.ftz.f32 %f238, %f237, %f402, %f397; |
|
fma.rn.ftz.f32 %f239, %f238, %f402, %f396; |
|
fma.rn.ftz.f32 %f240, %f239, %f402, %f395; |
|
neg.f32 %f241, %f402; |
|
selp.f32 %f242, %f241, %f13, %p14; |
|
fma.rn.ftz.f32 %f403, %f240, %f242, %f242; |
|
@%p15 bra $L__BB0_20; |
|
ex2.approx.ftz.f32 %f243, %f403; |
|
sub.f32 %f245, %f358, %f243; |
|
mov.b32 %r35, %f245; |
|
mov.b32 %r36, %f13; |
|
and.b32 %r37, %r36, -2147483648; |
|
or.b32 %r38, %r37, %r35; |
|
mov.b32 %f403, %r38; |
|
$L__BB0_20: |
|
.loc 1 0 0 |
|
{ .reg .b16 tmp; mov.b32 {tmp, %rs8}, %r5; } |
|
mov.b32 %f7, %r12; |
|
.loc 1 30 23 |
|
abs.ftz.f32 %f82, %f14; |
|
setp.ge.f32 %p17, %f82, 0f3F8060FE; |
|
mov.f32 %f410, 0f3789CA3C; |
|
mov.f32 %f409, 0fB9F560B9; |
|
mov.f32 %f408, 0f3BAC840B; |
|
mov.f32 %f407, 0fBD0C8162; |
|
mov.f32 %f406, 0f3E1CF906; |
|
mov.f32 %f405, 0f3F6A937E; |
|
mov.f32 %f404, 0f3F20D842; |
|
mov.f32 %f411, %f82; |
|
@%p17 bra $L__BB0_22; |
|
mul.f32 %f411, %f14, %f14; |
|
mov.f32 %f410, 0f38B1E96A; |
|
mov.f32 %f409, 0fBA574D20; |
|
mov.f32 %f408, 0f3BAAD5EA; |
|
mov.f32 %f407, 0fBCDC1BE7; |
|
mov.f32 %f406, 0f3DE718AF; |
|
mov.f32 %f405, 0fBEC093AC; |
|
mov.f32 %f404, 0f3E0375D3; |
|
$L__BB0_22: |
|
.loc 1 0 0 |
|
cvt.f32.bf16 %r13, %rs8; |
|
mul.f32 %f15, %f7, 0f3F3504F3; |
|
.loc 1 30 23 |
|
setp.ltu.f32 %p18, %f82, 0f3F8060FE; |
|
fma.rn.ftz.f32 %f260, %f410, %f411, %f409; |
|
fma.rn.ftz.f32 %f261, %f260, %f411, %f408; |
|
fma.rn.ftz.f32 %f262, %f261, %f411, %f407; |
|
fma.rn.ftz.f32 %f263, %f262, %f411, %f406; |
|
fma.rn.ftz.f32 %f264, %f263, %f411, %f405; |
|
fma.rn.ftz.f32 %f265, %f264, %f411, %f404; |
|
neg.f32 %f266, %f411; |
|
selp.f32 %f267, %f266, %f14, %p17; |
|
fma.rn.ftz.f32 %f412, %f265, %f267, %f267; |
|
@%p18 bra $L__BB0_24; |
|
ex2.approx.ftz.f32 %f268, %f412; |
|
sub.f32 %f270, %f358, %f268; |
|
mov.b32 %r39, %f270; |
|
mov.b32 %r40, %f14; |
|
and.b32 %r41, %r40, -2147483648; |
|
or.b32 %r42, %r41, %r39; |
|
mov.b32 %f412, %r42; |
|
$L__BB0_24: |
|
.loc 1 0 0 |
|
mov.b32 %f8, %r13; |
|
.loc 1 30 23 |
|
abs.ftz.f32 %f95, %f15; |
|
setp.ge.f32 %p20, %f95, 0f3F8060FE; |
|
mov.f32 %f419, 0f3789CA3C; |
|
mov.f32 %f418, 0fB9F560B9; |
|
mov.f32 %f417, 0f3BAC840B; |
|
mov.f32 %f416, 0fBD0C8162; |
|
mov.f32 %f415, 0f3E1CF906; |
|
mov.f32 %f414, 0f3F6A937E; |
|
mov.f32 %f413, 0f3F20D842; |
|
mov.f32 %f420, %f95; |
|
@%p20 bra $L__BB0_26; |
|
mul.f32 %f420, %f15, %f15; |
|
mov.f32 %f419, 0f38B1E96A; |
|
mov.f32 %f418, 0fBA574D20; |
|
mov.f32 %f417, 0f3BAAD5EA; |
|
mov.f32 %f416, 0fBCDC1BE7; |
|
mov.f32 %f415, 0f3DE718AF; |
|
mov.f32 %f414, 0fBEC093AC; |
|
mov.f32 %f413, 0f3E0375D3; |
|
$L__BB0_26: |
|
.loc 1 0 0 |
|
mul.f32 %f16, %f8, 0f3F3504F3; |
|
.loc 1 30 23 |
|
setp.ltu.f32 %p21, %f95, 0f3F8060FE; |
|
fma.rn.ftz.f32 %f285, %f419, %f420, %f418; |
|
fma.rn.ftz.f32 %f286, %f285, %f420, %f417; |
|
fma.rn.ftz.f32 %f287, %f286, %f420, %f416; |
|
fma.rn.ftz.f32 %f288, %f287, %f420, %f415; |
|
fma.rn.ftz.f32 %f289, %f288, %f420, %f414; |
|
fma.rn.ftz.f32 %f290, %f289, %f420, %f413; |
|
neg.f32 %f291, %f420; |
|
selp.f32 %f292, %f291, %f15, %p20; |
|
fma.rn.ftz.f32 %f421, %f290, %f292, %f292; |
|
@%p21 bra $L__BB0_28; |
|
ex2.approx.ftz.f32 %f293, %f421; |
|
sub.f32 %f295, %f358, %f293; |
|
mov.b32 %r43, %f295; |
|
mov.b32 %r44, %f15; |
|
and.b32 %r45, %r44, -2147483648; |
|
or.b32 %r46, %r45, %r43; |
|
mov.b32 %f421, %r46; |
|
$L__BB0_28: |
|
abs.ftz.f32 %f108, %f16; |
|
setp.ge.f32 %p23, %f108, 0f3F8060FE; |
|
mov.f32 %f428, 0f3789CA3C; |
|
mov.f32 %f427, 0fB9F560B9; |
|
mov.f32 %f426, 0f3BAC840B; |
|
mov.f32 %f425, 0fBD0C8162; |
|
mov.f32 %f424, 0f3E1CF906; |
|
mov.f32 %f423, 0f3F6A937E; |
|
mov.f32 %f422, 0f3F20D842; |
|
mov.f32 %f429, %f108; |
|
@%p23 bra $L__BB0_30; |
|
mul.f32 %f429, %f16, %f16; |
|
mov.f32 %f428, 0f38B1E96A; |
|
mov.f32 %f427, 0fBA574D20; |
|
mov.f32 %f426, 0f3BAAD5EA; |
|
mov.f32 %f425, 0fBCDC1BE7; |
|
mov.f32 %f424, 0f3DE718AF; |
|
mov.f32 %f423, 0fBEC093AC; |
|
mov.f32 %f422, 0f3E0375D3; |
|
$L__BB0_30: |
|
.loc 1 0 23 |
|
ld.param.u64 %rd2, [triton__0d1d2de_param_1]; |
|
cvt.s64.s32 %rd1, %r18; |
|
.loc 1 30 23 |
|
setp.ltu.f32 %p24, %f108, 0f3F8060FE; |
|
fma.rn.ftz.f32 %f310, %f428, %f429, %f427; |
|
fma.rn.ftz.f32 %f311, %f310, %f429, %f426; |
|
fma.rn.ftz.f32 %f312, %f311, %f429, %f425; |
|
fma.rn.ftz.f32 %f313, %f312, %f429, %f424; |
|
fma.rn.ftz.f32 %f314, %f313, %f429, %f423; |
|
fma.rn.ftz.f32 %f315, %f314, %f429, %f422; |
|
neg.f32 %f316, %f429; |
|
selp.f32 %f317, %f316, %f16, %p23; |
|
fma.rn.ftz.f32 %f430, %f315, %f317, %f317; |
|
@%p24 bra $L__BB0_32; |
|
ex2.approx.ftz.f32 %f318, %f430; |
|
sub.f32 %f320, %f358, %f318; |
|
mov.b32 %r47, %f320; |
|
mov.b32 %r48, %f16; |
|
and.b32 %r49, %r48, -2147483648; |
|
or.b32 %r50, %r49, %r47; |
|
mov.b32 %f430, %r50; |
|
$L__BB0_32: |
|
.loc 1 27 18 |
|
mul.f32 %f321, %f8, 0f3F000000; |
|
mul.f32 %f322, %f7, 0f3F000000; |
|
mul.f32 %f323, %f6, 0f3F000000; |
|
mul.f32 %f324, %f5, 0f3F000000; |
|
mul.f32 %f325, %f4, 0f3F000000; |
|
mul.f32 %f326, %f3, 0f3F000000; |
|
mul.f32 %f327, %f2, 0f3F000000; |
|
mul.f32 %f328, %f1, 0f3F000000; |
|
.loc 1 32 18 |
|
add.f32 %f329, %f367, 0f3F800000; |
|
add.f32 %f330, %f376, 0f3F800000; |
|
add.f32 %f331, %f385, 0f3F800000; |
|
add.f32 %f332, %f394, 0f3F800000; |
|
add.f32 %f333, %f403, 0f3F800000; |
|
add.f32 %f334, %f412, 0f3F800000; |
|
add.f32 %f335, %f421, 0f3F800000; |
|
add.f32 %f336, %f430, 0f3F800000; |
|
.loc 1 33 18 |
|
mul.f32 %f337, %f328, %f329; |
|
mul.f32 %f338, %f327, %f330; |
|
mul.f32 %f339, %f326, %f331; |
|
mul.f32 %f340, %f325, %f332; |
|
mul.f32 %f341, %f324, %f333; |
|
mul.f32 %f342, %f323, %f334; |
|
mul.f32 %f343, %f322, %f335; |
|
mul.f32 %f344, %f321, %f336; |
|
.loc 1 35 25 |
|
shl.b64 %rd7, %rd1, 1; |
|
add.s64 %rd6, %rd2, %rd7; |
|
.loc 1 35 37 |
|
mov.b32 %r51, %f337; |
|
cvt.rn.bf16.f32 %rs9, %r51; |
|
mov.b32 %r52, %f338; |
|
cvt.rn.bf16.f32 %rs10, %r52; |
|
mov.b32 %r53, %f339; |
|
cvt.rn.bf16.f32 %rs11, %r53; |
|
mov.b32 %r54, %f340; |
|
cvt.rn.bf16.f32 %rs12, %r54; |
|
mov.b32 %r55, %f341; |
|
cvt.rn.bf16.f32 %rs13, %r55; |
|
mov.b32 %r56, %f342; |
|
cvt.rn.bf16.f32 %rs14, %r56; |
|
mov.b32 %r57, %f343; |
|
cvt.rn.bf16.f32 %rs15, %r57; |
|
mov.b32 %r58, %f344; |
|
cvt.rn.bf16.f32 %rs16, %r58; |
|
mov.b32 %r63, {%rs9, %rs10}; |
|
mov.b32 %r64, {%rs11, %rs12}; |
|
mov.b32 %r65, {%rs13, %rs14}; |
|
mov.b32 %r66, {%rs15, %rs16}; |
|
@%p1 st.global.v4.b32 [ %rd6 + 0 ], { %r63, %r64, %r65, %r66 }; |
|
.loc 1 35 4 |
|
ret; |
|
$L__tmp1: |
|
$L__func_end0: |
|
|
|
} |
|
|
|
.visible .func (.param .b32 func_retval0) __nv_erff( |
|
.param .b32 __nv_erff_param_0 |
|
) |
|
{ |
|
.reg .pred %p<4>; |
|
.reg .b32 %r<5>; |
|
.reg .f32 %f<49>; |
|
$L__func_begin1: |
|
|
|
ld.param.f32 %f14, [__nv_erff_param_0]; |
|
abs.ftz.f32 %f1, %f14; |
|
setp.ge.f32 %p1, %f1, 0f3F8060FE; |
|
mov.f32 %f46, 0f3789CA3C; |
|
mov.f32 %f45, 0fB9F560B9; |
|
mov.f32 %f44, 0f3BAC840B; |
|
mov.f32 %f43, 0fBD0C8162; |
|
mov.f32 %f42, 0f3E1CF906; |
|
mov.f32 %f41, 0f3F6A937E; |
|
mov.f32 %f40, 0f3F20D842; |
|
mov.f32 %f47, %f1; |
|
@%p1 bra $L__BB1_2; |
|
mul.f32 %f47, %f14, %f14; |
|
mov.f32 %f46, 0f38B1E96A; |
|
mov.f32 %f45, 0fBA574D20; |
|
mov.f32 %f44, 0f3BAAD5EA; |
|
mov.f32 %f43, 0fBCDC1BE7; |
|
mov.f32 %f42, 0f3DE718AF; |
|
mov.f32 %f41, 0fBEC093AC; |
|
mov.f32 %f40, 0f3E0375D3; |
|
$L__BB1_2: |
|
setp.ltu.f32 %p2, %f1, 0f3F8060FE; |
|
fma.rn.ftz.f32 %f29, %f46, %f47, %f45; |
|
fma.rn.ftz.f32 %f30, %f29, %f47, %f44; |
|
fma.rn.ftz.f32 %f31, %f30, %f47, %f43; |
|
fma.rn.ftz.f32 %f32, %f31, %f47, %f42; |
|
fma.rn.ftz.f32 %f33, %f32, %f47, %f41; |
|
fma.rn.ftz.f32 %f34, %f33, %f47, %f40; |
|
neg.f32 %f35, %f47; |
|
selp.f32 %f36, %f35, %f14, %p1; |
|
fma.rn.ftz.f32 %f48, %f34, %f36, %f36; |
|
@%p2 bra $L__BB1_4; |
|
ex2.approx.ftz.f32 %f37, %f48; |
|
mov.f32 %f38, 0f3F800000; |
|
sub.f32 %f39, %f38, %f37; |
|
mov.b32 %r1, %f39; |
|
mov.b32 %r2, %f14; |
|
and.b32 %r3, %r2, -2147483648; |
|
or.b32 %r4, %r3, %r1; |
|
mov.b32 %f48, %r4; |
|
$L__BB1_4: |
|
st.param.f32 [func_retval0+0], %f48; |
|
ret; |
|
$L__func_end1: |
|
|
|
} |
|
.file 1 "/tmp/torchinductor_root/jf/cjfoqo3nutni5cmtw4brla34cz45fusadehkxfkr2fie2qgo7vwt.py" |
|
.section .debug_abbrev |
|
{ |
|
.b8 1 |
|
.b8 17 |
|
.b8 1 |
|
.b8 37 |
|
.b8 8 |
|
.b8 19 |
|
.b8 5 |
|
.b8 3 |
|
.b8 8 |
|
.b8 16 |
|
.b8 6 |
|
.b8 27 |
|
.b8 8 |
|
.b8 180 |
|
.b8 66 |
|
.b8 12 |
|
.b8 17 |
|
.b8 1 |
|
.b8 18 |
|
.b8 1 |
|
.b8 0 |
|
.b8 0 |
|
.b8 2 |
|
.b8 46 |
|
.b8 0 |
|
.b8 17 |
|
.b8 1 |
|
.b8 18 |
|
.b8 1 |
|
.b8 64 |
|
.b8 10 |
|
.b8 135 |
|
.b8 64 |
|
.b8 8 |
|
.b8 3 |
|
.b8 8 |
|
.b8 58 |
|
.b8 11 |
|
.b8 59 |
|
.b8 11 |
|
.b8 63 |
|
.b8 12 |
|
.b8 0 |
|
.b8 0 |
|
.b8 0 |
|
} |
|
.section .debug_info |
|
{ |
|
.b32 176 |
|
.b8 2 |
|
.b8 0 |
|
.b32 .debug_abbrev |
|
.b8 8 |
|
.b8 1 |
|
.b8 116 |
|
.b8 114 |
|
.b8 105 |
|
.b8 116 |
|
.b8 111 |
|
.b8 110 |
|
.b8 0 |
|
.b8 2 |
|
.b8 0 |
|
.b8 99 |
|
.b8 106 |
|
.b8 102 |
|
.b8 111 |
|
.b8 113 |
|
.b8 111 |
|
.b8 51 |
|
.b8 110 |
|
.b8 117 |
|
.b8 116 |
|
.b8 110 |
|
.b8 105 |
|
.b8 53 |
|
.b8 99 |
|
.b8 109 |
|
.b8 116 |
|
.b8 119 |
|
.b8 52 |
|
.b8 98 |
|
.b8 114 |
|
.b8 108 |
|
.b8 97 |
|
.b8 51 |
|
.b8 52 |
|
.b8 99 |
|
.b8 122 |
|
.b8 52 |
|
.b8 53 |
|
.b8 102 |
|
.b8 117 |
|
.b8 115 |
|
.b8 97 |
|
.b8 100 |
|
.b8 101 |
|
.b8 104 |
|
.b8 107 |
|
.b8 120 |
|
.b8 102 |
|
.b8 107 |
|
.b8 114 |
|
.b8 50 |
|
.b8 102 |
|
.b8 105 |
|
.b8 101 |
|
.b8 50 |
|
.b8 113 |
|
.b8 103 |
|
.b8 111 |
|
.b8 55 |
|
.b8 118 |
|
.b8 119 |
|
.b8 116 |
|
.b8 46 |
|
.b8 112 |
|
.b8 121 |
|
.b8 0 |
|
.b32 .debug_line |
|
.b8 47 |
|
.b8 116 |
|
.b8 109 |
|
.b8 112 |
|
.b8 47 |
|
.b8 116 |
|
.b8 111 |
|
.b8 114 |
|
.b8 99 |
|
.b8 104 |
|
.b8 105 |
|
.b8 110 |
|
.b8 100 |
|
.b8 117 |
|
.b8 99 |
|
.b8 116 |
|
.b8 111 |
|
.b8 114 |
|
.b8 95 |
|
.b8 114 |
|
.b8 111 |
|
.b8 111 |
|
.b8 116 |
|
.b8 47 |
|
.b8 106 |
|
.b8 102 |
|
.b8 0 |
|
.b8 1 |
|
.b64 $L__func_begin0 |
|
.b64 $L__func_end0 |
|
.b8 2 |
|
.b64 $L__func_begin0 |
|
.b64 $L__func_end0 |
|
.b8 1 |
|
.b8 156 |
|
.b8 116 |
|
.b8 114 |
|
.b8 105 |
|
.b8 116 |
|
.b8 111 |
|
.b8 110 |
|
.b8 95 |
|
.b8 95 |
|
.b8 48 |
|
.b8 100 |
|
.b8 49 |
|
.b8 100 |
|
.b8 50 |
|
.b8 100 |
|
.b8 101 |
|
.b8 0 |
|
.b8 116 |
|
.b8 114 |
|
.b8 105 |
|
.b8 116 |
|
.b8 111 |
|
.b8 110 |
|
.b8 95 |
|
.b8 95 |
|
.b8 48 |
|
.b8 100 |
|
.b8 49 |
|
.b8 100 |
|
.b8 50 |
|
.b8 100 |
|
.b8 101 |
|
.b8 0 |
|
.b8 1 |
|
.b8 18 |
|
.b8 1 |
|
.b8 0 |
|
} |
|
.section .debug_pubnames |
|
{ |
|
.b32 $L__pubNames_end0-$L__pubNames_start0 |
|
$L__pubNames_start0: |
|
.b8 2 |
|
.b8 0 |
|
.b32 .debug_info |
|
.b32 180 |
|
.b32 125 |
|
.b8 116 |
|
.b8 114 |
|
.b8 105 |
|
.b8 116 |
|
.b8 111 |
|
.b8 110 |
|
.b8 95 |
|
.b8 95 |
|
.b8 48 |
|
.b8 100 |
|
.b8 49 |
|
.b8 100 |
|
.b8 50 |
|
.b8 100 |
|
.b8 101 |
|
.b8 0 |
|
.b32 0 |
|
$L__pubNames_end0: |
|
} |
|
.section .debug_pubtypes |
|
{ |
|
.b32 $L__pubTypes_end0-$L__pubTypes_start0 |
|
$L__pubTypes_start0: |
|
.b8 2 |
|
.b8 0 |
|
.b32 .debug_info |
|
.b32 180 |
|
.b32 0 |
|
$L__pubTypes_end0: |
|
} |
|
.section .debug_loc { } |
|
|
