[llvm] [AMDGPU][MC] Adding symbolic name for the constant 1 over 2*pi (PR #160617)
Jun Wang via llvm-commits
llvm-commits at lists.llvm.org
Fri Nov 7 14:52:30 PST 2025
https://github.com/jwanggit86 updated https://github.com/llvm/llvm-project/pull/160617
>From f95cf322362f7d3df260fb5e76aa094d0424f60f Mon Sep 17 00:00:00 2001
From: Jun Wang <jwang86 at yahoo.com>
Date: Wed, 24 Sep 2025 16:33:03 -0700
Subject: [PATCH 1/3] [AMDGPU][MC] Adding symbolic name for the constant 1 over
2*pi
Adding a symbolic name, INV2PI, which can be used in assembly
instrucitons to represent 0.15915494 or 0.15915494309189532,
depending on data size.
---
.../Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp | 16 ++++++++++++++--
llvm/test/MC/AMDGPU/bf16_imm-fake16.s | 7 +++++++
llvm/test/MC/AMDGPU/bf16_imm.s | 3 +++
llvm/test/MC/AMDGPU/gfx1250_asm_valu_lit64.s | 7 +++++++
llvm/test/MC/AMDGPU/gfx1250_asm_vop3.s | 6 ++++++
llvm/test/MC/AMDGPU/gfx7_err_pos.s | 10 ++++++++++
llvm/test/MC/AMDGPU/gfx950_dlops.s | 3 +++
llvm/test/MC/AMDGPU/literal16.s | 6 ++++++
llvm/test/MC/AMDGPU/literals.s | 4 ++++
llvm/test/MC/AMDGPU/literalv216-err.s | 6 ++++++
llvm/test/MC/AMDGPU/literalv216.s | 4 ++++
llvm/test/MC/AMDGPU/mai-gfx950-err.s | 6 ++++++
llvm/test/MC/AMDGPU/mubuf.s | 4 ++++
13 files changed, 80 insertions(+), 2 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index 09338c533fdf2..59b3cd3ba0ab2 100644
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -3227,6 +3227,13 @@ AMDGPUAsmParser::parseRegister(bool RestoreOnFailure) {
return AMDGPUOperand::CreateReg(this, Reg, StartLoc, EndLoc);
}
+static bool isInv2PiToken(const AsmToken &Tok) {
+ if (!Tok.is(AsmToken::Identifier))
+ return false;
+ StringRef Str = Tok.getIdentifier();
+ return (Str.str() == "INV2PI" || Str.str() == "INV2PI64");
+}
+
ParseStatus AMDGPUAsmParser::parseImm(OperandVector &Operands,
bool HasSP3AbsModifier, LitModifier Lit) {
// TODO: add syntactic sugar for 1/(2*PI)
@@ -3253,11 +3260,12 @@ ParseStatus AMDGPUAsmParser::parseImm(OperandVector &Operands,
const auto& Tok = getToken();
const auto& NextTok = peekToken();
- bool IsReal = Tok.is(AsmToken::Real);
+ bool IsReal = Tok.is(AsmToken::Real) || isInv2PiToken(Tok);
SMLoc S = getLoc();
bool Negate = false;
- if (!IsReal && Tok.is(AsmToken::Minus) && NextTok.is(AsmToken::Real)) {
+ if (!IsReal && Tok.is(AsmToken::Minus) &&
+ (NextTok.is(AsmToken::Real) || isInv2PiToken(NextTok))) {
lex();
IsReal = true;
Negate = true;
@@ -3272,6 +3280,10 @@ ParseStatus AMDGPUAsmParser::parseImm(OperandVector &Operands,
// optional sign.
StringRef Num = getTokenStr();
+ if (Num.str() == "INV2PI")
+ Num = "0.15915494";
+ else if (Num.str() == "INV2PI64")
+ Num = "0.15915494309189532";
lex();
APFloat RealVal(APFloat::IEEEdouble());
diff --git a/llvm/test/MC/AMDGPU/bf16_imm-fake16.s b/llvm/test/MC/AMDGPU/bf16_imm-fake16.s
index ee697bee6ab2d..c8cdd2bf3588e 100644
--- a/llvm/test/MC/AMDGPU/bf16_imm-fake16.s
+++ b/llvm/test/MC/AMDGPU/bf16_imm-fake16.s
@@ -80,6 +80,13 @@ v_dot2_f32_bf16 v2, v1, -4.0, v2
v_dot2_f32_bf16 v2, v1, 0.15915494, v2
// CHECK: v_dot2_f32_bf16 v2, v1, 0.15915494, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xf1,0x09,0x1c]
+v_dot2_f32_bf16 v2, v1, INV2PI, v2
+// CHECK: v_dot2_f32_bf16 v2, v1, 0.15915494, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xf1,0x09,0x1c]
+
+// INV2PI64 is not interpreted as INV2PI
+v_dot2_f32_bf16 v2, v1, INV2PI64, v2
+// CHECK: v_dot2_f32_bf16 v2, v1, 0x3e23, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xff,0x09,0x1c,0x23,0x3e,0x00,0x00]
+
v_dot2_f32_bf16 v2, v1, 0x3e22, v2
// CHECK: v_dot2_f32_bf16 v2, v1, 0.15915494, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xf1,0x09,0x1c]
diff --git a/llvm/test/MC/AMDGPU/bf16_imm.s b/llvm/test/MC/AMDGPU/bf16_imm.s
index d79649073aa89..d734692c84b0e 100644
--- a/llvm/test/MC/AMDGPU/bf16_imm.s
+++ b/llvm/test/MC/AMDGPU/bf16_imm.s
@@ -50,6 +50,9 @@ v_dot2_bf16_bf16 v2.l, v0, 0x3e22, v2.l
v_dot2_bf16_bf16 v2.l, v0, v2, 0.15915494
// CHECK: v_dot2_bf16_bf16 v2.l, v0, v2, 0.15915494 ; encoding: [0x02,0x00,0x67,0xd6,0x00,0x05,0xe2,0x03]
+v_dot2_bf16_bf16 v2.l, v0, v2, INV2PI
+// CHECK: v_dot2_bf16_bf16 v2.l, v0, v2, 0.15915494 ; encoding: [0x02,0x00,0x67,0xd6,0x00,0x05,0xe2,0x03]
+
v_dot2_f32_bf16 v2, v1, 0, v2
// CHECK: v_dot2_f32_bf16 v2, v1, 0, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0x01,0x09,0x1c]
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_valu_lit64.s b/llvm/test/MC/AMDGPU/gfx1250_asm_valu_lit64.s
index 58da119f20e54..2ad2a90c96327 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_valu_lit64.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_valu_lit64.s
@@ -235,6 +235,13 @@ v_ceil_f64 v[254:255], 0x3fc45f306dc9c882
v_ceil_f64 v[254:255], 0.15915494309189532
// GFX1250: v_ceil_f64_e32 v[254:255], 0.15915494309189532 ; encoding: [0xf8,0x30,0xfc,0x7f]
+v_ceil_f64 v[254:255], INV2PI64
+// GFX1250: v_ceil_f64_e32 v[254:255], 0.15915494309189532 ; encoding: [0xf8,0x30,0xfc,0x7f]
+
+// using INV2PI produces a different result
+v_ceil_f64 v[254:255], INV2PI
+// GFX1250: v_ceil_f64_e32 v[254:255], lit64(0x3fc45f306725feed) ; encoding: [0xfe,0x30,0xfc,0x7f,0xed,0xfe,0x25,0x67,0x30,0x5f,0xc4,0x3f]
+
v_ceil_f64 v[254:255], -4.0
// GFX1250: v_ceil_f64_e32 v[254:255], -4.0 ; encoding: [0xf7,0x30,0xfc,0x7f]
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3.s b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3.s
index d913bd2db504b..4b6bf64267e53 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3.s
@@ -447,6 +447,9 @@ v_cvt_pk_bf8_f16 v1.l, 0x3118
v_cvt_pk_bf8_f16 v1.l, 0.15915494
// GFX1250: v_cvt_pk_bf8_f16 v1.l, 0x3118 ; encoding: [0x01,0x00,0x73,0xd7,0xff,0x00,0x00,0x00,0x18,0x31,0x00,0x00]
+v_cvt_pk_bf8_f16 v1.l, INV2PI
+// GFX1250: v_cvt_pk_bf8_f16 v1.l, 0x3118 ; encoding: [0x01,0x00,0x73,0xd7,0xff,0x00,0x00,0x00,0x18,0x31,0x00,0x00]
+
v_cvt_pk_fp8_f16 v1.l, v2
// GFX1250: v_cvt_pk_fp8_f16 v1.l, v2 ; encoding: [0x01,0x00,0x72,0xd7,0x02,0x01,0x00,0x00]
@@ -479,6 +482,9 @@ v_cvt_pk_fp8_f16 v1.l, 0x3118
v_cvt_pk_fp8_f16 v1.l, 0.15915494
// GFX1250: v_cvt_pk_fp8_f16 v1.l, 0x3118 ; encoding: [0x01,0x00,0x72,0xd7,0xff,0x00,0x00,0x00,0x18,0x31,0x00,0x00]
+v_cvt_pk_fp8_f16 v1.l, INV2PI
+// GFX1250: v_cvt_pk_fp8_f16 v1.l, 0x3118 ; encoding: [0x01,0x00,0x72,0xd7,0xff,0x00,0x00,0x00,0x18,0x31,0x00,0x00]
+
v_cvt_pk_f16_f32 v5, v1, v2
// GFX1250: v_cvt_pk_f16_f32 v5, v1, v2 ; encoding: [0x05,0x00,0x6f,0xd7,0x01,0x05,0x02,0x00]
diff --git a/llvm/test/MC/AMDGPU/gfx7_err_pos.s b/llvm/test/MC/AMDGPU/gfx7_err_pos.s
index 7b6b241e04707..f81b1e86fef04 100644
--- a/llvm/test/MC/AMDGPU/gfx7_err_pos.s
+++ b/llvm/test/MC/AMDGPU/gfx7_err_pos.s
@@ -37,6 +37,16 @@ v_and_b32_e64 v0, 0.159154943091895317852646485335, v1
// CHECK-NEXT:{{^}}v_and_b32_e64 v0, 0.159154943091895317852646485335, v1
// CHECK-NEXT:{{^}} ^
+v_and_b32_e64 v0, INV2PI, v1
+// CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: literal operands are not supported
+// CHECK-NEXT:{{^}}v_and_b32_e64 v0, INV2PI, v1
+// CHECK-NEXT:{{^}} ^
+
+v_and_b32_e64 v0, INV2PI64, v1
+// CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: literal operands are not supported
+// CHECK-NEXT:{{^}}v_and_b32_e64 v0, INV2PI64, v1
+// CHECK-NEXT:{{^}} ^
+
//==============================================================================
// cache policy is not supported for SMRD instructions
diff --git a/llvm/test/MC/AMDGPU/gfx950_dlops.s b/llvm/test/MC/AMDGPU/gfx950_dlops.s
index 4ae60ac785f49..52a4167f2fe0e 100644
--- a/llvm/test/MC/AMDGPU/gfx950_dlops.s
+++ b/llvm/test/MC/AMDGPU/gfx950_dlops.s
@@ -36,6 +36,9 @@ v_dot2_f32_bf16 v2, v1, -4.0, v2
v_dot2_f32_bf16 v2, v1, 0.15915494, v2
// GFX950: v_dot2_f32_bf16 v2, v1, 0.15915494, v2 ; encoding: [0x02,0x40,0x9a,0xd3,0x01,0xf1,0x09,0x1c]
+v_dot2_f32_bf16 v2, v1, INV2PI, v2
+// GFX950: v_dot2_f32_bf16 v2, v1, 0.15915494, v2 ; encoding: [0x02,0x40,0x9a,0xd3,0x01,0xf1,0x09,0x1c]
+
v_dot2_f32_bf16 v2, 0.5, v1, v2
// GFX950: v_dot2_f32_bf16 v2, 0.5, v1, v2 ; encoding: [0x02,0x40,0x9a,0xd3,0xf0,0x02,0x0a,0x1c]
diff --git a/llvm/test/MC/AMDGPU/literal16.s b/llvm/test/MC/AMDGPU/literal16.s
index fd8ab05ec3ab4..bf5028c17e077 100644
--- a/llvm/test/MC/AMDGPU/literal16.s
+++ b/llvm/test/MC/AMDGPU/literal16.s
@@ -42,9 +42,15 @@ v_add_f16 v1, -4.0, v2
v_add_f16 v1, 0.15915494, v2
// VI: v_add_f16_e32 v1, 0.15915494, v2 ; encoding: [0xf8,0x04,0x02,0x3e]
+v_add_f16 v1, INV2PI, v2
+// VI: v_add_f16_e32 v1, 0.15915494, v2 ; encoding: [0xf8,0x04,0x02,0x3e]
+
v_add_f16 v1, -0.15915494, v2
// VI: v_add_f16_e32 v1, 0xb118, v2 ; encoding: [0xff,0x04,0x02,0x3e,0x18,0xb1,0x00,0x00]
+v_add_f16 v1, -INV2PI, v2
+// VI: v_add_f16_e32 v1, 0xb118, v2 ; encoding: [0xff,0x04,0x02,0x3e,0x18,0xb1,0x00,0x00]
+
v_add_f16 v1, -1, v2
// VI: v_add_f16_e32 v1, -1, v2 ; encoding: [0xc1,0x04,0x02,0x3e]
diff --git a/llvm/test/MC/AMDGPU/literals.s b/llvm/test/MC/AMDGPU/literals.s
index be4e0defa5760..bcb2c59bc56e9 100644
--- a/llvm/test/MC/AMDGPU/literals.s
+++ b/llvm/test/MC/AMDGPU/literals.s
@@ -1037,6 +1037,10 @@ v_trunc_f32 v0, lit(0.159154943091895317852646485335)
// GFX89: v_trunc_f32_e32 v0, lit(0x3e22f983) ; encoding: [0xff,0x38,0x00,0x7e,0x83,0xf9,0x22,0x3e]
// SICI: v_trunc_f32_e32 v0, lit(0x3e22f983) ; encoding: [0xff,0x42,0x00,0x7e,0x83,0xf9,0x22,0x3e]
+// SICI: v_trunc_f32_e32 v0, 0x3e22f983 ; encoding: [0xff,0x42,0x00,0x7e,0x83,0xf9,0x22,0x3e]
+// GFX89: v_trunc_f32_e32 v0, 0.15915494 ; encoding: [0xf8,0x38,0x00,0x7e]
+v_trunc_f32 v0, INV2PI64
+
//---------------------------------------------------------------------------//
// integer literal truncation checks
//---------------------------------------------------------------------------//
diff --git a/llvm/test/MC/AMDGPU/literalv216-err.s b/llvm/test/MC/AMDGPU/literalv216-err.s
index 73148711d0425..a52f72d28e9ce 100644
--- a/llvm/test/MC/AMDGPU/literalv216-err.s
+++ b/llvm/test/MC/AMDGPU/literalv216-err.s
@@ -13,6 +13,12 @@ v_pk_add_f16 v1, 64.0, v2
v_pk_add_f16 v1, -0.15915494, v2
// GFX9: :[[@LINE-1]]:{{[0-9]+}}: error: literal operands are not supported
+v_pk_add_f16 v1, -INV2PI, v2
+// GFX9: :[[@LINE-1]]:{{[0-9]+}}: error: literal operands are not supported
+
+v_pk_add_f16 v1, -INV2PI64, v2
+// GFX9: :[[@LINE-1]]:{{[0-9]+}}: error: literal operands are not supported
+
v_pk_add_f16 v1, -0.0, v2
// GFX9: :[[@LINE-1]]:{{[0-9]+}}: error: literal operands are not supported
diff --git a/llvm/test/MC/AMDGPU/literalv216.s b/llvm/test/MC/AMDGPU/literalv216.s
index c695bc3600c38..20c3d34c0bb86 100644
--- a/llvm/test/MC/AMDGPU/literalv216.s
+++ b/llvm/test/MC/AMDGPU/literalv216.s
@@ -60,6 +60,10 @@ v_pk_add_f16 v1, 0.15915494, v2
// GFX9: v_pk_add_f16 v1, 0.15915494, v2 ; encoding: [0x01,0x40,0x8f,0xd3,0xf8,0x04,0x02,0x18]
// GFX10: v_pk_add_f16 v1, 0.15915494, v2 ; encoding: [0x01,0x40,0x0f,0xcc,0xf8,0x04,0x02,0x18]
+v_pk_add_f16 v1, INV2PI, v2
+// GFX9: v_pk_add_f16 v1, 0.15915494, v2 ; encoding: [0x01,0x40,0x8f,0xd3,0xf8,0x04,0x02,0x18]
+// GFX10: v_pk_add_f16 v1, 0.15915494, v2 ; encoding: [0x01,0x40,0x0f,0xcc,0xf8,0x04,0x02,0x18]
+
v_pk_add_f16 v1, -1, v2
// GFX9: v_pk_add_f16 v1, -1, v2 ; encoding: [0x01,0x40,0x8f,0xd3,0xc1,0x04,0x02,0x18]
// GFX10: v_pk_add_f16 v1, -1, v2 ; encoding: [0x01,0x40,0x0f,0xcc,0xc1,0x04,0x02,0x18]
diff --git a/llvm/test/MC/AMDGPU/mai-gfx950-err.s b/llvm/test/MC/AMDGPU/mai-gfx950-err.s
index 747deab3bfcae..9848c75a66937 100644
--- a/llvm/test/MC/AMDGPU/mai-gfx950-err.s
+++ b/llvm/test/MC/AMDGPU/mai-gfx950-err.s
@@ -195,6 +195,9 @@ v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], -4.0, v24
v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], 0.15915494, v24
// CHECK: :[[@LINE-1]]:72: error: invalid operand for instruction
+v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], INV2PI, v24
+// CHECK: :[[@LINE-1]]:72: error: invalid operand for instruction
+
v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], 16, v49
// CHECK: :[[@LINE-1]]:73: error: invalid operand for instruction
@@ -206,3 +209,6 @@ v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], 4.0, v24
v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], 0.15915494, v24
// CHECK: :[[@LINE-1]]:73: error: invalid operand for instruction
+
+v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], INV2PI, v24
+// CHECK: :[[@LINE-1]]:73: error: invalid operand for instruction
diff --git a/llvm/test/MC/AMDGPU/mubuf.s b/llvm/test/MC/AMDGPU/mubuf.s
index 600580ac51556..2ea6a3b5150a7 100644
--- a/llvm/test/MC/AMDGPU/mubuf.s
+++ b/llvm/test/MC/AMDGPU/mubuf.s
@@ -647,6 +647,10 @@ buffer_atomic_add v5, off, s[8:11], 0.15915494 offset:4095 glc
// NOSICI: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction
// VI: buffer_atomic_add v5, off, s[8:11], 0.15915494 offset:4095 glc ; encoding: [0xff,0x4f,0x08,0xe1,0x00,0x05,0x02,0xf8]
+buffer_atomic_add v5, off, s[8:11], INV2PI offset:4095 glc
+// NOSICI: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction
+// VI: buffer_atomic_add v5, off, s[8:11], 0.15915494 offset:4095 glc ; encoding: [0xff,0x4f,0x08,0xe1,0x00,0x05,0x02,0xf8]
+
buffer_atomic_fcmpswap v[0:1], off, s[0:3], s0 offset:4095
// SICI: buffer_atomic_fcmpswap v[0:1], off, s[0:3], s0 offset:4095 ; encoding: [0xff,0x0f,0xf8,0xe0,0x00,0x00,0x00,0x00]
// NOVI: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
>From 2c19357ae2f0aeaf5723ba5abf3533f6ea84ab31 Mon Sep 17 00:00:00 2001
From: Jun Wang <jwang86 at yahoo.com>
Date: Wed, 22 Oct 2025 16:11:55 -0700
Subject: [PATCH 2/3] Changing two symbolic names to one
---
.../Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp | 15 +++++++--------
llvm/test/MC/AMDGPU/bf16_imm-fake16.s | 4 ----
llvm/test/MC/AMDGPU/gfx1250_asm_valu_lit64.s | 6 +-----
llvm/test/MC/AMDGPU/gfx7_err_pos.s | 5 -----
llvm/test/MC/AMDGPU/literals.s | 2 +-
llvm/test/MC/AMDGPU/literalv216-err.s | 3 ---
6 files changed, 9 insertions(+), 26 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index 59b3cd3ba0ab2..55e8e9121d6b2 100644
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -2408,8 +2408,10 @@ void AMDGPUOperand::addLiteralImmOperand(MCInst &Inst, int64_t Val, bool ApplyMo
case AMDGPU::OPERAND_REG_INLINE_C_BF16:
case AMDGPU::OPERAND_REG_INLINE_C_V2BF16:
case AMDGPU::OPERAND_REG_IMM_V2BF16:
+ // If the symbol INV2PI is used as the operand, the value is set to
+ // 0x3fc45f306dc9c882 in parseImm().
if (Lit == LitModifier::None && AsmParser->hasInv2PiInlineImm() &&
- Literal == 0x3fc45f306725feed) {
+ (Literal == 0x3fc45f306725feed || Literal == 0x3fc45f306dc9c882)) {
// This is the 1/(2*pi) which is going to be truncated to bf16 with the
// loss of precision. The constant represents ideomatic fp32 value of
// 1/(2*pi) = 0.15915494 since bf16 is in fact fp32 with cleared low 16
@@ -3228,10 +3230,7 @@ AMDGPUAsmParser::parseRegister(bool RestoreOnFailure) {
}
static bool isInv2PiToken(const AsmToken &Tok) {
- if (!Tok.is(AsmToken::Identifier))
- return false;
- StringRef Str = Tok.getIdentifier();
- return (Str.str() == "INV2PI" || Str.str() == "INV2PI64");
+ return Tok.is(AsmToken::Identifier) && Tok.getIdentifier() == "INV2PI";
}
ParseStatus AMDGPUAsmParser::parseImm(OperandVector &Operands,
@@ -3280,9 +3279,9 @@ ParseStatus AMDGPUAsmParser::parseImm(OperandVector &Operands,
// optional sign.
StringRef Num = getTokenStr();
- if (Num.str() == "INV2PI")
- Num = "0.15915494";
- else if (Num.str() == "INV2PI64")
+ if (Num == "INV2PI")
+ // Setting the imm to this for INV2PI works for all types except bf16.
+ // In addLiteralImmOperand() we specifically check for this.
Num = "0.15915494309189532";
lex();
diff --git a/llvm/test/MC/AMDGPU/bf16_imm-fake16.s b/llvm/test/MC/AMDGPU/bf16_imm-fake16.s
index c8cdd2bf3588e..482bb3c87ce0b 100644
--- a/llvm/test/MC/AMDGPU/bf16_imm-fake16.s
+++ b/llvm/test/MC/AMDGPU/bf16_imm-fake16.s
@@ -83,10 +83,6 @@ v_dot2_f32_bf16 v2, v1, 0.15915494, v2
v_dot2_f32_bf16 v2, v1, INV2PI, v2
// CHECK: v_dot2_f32_bf16 v2, v1, 0.15915494, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xf1,0x09,0x1c]
-// INV2PI64 is not interpreted as INV2PI
-v_dot2_f32_bf16 v2, v1, INV2PI64, v2
-// CHECK: v_dot2_f32_bf16 v2, v1, 0x3e23, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xff,0x09,0x1c,0x23,0x3e,0x00,0x00]
-
v_dot2_f32_bf16 v2, v1, 0x3e22, v2
// CHECK: v_dot2_f32_bf16 v2, v1, 0.15915494, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xf1,0x09,0x1c]
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_valu_lit64.s b/llvm/test/MC/AMDGPU/gfx1250_asm_valu_lit64.s
index 2ad2a90c96327..4c0d5a69c2907 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_valu_lit64.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_valu_lit64.s
@@ -235,12 +235,8 @@ v_ceil_f64 v[254:255], 0x3fc45f306dc9c882
v_ceil_f64 v[254:255], 0.15915494309189532
// GFX1250: v_ceil_f64_e32 v[254:255], 0.15915494309189532 ; encoding: [0xf8,0x30,0xfc,0x7f]
-v_ceil_f64 v[254:255], INV2PI64
-// GFX1250: v_ceil_f64_e32 v[254:255], 0.15915494309189532 ; encoding: [0xf8,0x30,0xfc,0x7f]
-
-// using INV2PI produces a different result
v_ceil_f64 v[254:255], INV2PI
-// GFX1250: v_ceil_f64_e32 v[254:255], lit64(0x3fc45f306725feed) ; encoding: [0xfe,0x30,0xfc,0x7f,0xed,0xfe,0x25,0x67,0x30,0x5f,0xc4,0x3f]
+// GFX1250: v_ceil_f64_e32 v[254:255], 0.15915494309189532 ; encoding: [0xf8,0x30,0xfc,0x7f]
v_ceil_f64 v[254:255], -4.0
// GFX1250: v_ceil_f64_e32 v[254:255], -4.0 ; encoding: [0xf7,0x30,0xfc,0x7f]
diff --git a/llvm/test/MC/AMDGPU/gfx7_err_pos.s b/llvm/test/MC/AMDGPU/gfx7_err_pos.s
index f81b1e86fef04..d6ce10d6fac72 100644
--- a/llvm/test/MC/AMDGPU/gfx7_err_pos.s
+++ b/llvm/test/MC/AMDGPU/gfx7_err_pos.s
@@ -42,11 +42,6 @@ v_and_b32_e64 v0, INV2PI, v1
// CHECK-NEXT:{{^}}v_and_b32_e64 v0, INV2PI, v1
// CHECK-NEXT:{{^}} ^
-v_and_b32_e64 v0, INV2PI64, v1
-// CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: literal operands are not supported
-// CHECK-NEXT:{{^}}v_and_b32_e64 v0, INV2PI64, v1
-// CHECK-NEXT:{{^}} ^
-
//==============================================================================
// cache policy is not supported for SMRD instructions
diff --git a/llvm/test/MC/AMDGPU/literals.s b/llvm/test/MC/AMDGPU/literals.s
index bcb2c59bc56e9..f749f6e16fb03 100644
--- a/llvm/test/MC/AMDGPU/literals.s
+++ b/llvm/test/MC/AMDGPU/literals.s
@@ -1039,7 +1039,7 @@ v_trunc_f32 v0, lit(0.159154943091895317852646485335)
// SICI: v_trunc_f32_e32 v0, 0x3e22f983 ; encoding: [0xff,0x42,0x00,0x7e,0x83,0xf9,0x22,0x3e]
// GFX89: v_trunc_f32_e32 v0, 0.15915494 ; encoding: [0xf8,0x38,0x00,0x7e]
-v_trunc_f32 v0, INV2PI64
+v_trunc_f32 v0, INV2PI
//---------------------------------------------------------------------------//
// integer literal truncation checks
diff --git a/llvm/test/MC/AMDGPU/literalv216-err.s b/llvm/test/MC/AMDGPU/literalv216-err.s
index a52f72d28e9ce..fdb262db40ee5 100644
--- a/llvm/test/MC/AMDGPU/literalv216-err.s
+++ b/llvm/test/MC/AMDGPU/literalv216-err.s
@@ -16,9 +16,6 @@ v_pk_add_f16 v1, -0.15915494, v2
v_pk_add_f16 v1, -INV2PI, v2
// GFX9: :[[@LINE-1]]:{{[0-9]+}}: error: literal operands are not supported
-v_pk_add_f16 v1, -INV2PI64, v2
-// GFX9: :[[@LINE-1]]:{{[0-9]+}}: error: literal operands are not supported
-
v_pk_add_f16 v1, -0.0, v2
// GFX9: :[[@LINE-1]]:{{[0-9]+}}: error: literal operands are not supported
>From 3a3ef24b6f1db765048977f8f33f0f15c21a393b Mon Sep 17 00:00:00 2001
From: Jun Wang <jwang86 at yahoo.com>
Date: Fri, 7 Nov 2025 14:50:57 -0800
Subject: [PATCH 3/3] Make printer print INV2PI instead of of literal numbers
to support assembler-disassembler round-trip.
---
.../AMDGPU/MCTargetDesc/AMDGPUInstPrinter.cpp | 8 +--
.../GlobalISel/llvm.amdgcn.writelane.ll | 4 +-
.../AMDGPU/bitreverse-inline-immediates.ll | 2 +-
llvm/test/CodeGen/AMDGPU/fneg-combines.f16.ll | 24 ++++-----
llvm/test/CodeGen/AMDGPU/fneg-combines.ll | 14 ++---
llvm/test/CodeGen/AMDGPU/fneg-combines.new.ll | 42 +++++++--------
llvm/test/CodeGen/AMDGPU/imm.ll | 8 +--
llvm/test/CodeGen/AMDGPU/known-never-snan.ll | 4 +-
....amdgcn.mfma.scale.f32.16x16x128.f8f6f4.ll | 4 +-
...m.amdgcn.mfma.scale.f32.32x32x64.f8f6f4.ll | 4 +-
llvm/test/CodeGen/AMDGPU/llvm.cos.f16.ll | 36 ++++++-------
llvm/test/CodeGen/AMDGPU/llvm.sin.f16.ll | 36 ++++++-------
llvm/test/CodeGen/AMDGPU/mad-mix-bf16.ll | 4 +-
llvm/test/CodeGen/AMDGPU/mad-mix.ll | 44 ++++++++--------
.../AMDGPU/select-fabs-fneg-extract.ll | 6 +--
llvm/test/MC/AMDGPU/bf16_imm-fake16.s | 14 ++---
llvm/test/MC/AMDGPU/bf16_imm.s | 14 ++---
llvm/test/MC/AMDGPU/gfx1250_asm_valu_lit64.s | 6 +--
llvm/test/MC/AMDGPU/gfx950_dlops.s | 4 +-
llvm/test/MC/AMDGPU/inline-imm-inv2pi.s | 2 +-
llvm/test/MC/AMDGPU/literal16.s | 6 +--
llvm/test/MC/AMDGPU/literals.s | 52 +++++++++----------
llvm/test/MC/AMDGPU/literalv216.s | 12 ++---
llvm/test/MC/AMDGPU/mubuf.s | 4 +-
llvm/test/MC/Disassembler/AMDGPU/bf16_imm.txt | 8 +--
.../AMDGPU/gfx10_vop3p_literalv216.txt | 2 +-
.../AMDGPU/gfx1250_dasm_valu_lit64.txt | 2 +-
.../MC/Disassembler/AMDGPU/gfx8-literal.txt | 4 +-
.../MC/Disassembler/AMDGPU/gfx8-literal16.txt | 2 +-
.../Disassembler/AMDGPU/gfx950_dasm_vop3.txt | 2 +-
30 files changed, 187 insertions(+), 187 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUInstPrinter.cpp b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUInstPrinter.cpp
index 703ec0a4befa5..5eaed83a544e6 100644
--- a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUInstPrinter.cpp
+++ b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUInstPrinter.cpp
@@ -524,7 +524,7 @@ static bool printImmediateFP16(uint32_t Imm, const MCSubtargetInfo &STI,
else if (Imm == 0xC400)
O << "-4.0";
else if (Imm == 0x3118 && STI.hasFeature(AMDGPU::FeatureInv2PiInlineImm))
- O << "0.15915494";
+ O << "INV2PI";
else
return false;
@@ -550,7 +550,7 @@ static bool printImmediateBFloat16(uint32_t Imm, const MCSubtargetInfo &STI,
else if (Imm == 0xC080)
O << "-4.0";
else if (Imm == 0x3E22 && STI.hasFeature(AMDGPU::FeatureInv2PiInlineImm))
- O << "0.15915494";
+ O << "INV2PI";
else
return false;
@@ -648,7 +648,7 @@ bool AMDGPUInstPrinter::printImmediateFloat32(uint32_t Imm,
O << "-4.0";
else if (Imm == 0x3e22f983 &&
STI.hasFeature(AMDGPU::FeatureInv2PiInlineImm))
- O << "0.15915494";
+ O << "INV2PI";
else
return false;
@@ -699,7 +699,7 @@ void AMDGPUInstPrinter::printImmediate64(uint64_t Imm,
O << "-4.0";
else if (Imm == 0x3fc45f306dc9c882 &&
STI.hasFeature(AMDGPU::FeatureInv2PiInlineImm))
- O << "0.15915494309189532";
+ O << "INV2PI";
else
printLiteral64(Imm, O, IsFP);
}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.writelane.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.writelane.ll
index 521300bdd477b..b732e3a22868a 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.writelane.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.writelane.ll
@@ -112,12 +112,12 @@ define amdgpu_ps float @test_writelane_imminv2pi_s_v(i32 inreg %lane, i32 %vdst.
;
; GFX8-LABEL: test_writelane_imminv2pi_s_v:
; GFX8: ; %bb.0:
-; GFX8-NEXT: v_writelane_b32 v0, 0.15915494, s2
+; GFX8-NEXT: v_writelane_b32 v0, INV2PI, s2
; GFX8-NEXT: ; return to shader part epilog
;
; GFX10-LABEL: test_writelane_imminv2pi_s_v:
; GFX10: ; %bb.0:
-; GFX10-NEXT: v_writelane_b32 v0, 0.15915494, s2
+; GFX10-NEXT: v_writelane_b32 v0, INV2PI, s2
; GFX10-NEXT: ; return to shader part epilog
%writelane = call i32 @llvm.amdgcn.writelane(i32 bitcast (float 0x3FC45F3060000000 to i32), i32 %lane, i32 %vdst.in)
%writelane.cast = bitcast i32 %writelane to float
diff --git a/llvm/test/CodeGen/AMDGPU/bitreverse-inline-immediates.ll b/llvm/test/CodeGen/AMDGPU/bitreverse-inline-immediates.ll
index 58a4a22f049c1..8dd56beb7b01c 100644
--- a/llvm/test/CodeGen/AMDGPU/bitreverse-inline-immediates.ll
+++ b/llvm/test/CodeGen/AMDGPU/bitreverse-inline-immediates.ll
@@ -364,7 +364,7 @@ define void @materialize_not_neg4.0_i32(ptr addrspace(1) %out) {
; GCN-LABEL: {{^}}materialize_not_inv2pi_i32:
; SI: v_mov_b32_e32 v{{[0-9]+}}, 0xc1dd067c
-; VI: v_not_b32_e32 v{{[0-9]+}}, 0.15915494
+; VI: v_not_b32_e32 v{{[0-9]+}}, INV2PI
define void @materialize_not_inv2pi_i32(ptr addrspace(1) %out) {
store i32 -1042479492, ptr addrspace(1) %out
ret void
diff --git a/llvm/test/CodeGen/AMDGPU/fneg-combines.f16.ll b/llvm/test/CodeGen/AMDGPU/fneg-combines.f16.ll
index b14e8c44ffcce..e78798b2496e4 100644
--- a/llvm/test/CodeGen/AMDGPU/fneg-combines.f16.ll
+++ b/llvm/test/CodeGen/AMDGPU/fneg-combines.f16.ll
@@ -1470,7 +1470,7 @@ define half @v_fneg_inv2pi_minnum_f16(half %a) #0 {
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; VI-NEXT: v_max_f16_e32 v0, v0, v0
-; VI-NEXT: v_min_f16_e32 v0, 0.15915494, v0
+; VI-NEXT: v_min_f16_e32 v0, INV2PI, v0
; VI-NEXT: v_xor_b32_e32 v0, 0x8000, v0
; VI-NEXT: s_setpc_b64 s[30:31]
;
@@ -1479,7 +1479,7 @@ define half @v_fneg_inv2pi_minnum_f16(half %a) #0 {
; GFX11-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; GFX11-NEXT: v_max_f16_e32 v0, v0, v0
; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
-; GFX11-NEXT: v_min_f16_e32 v0, 0.15915494, v0
+; GFX11-NEXT: v_min_f16_e32 v0, INV2PI, v0
; GFX11-NEXT: v_xor_b32_e32 v0, 0x8000, v0
; GFX11-NEXT: s_setpc_b64 s[30:31]
; GFX11-SAFE-TRUE16-LABEL: v_fneg_inv2pi_minnum_f16:
@@ -1487,7 +1487,7 @@ define half @v_fneg_inv2pi_minnum_f16(half %a) #0 {
; GFX11-SAFE-TRUE16-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; GFX11-SAFE-TRUE16-NEXT: v_max_f16_e32 v0.l, v0.l, v0.l
; GFX11-SAFE-TRUE16-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
-; GFX11-SAFE-TRUE16-NEXT: v_min_f16_e32 v0.l, 0.15915494, v0.l
+; GFX11-SAFE-TRUE16-NEXT: v_min_f16_e32 v0.l, INV2PI, v0.l
; GFX11-SAFE-TRUE16-NEXT: v_xor_b16 v0.l, 0x8000, v0.l
; GFX11-SAFE-TRUE16-NEXT: s_setpc_b64 s[30:31]
; GFX11-NSZ-TRUE16-LABEL: v_fneg_inv2pi_minnum_f16:
@@ -1495,7 +1495,7 @@ define half @v_fneg_inv2pi_minnum_f16(half %a) #0 {
; GFX11-NSZ-TRUE16-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; GFX11-NSZ-TRUE16-NEXT: v_max_f16_e32 v0.l, v0.l, v0.l
; GFX11-NSZ-TRUE16-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
-; GFX11-NSZ-TRUE16-NEXT: v_min_f16_e32 v0.l, 0.15915494, v0.l
+; GFX11-NSZ-TRUE16-NEXT: v_min_f16_e32 v0.l, INV2PI, v0.l
; GFX11-NSZ-TRUE16-NEXT: v_xor_b16 v0.l, 0x8000, v0.l
; GFX11-NSZ-TRUE16-NEXT: s_setpc_b64 s[30:31]
%min = call half @llvm.minnum.f16(half 0xH3118, half %a)
@@ -1516,7 +1516,7 @@ define half @v_fneg_neg_inv2pi_minnum_f16(half %a) #0 {
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; VI-NEXT: v_max_f16_e32 v0, v0, v0
-; VI-NEXT: v_min_f16_e32 v0, 0.15915494, v0
+; VI-NEXT: v_min_f16_e32 v0, INV2PI, v0
; VI-NEXT: v_xor_b32_e32 v0, 0x8000, v0
; VI-NEXT: s_setpc_b64 s[30:31]
;
@@ -1525,7 +1525,7 @@ define half @v_fneg_neg_inv2pi_minnum_f16(half %a) #0 {
; GFX11-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; GFX11-NEXT: v_max_f16_e32 v0, v0, v0
; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
-; GFX11-NEXT: v_min_f16_e32 v0, 0.15915494, v0
+; GFX11-NEXT: v_min_f16_e32 v0, INV2PI, v0
; GFX11-NEXT: v_xor_b32_e32 v0, 0x8000, v0
; GFX11-NEXT: s_setpc_b64 s[30:31]
; GFX11-SAFE-TRUE16-LABEL: v_fneg_neg_inv2pi_minnum_f16:
@@ -1533,7 +1533,7 @@ define half @v_fneg_neg_inv2pi_minnum_f16(half %a) #0 {
; GFX11-SAFE-TRUE16-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; GFX11-SAFE-TRUE16-NEXT: v_max_f16_e32 v0.l, v0.l, v0.l
; GFX11-SAFE-TRUE16-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
-; GFX11-SAFE-TRUE16-NEXT: v_min_f16_e32 v0.l, 0.15915494, v0.l
+; GFX11-SAFE-TRUE16-NEXT: v_min_f16_e32 v0.l, INV2PI, v0.l
; GFX11-SAFE-TRUE16-NEXT: v_xor_b16 v0.l, 0x8000, v0.l
; GFX11-SAFE-TRUE16-NEXT: s_setpc_b64 s[30:31]
; GFX11-NSZ-TRUE16-LABEL: v_fneg_neg_inv2pi_minnum_f16:
@@ -1541,7 +1541,7 @@ define half @v_fneg_neg_inv2pi_minnum_f16(half %a) #0 {
; GFX11-NSZ-TRUE16-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; GFX11-NSZ-TRUE16-NEXT: v_max_f16_e32 v0.l, v0.l, v0.l
; GFX11-NSZ-TRUE16-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
-; GFX11-NSZ-TRUE16-NEXT: v_min_f16_e32 v0.l, 0.15915494, v0.l
+; GFX11-NSZ-TRUE16-NEXT: v_min_f16_e32 v0.l, INV2PI, v0.l
; GFX11-NSZ-TRUE16-NEXT: v_xor_b16 v0.l, 0x8000, v0.l
; GFX11-NSZ-TRUE16-NEXT: s_setpc_b64 s[30:31]
%min = call half @llvm.minnum.f16(half 0xH3118, half %a)
@@ -1650,7 +1650,7 @@ define half @v_fneg_inv2pi_minnum_foldable_use_f16(half %a, half %b) #0 {
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; VI-NEXT: v_max_f16_e32 v0, v0, v0
-; VI-NEXT: v_min_f16_e32 v0, 0.15915494, v0
+; VI-NEXT: v_min_f16_e32 v0, INV2PI, v0
; VI-NEXT: v_mul_f16_e64 v0, -v0, v1
; VI-NEXT: s_setpc_b64 s[30:31]
;
@@ -1659,7 +1659,7 @@ define half @v_fneg_inv2pi_minnum_foldable_use_f16(half %a, half %b) #0 {
; GFX11-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; GFX11-NEXT: v_max_f16_e32 v0, v0, v0
; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
-; GFX11-NEXT: v_min_f16_e32 v0, 0.15915494, v0
+; GFX11-NEXT: v_min_f16_e32 v0, INV2PI, v0
; GFX11-NEXT: v_mul_f16_e64 v0, -v0, v1
; GFX11-NEXT: s_setpc_b64 s[30:31]
; GFX11-SAFE-TRUE16-LABEL: v_fneg_inv2pi_minnum_foldable_use_f16:
@@ -1667,7 +1667,7 @@ define half @v_fneg_inv2pi_minnum_foldable_use_f16(half %a, half %b) #0 {
; GFX11-SAFE-TRUE16-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; GFX11-SAFE-TRUE16-NEXT: v_max_f16_e32 v0.l, v0.l, v0.l
; GFX11-SAFE-TRUE16-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
-; GFX11-SAFE-TRUE16-NEXT: v_min_f16_e32 v0.l, 0.15915494, v0.l
+; GFX11-SAFE-TRUE16-NEXT: v_min_f16_e32 v0.l, INV2PI, v0.l
; GFX11-SAFE-TRUE16-NEXT: v_mul_f16_e64 v0.l, -v0.l, v1.l
; GFX11-SAFE-TRUE16-NEXT: s_setpc_b64 s[30:31]
; GFX11-NSZ-TRUE16-LABEL: v_fneg_inv2pi_minnum_foldable_use_f16:
@@ -1675,7 +1675,7 @@ define half @v_fneg_inv2pi_minnum_foldable_use_f16(half %a, half %b) #0 {
; GFX11-NSZ-TRUE16-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; GFX11-NSZ-TRUE16-NEXT: v_max_f16_e32 v0.l, v0.l, v0.l
; GFX11-NSZ-TRUE16-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
-; GFX11-NSZ-TRUE16-NEXT: v_min_f16_e32 v0.l, 0.15915494, v0.l
+; GFX11-NSZ-TRUE16-NEXT: v_min_f16_e32 v0.l, INV2PI, v0.l
; GFX11-NSZ-TRUE16-NEXT: v_mul_f16_e64 v0.l, -v0.l, v1.l
; GFX11-NSZ-TRUE16-NEXT: s_setpc_b64 s[30:31]
%min = call half @llvm.minnum.f16(half 0xH3118, half %a)
diff --git a/llvm/test/CodeGen/AMDGPU/fneg-combines.ll b/llvm/test/CodeGen/AMDGPU/fneg-combines.ll
index b3202cbe30d0b..9538ae10979b3 100644
--- a/llvm/test/CodeGen/AMDGPU/fneg-combines.ll
+++ b/llvm/test/CodeGen/AMDGPU/fneg-combines.ll
@@ -1857,7 +1857,7 @@ define amdgpu_kernel void @v_fneg_inv2pi_minnum_f32(ptr addrspace(1) %out, ptr a
; VI-NEXT: v_mov_b32_e32 v1, s1
; VI-NEXT: v_addc_u32_e32 v1, vcc, 0, v1, vcc
; VI-NEXT: v_mul_f32_e32 v2, 1.0, v3
-; VI-NEXT: v_min_f32_e32 v2, 0.15915494, v2
+; VI-NEXT: v_min_f32_e32 v2, INV2PI, v2
; VI-NEXT: v_xor_b32_e32 v2, 0x80000000, v2
; VI-NEXT: flat_store_dword v[0:1], v2
; VI-NEXT: s_endpgm
@@ -1905,7 +1905,7 @@ define amdgpu_kernel void @v_fneg_neg_inv2pi_minnum_f32(ptr addrspace(1) %out, p
; VI-NEXT: v_add_u32_e32 v0, vcc, s0, v2
; VI-NEXT: v_addc_u32_e32 v1, vcc, 0, v1, vcc
; VI-NEXT: v_mul_f32_e32 v2, -1.0, v3
-; VI-NEXT: v_max_f32_e32 v2, 0.15915494, v2
+; VI-NEXT: v_max_f32_e32 v2, INV2PI, v2
; VI-NEXT: flat_store_dword v[0:1], v2
; VI-NEXT: s_endpgm
%tid = call i32 @llvm.amdgcn.workitem.id.x()
@@ -1953,7 +1953,7 @@ define amdgpu_kernel void @v_fneg_inv2pi_minnum_f16(ptr addrspace(1) %out, ptr a
; VI-NEXT: v_mov_b32_e32 v1, s1
; VI-NEXT: v_addc_u32_e32 v1, vcc, 0, v1, vcc
; VI-NEXT: v_max_f16_e32 v2, v3, v3
-; VI-NEXT: v_min_f16_e32 v2, 0.15915494, v2
+; VI-NEXT: v_min_f16_e32 v2, INV2PI, v2
; VI-NEXT: v_xor_b32_e32 v2, 0x8000, v2
; VI-NEXT: flat_store_short v[0:1], v2
; VI-NEXT: s_endpgm
@@ -2002,7 +2002,7 @@ define amdgpu_kernel void @v_fneg_neg_inv2pi_minnum_f16(ptr addrspace(1) %out, p
; VI-NEXT: v_add_u32_e32 v0, vcc, s0, v2
; VI-NEXT: v_addc_u32_e32 v1, vcc, 0, v1, vcc
; VI-NEXT: v_max_f16_e64 v2, -v3, -v3
-; VI-NEXT: v_max_f16_e32 v2, 0.15915494, v2
+; VI-NEXT: v_max_f16_e32 v2, INV2PI, v2
; VI-NEXT: flat_store_short v[0:1], v2
; VI-NEXT: s_endpgm
%tid = call i32 @llvm.amdgcn.workitem.id.x()
@@ -2051,7 +2051,7 @@ define amdgpu_kernel void @v_fneg_inv2pi_minnum_f64(ptr addrspace(1) %out, ptr a
; VI-NEXT: v_add_u32_e32 v2, vcc, s0, v2
; VI-NEXT: v_addc_u32_e32 v3, vcc, 0, v3, vcc
; VI-NEXT: v_max_f64 v[0:1], v[0:1], v[0:1]
-; VI-NEXT: v_min_f64 v[0:1], v[0:1], 0.15915494309189532
+; VI-NEXT: v_min_f64 v[0:1], v[0:1], INV2PI
; VI-NEXT: v_xor_b32_e32 v1, 0x80000000, v1
; VI-NEXT: flat_store_dwordx2 v[2:3], v[0:1]
; VI-NEXT: s_endpgm
@@ -2101,7 +2101,7 @@ define amdgpu_kernel void @v_fneg_neg_inv2pi_minnum_f64(ptr addrspace(1) %out, p
; VI-NEXT: v_add_u32_e32 v2, vcc, s0, v2
; VI-NEXT: v_addc_u32_e32 v3, vcc, 0, v3, vcc
; VI-NEXT: v_max_f64 v[0:1], -v[0:1], -v[0:1]
-; VI-NEXT: v_max_f64 v[0:1], v[0:1], 0.15915494309189532
+; VI-NEXT: v_max_f64 v[0:1], v[0:1], INV2PI
; VI-NEXT: flat_store_dwordx2 v[2:3], v[0:1]
; VI-NEXT: s_endpgm
%tid = call i32 @llvm.amdgcn.workitem.id.x()
@@ -2235,7 +2235,7 @@ define amdgpu_kernel void @v_fneg_inv2pi_minnum_foldable_use_f32(ptr addrspace(1
; VI-NEXT: v_mov_b32_e32 v1, s1
; VI-NEXT: v_addc_u32_e32 v1, vcc, 0, v1, vcc
; VI-NEXT: v_mul_f32_e32 v2, 1.0, v4
-; VI-NEXT: v_min_f32_e32 v2, 0.15915494, v2
+; VI-NEXT: v_min_f32_e32 v2, INV2PI, v2
; VI-NEXT: v_mul_f32_e64 v2, -v2, v3
; VI-NEXT: flat_store_dword v[0:1], v2
; VI-NEXT: s_endpgm
diff --git a/llvm/test/CodeGen/AMDGPU/fneg-combines.new.ll b/llvm/test/CodeGen/AMDGPU/fneg-combines.new.ll
index 833be2066cd54..977b0131e6be4 100644
--- a/llvm/test/CodeGen/AMDGPU/fneg-combines.new.ll
+++ b/llvm/test/CodeGen/AMDGPU/fneg-combines.new.ll
@@ -907,7 +907,7 @@ define float @v_fneg_inv2pi_minnum_f32(float %a) #0 {
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; VI-NEXT: v_mul_f32_e32 v0, 1.0, v0
-; VI-NEXT: v_min_f32_e32 v0, 0.15915494, v0
+; VI-NEXT: v_min_f32_e32 v0, INV2PI, v0
; VI-NEXT: v_xor_b32_e32 v0, 0x80000000, v0
; VI-NEXT: s_setpc_b64 s[30:31]
%min = call float @llvm.minnum.f32(float 0x3FC45F3060000000, float %a)
@@ -927,7 +927,7 @@ define float @v_fneg_neg_inv2pi_minnum_f32(float %a) #0 {
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; VI-NEXT: v_mul_f32_e32 v0, -1.0, v0
-; VI-NEXT: v_max_f32_e32 v0, 0.15915494, v0
+; VI-NEXT: v_max_f32_e32 v0, INV2PI, v0
; VI-NEXT: s_setpc_b64 s[30:31]
%min = call float @llvm.minnum.f32(float 0xBFC45F3060000000, float %a)
%fneg = fneg float %min
@@ -947,7 +947,7 @@ define half @v_fneg_inv2pi_minnum_f16(half %a) #0 {
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; VI-NEXT: v_max_f16_e32 v0, v0, v0
-; VI-NEXT: v_min_f16_e32 v0, 0.15915494, v0
+; VI-NEXT: v_min_f16_e32 v0, INV2PI, v0
; VI-NEXT: v_xor_b32_e32 v0, 0x8000, v0
; VI-NEXT: s_setpc_b64 s[30:31]
%min = call half @llvm.minnum.f16(half 0xH3118, half %a)
@@ -968,7 +968,7 @@ define half @v_fneg_neg_inv2pi_minnum_f16(half %a) #0 {
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; VI-NEXT: v_max_f16_e64 v0, -v0, -v0
-; VI-NEXT: v_max_f16_e32 v0, 0.15915494, v0
+; VI-NEXT: v_max_f16_e32 v0, INV2PI, v0
; VI-NEXT: s_setpc_b64 s[30:31]
%min = call half @llvm.minnum.f16(half 0xHB118, half %a)
%fneg = fneg half %min
@@ -989,7 +989,7 @@ define double @v_fneg_inv2pi_minnum_f64(double %a) #0 {
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; VI-NEXT: v_max_f64 v[0:1], v[0:1], v[0:1]
-; VI-NEXT: v_min_f64 v[0:1], v[0:1], 0.15915494309189532
+; VI-NEXT: v_min_f64 v[0:1], v[0:1], INV2PI
; VI-NEXT: v_xor_b32_e32 v1, 0x80000000, v1
; VI-NEXT: s_setpc_b64 s[30:31]
%min = call double @llvm.minnum.f64(double 0x3fc45f306dc9c882, double %a)
@@ -1011,7 +1011,7 @@ define double @v_fneg_neg_inv2pi_minnum_f64(double %a) #0 {
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; VI-NEXT: v_max_f64 v[0:1], -v[0:1], -v[0:1]
-; VI-NEXT: v_max_f64 v[0:1], v[0:1], 0.15915494309189532
+; VI-NEXT: v_max_f64 v[0:1], v[0:1], INV2PI
; VI-NEXT: s_setpc_b64 s[30:31]
%min = call double @llvm.minnum.f64(double 0xbfc45f306dc9c882, double %a)
%fneg = fneg double %min
@@ -1056,7 +1056,7 @@ define float @v_fneg_inv2pi_minnum_foldable_use_f32(float %a, float %b) #0 {
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; VI-NEXT: v_mul_f32_e32 v0, 1.0, v0
-; VI-NEXT: v_min_f32_e32 v0, 0.15915494, v0
+; VI-NEXT: v_min_f32_e32 v0, INV2PI, v0
; VI-NEXT: v_mul_f32_e64 v0, -v0, v1
; VI-NEXT: s_setpc_b64 s[30:31]
%min = call float @llvm.minnum.f32(float 0x3FC45F3060000000, float %a)
@@ -1450,7 +1450,7 @@ define float @v_fneg_inv2pi_minimum_f32(float %a) #0 {
; VI-LABEL: v_fneg_inv2pi_minimum_f32:
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; VI-NEXT: v_min_f32_e32 v1, 0.15915494, v0
+; VI-NEXT: v_min_f32_e32 v1, INV2PI, v0
; VI-NEXT: v_mov_b32_e32 v2, 0x7fc00000
; VI-NEXT: v_cmp_o_f32_e32 vcc, v0, v0
; VI-NEXT: v_cndmask_b32_e32 v0, v2, v1, vcc
@@ -1475,7 +1475,7 @@ define float @v_fneg_neg_inv2pi_minimum_f32(float %a) #0 {
; VI-LABEL: v_fneg_neg_inv2pi_minimum_f32:
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; VI-NEXT: v_max_f32_e64 v1, -v0, 0.15915494
+; VI-NEXT: v_max_f32_e64 v1, -v0, INV2PI
; VI-NEXT: v_mov_b32_e32 v2, 0x7fc00000
; VI-NEXT: v_cmp_o_f32_e64 vcc, -v0, -v0
; VI-NEXT: v_cndmask_b32_e32 v0, v2, v1, vcc
@@ -1500,7 +1500,7 @@ define half @v_fneg_inv2pi_minimum_f16(half %a) #0 {
; VI-LABEL: v_fneg_inv2pi_minimum_f16:
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; VI-NEXT: v_min_f16_e32 v1, 0.15915494, v0
+; VI-NEXT: v_min_f16_e32 v1, INV2PI, v0
; VI-NEXT: v_mov_b32_e32 v2, 0x7e00
; VI-NEXT: v_cmp_o_f16_e32 vcc, v0, v0
; VI-NEXT: v_cndmask_b32_e32 v0, v2, v1, vcc
@@ -1526,7 +1526,7 @@ define half @v_fneg_neg_inv2pi_minimum_f16(half %a) #0 {
; VI-LABEL: v_fneg_neg_inv2pi_minimum_f16:
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; VI-NEXT: v_max_f16_e64 v1, -v0, 0.15915494
+; VI-NEXT: v_max_f16_e64 v1, -v0, INV2PI
; VI-NEXT: v_mov_b32_e32 v2, 0x7e00
; VI-NEXT: v_cmp_o_f16_e64 vcc, -v0, -v0
; VI-NEXT: v_cndmask_b32_e32 v0, v2, v1, vcc
@@ -1552,7 +1552,7 @@ define double @v_fneg_inv2pi_minimum_f64(double %a) #0 {
; VI-LABEL: v_fneg_inv2pi_minimum_f64:
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; VI-NEXT: v_min_f64 v[2:3], v[0:1], 0.15915494309189532
+; VI-NEXT: v_min_f64 v[2:3], v[0:1], INV2PI
; VI-NEXT: v_cmp_u_f64_e32 vcc, v[0:1], v[0:1]
; VI-NEXT: v_mov_b32_e32 v1, 0xfff80000
; VI-NEXT: v_cndmask_b32_e64 v0, v2, 0, vcc
@@ -1579,7 +1579,7 @@ define double @v_fneg_neg_inv2pi_minimum_f64(double %a) #0 {
; VI-LABEL: v_fneg_neg_inv2pi_minimum_f64:
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; VI-NEXT: v_max_f64 v[2:3], -v[0:1], 0.15915494309189532
+; VI-NEXT: v_max_f64 v[2:3], -v[0:1], INV2PI
; VI-NEXT: v_cmp_u_f64_e64 vcc, -v[0:1], -v[0:1]
; VI-NEXT: v_mov_b32_e32 v1, 0x7ff80000
; VI-NEXT: v_cndmask_b32_e64 v0, v2, 0, vcc
@@ -1635,7 +1635,7 @@ define float @v_fneg_inv2pi_minimum_foldable_use_f32(float %a, float %b) #0 {
; VI-LABEL: v_fneg_inv2pi_minimum_foldable_use_f32:
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; VI-NEXT: v_min_f32_e32 v2, 0.15915494, v0
+; VI-NEXT: v_min_f32_e32 v2, INV2PI, v0
; VI-NEXT: v_mov_b32_e32 v3, 0x7fc00000
; VI-NEXT: v_cmp_o_f32_e32 vcc, v0, v0
; VI-NEXT: v_cndmask_b32_e32 v0, v3, v2, vcc
@@ -2049,7 +2049,7 @@ define float @v_fneg_inv2pi_minimumnum_f32(float %a) #0 {
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; VI-NEXT: v_mul_f32_e32 v0, 1.0, v0
-; VI-NEXT: v_min_f32_e32 v0, 0.15915494, v0
+; VI-NEXT: v_min_f32_e32 v0, INV2PI, v0
; VI-NEXT: v_xor_b32_e32 v0, 0x80000000, v0
; VI-NEXT: s_setpc_b64 s[30:31]
%min = call float @llvm.minimumnum.f32(float 0x3FC45F3060000000, float %a)
@@ -2069,7 +2069,7 @@ define float @v_fneg_neg_inv2pi_minimumnum_f32(float %a) #0 {
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; VI-NEXT: v_mul_f32_e32 v0, -1.0, v0
-; VI-NEXT: v_max_f32_e32 v0, 0.15915494, v0
+; VI-NEXT: v_max_f32_e32 v0, INV2PI, v0
; VI-NEXT: s_setpc_b64 s[30:31]
%min = call float @llvm.minimumnum.f32(float 0xBFC45F3060000000, float %a)
%fneg = fneg float %min
@@ -2089,7 +2089,7 @@ define half @v_fneg_inv2pi_minimumnum_f16(half %a) #0 {
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; VI-NEXT: v_max_f16_e32 v0, v0, v0
-; VI-NEXT: v_min_f16_e32 v0, 0.15915494, v0
+; VI-NEXT: v_min_f16_e32 v0, INV2PI, v0
; VI-NEXT: v_xor_b32_e32 v0, 0x8000, v0
; VI-NEXT: s_setpc_b64 s[30:31]
%min = call half @llvm.minimumnum.f16(half 0xH3118, half %a)
@@ -2110,7 +2110,7 @@ define half @v_fneg_neg_inv2pi_minimumnum_f16(half %a) #0 {
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; VI-NEXT: v_max_f16_e64 v0, -v0, -v0
-; VI-NEXT: v_max_f16_e32 v0, 0.15915494, v0
+; VI-NEXT: v_max_f16_e32 v0, INV2PI, v0
; VI-NEXT: s_setpc_b64 s[30:31]
%min = call half @llvm.minimumnum.f16(half 0xHB118, half %a)
%fneg = fneg half %min
@@ -2131,7 +2131,7 @@ define double @v_fneg_inv2pi_minimumnum_f64(double %a) #0 {
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; VI-NEXT: v_max_f64 v[0:1], v[0:1], v[0:1]
-; VI-NEXT: v_min_f64 v[0:1], v[0:1], 0.15915494309189532
+; VI-NEXT: v_min_f64 v[0:1], v[0:1], INV2PI
; VI-NEXT: v_xor_b32_e32 v1, 0x80000000, v1
; VI-NEXT: s_setpc_b64 s[30:31]
%min = call double @llvm.minimumnum.f64(double 0x3fc45f306dc9c882, double %a)
@@ -2153,7 +2153,7 @@ define double @v_fneg_neg_inv2pi_minimumnum_f64(double %a) #0 {
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; VI-NEXT: v_max_f64 v[0:1], -v[0:1], -v[0:1]
-; VI-NEXT: v_max_f64 v[0:1], v[0:1], 0.15915494309189532
+; VI-NEXT: v_max_f64 v[0:1], v[0:1], INV2PI
; VI-NEXT: s_setpc_b64 s[30:31]
%min = call double @llvm.minimumnum.f64(double 0xbfc45f306dc9c882, double %a)
%fneg = fneg double %min
@@ -2198,7 +2198,7 @@ define float @v_fneg_inv2pi_minimumnum_foldable_use_f32(float %a, float %b) #0 {
; VI: ; %bb.0:
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; VI-NEXT: v_mul_f32_e32 v0, 1.0, v0
-; VI-NEXT: v_min_f32_e32 v0, 0.15915494, v0
+; VI-NEXT: v_min_f32_e32 v0, INV2PI, v0
; VI-NEXT: v_mul_f32_e64 v0, -v0, v1
; VI-NEXT: s_setpc_b64 s[30:31]
%min = call float @llvm.minimumnum.f32(float 0x3FC45F3060000000, float %a)
diff --git a/llvm/test/CodeGen/AMDGPU/imm.ll b/llvm/test/CodeGen/AMDGPU/imm.ll
index 21390003ee565..4495241b582db 100644
--- a/llvm/test/CodeGen/AMDGPU/imm.ll
+++ b/llvm/test/CodeGen/AMDGPU/imm.ll
@@ -508,7 +508,7 @@ define amdgpu_kernel void @store_inline_imm_inv_2pi_f32(ptr addrspace(1) %out) {
; VI-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24
; VI-NEXT: s_mov_b32 s3, 0xf000
; VI-NEXT: s_mov_b32 s2, -1
-; VI-NEXT: v_mov_b32_e32 v0, 0.15915494
+; VI-NEXT: v_mov_b32_e32 v0, INV2PI
; VI-NEXT: s_waitcnt lgkmcnt(0)
; VI-NEXT: buffer_store_dword v0, off, s[0:3], 0
; VI-NEXT: s_endpgm
@@ -518,7 +518,7 @@ define amdgpu_kernel void @store_inline_imm_inv_2pi_f32(ptr addrspace(1) %out) {
; GFX942-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24
; GFX942-NEXT: s_mov_b32 s3, 0xf000
; GFX942-NEXT: s_mov_b32 s2, -1
-; GFX942-NEXT: v_mov_b32_e32 v0, 0.15915494
+; GFX942-NEXT: v_mov_b32_e32 v0, INV2PI
; GFX942-NEXT: s_waitcnt lgkmcnt(0)
; GFX942-NEXT: buffer_store_dword v0, off, s[0:3], 0
; GFX942-NEXT: s_endpgm
@@ -1756,7 +1756,7 @@ define amdgpu_kernel void @add_inline_imm_inv_2pi_f64(ptr addrspace(1) %out, [8
; VI-NEXT: s_mov_b32 s3, 0xf000
; VI-NEXT: s_mov_b32 s2, -1
; VI-NEXT: s_waitcnt lgkmcnt(0)
-; VI-NEXT: v_add_f64 v[0:1], s[0:1], 0.15915494309189532
+; VI-NEXT: v_add_f64 v[0:1], s[0:1], INV2PI
; VI-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24
; VI-NEXT: s_waitcnt lgkmcnt(0)
; VI-NEXT: buffer_store_dwordx2 v[0:1], off, s[0:3], 0
@@ -1769,7 +1769,7 @@ define amdgpu_kernel void @add_inline_imm_inv_2pi_f64(ptr addrspace(1) %out, [8
; GFX942-NEXT: s_mov_b32 s3, 0xf000
; GFX942-NEXT: s_mov_b32 s2, -1
; GFX942-NEXT: s_waitcnt lgkmcnt(0)
-; GFX942-NEXT: v_add_f64 v[0:1], s[6:7], 0.15915494309189532
+; GFX942-NEXT: v_add_f64 v[0:1], s[6:7], INV2PI
; GFX942-NEXT: buffer_store_dwordx2 v[0:1], off, s[0:3], 0
; GFX942-NEXT: s_endpgm
%y = fadd double %x, 0x3fc45f306dc9c882
diff --git a/llvm/test/CodeGen/AMDGPU/known-never-snan.ll b/llvm/test/CodeGen/AMDGPU/known-never-snan.ll
index 5691fc8740a6b..0aac440e94596 100644
--- a/llvm/test/CodeGen/AMDGPU/known-never-snan.ll
+++ b/llvm/test/CodeGen/AMDGPU/known-never-snan.ll
@@ -373,7 +373,7 @@ define float @v_test_known_not_snan_sin_input_fmed3_r_i_i_f32(float %a) #0 {
; GCN-LABEL: v_test_known_not_snan_sin_input_fmed3_r_i_i_f32:
; GCN: ; %bb.0:
; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GCN-NEXT: v_mul_f32_e32 v0, 0.15915494, v0
+; GCN-NEXT: v_mul_f32_e32 v0, INV2PI, v0
; GCN-NEXT: v_fract_f32_e32 v0, v0
; GCN-NEXT: v_sin_f32_e32 v0, v0
; GCN-NEXT: v_med3_f32 v0, v0, 2.0, 4.0
@@ -388,7 +388,7 @@ define float @v_test_known_not_snan_cos_input_fmed3_r_i_i_f32(float %a) #0 {
; GCN-LABEL: v_test_known_not_snan_cos_input_fmed3_r_i_i_f32:
; GCN: ; %bb.0:
; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GCN-NEXT: v_mul_f32_e32 v0, 0.15915494, v0
+; GCN-NEXT: v_mul_f32_e32 v0, INV2PI, v0
; GCN-NEXT: v_fract_f32_e32 v0, v0
; GCN-NEXT: v_cos_f32_e32 v0, v0
; GCN-NEXT: v_med3_f32 v0, v0, 2.0, 4.0
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.ll
index aae14c8cc87b3..32154561494ba 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.ll
@@ -2126,7 +2126,7 @@ define amdgpu_kernel void @test_mfma_scale_f32_16x16x128_f8f6f4__vgprcd___scaleA
; SDAG-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x0
; SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x40
; SDAG-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x50
-; SDAG-NEXT: v_mov_b32_e32 v21, 0.15915494
+; SDAG-NEXT: v_mov_b32_e32 v21, INV2PI
; SDAG-NEXT: v_mov_b32_e32 v22, 1.0
; SDAG-NEXT: v_mov_b32_e32 v20, 0
; SDAG-NEXT: s_waitcnt lgkmcnt(0)
@@ -2160,7 +2160,7 @@ define amdgpu_kernel void @test_mfma_scale_f32_16x16x128_f8f6f4__vgprcd___scaleA
; GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x40
; GISEL-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x50
; GISEL-NEXT: v_mov_b32_e32 v20, 1.0
-; GISEL-NEXT: v_mov_b32_e32 v21, 0.15915494
+; GISEL-NEXT: v_mov_b32_e32 v21, INV2PI
; GISEL-NEXT: s_waitcnt lgkmcnt(0)
; GISEL-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
; GISEL-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.scale.f32.32x32x64.f8f6f4.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.scale.f32.32x32x64.f8f6f4.ll
index f0205a3a788ed..d88bfb0ec6aa5 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.scale.f32.32x32x64.f8f6f4.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.scale.f32.32x32x64.f8f6f4.ll
@@ -4415,7 +4415,7 @@ define <16 x float> @test_mfma_scale_f32_32x32x64_f8f6f4_0_0__scaleA_FP_literal_
; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; SDAG-NEXT: scratch_load_dword a15, off, s32
; SDAG-NEXT: v_mov_b32_e32 v31, 1.0
-; SDAG-NEXT: v_mov_b32_e32 v32, 0.15915494
+; SDAG-NEXT: v_mov_b32_e32 v32, INV2PI
; SDAG-NEXT: v_accvgpr_write_b32 a0, v16
; SDAG-NEXT: v_accvgpr_write_b32 a1, v17
; SDAG-NEXT: v_accvgpr_write_b32 a2, v18
@@ -4458,7 +4458,7 @@ define <16 x float> @test_mfma_scale_f32_32x32x64_f8f6f4_0_0__scaleA_FP_literal_
; GISEL: ; %bb.0:
; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; GISEL-NEXT: scratch_load_dword a15, off, s32
-; GISEL-NEXT: v_mov_b32_e32 v31, 0.15915494
+; GISEL-NEXT: v_mov_b32_e32 v31, INV2PI
; GISEL-NEXT: v_mov_b32_e32 v32, 1.0
; GISEL-NEXT: v_accvgpr_write_b32 a0, v16
; GISEL-NEXT: v_accvgpr_write_b32 a1, v17
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.cos.f16.ll b/llvm/test/CodeGen/AMDGPU/llvm.cos.f16.ll
index 7d63e22d84b72..a39325beb5037 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.cos.f16.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.cos.f16.ll
@@ -40,7 +40,7 @@ define amdgpu_kernel void @cos_f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX8-NEXT: flat_load_ushort v0, v[0:1]
; GFX8-NEXT: v_mov_b32_e32 v1, s1
; GFX8-NEXT: s_waitcnt vmcnt(0)
-; GFX8-NEXT: v_mul_f16_e32 v0, 0.15915494, v0
+; GFX8-NEXT: v_mul_f16_e32 v0, INV2PI, v0
; GFX8-NEXT: v_fract_f16_e32 v0, v0
; GFX8-NEXT: v_cos_f16_e32 v2, v0
; GFX8-NEXT: v_mov_b32_e32 v0, s0
@@ -54,7 +54,7 @@ define amdgpu_kernel void @cos_f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX9-NEXT: s_waitcnt lgkmcnt(0)
; GFX9-NEXT: global_load_ushort v1, v0, s[2:3]
; GFX9-NEXT: s_waitcnt vmcnt(0)
-; GFX9-NEXT: v_mul_f16_e32 v1, 0.15915494, v1
+; GFX9-NEXT: v_mul_f16_e32 v1, INV2PI, v1
; GFX9-NEXT: v_cos_f16_e32 v1, v1
; GFX9-NEXT: global_store_short v0, v1, s[0:1]
; GFX9-NEXT: s_endpgm
@@ -66,7 +66,7 @@ define amdgpu_kernel void @cos_f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX10-NEXT: s_waitcnt lgkmcnt(0)
; GFX10-NEXT: global_load_ushort v1, v0, s[2:3]
; GFX10-NEXT: s_waitcnt vmcnt(0)
-; GFX10-NEXT: v_mul_f16_e32 v1, 0.15915494, v1
+; GFX10-NEXT: v_mul_f16_e32 v1, INV2PI, v1
; GFX10-NEXT: v_cos_f16_e32 v1, v1
; GFX10-NEXT: global_store_short v0, v1, s[0:1]
; GFX10-NEXT: s_endpgm
@@ -78,7 +78,7 @@ define amdgpu_kernel void @cos_f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX11-TRUE16-NEXT: s_waitcnt lgkmcnt(0)
; GFX11-TRUE16-NEXT: global_load_d16_b16 v0, v1, s[2:3]
; GFX11-TRUE16-NEXT: s_waitcnt vmcnt(0)
-; GFX11-TRUE16-NEXT: v_mul_f16_e32 v0.l, 0.15915494, v0.l
+; GFX11-TRUE16-NEXT: v_mul_f16_e32 v0.l, INV2PI, v0.l
; GFX11-TRUE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
; GFX11-TRUE16-NEXT: v_cos_f16_e32 v0.l, v0.l
; GFX11-TRUE16-NEXT: global_store_b16 v1, v0, s[0:1]
@@ -91,7 +91,7 @@ define amdgpu_kernel void @cos_f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX11-FAKE16-NEXT: s_waitcnt lgkmcnt(0)
; GFX11-FAKE16-NEXT: global_load_u16 v1, v0, s[2:3]
; GFX11-FAKE16-NEXT: s_waitcnt vmcnt(0)
-; GFX11-FAKE16-NEXT: v_mul_f16_e32 v1, 0.15915494, v1
+; GFX11-FAKE16-NEXT: v_mul_f16_e32 v1, INV2PI, v1
; GFX11-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
; GFX11-FAKE16-NEXT: v_cos_f16_e32 v1, v1
; GFX11-FAKE16-NEXT: global_store_b16 v0, v1, s[0:1]
@@ -104,7 +104,7 @@ define amdgpu_kernel void @cos_f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX12-TRUE16-NEXT: s_wait_kmcnt 0x0
; GFX12-TRUE16-NEXT: global_load_d16_b16 v0, v1, s[2:3]
; GFX12-TRUE16-NEXT: s_wait_loadcnt 0x0
-; GFX12-TRUE16-NEXT: v_mul_f16_e32 v0.l, 0.15915494, v0.l
+; GFX12-TRUE16-NEXT: v_mul_f16_e32 v0.l, INV2PI, v0.l
; GFX12-TRUE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
; GFX12-TRUE16-NEXT: v_cos_f16_e32 v0.l, v0.l
; GFX12-TRUE16-NEXT: global_store_b16 v1, v0, s[0:1]
@@ -117,7 +117,7 @@ define amdgpu_kernel void @cos_f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX12-FAKE16-NEXT: s_wait_kmcnt 0x0
; GFX12-FAKE16-NEXT: global_load_u16 v1, v0, s[2:3]
; GFX12-FAKE16-NEXT: s_wait_loadcnt 0x0
-; GFX12-FAKE16-NEXT: v_mul_f16_e32 v1, 0.15915494, v1
+; GFX12-FAKE16-NEXT: v_mul_f16_e32 v1, INV2PI, v1
; GFX12-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
; GFX12-FAKE16-NEXT: v_cos_f16_e32 v1, v1
; GFX12-FAKE16-NEXT: global_store_b16 v0, v1, s[0:1]
@@ -169,7 +169,7 @@ define amdgpu_kernel void @cos_v2f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX8-NEXT: v_mov_b32_e32 v1, 0x3118
; GFX8-NEXT: s_waitcnt vmcnt(0)
; GFX8-NEXT: v_mul_f16_sdwa v1, v0, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:DWORD
-; GFX8-NEXT: v_mul_f16_e32 v0, 0.15915494, v0
+; GFX8-NEXT: v_mul_f16_e32 v0, INV2PI, v0
; GFX8-NEXT: v_fract_f16_e32 v1, v1
; GFX8-NEXT: v_fract_f16_e32 v0, v0
; GFX8-NEXT: v_cos_f16_sdwa v2, v1 dst_sel:WORD_1 dst_unused:UNUSED_PAD src0_sel:DWORD
@@ -188,7 +188,7 @@ define amdgpu_kernel void @cos_v2f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX9-NEXT: s_waitcnt lgkmcnt(0)
; GFX9-NEXT: global_load_dword v1, v0, s[2:3]
; GFX9-NEXT: s_waitcnt vmcnt(0)
-; GFX9-NEXT: v_mul_f16_e32 v3, 0.15915494, v1
+; GFX9-NEXT: v_mul_f16_e32 v3, INV2PI, v1
; GFX9-NEXT: v_mul_f16_sdwa v1, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:DWORD
; GFX9-NEXT: v_cos_f16_e32 v2, v3
; GFX9-NEXT: v_cos_f16_e32 v1, v1
@@ -204,7 +204,7 @@ define amdgpu_kernel void @cos_v2f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX10-NEXT: s_waitcnt lgkmcnt(0)
; GFX10-NEXT: global_load_dword v1, v0, s[2:3]
; GFX10-NEXT: s_waitcnt vmcnt(0)
-; GFX10-NEXT: v_mul_f16_e32 v3, 0.15915494, v1
+; GFX10-NEXT: v_mul_f16_e32 v3, INV2PI, v1
; GFX10-NEXT: v_mul_f16_sdwa v1, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:DWORD
; GFX10-NEXT: v_cos_f16_e32 v2, v3
; GFX10-NEXT: v_cos_f16_e32 v1, v1
@@ -220,9 +220,9 @@ define amdgpu_kernel void @cos_v2f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX11-TRUE16-NEXT: global_load_b32 v0, v1, s[2:3]
; GFX11-TRUE16-NEXT: s_waitcnt vmcnt(0)
; GFX11-TRUE16-NEXT: v_lshrrev_b32_e32 v2, 16, v0
-; GFX11-TRUE16-NEXT: v_mul_f16_e32 v0.l, 0.15915494, v0.l
+; GFX11-TRUE16-NEXT: v_mul_f16_e32 v0.l, INV2PI, v0.l
; GFX11-TRUE16-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
-; GFX11-TRUE16-NEXT: v_mul_f16_e32 v0.h, 0.15915494, v2.l
+; GFX11-TRUE16-NEXT: v_mul_f16_e32 v0.h, INV2PI, v2.l
; GFX11-TRUE16-NEXT: v_cos_f16_e32 v0.l, v0.l
; GFX11-TRUE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
; GFX11-TRUE16-NEXT: v_cos_f16_e32 v0.h, v0.h
@@ -239,9 +239,9 @@ define amdgpu_kernel void @cos_v2f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX11-FAKE16-NEXT: global_load_b32 v1, v0, s[2:3]
; GFX11-FAKE16-NEXT: s_waitcnt vmcnt(0)
; GFX11-FAKE16-NEXT: v_lshrrev_b32_e32 v2, 16, v1
-; GFX11-FAKE16-NEXT: v_mul_f16_e32 v1, 0.15915494, v1
+; GFX11-FAKE16-NEXT: v_mul_f16_e32 v1, INV2PI, v1
; GFX11-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
-; GFX11-FAKE16-NEXT: v_mul_f16_e32 v2, 0.15915494, v2
+; GFX11-FAKE16-NEXT: v_mul_f16_e32 v2, INV2PI, v2
; GFX11-FAKE16-NEXT: v_cos_f16_e32 v1, v1
; GFX11-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
; GFX11-FAKE16-NEXT: v_cos_f16_e32 v2, v2
@@ -258,9 +258,9 @@ define amdgpu_kernel void @cos_v2f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX12-TRUE16-NEXT: global_load_b32 v0, v1, s[2:3]
; GFX12-TRUE16-NEXT: s_wait_loadcnt 0x0
; GFX12-TRUE16-NEXT: v_lshrrev_b32_e32 v2, 16, v0
-; GFX12-TRUE16-NEXT: v_mul_f16_e32 v0.l, 0.15915494, v0.l
+; GFX12-TRUE16-NEXT: v_mul_f16_e32 v0.l, INV2PI, v0.l
; GFX12-TRUE16-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
-; GFX12-TRUE16-NEXT: v_mul_f16_e32 v0.h, 0.15915494, v2.l
+; GFX12-TRUE16-NEXT: v_mul_f16_e32 v0.h, INV2PI, v2.l
; GFX12-TRUE16-NEXT: v_cos_f16_e32 v0.l, v0.l
; GFX12-TRUE16-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(TRANS32_DEP_1)
; GFX12-TRUE16-NEXT: v_cos_f16_e32 v0.h, v0.h
@@ -276,9 +276,9 @@ define amdgpu_kernel void @cos_v2f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX12-FAKE16-NEXT: global_load_b32 v1, v0, s[2:3]
; GFX12-FAKE16-NEXT: s_wait_loadcnt 0x0
; GFX12-FAKE16-NEXT: v_lshrrev_b32_e32 v2, 16, v1
-; GFX12-FAKE16-NEXT: v_mul_f16_e32 v1, 0.15915494, v1
+; GFX12-FAKE16-NEXT: v_mul_f16_e32 v1, INV2PI, v1
; GFX12-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
-; GFX12-FAKE16-NEXT: v_mul_f16_e32 v2, 0.15915494, v2
+; GFX12-FAKE16-NEXT: v_mul_f16_e32 v2, INV2PI, v2
; GFX12-FAKE16-NEXT: v_cos_f16_e32 v1, v1
; GFX12-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(TRANS32_DEP_1)
; GFX12-FAKE16-NEXT: v_cos_f16_e32 v2, v2
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.sin.f16.ll b/llvm/test/CodeGen/AMDGPU/llvm.sin.f16.ll
index ba03115c51536..2dd45192be9dd 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.sin.f16.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.sin.f16.ll
@@ -40,7 +40,7 @@ define amdgpu_kernel void @sin_f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX8-NEXT: flat_load_ushort v0, v[0:1]
; GFX8-NEXT: v_mov_b32_e32 v1, s1
; GFX8-NEXT: s_waitcnt vmcnt(0)
-; GFX8-NEXT: v_mul_f16_e32 v0, 0.15915494, v0
+; GFX8-NEXT: v_mul_f16_e32 v0, INV2PI, v0
; GFX8-NEXT: v_fract_f16_e32 v0, v0
; GFX8-NEXT: v_sin_f16_e32 v2, v0
; GFX8-NEXT: v_mov_b32_e32 v0, s0
@@ -54,7 +54,7 @@ define amdgpu_kernel void @sin_f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX9-NEXT: s_waitcnt lgkmcnt(0)
; GFX9-NEXT: global_load_ushort v1, v0, s[2:3]
; GFX9-NEXT: s_waitcnt vmcnt(0)
-; GFX9-NEXT: v_mul_f16_e32 v1, 0.15915494, v1
+; GFX9-NEXT: v_mul_f16_e32 v1, INV2PI, v1
; GFX9-NEXT: v_sin_f16_e32 v1, v1
; GFX9-NEXT: global_store_short v0, v1, s[0:1]
; GFX9-NEXT: s_endpgm
@@ -66,7 +66,7 @@ define amdgpu_kernel void @sin_f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX10-NEXT: s_waitcnt lgkmcnt(0)
; GFX10-NEXT: global_load_ushort v1, v0, s[2:3]
; GFX10-NEXT: s_waitcnt vmcnt(0)
-; GFX10-NEXT: v_mul_f16_e32 v1, 0.15915494, v1
+; GFX10-NEXT: v_mul_f16_e32 v1, INV2PI, v1
; GFX10-NEXT: v_sin_f16_e32 v1, v1
; GFX10-NEXT: global_store_short v0, v1, s[0:1]
; GFX10-NEXT: s_endpgm
@@ -78,7 +78,7 @@ define amdgpu_kernel void @sin_f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX11-TRUE16-NEXT: s_waitcnt lgkmcnt(0)
; GFX11-TRUE16-NEXT: global_load_d16_b16 v0, v1, s[2:3]
; GFX11-TRUE16-NEXT: s_waitcnt vmcnt(0)
-; GFX11-TRUE16-NEXT: v_mul_f16_e32 v0.l, 0.15915494, v0.l
+; GFX11-TRUE16-NEXT: v_mul_f16_e32 v0.l, INV2PI, v0.l
; GFX11-TRUE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
; GFX11-TRUE16-NEXT: v_sin_f16_e32 v0.l, v0.l
; GFX11-TRUE16-NEXT: global_store_b16 v1, v0, s[0:1]
@@ -91,7 +91,7 @@ define amdgpu_kernel void @sin_f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX11-FAKE16-NEXT: s_waitcnt lgkmcnt(0)
; GFX11-FAKE16-NEXT: global_load_u16 v1, v0, s[2:3]
; GFX11-FAKE16-NEXT: s_waitcnt vmcnt(0)
-; GFX11-FAKE16-NEXT: v_mul_f16_e32 v1, 0.15915494, v1
+; GFX11-FAKE16-NEXT: v_mul_f16_e32 v1, INV2PI, v1
; GFX11-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
; GFX11-FAKE16-NEXT: v_sin_f16_e32 v1, v1
; GFX11-FAKE16-NEXT: global_store_b16 v0, v1, s[0:1]
@@ -104,7 +104,7 @@ define amdgpu_kernel void @sin_f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX12-TRUE16-NEXT: s_wait_kmcnt 0x0
; GFX12-TRUE16-NEXT: global_load_d16_b16 v0, v1, s[2:3]
; GFX12-TRUE16-NEXT: s_wait_loadcnt 0x0
-; GFX12-TRUE16-NEXT: v_mul_f16_e32 v0.l, 0.15915494, v0.l
+; GFX12-TRUE16-NEXT: v_mul_f16_e32 v0.l, INV2PI, v0.l
; GFX12-TRUE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
; GFX12-TRUE16-NEXT: v_sin_f16_e32 v0.l, v0.l
; GFX12-TRUE16-NEXT: global_store_b16 v1, v0, s[0:1]
@@ -117,7 +117,7 @@ define amdgpu_kernel void @sin_f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX12-FAKE16-NEXT: s_wait_kmcnt 0x0
; GFX12-FAKE16-NEXT: global_load_u16 v1, v0, s[2:3]
; GFX12-FAKE16-NEXT: s_wait_loadcnt 0x0
-; GFX12-FAKE16-NEXT: v_mul_f16_e32 v1, 0.15915494, v1
+; GFX12-FAKE16-NEXT: v_mul_f16_e32 v1, INV2PI, v1
; GFX12-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
; GFX12-FAKE16-NEXT: v_sin_f16_e32 v1, v1
; GFX12-FAKE16-NEXT: global_store_b16 v0, v1, s[0:1]
@@ -169,7 +169,7 @@ define amdgpu_kernel void @sin_v2f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX8-NEXT: v_mov_b32_e32 v1, 0x3118
; GFX8-NEXT: s_waitcnt vmcnt(0)
; GFX8-NEXT: v_mul_f16_sdwa v1, v0, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:DWORD
-; GFX8-NEXT: v_mul_f16_e32 v0, 0.15915494, v0
+; GFX8-NEXT: v_mul_f16_e32 v0, INV2PI, v0
; GFX8-NEXT: v_fract_f16_e32 v1, v1
; GFX8-NEXT: v_fract_f16_e32 v0, v0
; GFX8-NEXT: v_sin_f16_sdwa v2, v1 dst_sel:WORD_1 dst_unused:UNUSED_PAD src0_sel:DWORD
@@ -188,7 +188,7 @@ define amdgpu_kernel void @sin_v2f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX9-NEXT: s_waitcnt lgkmcnt(0)
; GFX9-NEXT: global_load_dword v1, v0, s[2:3]
; GFX9-NEXT: s_waitcnt vmcnt(0)
-; GFX9-NEXT: v_mul_f16_e32 v3, 0.15915494, v1
+; GFX9-NEXT: v_mul_f16_e32 v3, INV2PI, v1
; GFX9-NEXT: v_mul_f16_sdwa v1, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:DWORD
; GFX9-NEXT: v_sin_f16_e32 v2, v3
; GFX9-NEXT: v_sin_f16_e32 v1, v1
@@ -204,7 +204,7 @@ define amdgpu_kernel void @sin_v2f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX10-NEXT: s_waitcnt lgkmcnt(0)
; GFX10-NEXT: global_load_dword v1, v0, s[2:3]
; GFX10-NEXT: s_waitcnt vmcnt(0)
-; GFX10-NEXT: v_mul_f16_e32 v3, 0.15915494, v1
+; GFX10-NEXT: v_mul_f16_e32 v3, INV2PI, v1
; GFX10-NEXT: v_mul_f16_sdwa v1, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:DWORD
; GFX10-NEXT: v_sin_f16_e32 v2, v3
; GFX10-NEXT: v_sin_f16_e32 v1, v1
@@ -220,9 +220,9 @@ define amdgpu_kernel void @sin_v2f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX11-TRUE16-NEXT: global_load_b32 v0, v1, s[2:3]
; GFX11-TRUE16-NEXT: s_waitcnt vmcnt(0)
; GFX11-TRUE16-NEXT: v_lshrrev_b32_e32 v2, 16, v0
-; GFX11-TRUE16-NEXT: v_mul_f16_e32 v0.l, 0.15915494, v0.l
+; GFX11-TRUE16-NEXT: v_mul_f16_e32 v0.l, INV2PI, v0.l
; GFX11-TRUE16-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
-; GFX11-TRUE16-NEXT: v_mul_f16_e32 v0.h, 0.15915494, v2.l
+; GFX11-TRUE16-NEXT: v_mul_f16_e32 v0.h, INV2PI, v2.l
; GFX11-TRUE16-NEXT: v_sin_f16_e32 v0.l, v0.l
; GFX11-TRUE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
; GFX11-TRUE16-NEXT: v_sin_f16_e32 v0.h, v0.h
@@ -239,9 +239,9 @@ define amdgpu_kernel void @sin_v2f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX11-FAKE16-NEXT: global_load_b32 v1, v0, s[2:3]
; GFX11-FAKE16-NEXT: s_waitcnt vmcnt(0)
; GFX11-FAKE16-NEXT: v_lshrrev_b32_e32 v2, 16, v1
-; GFX11-FAKE16-NEXT: v_mul_f16_e32 v1, 0.15915494, v1
+; GFX11-FAKE16-NEXT: v_mul_f16_e32 v1, INV2PI, v1
; GFX11-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
-; GFX11-FAKE16-NEXT: v_mul_f16_e32 v2, 0.15915494, v2
+; GFX11-FAKE16-NEXT: v_mul_f16_e32 v2, INV2PI, v2
; GFX11-FAKE16-NEXT: v_sin_f16_e32 v1, v1
; GFX11-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
; GFX11-FAKE16-NEXT: v_sin_f16_e32 v2, v2
@@ -258,9 +258,9 @@ define amdgpu_kernel void @sin_v2f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX12-TRUE16-NEXT: global_load_b32 v0, v1, s[2:3]
; GFX12-TRUE16-NEXT: s_wait_loadcnt 0x0
; GFX12-TRUE16-NEXT: v_lshrrev_b32_e32 v2, 16, v0
-; GFX12-TRUE16-NEXT: v_mul_f16_e32 v0.l, 0.15915494, v0.l
+; GFX12-TRUE16-NEXT: v_mul_f16_e32 v0.l, INV2PI, v0.l
; GFX12-TRUE16-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
-; GFX12-TRUE16-NEXT: v_mul_f16_e32 v0.h, 0.15915494, v2.l
+; GFX12-TRUE16-NEXT: v_mul_f16_e32 v0.h, INV2PI, v2.l
; GFX12-TRUE16-NEXT: v_sin_f16_e32 v0.l, v0.l
; GFX12-TRUE16-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(TRANS32_DEP_1)
; GFX12-TRUE16-NEXT: v_sin_f16_e32 v0.h, v0.h
@@ -276,9 +276,9 @@ define amdgpu_kernel void @sin_v2f16(ptr addrspace(1) %r, ptr addrspace(1) %a) {
; GFX12-FAKE16-NEXT: global_load_b32 v1, v0, s[2:3]
; GFX12-FAKE16-NEXT: s_wait_loadcnt 0x0
; GFX12-FAKE16-NEXT: v_lshrrev_b32_e32 v2, 16, v1
-; GFX12-FAKE16-NEXT: v_mul_f16_e32 v1, 0.15915494, v1
+; GFX12-FAKE16-NEXT: v_mul_f16_e32 v1, INV2PI, v1
; GFX12-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
-; GFX12-FAKE16-NEXT: v_mul_f16_e32 v2, 0.15915494, v2
+; GFX12-FAKE16-NEXT: v_mul_f16_e32 v2, INV2PI, v2
; GFX12-FAKE16-NEXT: v_sin_f16_e32 v1, v1
; GFX12-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(TRANS32_DEP_1)
; GFX12-FAKE16-NEXT: v_sin_f16_e32 v2, v2
diff --git a/llvm/test/CodeGen/AMDGPU/mad-mix-bf16.ll b/llvm/test/CodeGen/AMDGPU/mad-mix-bf16.ll
index c96ba754c0811..b3f3670fcbc05 100644
--- a/llvm/test/CodeGen/AMDGPU/mad-mix-bf16.ll
+++ b/llvm/test/CodeGen/AMDGPU/mad-mix-bf16.ll
@@ -219,7 +219,7 @@ define float @v_mad_mix_f32_bf16lo_bf16lo_f32imminv2pi(bfloat %src0, bfloat %src
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_wait_loadcnt_dscnt 0x0
; GFX1250-NEXT: s_wait_kmcnt 0x0
-; GFX1250-NEXT: s_mov_b32 s0, 0.15915494
+; GFX1250-NEXT: s_mov_b32 s0, INV2PI
; GFX1250-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; GFX1250-NEXT: v_fma_mix_f32_bf16 v0, v0, v1, s0 op_sel_hi:[1,1,0]
; GFX1250-NEXT: s_set_pc_i64 s[30:31]
@@ -306,7 +306,7 @@ define <2 x float> @v_mad_mix_v2f32_f32imminv2pi(<2 x bfloat> %src0, <2 x bfloat
; GFX1250-NEXT: v_dual_lshlrev_b32 v2, 16, v0 :: v_dual_lshlrev_b32 v4, 16, v1
; GFX1250-NEXT: v_and_b32_e32 v5, 0xffff0000, v1
; GFX1250-NEXT: s_delay_alu instid0(VALU_DEP_1)
-; GFX1250-NEXT: v_pk_fma_f32 v[0:1], v[2:3], v[4:5], 0.15915494 op_sel_hi:[1,1,0]
+; GFX1250-NEXT: v_pk_fma_f32 v[0:1], v[2:3], v[4:5], INV2PI op_sel_hi:[1,1,0]
; GFX1250-NEXT: s_set_pc_i64 s[30:31]
%src0.ext = fpext <2 x bfloat> %src0 to <2 x float>
%src1.ext = fpext <2 x bfloat> %src1 to <2 x float>
diff --git a/llvm/test/CodeGen/AMDGPU/mad-mix.ll b/llvm/test/CodeGen/AMDGPU/mad-mix.ll
index 95df131e21358..2c8fc419c9861 100644
--- a/llvm/test/CodeGen/AMDGPU/mad-mix.ll
+++ b/llvm/test/CodeGen/AMDGPU/mad-mix.ll
@@ -908,7 +908,7 @@ define float @v_mad_mix_f32_f16lo_f16lo_f32imminv2pi(half %src0, half %src1) #0
; SDAG-GFX1100-LABEL: v_mad_mix_f32_f16lo_f16lo_f32imminv2pi:
; SDAG-GFX1100: ; %bb.0:
; SDAG-GFX1100-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; SDAG-GFX1100-NEXT: s_mov_b32 s0, 0.15915494
+; SDAG-GFX1100-NEXT: s_mov_b32 s0, INV2PI
; SDAG-GFX1100-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; SDAG-GFX1100-NEXT: v_fma_mix_f32 v0, v0, v1, s0 op_sel_hi:[1,1,0]
; SDAG-GFX1100-NEXT: s_setpc_b64 s[30:31]
@@ -916,14 +916,14 @@ define float @v_mad_mix_f32_f16lo_f16lo_f32imminv2pi(half %src0, half %src1) #0
; SDAG-GFX900-LABEL: v_mad_mix_f32_f16lo_f16lo_f32imminv2pi:
; SDAG-GFX900: ; %bb.0:
; SDAG-GFX900-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; SDAG-GFX900-NEXT: s_mov_b32 s4, 0.15915494
+; SDAG-GFX900-NEXT: s_mov_b32 s4, INV2PI
; SDAG-GFX900-NEXT: v_mad_mix_f32 v0, v0, v1, s4 op_sel_hi:[1,1,0]
; SDAG-GFX900-NEXT: s_setpc_b64 s[30:31]
;
; SDAG-GFX906-LABEL: v_mad_mix_f32_f16lo_f16lo_f32imminv2pi:
; SDAG-GFX906: ; %bb.0:
; SDAG-GFX906-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; SDAG-GFX906-NEXT: s_mov_b32 s4, 0.15915494
+; SDAG-GFX906-NEXT: s_mov_b32 s4, INV2PI
; SDAG-GFX906-NEXT: v_fma_mix_f32 v0, v0, v1, s4 op_sel_hi:[1,1,0]
; SDAG-GFX906-NEXT: s_setpc_b64 s[30:31]
;
@@ -932,7 +932,7 @@ define float @v_mad_mix_f32_f16lo_f16lo_f32imminv2pi(half %src0, half %src1) #0
; GFX9GEN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; GFX9GEN-NEXT: v_cvt_f32_f16_e32 v0, v0
; GFX9GEN-NEXT: v_cvt_f32_f16_e32 v1, v1
-; GFX9GEN-NEXT: v_mad_f32 v0, v0, v1, 0.15915494
+; GFX9GEN-NEXT: v_mad_f32 v0, v0, v1, INV2PI
; GFX9GEN-NEXT: s_setpc_b64 s[30:31]
;
; VI-LABEL: v_mad_mix_f32_f16lo_f16lo_f32imminv2pi:
@@ -940,7 +940,7 @@ define float @v_mad_mix_f32_f16lo_f16lo_f32imminv2pi(half %src0, half %src1) #0
; VI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; VI-NEXT: v_cvt_f32_f16_e32 v0, v0
; VI-NEXT: v_cvt_f32_f16_e32 v1, v1
-; VI-NEXT: v_mad_f32 v0, v0, v1, 0.15915494
+; VI-NEXT: v_mad_f32 v0, v0, v1, INV2PI
; VI-NEXT: s_setpc_b64 s[30:31]
;
; SDAG-CI-LABEL: v_mad_mix_f32_f16lo_f16lo_f32imminv2pi:
@@ -952,7 +952,7 @@ define float @v_mad_mix_f32_f16lo_f16lo_f32imminv2pi(half %src0, half %src1) #0
; GISEL-GFX1100-LABEL: v_mad_mix_f32_f16lo_f16lo_f32imminv2pi:
; GISEL-GFX1100: ; %bb.0:
; GISEL-GFX1100-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GISEL-GFX1100-NEXT: v_mov_b32_e32 v2, 0.15915494
+; GISEL-GFX1100-NEXT: v_mov_b32_e32 v2, INV2PI
; GISEL-GFX1100-NEXT: s_delay_alu instid0(VALU_DEP_1)
; GISEL-GFX1100-NEXT: v_fma_mix_f32 v0, v0, v1, v2 op_sel_hi:[1,1,0]
; GISEL-GFX1100-NEXT: s_setpc_b64 s[30:31]
@@ -960,14 +960,14 @@ define float @v_mad_mix_f32_f16lo_f16lo_f32imminv2pi(half %src0, half %src1) #0
; GISEL-GFX900-LABEL: v_mad_mix_f32_f16lo_f16lo_f32imminv2pi:
; GISEL-GFX900: ; %bb.0:
; GISEL-GFX900-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GISEL-GFX900-NEXT: v_mov_b32_e32 v2, 0.15915494
+; GISEL-GFX900-NEXT: v_mov_b32_e32 v2, INV2PI
; GISEL-GFX900-NEXT: v_mad_mix_f32 v0, v0, v1, v2 op_sel_hi:[1,1,0]
; GISEL-GFX900-NEXT: s_setpc_b64 s[30:31]
;
; GISEL-GFX906-LABEL: v_mad_mix_f32_f16lo_f16lo_f32imminv2pi:
; GISEL-GFX906: ; %bb.0:
; GISEL-GFX906-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GISEL-GFX906-NEXT: v_mov_b32_e32 v2, 0.15915494
+; GISEL-GFX906-NEXT: v_mov_b32_e32 v2, INV2PI
; GISEL-GFX906-NEXT: v_fma_mix_f32 v0, v0, v1, v2 op_sel_hi:[1,1,0]
; GISEL-GFX906-NEXT: s_setpc_b64 s[30:31]
;
@@ -1467,7 +1467,7 @@ define <2 x float> @v_mad_mix_v2f32_f32imminv2pi(<2 x half> %src0, <2 x half> %s
; SDAG-GFX1100-LABEL: v_mad_mix_v2f32_f32imminv2pi:
; SDAG-GFX1100: ; %bb.0:
; SDAG-GFX1100-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; SDAG-GFX1100-NEXT: s_mov_b32 s0, 0.15915494
+; SDAG-GFX1100-NEXT: s_mov_b32 s0, INV2PI
; SDAG-GFX1100-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
; SDAG-GFX1100-NEXT: v_fma_mix_f32 v2, v0, v1, s0 op_sel_hi:[1,1,0]
; SDAG-GFX1100-NEXT: v_fma_mix_f32 v1, v0, v1, s0 op_sel:[1,1,0] op_sel_hi:[1,1,0]
@@ -1477,7 +1477,7 @@ define <2 x float> @v_mad_mix_v2f32_f32imminv2pi(<2 x half> %src0, <2 x half> %s
; SDAG-GFX900-LABEL: v_mad_mix_v2f32_f32imminv2pi:
; SDAG-GFX900: ; %bb.0:
; SDAG-GFX900-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; SDAG-GFX900-NEXT: s_mov_b32 s4, 0.15915494
+; SDAG-GFX900-NEXT: s_mov_b32 s4, INV2PI
; SDAG-GFX900-NEXT: v_mad_mix_f32 v2, v0, v1, s4 op_sel_hi:[1,1,0]
; SDAG-GFX900-NEXT: v_mad_mix_f32 v1, v0, v1, s4 op_sel:[1,1,0] op_sel_hi:[1,1,0]
; SDAG-GFX900-NEXT: v_mov_b32_e32 v0, v2
@@ -1486,7 +1486,7 @@ define <2 x float> @v_mad_mix_v2f32_f32imminv2pi(<2 x half> %src0, <2 x half> %s
; SDAG-GFX906-LABEL: v_mad_mix_v2f32_f32imminv2pi:
; SDAG-GFX906: ; %bb.0:
; SDAG-GFX906-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; SDAG-GFX906-NEXT: s_mov_b32 s4, 0.15915494
+; SDAG-GFX906-NEXT: s_mov_b32 s4, INV2PI
; SDAG-GFX906-NEXT: v_fma_mix_f32 v2, v0, v1, s4 op_sel_hi:[1,1,0]
; SDAG-GFX906-NEXT: v_fma_mix_f32 v1, v0, v1, s4 op_sel:[1,1,0] op_sel_hi:[1,1,0]
; SDAG-GFX906-NEXT: v_mov_b32_e32 v0, v2
@@ -1499,8 +1499,8 @@ define <2 x float> @v_mad_mix_v2f32_f32imminv2pi(<2 x half> %src0, <2 x half> %s
; SDAG-GFX9GEN-NEXT: v_cvt_f32_f16_e32 v0, v0
; SDAG-GFX9GEN-NEXT: v_cvt_f32_f16_e32 v3, v1
; SDAG-GFX9GEN-NEXT: v_cvt_f32_f16_sdwa v1, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
-; SDAG-GFX9GEN-NEXT: v_mad_f32 v0, v0, v3, 0.15915494
-; SDAG-GFX9GEN-NEXT: v_mad_f32 v1, v2, v1, 0.15915494
+; SDAG-GFX9GEN-NEXT: v_mad_f32 v0, v0, v3, INV2PI
+; SDAG-GFX9GEN-NEXT: v_mad_f32 v1, v2, v1, INV2PI
; SDAG-GFX9GEN-NEXT: s_setpc_b64 s[30:31]
;
; SDAG-VI-LABEL: v_mad_mix_v2f32_f32imminv2pi:
@@ -1510,8 +1510,8 @@ define <2 x float> @v_mad_mix_v2f32_f32imminv2pi(<2 x half> %src0, <2 x half> %s
; SDAG-VI-NEXT: v_cvt_f32_f16_e32 v0, v0
; SDAG-VI-NEXT: v_cvt_f32_f16_e32 v3, v1
; SDAG-VI-NEXT: v_cvt_f32_f16_sdwa v1, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
-; SDAG-VI-NEXT: v_mad_f32 v0, v0, v3, 0.15915494
-; SDAG-VI-NEXT: v_mad_f32 v1, v2, v1, 0.15915494
+; SDAG-VI-NEXT: v_mad_f32 v0, v0, v3, INV2PI
+; SDAG-VI-NEXT: v_mad_f32 v1, v2, v1, INV2PI
; SDAG-VI-NEXT: s_setpc_b64 s[30:31]
;
; SDAG-CI-LABEL: v_mad_mix_v2f32_f32imminv2pi:
@@ -1533,7 +1533,7 @@ define <2 x float> @v_mad_mix_v2f32_f32imminv2pi(<2 x half> %src0, <2 x half> %s
; GISEL-GFX1100-LABEL: v_mad_mix_v2f32_f32imminv2pi:
; GISEL-GFX1100: ; %bb.0:
; GISEL-GFX1100-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GISEL-GFX1100-NEXT: v_mov_b32_e32 v3, 0.15915494
+; GISEL-GFX1100-NEXT: v_mov_b32_e32 v3, INV2PI
; GISEL-GFX1100-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
; GISEL-GFX1100-NEXT: v_fma_mix_f32 v2, v0, v1, v3 op_sel_hi:[1,1,0]
; GISEL-GFX1100-NEXT: v_fma_mix_f32 v1, v0, v1, v3 op_sel:[1,1,0] op_sel_hi:[1,1,0]
@@ -1543,7 +1543,7 @@ define <2 x float> @v_mad_mix_v2f32_f32imminv2pi(<2 x half> %src0, <2 x half> %s
; GISEL-GFX900-LABEL: v_mad_mix_v2f32_f32imminv2pi:
; GISEL-GFX900: ; %bb.0:
; GISEL-GFX900-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GISEL-GFX900-NEXT: v_mov_b32_e32 v3, 0.15915494
+; GISEL-GFX900-NEXT: v_mov_b32_e32 v3, INV2PI
; GISEL-GFX900-NEXT: v_mad_mix_f32 v2, v0, v1, v3 op_sel_hi:[1,1,0]
; GISEL-GFX900-NEXT: v_mad_mix_f32 v1, v0, v1, v3 op_sel:[1,1,0] op_sel_hi:[1,1,0]
; GISEL-GFX900-NEXT: v_mov_b32_e32 v0, v2
@@ -1552,7 +1552,7 @@ define <2 x float> @v_mad_mix_v2f32_f32imminv2pi(<2 x half> %src0, <2 x half> %s
; GISEL-GFX906-LABEL: v_mad_mix_v2f32_f32imminv2pi:
; GISEL-GFX906: ; %bb.0:
; GISEL-GFX906-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GISEL-GFX906-NEXT: v_mov_b32_e32 v3, 0.15915494
+; GISEL-GFX906-NEXT: v_mov_b32_e32 v3, INV2PI
; GISEL-GFX906-NEXT: v_fma_mix_f32 v2, v0, v1, v3 op_sel_hi:[1,1,0]
; GISEL-GFX906-NEXT: v_fma_mix_f32 v1, v0, v1, v3 op_sel:[1,1,0] op_sel_hi:[1,1,0]
; GISEL-GFX906-NEXT: v_mov_b32_e32 v0, v2
@@ -1565,8 +1565,8 @@ define <2 x float> @v_mad_mix_v2f32_f32imminv2pi(<2 x half> %src0, <2 x half> %s
; GISEL-GFX9GEN-NEXT: v_cvt_f32_f16_sdwa v3, v0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
; GISEL-GFX9GEN-NEXT: v_cvt_f32_f16_e32 v0, v1
; GISEL-GFX9GEN-NEXT: v_cvt_f32_f16_sdwa v1, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
-; GISEL-GFX9GEN-NEXT: v_mad_f32 v0, v2, v0, 0.15915494
-; GISEL-GFX9GEN-NEXT: v_mad_f32 v1, v3, v1, 0.15915494
+; GISEL-GFX9GEN-NEXT: v_mad_f32 v0, v2, v0, INV2PI
+; GISEL-GFX9GEN-NEXT: v_mad_f32 v1, v3, v1, INV2PI
; GISEL-GFX9GEN-NEXT: s_setpc_b64 s[30:31]
;
; GISEL-VI-LABEL: v_mad_mix_v2f32_f32imminv2pi:
@@ -1576,8 +1576,8 @@ define <2 x float> @v_mad_mix_v2f32_f32imminv2pi(<2 x half> %src0, <2 x half> %s
; GISEL-VI-NEXT: v_cvt_f32_f16_sdwa v3, v0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
; GISEL-VI-NEXT: v_cvt_f32_f16_e32 v0, v1
; GISEL-VI-NEXT: v_cvt_f32_f16_sdwa v1, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
-; GISEL-VI-NEXT: v_mad_f32 v0, v2, v0, 0.15915494
-; GISEL-VI-NEXT: v_mad_f32 v1, v3, v1, 0.15915494
+; GISEL-VI-NEXT: v_mad_f32 v0, v2, v0, INV2PI
+; GISEL-VI-NEXT: v_mad_f32 v1, v3, v1, INV2PI
; GISEL-VI-NEXT: s_setpc_b64 s[30:31]
;
; GISEL-CI-LABEL: v_mad_mix_v2f32_f32imminv2pi:
diff --git a/llvm/test/CodeGen/AMDGPU/select-fabs-fneg-extract.ll b/llvm/test/CodeGen/AMDGPU/select-fabs-fneg-extract.ll
index c402b692f797f..cc11eb04fb998 100644
--- a/llvm/test/CodeGen/AMDGPU/select-fabs-fneg-extract.ll
+++ b/llvm/test/CodeGen/AMDGPU/select-fabs-fneg-extract.ll
@@ -392,7 +392,7 @@ define amdgpu_kernel void @add_select_fneg_inv2pi_f32(i32 %c) #0 {
; SI-DAG: v_mov_b32_e32 [[K:v[0-9]+]], 0x3e22f983
; SI: v_cndmask_b32_e32 [[SELECT:v[0-9]+]], [[K]], [[X]], vcc
-; VI: v_cndmask_b32_e32 [[SELECT:v[0-9]+]], 0.15915494, [[X]], vcc
+; VI: v_cndmask_b32_e32 [[SELECT:v[0-9]+]], INV2PI, [[X]], vcc
; GCN: v_sub_f32_e32 v{{[0-9]+}}, [[Y]], [[SELECT]]
define amdgpu_kernel void @add_select_fneg_neginv2pi_f32(i32 %c) #0 {
@@ -828,7 +828,7 @@ define amdgpu_kernel void @select_fneg_posk_src_rcp_f32(i32 %c) #0 {
; SI: v_cndmask_b32_e64 [[SELECT:v[0-9]+]], [[K]], -|[[X]]|, [[VCC]]
; SI: v_mul_f32_e32 v{{[0-9]+}}, [[SELECT]], [[Y]]
-; VI: v_cndmask_b32_e64 [[SELECT:v[0-9]+]], 0.15915494, -|[[X]]|, [[VCC]]
+; VI: v_cndmask_b32_e64 [[SELECT:v[0-9]+]], INV2PI, -|[[X]]|, [[VCC]]
; VI: v_mul_f32_e32 v{{[0-9]+}}, [[SELECT]], [[Y]]
define amdgpu_kernel void @mul_select_negfabs_posk_inv2pi_f32(i32 %c) #0 {
%x = load volatile float, ptr addrspace(1) poison
@@ -855,7 +855,7 @@ define amdgpu_kernel void @mul_select_negfabs_posk_inv2pi_f32(i32 %c) #0 {
; SI: v_mul_f32_e32 v{{[0-9]+}}, [[SELECT]], [[Y]]
-; VI: v_cndmask_b32_e64 [[SELECT:v[0-9]+]], 0.15915494, -|[[X]]|, [[VCC]]
+; VI: v_cndmask_b32_e64 [[SELECT:v[0-9]+]], INV2PI, -|[[X]]|, [[VCC]]
; VI: v_mul_f32_e32 v{{[0-9]+}}, [[SELECT]], [[Y]]
define amdgpu_kernel void @mul_select_posk_inv2pi_negfabs_f32(i32 %c) #0 {
%x = load volatile float, ptr addrspace(1) poison
diff --git a/llvm/test/MC/AMDGPU/bf16_imm-fake16.s b/llvm/test/MC/AMDGPU/bf16_imm-fake16.s
index 482bb3c87ce0b..9b210ddf17603 100644
--- a/llvm/test/MC/AMDGPU/bf16_imm-fake16.s
+++ b/llvm/test/MC/AMDGPU/bf16_imm-fake16.s
@@ -39,16 +39,16 @@ v_dot2_bf16_bf16 v2, v0, -4.0, v2
// which cannot be accurately represented in bf16.
v_dot2_bf16_bf16 v2, v0, 0.158203125, v2
-// CHECK: v_dot2_bf16_bf16 v2, v0, 0.15915494, v2 ; encoding: [0x02,0x00,0x67,0xd6,0x00,0xf1,0x09,0x04]
+// CHECK: v_dot2_bf16_bf16 v2, v0, INV2PI, v2 ; encoding: [0x02,0x00,0x67,0xd6,0x00,0xf1,0x09,0x04]
v_dot2_bf16_bf16 v2, v0, 0.15915494, v2
-// CHECK: v_dot2_bf16_bf16 v2, v0, 0.15915494, v2 ; encoding: [0x02,0x00,0x67,0xd6,0x00,0xf1,0x09,0x04]
+// CHECK: v_dot2_bf16_bf16 v2, v0, INV2PI, v2 ; encoding: [0x02,0x00,0x67,0xd6,0x00,0xf1,0x09,0x04]
v_dot2_bf16_bf16 v2, v0, 0x3e22, v2
-// CHECK: v_dot2_bf16_bf16 v2, v0, 0.15915494, v2 ; encoding: [0x02,0x00,0x67,0xd6,0x00,0xf1,0x09,0x04]
+// CHECK: v_dot2_bf16_bf16 v2, v0, INV2PI, v2 ; encoding: [0x02,0x00,0x67,0xd6,0x00,0xf1,0x09,0x04]
v_dot2_bf16_bf16 v2, v0, v2, 0.15915494
-// CHECK: v_dot2_bf16_bf16 v2, v0, v2, 0.15915494 ; encoding: [0x02,0x00,0x67,0xd6,0x00,0x05,0xe2,0x03]
+// CHECK: v_dot2_bf16_bf16 v2, v0, v2, INV2PI ; encoding: [0x02,0x00,0x67,0xd6,0x00,0x05,0xe2,0x03]
v_dot2_f32_bf16 v2, v1, 0, v2
// CHECK: v_dot2_f32_bf16 v2, v1, 0, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0x01,0x09,0x1c]
@@ -78,13 +78,13 @@ v_dot2_f32_bf16 v2, v1, -4.0, v2
// CHECK: v_dot2_f32_bf16 v2, v1, -4.0, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xef,0x09,0x1c]
v_dot2_f32_bf16 v2, v1, 0.15915494, v2
-// CHECK: v_dot2_f32_bf16 v2, v1, 0.15915494, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xf1,0x09,0x1c]
+// CHECK: v_dot2_f32_bf16 v2, v1, INV2PI, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xf1,0x09,0x1c]
v_dot2_f32_bf16 v2, v1, INV2PI, v2
-// CHECK: v_dot2_f32_bf16 v2, v1, 0.15915494, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xf1,0x09,0x1c]
+// CHECK: v_dot2_f32_bf16 v2, v1, INV2PI, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xf1,0x09,0x1c]
v_dot2_f32_bf16 v2, v1, 0x3e22, v2
-// CHECK: v_dot2_f32_bf16 v2, v1, 0.15915494, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xf1,0x09,0x1c]
+// CHECK: v_dot2_f32_bf16 v2, v1, INV2PI, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xf1,0x09,0x1c]
v_dot2_f32_bf16 v2, 0.5, v1, v2
// CHECK: v_dot2_f32_bf16 v2, 0.5, v1, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0xf0,0x02,0x0a,0x1c]
diff --git a/llvm/test/MC/AMDGPU/bf16_imm.s b/llvm/test/MC/AMDGPU/bf16_imm.s
index d734692c84b0e..f09d6c20af7b6 100644
--- a/llvm/test/MC/AMDGPU/bf16_imm.s
+++ b/llvm/test/MC/AMDGPU/bf16_imm.s
@@ -39,19 +39,19 @@ v_dot2_bf16_bf16 v2.l, v0, -4.0, v2.l
// which cannot be accurately represented in bf16.
v_dot2_bf16_bf16 v2.l, v0, 0.158203125, v2.l
-// CHECK: v_dot2_bf16_bf16 v2.l, v0, 0.15915494, v2.l ; encoding: [0x02,0x00,0x67,0xd6,0x00,0xf1,0x09,0x04]
+// CHECK: v_dot2_bf16_bf16 v2.l, v0, INV2PI, v2.l ; encoding: [0x02,0x00,0x67,0xd6,0x00,0xf1,0x09,0x04]
v_dot2_bf16_bf16 v2.l, v0, 0.15915494, v2.l
-// CHECK: v_dot2_bf16_bf16 v2.l, v0, 0.15915494, v2.l ; encoding: [0x02,0x00,0x67,0xd6,0x00,0xf1,0x09,0x04]
+// CHECK: v_dot2_bf16_bf16 v2.l, v0, INV2PI, v2.l ; encoding: [0x02,0x00,0x67,0xd6,0x00,0xf1,0x09,0x04]
v_dot2_bf16_bf16 v2.l, v0, 0x3e22, v2.l
-// CHECK: v_dot2_bf16_bf16 v2.l, v0, 0.15915494, v2.l ; encoding: [0x02,0x00,0x67,0xd6,0x00,0xf1,0x09,0x04]
+// CHECK: v_dot2_bf16_bf16 v2.l, v0, INV2PI, v2.l ; encoding: [0x02,0x00,0x67,0xd6,0x00,0xf1,0x09,0x04]
v_dot2_bf16_bf16 v2.l, v0, v2, 0.15915494
-// CHECK: v_dot2_bf16_bf16 v2.l, v0, v2, 0.15915494 ; encoding: [0x02,0x00,0x67,0xd6,0x00,0x05,0xe2,0x03]
+// CHECK: v_dot2_bf16_bf16 v2.l, v0, v2, INV2PI ; encoding: [0x02,0x00,0x67,0xd6,0x00,0x05,0xe2,0x03]
v_dot2_bf16_bf16 v2.l, v0, v2, INV2PI
-// CHECK: v_dot2_bf16_bf16 v2.l, v0, v2, 0.15915494 ; encoding: [0x02,0x00,0x67,0xd6,0x00,0x05,0xe2,0x03]
+// CHECK: v_dot2_bf16_bf16 v2.l, v0, v2, INV2PI ; encoding: [0x02,0x00,0x67,0xd6,0x00,0x05,0xe2,0x03]
v_dot2_f32_bf16 v2, v1, 0, v2
// CHECK: v_dot2_f32_bf16 v2, v1, 0, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0x01,0x09,0x1c]
@@ -81,10 +81,10 @@ v_dot2_f32_bf16 v2, v1, -4.0, v2
// CHECK: v_dot2_f32_bf16 v2, v1, -4.0, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xef,0x09,0x1c]
v_dot2_f32_bf16 v2, v1, 0.15915494, v2
-// CHECK: v_dot2_f32_bf16 v2, v1, 0.15915494, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xf1,0x09,0x1c]
+// CHECK: v_dot2_f32_bf16 v2, v1, INV2PI, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xf1,0x09,0x1c]
v_dot2_f32_bf16 v2, v1, 0x3e22, v2
-// CHECK: v_dot2_f32_bf16 v2, v1, 0.15915494, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xf1,0x09,0x1c]
+// CHECK: v_dot2_f32_bf16 v2, v1, INV2PI, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xf1,0x09,0x1c]
v_dot2_f32_bf16 v2, 0.5, v1, v2
// CHECK: v_dot2_f32_bf16 v2, 0.5, v1, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0xf0,0x02,0x0a,0x1c]
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_valu_lit64.s b/llvm/test/MC/AMDGPU/gfx1250_asm_valu_lit64.s
index 4c0d5a69c2907..a8cbcd6dd78da 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_valu_lit64.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_valu_lit64.s
@@ -230,13 +230,13 @@ v_mov_b64 v[0:1], 0x12345678
// 1.0 / (2.0 * pi)
v_ceil_f64 v[254:255], 0x3fc45f306dc9c882
-// GFX1250: v_ceil_f64_e32 v[254:255], 0.15915494309189532 ; encoding: [0xf8,0x30,0xfc,0x7f]
+// GFX1250: v_ceil_f64_e32 v[254:255], INV2PI ; encoding: [0xf8,0x30,0xfc,0x7f]
v_ceil_f64 v[254:255], 0.15915494309189532
-// GFX1250: v_ceil_f64_e32 v[254:255], 0.15915494309189532 ; encoding: [0xf8,0x30,0xfc,0x7f]
+// GFX1250: v_ceil_f64_e32 v[254:255], INV2PI ; encoding: [0xf8,0x30,0xfc,0x7f]
v_ceil_f64 v[254:255], INV2PI
-// GFX1250: v_ceil_f64_e32 v[254:255], 0.15915494309189532 ; encoding: [0xf8,0x30,0xfc,0x7f]
+// GFX1250: v_ceil_f64_e32 v[254:255], INV2PI ; encoding: [0xf8,0x30,0xfc,0x7f]
v_ceil_f64 v[254:255], -4.0
// GFX1250: v_ceil_f64_e32 v[254:255], -4.0 ; encoding: [0xf7,0x30,0xfc,0x7f]
diff --git a/llvm/test/MC/AMDGPU/gfx950_dlops.s b/llvm/test/MC/AMDGPU/gfx950_dlops.s
index 52a4167f2fe0e..674de4f0958ad 100644
--- a/llvm/test/MC/AMDGPU/gfx950_dlops.s
+++ b/llvm/test/MC/AMDGPU/gfx950_dlops.s
@@ -34,10 +34,10 @@ v_dot2_f32_bf16 v2, v1, -4.0, v2
// GFX950: v_dot2_f32_bf16 v2, v1, -4.0, v2 ; encoding: [0x02,0x40,0x9a,0xd3,0x01,0xef,0x09,0x1c]
v_dot2_f32_bf16 v2, v1, 0.15915494, v2
-// GFX950: v_dot2_f32_bf16 v2, v1, 0.15915494, v2 ; encoding: [0x02,0x40,0x9a,0xd3,0x01,0xf1,0x09,0x1c]
+// GFX950: v_dot2_f32_bf16 v2, v1, INV2PI, v2 ; encoding: [0x02,0x40,0x9a,0xd3,0x01,0xf1,0x09,0x1c]
v_dot2_f32_bf16 v2, v1, INV2PI, v2
-// GFX950: v_dot2_f32_bf16 v2, v1, 0.15915494, v2 ; encoding: [0x02,0x40,0x9a,0xd3,0x01,0xf1,0x09,0x1c]
+// GFX950: v_dot2_f32_bf16 v2, v1, INV2PI, v2 ; encoding: [0x02,0x40,0x9a,0xd3,0x01,0xf1,0x09,0x1c]
v_dot2_f32_bf16 v2, 0.5, v1, v2
// GFX950: v_dot2_f32_bf16 v2, 0.5, v1, v2 ; encoding: [0x02,0x40,0x9a,0xd3,0xf0,0x02,0x0a,0x1c]
diff --git a/llvm/test/MC/AMDGPU/inline-imm-inv2pi.s b/llvm/test/MC/AMDGPU/inline-imm-inv2pi.s
index d730bb32c9f13..631307c2a60ae 100644
--- a/llvm/test/MC/AMDGPU/inline-imm-inv2pi.s
+++ b/llvm/test/MC/AMDGPU/inline-imm-inv2pi.s
@@ -6,5 +6,5 @@
// immediate or not.
// SI: v_cvt_f32_f16_e32 v0, 0x3118 ; encoding: [0xff,0x16,0x00,0x7e,0x18,0x31,0x00,0x00]
-// VI: v_cvt_f32_f16_e32 v0, 0.15915494 ; encoding: [0xf8,0x16,0x00,0x7e]
+// VI: v_cvt_f32_f16_e32 v0, INV2PI ; encoding: [0xf8,0x16,0x00,0x7e]
v_cvt_f32_f16_e32 v0, 0x3118
diff --git a/llvm/test/MC/AMDGPU/literal16.s b/llvm/test/MC/AMDGPU/literal16.s
index bf5028c17e077..af1d3a81ac052 100644
--- a/llvm/test/MC/AMDGPU/literal16.s
+++ b/llvm/test/MC/AMDGPU/literal16.s
@@ -40,10 +40,10 @@ v_add_f16 v1, -4.0, v2
// VI: v_add_f16_e32 v1, -4.0, v2 ; encoding: [0xf7,0x04,0x02,0x3e]
v_add_f16 v1, 0.15915494, v2
-// VI: v_add_f16_e32 v1, 0.15915494, v2 ; encoding: [0xf8,0x04,0x02,0x3e]
+// VI: v_add_f16_e32 v1, INV2PI, v2 ; encoding: [0xf8,0x04,0x02,0x3e]
v_add_f16 v1, INV2PI, v2
-// VI: v_add_f16_e32 v1, 0.15915494, v2 ; encoding: [0xf8,0x04,0x02,0x3e]
+// VI: v_add_f16_e32 v1, INV2PI, v2 ; encoding: [0xf8,0x04,0x02,0x3e]
v_add_f16 v1, -0.15915494, v2
// VI: v_add_f16_e32 v1, 0xb118, v2 ; encoding: [0xff,0x04,0x02,0x3e,0x18,0xb1,0x00,0x00]
@@ -125,7 +125,7 @@ v_add_f16 v1, 0xc400, v2
// VI: v_add_f16_e32 v1, -4.0, v2 ; encoding: [0xf7,0x04,0x02,0x3e]
v_add_f16 v1, 0x3118, v2
-// VI: v_add_f16_e32 v1, 0.15915494, v2 ; encoding: [0xf8,0x04,0x02,0x3e]
+// VI: v_add_f16_e32 v1, INV2PI, v2 ; encoding: [0xf8,0x04,0x02,0x3e]
v_add_f16 v1, -32768, v2
// VI: v_add_f16_e32 v1, 0x8000, v2 ; encoding: [0xff,0x04,0x02,0x3e,0x00,0x80,0x00,0x00]
diff --git a/llvm/test/MC/AMDGPU/literals.s b/llvm/test/MC/AMDGPU/literals.s
index f749f6e16fb03..9658712b270eb 100644
--- a/llvm/test/MC/AMDGPU/literals.s
+++ b/llvm/test/MC/AMDGPU/literals.s
@@ -957,16 +957,16 @@ v_trunc_f32_e32 v0, 0x3fc45f306dc9c882
// NOGCN: :[[@LINE-1]]:21: error: invalid operand for instruction
v_fract_f64_e32 v[0:1], 0x3fc45f306dc9c882
-// GFX11: v_fract_f64_e32 v[0:1], 0.15915494309189532 ; encoding: [0xf8,0x7c,0x00,0x7e]
-// GFX12XX: v_fract_f64_e32 v[0:1], 0.15915494309189532 ; encoding: [0xf8,0x7c,0x00,0x7e]
-// GFX89: v_fract_f64_e32 v[0:1], 0.15915494309189532 ; encoding: [0xf8,0x64,0x00,0x7e]
+// GFX11: v_fract_f64_e32 v[0:1], INV2PI ; encoding: [0xf8,0x7c,0x00,0x7e]
+// GFX12XX: v_fract_f64_e32 v[0:1], INV2PI ; encoding: [0xf8,0x7c,0x00,0x7e]
+// GFX89: v_fract_f64_e32 v[0:1], INV2PI ; encoding: [0xf8,0x64,0x00,0x7e]
// NOSICI: :[[@LINE-4]]:25: error: invalid operand for instruction
// NOSICIVI: :[[@LINE-2]]:25: error: invalid operand for instruction
v_trunc_f32_e32 v0, 0x3e22f983
-// GFX11: v_trunc_f32_e32 v0, 0.15915494 ; encoding: [0xf8,0x42,0x00,0x7e]
-// GFX12XX: v_trunc_f32_e32 v0, 0.15915494 ; encoding: [0xf8,0x42,0x00,0x7e]
-// GFX89: v_trunc_f32_e32 v0, 0.15915494 ; encoding: [0xf8,0x38,0x00,0x7e]
+// GFX11: v_trunc_f32_e32 v0, INV2PI ; encoding: [0xf8,0x42,0x00,0x7e]
+// GFX12XX: v_trunc_f32_e32 v0, INV2PI ; encoding: [0xf8,0x42,0x00,0x7e]
+// GFX89: v_trunc_f32_e32 v0, INV2PI ; encoding: [0xf8,0x38,0x00,0x7e]
// SICI: v_trunc_f32_e32 v0, 0x3e22f983 ; encoding: [0xff,0x42,0x00,0x7e,0x83,0xf9,0x22,0x3e]
v_fract_f64_e32 v[0:1], 0x3e22f983
@@ -979,16 +979,16 @@ v_trunc_f32_e64 v0, 0x3fc45f306dc9c882
// NOGCN: :[[@LINE-1]]:21: error: invalid operand for instruction
v_fract_f64_e64 v[0:1], 0x3fc45f306dc9c882
-// GFX11: v_fract_f64_e64 v[0:1], 0.15915494309189532 ; encoding: [0x00,0x00,0xbe,0xd5,0xf8,0x00,0x00,0x00]
-// GFX12XX: v_fract_f64_e64 v[0:1], 0.15915494309189532 ; encoding: [0x00,0x00,0xbe,0xd5,0xf8,0x00,0x00,0x00]
-// GFX89: v_fract_f64_e64 v[0:1], 0.15915494309189532 ; encoding: [0x00,0x00,0x72,0xd1,0xf8,0x00,0x00,0x00]
+// GFX11: v_fract_f64_e64 v[0:1], INV2PI ; encoding: [0x00,0x00,0xbe,0xd5,0xf8,0x00,0x00,0x00]
+// GFX12XX: v_fract_f64_e64 v[0:1], INV2PI ; encoding: [0x00,0x00,0xbe,0xd5,0xf8,0x00,0x00,0x00]
+// GFX89: v_fract_f64_e64 v[0:1], INV2PI ; encoding: [0x00,0x00,0x72,0xd1,0xf8,0x00,0x00,0x00]
// NOSICI: :[[@LINE-4]]:25: error: invalid operand for instruction
// NOSICIVI: :[[@LINE-2]]:25: error: invalid operand for instruction
v_trunc_f32_e64 v0, 0x3e22f983
-// GFX11: v_trunc_f32_e64 v0, 0.15915494 ; encoding: [0x00,0x00,0xa1,0xd5,0xf8,0x00,0x00,0x00]
-// GFX12XX: v_trunc_f32_e64 v0, 0.15915494 ; encoding: [0x00,0x00,0xa1,0xd5,0xf8,0x00,0x00,0x00]
-// GFX89: v_trunc_f32_e64 v0, 0.15915494 ; encoding: [0x00,0x00,0x5c,0xd1,0xf8,0x00,0x00,0x00]
+// GFX11: v_trunc_f32_e64 v0, INV2PI ; encoding: [0x00,0x00,0xa1,0xd5,0xf8,0x00,0x00,0x00]
+// GFX12XX: v_trunc_f32_e64 v0, INV2PI ; encoding: [0x00,0x00,0xa1,0xd5,0xf8,0x00,0x00,0x00]
+// GFX89: v_trunc_f32_e64 v0, INV2PI ; encoding: [0x00,0x00,0x5c,0xd1,0xf8,0x00,0x00,0x00]
// NOSICI: :[[@LINE-4]]:21: error: literal operands are not supported
// NOSICIVI: :[[@LINE-2]]:21: error: literal operands are not supported
@@ -1000,35 +1000,35 @@ v_fract_f64_e64 v[0:1], 0x3e22f983
// NOSICIVI: :[[@LINE-1]]:25: error: literal operands are not supported
s_mov_b64_e32 s[0:1], 0.159154943091895317852646485335
-// GFX8PLUS: s_mov_b64 s[0:1], 0.15915494309189532 ; encoding: [0xf8,0x01,0x80,0xbe]
+// GFX8PLUS: s_mov_b64 s[0:1], INV2PI ; encoding: [0xf8,0x01,0x80,0xbe]
// NOSICI: :[[@LINE-2]]:23: error: invalid operand for instruction
// NOSICIVI: :[[@LINE-2]]:23: error: invalid operand for instruction
v_and_b32_e32 v0, 0.159154943091895317852646485335, v1
-// GFX11: v_and_b32_e32 v0, 0.15915494, v1 ; encoding: [0xf8,0x02,0x00,0x36]
-// GFX12XX: v_and_b32_e32 v0, 0.15915494, v1 ; encoding: [0xf8,0x02,0x00,0x36]
-// GFX89: v_and_b32_e32 v0, 0.15915494, v1 ; encoding: [0xf8,0x02,0x00,0x26]
+// GFX11: v_and_b32_e32 v0, INV2PI, v1 ; encoding: [0xf8,0x02,0x00,0x36]
+// GFX12XX: v_and_b32_e32 v0, INV2PI, v1 ; encoding: [0xf8,0x02,0x00,0x36]
+// GFX89: v_and_b32_e32 v0, INV2PI, v1 ; encoding: [0xf8,0x02,0x00,0x26]
// SICI: v_and_b32_e32 v0, 0x3e22f983, v1 ; encoding: [0xff,0x02,0x00,0x36,0x83,0xf9,0x22,0x3e]
v_and_b32_e64 v0, 0.159154943091895317852646485335, v1
-// GFX11: v_and_b32_e64 v0, 0.15915494, v1 ; encoding: [0x00,0x00,0x1b,0xd5,0xf8,0x02,0x02,0x00]
-// GFX12XX: v_and_b32_e64 v0, 0.15915494, v1 ; encoding: [0x00,0x00,0x1b,0xd5,0xf8,0x02,0x02,0x00]
-// GFX89: v_and_b32_e64 v0, 0.15915494, v1 ; encoding: [0x00,0x00,0x13,0xd1,0xf8,0x02,0x02,0x00]
+// GFX11: v_and_b32_e64 v0, INV2PI, v1 ; encoding: [0x00,0x00,0x1b,0xd5,0xf8,0x02,0x02,0x00]
+// GFX12XX: v_and_b32_e64 v0, INV2PI, v1 ; encoding: [0x00,0x00,0x1b,0xd5,0xf8,0x02,0x02,0x00]
+// GFX89: v_and_b32_e64 v0, INV2PI, v1 ; encoding: [0x00,0x00,0x13,0xd1,0xf8,0x02,0x02,0x00]
// NOSICI: :[[@LINE-4]]:19: error: literal operands are not supported
// NOSICIVI: :[[@LINE-2]]:19: error: literal operands are not supported
v_fract_f64 v[0:1], 0.159154943091895317852646485335
-// GFX11: v_fract_f64_e32 v[0:1], 0.15915494309189532 ; encoding: [0xf8,0x7c,0x00,0x7e]
-// GFX12XX: v_fract_f64_e32 v[0:1], 0.15915494309189532 ; encoding: [0xf8,0x7c,0x00,0x7e]
-// GFX89: v_fract_f64_e32 v[0:1], 0.15915494309189532 ; encoding: [0xf8,0x64,0x00,0x7e]
+// GFX11: v_fract_f64_e32 v[0:1], INV2PI ; encoding: [0xf8,0x7c,0x00,0x7e]
+// GFX12XX: v_fract_f64_e32 v[0:1], INV2PI ; encoding: [0xf8,0x7c,0x00,0x7e]
+// GFX89: v_fract_f64_e32 v[0:1], INV2PI ; encoding: [0xf8,0x64,0x00,0x7e]
// NOSICI: :[[@LINE-4]]:1: warning: Can't encode literal as exact 64-bit floating-point operand. Low 32-bits will be set to zero
// SICI: v_fract_f64_e32 v[0:1], 0x3fc45f30 ; encoding: [0xff,0x7c,0x00,0x7e,0x30,0x5f,0xc4,0x3f]
// NOSICIVI: :[[@LINE-3]]:1: warning: Can't encode literal as exact 64-bit floating-point operand. Low 32-bits will be set to zero
v_trunc_f32 v0, 0.159154943091895317852646485335
-// GFX11: v_trunc_f32_e32 v0, 0.15915494 ; encoding: [0xf8,0x42,0x00,0x7e]
-// GFX12XX: v_trunc_f32_e32 v0, 0.15915494 ; encoding: [0xf8,0x42,0x00,0x7e]
-// GFX89: v_trunc_f32_e32 v0, 0.15915494 ; encoding: [0xf8,0x38,0x00,0x7e]
+// GFX11: v_trunc_f32_e32 v0, INV2PI ; encoding: [0xf8,0x42,0x00,0x7e]
+// GFX12XX: v_trunc_f32_e32 v0, INV2PI ; encoding: [0xf8,0x42,0x00,0x7e]
+// GFX89: v_trunc_f32_e32 v0, INV2PI ; encoding: [0xf8,0x38,0x00,0x7e]
// SICI: v_trunc_f32_e32 v0, 0x3e22f983 ; encoding: [0xff,0x42,0x00,0x7e,0x83,0xf9,0x22,0x3e]
v_trunc_f32 v0, lit(0.159154943091895317852646485335)
@@ -1038,7 +1038,7 @@ v_trunc_f32 v0, lit(0.159154943091895317852646485335)
// SICI: v_trunc_f32_e32 v0, lit(0x3e22f983) ; encoding: [0xff,0x42,0x00,0x7e,0x83,0xf9,0x22,0x3e]
// SICI: v_trunc_f32_e32 v0, 0x3e22f983 ; encoding: [0xff,0x42,0x00,0x7e,0x83,0xf9,0x22,0x3e]
-// GFX89: v_trunc_f32_e32 v0, 0.15915494 ; encoding: [0xf8,0x38,0x00,0x7e]
+// GFX89: v_trunc_f32_e32 v0, INV2PI ; encoding: [0xf8,0x38,0x00,0x7e]
v_trunc_f32 v0, INV2PI
//---------------------------------------------------------------------------//
diff --git a/llvm/test/MC/AMDGPU/literalv216.s b/llvm/test/MC/AMDGPU/literalv216.s
index 20c3d34c0bb86..74c693f1b3fef 100644
--- a/llvm/test/MC/AMDGPU/literalv216.s
+++ b/llvm/test/MC/AMDGPU/literalv216.s
@@ -57,12 +57,12 @@ v_pk_add_f16 v1, -4.0, v2
// GFX10: v_pk_add_f16 v1, -4.0, v2 ; encoding: [0x01,0x40,0x0f,0xcc,0xf7,0x04,0x02,0x18]
v_pk_add_f16 v1, 0.15915494, v2
-// GFX9: v_pk_add_f16 v1, 0.15915494, v2 ; encoding: [0x01,0x40,0x8f,0xd3,0xf8,0x04,0x02,0x18]
-// GFX10: v_pk_add_f16 v1, 0.15915494, v2 ; encoding: [0x01,0x40,0x0f,0xcc,0xf8,0x04,0x02,0x18]
+// GFX9: v_pk_add_f16 v1, INV2PI, v2 ; encoding: [0x01,0x40,0x8f,0xd3,0xf8,0x04,0x02,0x18]
+// GFX10: v_pk_add_f16 v1, INV2PI, v2 ; encoding: [0x01,0x40,0x0f,0xcc,0xf8,0x04,0x02,0x18]
v_pk_add_f16 v1, INV2PI, v2
-// GFX9: v_pk_add_f16 v1, 0.15915494, v2 ; encoding: [0x01,0x40,0x8f,0xd3,0xf8,0x04,0x02,0x18]
-// GFX10: v_pk_add_f16 v1, 0.15915494, v2 ; encoding: [0x01,0x40,0x0f,0xcc,0xf8,0x04,0x02,0x18]
+// GFX9: v_pk_add_f16 v1, INV2PI, v2 ; encoding: [0x01,0x40,0x8f,0xd3,0xf8,0x04,0x02,0x18]
+// GFX10: v_pk_add_f16 v1, INV2PI, v2 ; encoding: [0x01,0x40,0x0f,0xcc,0xf8,0x04,0x02,0x18]
v_pk_add_f16 v1, -1, v2
// GFX9: v_pk_add_f16 v1, -1, v2 ; encoding: [0x01,0x40,0x8f,0xd3,0xc1,0x04,0x02,0x18]
@@ -157,8 +157,8 @@ v_pk_add_f16 v1, 0xc400, v2
// GFX10: v_pk_add_f16 v1, -4.0, v2 ; encoding: [0x01,0x40,0x0f,0xcc,0xf7,0x04,0x02,0x18]
v_pk_add_f16 v1, 0x3118, v2
-// GFX9: v_pk_add_f16 v1, 0.15915494, v2 ; encoding: [0x01,0x40,0x8f,0xd3,0xf8,0x04,0x02,0x18]
-// GFX10: v_pk_add_f16 v1, 0.15915494, v2 ; encoding: [0x01,0x40,0x0f,0xcc,0xf8,0x04,0x02,0x18]
+// GFX9: v_pk_add_f16 v1, INV2PI, v2 ; encoding: [0x01,0x40,0x8f,0xd3,0xf8,0x04,0x02,0x18]
+// GFX10: v_pk_add_f16 v1, INV2PI, v2 ; encoding: [0x01,0x40,0x0f,0xcc,0xf8,0x04,0x02,0x18]
v_pk_add_f16 v1, 65535, v2
// NOGFX9: :[[@LINE-1]]:{{[0-9]+}}: error: literal operands are not supported
diff --git a/llvm/test/MC/AMDGPU/mubuf.s b/llvm/test/MC/AMDGPU/mubuf.s
index 2ea6a3b5150a7..e7e2dc7522894 100644
--- a/llvm/test/MC/AMDGPU/mubuf.s
+++ b/llvm/test/MC/AMDGPU/mubuf.s
@@ -645,11 +645,11 @@ buffer_atomic_add v5, off, s[8:11], 0.5 offset:4095 glc
buffer_atomic_add v5, off, s[8:11], 0.15915494 offset:4095 glc
// NOSICI: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction
-// VI: buffer_atomic_add v5, off, s[8:11], 0.15915494 offset:4095 glc ; encoding: [0xff,0x4f,0x08,0xe1,0x00,0x05,0x02,0xf8]
+// VI: buffer_atomic_add v5, off, s[8:11], INV2PI offset:4095 glc ; encoding: [0xff,0x4f,0x08,0xe1,0x00,0x05,0x02,0xf8]
buffer_atomic_add v5, off, s[8:11], INV2PI offset:4095 glc
// NOSICI: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction
-// VI: buffer_atomic_add v5, off, s[8:11], 0.15915494 offset:4095 glc ; encoding: [0xff,0x4f,0x08,0xe1,0x00,0x05,0x02,0xf8]
+// VI: buffer_atomic_add v5, off, s[8:11], INV2PI offset:4095 glc ; encoding: [0xff,0x4f,0x08,0xe1,0x00,0x05,0x02,0xf8]
buffer_atomic_fcmpswap v[0:1], off, s[0:3], s0 offset:4095
// SICI: buffer_atomic_fcmpswap v[0:1], off, s[0:3], s0 offset:4095 ; encoding: [0xff,0x0f,0xf8,0xe0,0x00,0x00,0x00,0x00]
diff --git a/llvm/test/MC/Disassembler/AMDGPU/bf16_imm.txt b/llvm/test/MC/Disassembler/AMDGPU/bf16_imm.txt
index 140e52425fc19..ca7c7bf4bc83b 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/bf16_imm.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/bf16_imm.txt
@@ -49,8 +49,8 @@
# CHECK-FAKE16: v_dot2_bf16_bf16 v2, v0, -4.0, v2 ; encoding: [0x02,0x00,0x67,0xd6,0x00,0xef,0x09,0x04]
0x02,0x00,0x67,0xd6,0x00,0xf1,0x09,0x04
-# CHECK-REAL16: v_dot2_bf16_bf16 v2.l, v0, 0.15915494, v2.l ; encoding: [0x02,0x00,0x67,0xd6,0x00,0xf1,0x09,0x04]
-# CHECK-FAKE16: v_dot2_bf16_bf16 v2, v0, 0.15915494, v2 ; encoding: [0x02,0x00,0x67,0xd6,0x00,0xf1,0x09,0x04]
+# CHECK-REAL16: v_dot2_bf16_bf16 v2.l, v0, INV2PI, v2.l ; encoding: [0x02,0x00,0x67,0xd6,0x00,0xf1,0x09,0x04]
+# CHECK-FAKE16: v_dot2_bf16_bf16 v2, v0, INV2PI, v2 ; encoding: [0x02,0x00,0x67,0xd6,0x00,0xf1,0x09,0x04]
0x02,0x40,0x1a,0xcc,0x01,0x01,0x09,0x1c
# CHECK: v_dot2_f32_bf16 v2, v1, 0, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0x01,0x09,0x1c]
@@ -80,10 +80,10 @@
# CHECK: v_dot2_f32_bf16 v2, v1, -4.0, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xef,0x09,0x1c]
0x02,0x40,0x1a,0xcc,0x01,0xf1,0x09,0x1c
-# CHECK: v_dot2_f32_bf16 v2, v1, 0.15915494, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xf1,0x09,0x1c]
+# CHECK: v_dot2_f32_bf16 v2, v1, INV2PI, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xf1,0x09,0x1c]
0x02,0x40,0x1a,0xcc,0x01,0xf1,0x09,0x1c
-# CHECK: v_dot2_f32_bf16 v2, v1, 0.15915494, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xf1,0x09,0x1c]
+# CHECK: v_dot2_f32_bf16 v2, v1, INV2PI, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0x01,0xf1,0x09,0x1c]
0x02,0x40,0x1a,0xcc,0xf0,0x02,0x0a,0x1c
# CHECK: v_dot2_f32_bf16 v2, 0.5, v1, v2 ; encoding: [0x02,0x40,0x1a,0xcc,0xf0,0x02,0x0a,0x1c]
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx10_vop3p_literalv216.txt b/llvm/test/MC/Disassembler/AMDGPU/gfx10_vop3p_literalv216.txt
index a022c79fe97e6..a8f5246eba750 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx10_vop3p_literalv216.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx10_vop3p_literalv216.txt
@@ -34,7 +34,7 @@
# GFX10: v_pk_add_f16 v1, -4.0, v2 ; encoding: [0x01,0x40,0x0f,0xcc,0xf7,0x04,0x02,0x18]
0x01,0x00,0x0f,0xcc,0xf7,0x04,0x02,0x18
-# GFX10: v_pk_add_f16 v1, 0.15915494, v2 ; encoding: [0x01,0x40,0x0f,0xcc,0xf8,0x04,0x02,0x18]
+# GFX10: v_pk_add_f16 v1, INV2PI, v2 ; encoding: [0x01,0x40,0x0f,0xcc,0xf8,0x04,0x02,0x18]
0x01,0x00,0x0f,0xcc,0xf8,0x04,0x02,0x18
# GFX10: v_pk_add_f16 v1, -1, v2 ; encoding: [0x01,0x40,0x0f,0xcc,0xc1,0x04,0x02,0x18]
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_valu_lit64.txt b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_valu_lit64.txt
index cce6a74bf64a3..9f296dffb605e 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_valu_lit64.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_valu_lit64.txt
@@ -215,7 +215,7 @@
# GFX1250: v_mov_b64_e32 v[0:1], 0x12345678 ; encoding: [0xff,0x3a,0x00,0x7e,0x78,0x56,0x34,0x12]
0xf8,0x30,0xfc,0x7f
-# GFX1250: v_ceil_f64_e32 v[254:255], 0.15915494309189532 ; encoding: [0xf8,0x30,0xfc,0x7f]
+# GFX1250: v_ceil_f64_e32 v[254:255], INV2PI ; encoding: [0xf8,0x30,0xfc,0x7f]
0xf7,0x30,0xfc,0x7f
# GFX1250: v_ceil_f64_e32 v[254:255], -4.0 ; encoding: [0xf7,0x30,0xfc,0x7f]
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx8-literal.txt b/llvm/test/MC/Disassembler/AMDGPU/gfx8-literal.txt
index d2c24f3fbaef5..7224d6c273789 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx8-literal.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx8-literal.txt
@@ -1,7 +1,7 @@
# RUN: llvm-mc -triple=amdgcn -mcpu=tonga -disassemble -show-encoding %s | FileCheck -check-prefix=VI %s
-# VI: v_fract_f64_e32 v[0:1], 0.15915494309189532 ; encoding: [0xf8,0x64,0x00,0x7e]
+# VI: v_fract_f64_e32 v[0:1], INV2PI ; encoding: [0xf8,0x64,0x00,0x7e]
0xf8,0x64,0x00,0x7e
-# VI: v_trunc_f32_e32 v0, 0.15915494 ; encoding: [0xf8,0x38,0x00,0x7e]
+# VI: v_trunc_f32_e32 v0, INV2PI ; encoding: [0xf8,0x38,0x00,0x7e]
0xf8,0x38,0x00,0x7e
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx8-literal16.txt b/llvm/test/MC/Disassembler/AMDGPU/gfx8-literal16.txt
index 856d7c22177ff..0df3b87cbe23e 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx8-literal16.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx8-literal16.txt
@@ -24,7 +24,7 @@
# VI: v_add_f16_e32 v1, -4.0, v3 ; encoding: [0xf7,0x06,0x02,0x3e]
0xf7 0x06 0x02 0x3e
-# VI: v_add_f16_e32 v1, 0.15915494, v3 ; encoding: [0xf8,0x06,0x02,0x3e]
+# VI: v_add_f16_e32 v1, INV2PI, v3 ; encoding: [0xf8,0x06,0x02,0x3e]
0xf8 0x06 0x02 0x3e
# VI: v_add_f16_e32 v1, 0x41, v3 ; encoding: [0xff,0x06,0x02,0x3e,0x41,0x00,0x00,0x00]
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx950_dasm_vop3.txt b/llvm/test/MC/Disassembler/AMDGPU/gfx950_dasm_vop3.txt
index 97bc68b0774b1..f13b768f25650 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx950_dasm_vop3.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx950_dasm_vop3.txt
@@ -813,7 +813,7 @@
# GFX950: v_dot2_f32_bf16 v2, v1, -4.0, v2 ; encoding: [0x02,0x40,0x9a,0xd3,0x01,0xef,0x09,0x1c]
0x02,0x40,0x9a,0xd3,0x01,0xef,0x09,0x1c
-# GFX950: v_dot2_f32_bf16 v2, v1, 0.15915494, v2 ; encoding: [0x02,0x40,0x9a,0xd3,0x01,0xf1,0x09,0x1c]
+# GFX950: v_dot2_f32_bf16 v2, v1, INV2PI, v2 ; encoding: [0x02,0x40,0x9a,0xd3,0x01,0xf1,0x09,0x1c]
0x02,0x40,0x9a,0xd3,0x01,0xf1,0x09,0x1c
# GFX950: v_dot2_f32_bf16 v2, 0.5, v1, v2 ; encoding: [0x02,0x40,0x9a,0xd3,0xf0,0x02,0x0a,0x1c]
More information about the llvm-commits
mailing list