diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index 2ced4d6813766..0fff784c350e4 100644
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -3166,6 +3166,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, bool HasLit,
                                       bool HasLit64) {
@@ -3190,11 +3197,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;
@@ -3210,6 +3218,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 7395a51250490..edf1aa0f60f8e 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_valu_lit64.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_valu_lit64.s
@@ -233,6 +233,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 b4d4e365d0453..081ce520d7d27 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3.s
@@ -435,6 +435,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]
 
@@ -467,6 +470,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 783947544d221..b38327cd7f0c1 100644
--- a/llvm/test/MC/AMDGPU/literals.s
+++ b/llvm/test/MC/AMDGPU/literals.s
@@ -460,22 +460,58 @@ v_fract_f64_e64 v[0:1], 0x3e22f983
 // GFX89: s_mov_b64 s[0:1], 0.15915494309189532 ; encoding: [0xf8,0x01,0x80,0xbe]
 s_mov_b64_e32 s[0:1], 0.159154943091895317852646485335
 
+// NOSICI: :[[@LINE+2]]:{{[0-9]+}}: error: invalid operand for instruction
+// GFX89: s_mov_b64 s[0:1], 0.15915494309189532 ; encoding: [0xf8,0x01,0x80,0xbe]
+s_mov_b64_e32 s[0:1], INV2PI64
+
+// NOSICI: :[[@LINE+2]]:{{[0-9]+}}: error: invalid operand for instruction
+// NOGFX89: :[[@LINE+1]]:{{[0-9]+}}: error: invalid operand for instruction
+s_mov_b64_e32 s[0:1], 0.15915494
+
+// NOSICI: :[[@LINE+2]]:{{[0-9]+}}: error: invalid operand for instruction
+// NOGFX89: :[[@LINE+1]]:{{[0-9]+}}: error: invalid operand for instruction
+s_mov_b64_e32 s[0:1], INV2PI
+
 // SICI: v_and_b32_e32 v0, 0x3e22f983, v1 ; encoding: [0xff,0x02,0x00,0x36,0x83,0xf9,0x22,0x3e]
 // GFX89: v_and_b32_e32 v0, 0.15915494, v1 ; encoding: [0xf8,0x02,0x00,0x26]
 v_and_b32_e32 v0, 0.159154943091895317852646485335, v1
 
+// SICI: v_and_b32_e32 v0, 0x3e22f983, v1 ; encoding: [0xff,0x02,0x00,0x36,0x83,0xf9,0x22,0x3e]
+// GFX89: v_and_b32_e32 v0, 0.15915494, v1 ; encoding: [0xf8,0x02,0x00,0x26]
+v_and_b32_e32 v0, INV2PI, v1
+
+// SICI: v_and_b32_e32 v0, 0x3e22f983, v1 ; encoding: [0xff,0x02,0x00,0x36,0x83,0xf9,0x22,0x3e]
+// GFX89: v_and_b32_e32 v0, 0.15915494, v1 ; encoding: [0xf8,0x02,0x00,0x26]
+v_and_b32_e32 v0, INV2PI64, v1
+
 // NOSICI: :[[@LINE+2]]:{{[0-9]+}}: error: literal operands are not supported
 // GFX89: v_and_b32_e64 v0, 0.15915494, v1 ; encoding: [0x00,0x00,0x13,0xd1,0xf8,0x02,0x02,0x00]
 v_and_b32_e64 v0, 0.159154943091895317852646485335, v1
 
+// NOSICI: :[[@LINE+2]]:{{[0-9]+}}: error: literal operands are not supported
+// GFX89: v_and_b32_e64 v0, 0.15915494, v1 ; encoding: [0x00,0x00,0x13,0xd1,0xf8,0x02,0x02,0x00]
+v_and_b32_e64 v0, INV2PI, v1
+
+// NOSICI: :[[@LINE+2]]:{{[0-9]+}}: error: literal operands are not supported
+// GFX89: v_and_b32_e64 v0, 0.15915494, v1 ; encoding: [0x00,0x00,0x13,0xd1,0xf8,0x02,0x02,0x00]
+v_and_b32_e64 v0, INV2PI64, v1
+
 // SICI: v_fract_f64_e32 v[0:1], 0x3fc45f30 ; encoding: [0xff,0x7c,0x00,0x7e,0x30,0x5f,0xc4,0x3f]
 // GFX89: v_fract_f64_e32 v[0:1], 0.15915494309189532 ; encoding: [0xf8,0x64,0x00,0x7e]
 v_fract_f64 v[0:1], 0.159154943091895317852646485335
 
+// SICI: v_fract_f64_e32 v[0:1], 0x3fc45f30 ; encoding: [0xff,0x7c,0x00,0x7e,0x30,0x5f,0xc4,0x3f]
+// GFX89: v_fract_f64_e32 v[0:1], 0.15915494309189532 ; encoding: [0xf8,0x64,0x00,0x7e]
+v_fract_f64 v[0:1], INV2PI64
+
 // 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, 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
+
 //---------------------------------------------------------------------------//
 // 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