[AMDGPU][MC] Adding symbolic name for the constant 1 over 2*pi #160617

jwanggit86 · 2025-09-24T23:40:56Z

Two symbolic names are added: INV2PI and INV2PI64, which can be used in assembly instrucitons to respectively represent 0.15915494 and 0.15915494309189532.

Fixes #38644.

Two symbolic names are added: INV2PI and INV2PI64, which can be used in assembly instrucitons to respectively represent 0.15915494 and 0.15915494309189532.

llvmbot · 2025-09-24T23:41:22Z

@llvm/pr-subscribers-backend-amdgpu

Author: Jun Wang (jwanggit86)

Changes

Two symbolic names are added: INV2PI and INV2PI64, which can be used in assembly instrucitons to respectively represent 0.15915494 and 0.15915494309189532.

Full diff: https://github.com/llvm/llvm-project/pull/160617.diff

13 Files Affected:

(modified) llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp (+14-2)
(modified) llvm/test/MC/AMDGPU/bf16_imm-fake16.s (+7)
(modified) llvm/test/MC/AMDGPU/bf16_imm.s (+3)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_valu_lit64.s (+7)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3.s (+6)
(modified) llvm/test/MC/AMDGPU/gfx7_err_pos.s (+10)
(modified) llvm/test/MC/AMDGPU/gfx950_dlops.s (+3)
(modified) llvm/test/MC/AMDGPU/literal16.s (+6)
(modified) llvm/test/MC/AMDGPU/literals.s (+36)
(modified) llvm/test/MC/AMDGPU/literalv216-err.s (+6)
(modified) llvm/test/MC/AMDGPU/literalv216.s (+4)
(modified) llvm/test/MC/AMDGPU/mai-gfx950-err.s (+6)
(modified) llvm/test/MC/AMDGPU/mubuf.s (+4)

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

jwanggit86 · 2025-09-24T23:43:30Z

For issue 38644.

arsenm

Does SP3 support this?

arsenm · 2025-09-25T01:19:52Z

llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

+  if (!Tok.is(AsmToken::Identifier))
+    return false;
+  StringRef Str = Tok.getIdentifier();
+  return (Str.str() == "INV2PI" || Str.str() == "INV2PI64");


Suggested change

return (Str.str() == "INV2PI" || Str.str() == "INV2PI64");

return Str == "INV2PI" || Str. == "INV2PI64";

Don't need all these conversions to std::string

arsenm · 2025-09-25T01:20:05Z

llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

    // optional sign.

    StringRef Num = getTokenStr();
+    if (Num.str() == "INV2PI")


Suggested change

if (Num.str() == "INV2PI")

if (Num == "INV2PI")

arsenm · 2025-09-25T01:20:12Z

llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

    StringRef Num = getTokenStr();
+    if (Num.str() == "INV2PI")
+      Num = "0.15915494";
+    else if (Num.str() == "INV2PI64")


Suggested change

else if (Num.str() == "INV2PI64")

else if (Num == "INV2PI64")

rampitec · 2025-09-25T05:56:52Z

This is also yet another thing which does not roundtrip on disasm.

rampitec · 2025-09-25T06:04:32Z

This is inline literal. If we really want it, it probably needs to be MCExpr.

rampitec · 2025-09-25T06:30:59Z

And then conceptually this is a fundamental constant. It does not change depending on the number of bits, only how we round it changes. So, there shall be no 2 separate symbolic names, we always know the operand size from the mnemonic. This is not really a some number which we can represent as a decimal fraction, but really an inline idiom. Basically, it means it then shall accept INV2PI in any context and assemble it to a corresponding inline literal constant. Ideally you do not even have a number anywhere. And then ideally lit(INV2PI) or lit64(INV2PI) shall create a real literal instead of inline for that, that is where you need to know the number, but we are parsecs away from it.

jayfoad · 2025-09-25T08:48:10Z

Does SP3 support this?

I don't think SP3 has anything like this. Personally I don't think it is a good idea.

kosarev · 2025-09-25T14:28:24Z

For what the change does, it seems a couple of createConstantSymbol() calls would suffice.

But as Stanislav said, having a proper universal designation for the inline constant looks more appealing.

[AMDGPU][MC] Adding symbolic name for the constant 1 over 2*pi

98923c5

Two symbolic names are added: INV2PI and INV2PI64, which can be used in assembly instrucitons to respectively represent 0.15915494 and 0.15915494309189532.

jwanggit86 requested review from Sisyph and rampitec September 24, 2025 23:40

jwanggit86 added backend:AMDGPU mc labels Sep 24, 2025

arsenm reviewed Sep 25, 2025

View reviewed changes

arsenm requested review from jayfoad and kosarev September 25, 2025 01:21

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

[AMDGPU][MC] Adding symbolic name for the constant 1 over 2*pi #160617

[AMDGPU][MC] Adding symbolic name for the constant 1 over 2*pi #160617

Uh oh!

jwanggit86 commented Sep 24, 2025 •

edited by arsenm

Loading

Uh oh!

llvmbot commented Sep 24, 2025

Uh oh!

jwanggit86 commented Sep 24, 2025 •

edited

Loading

Uh oh!

arsenm left a comment

Uh oh!

arsenm Sep 25, 2025

Uh oh!

arsenm Sep 25, 2025

Uh oh!

arsenm Sep 25, 2025

Uh oh!

rampitec commented Sep 25, 2025

Uh oh!

rampitec commented Sep 25, 2025

Uh oh!

rampitec commented Sep 25, 2025

Uh oh!

jayfoad commented Sep 25, 2025

Uh oh!

kosarev commented Sep 25, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants

	return (Str.str() == "INV2PI" \|\| Str.str() == "INV2PI64");
	return Str == "INV2PI" \|\| Str. == "INV2PI64";

	else if (Num.str() == "INV2PI64")
	else if (Num == "INV2PI64")

[AMDGPU][MC] Adding symbolic name for the constant 1 over 2*pi #160617

Are you sure you want to change the base?

[AMDGPU][MC] Adding symbolic name for the constant 1 over 2*pi #160617

Uh oh!

Conversation

jwanggit86 commented Sep 24, 2025 • edited by arsenm Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

llvmbot commented Sep 24, 2025

Uh oh!

jwanggit86 commented Sep 24, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

arsenm left a comment

Choose a reason for hiding this comment

Uh oh!

arsenm Sep 25, 2025

Choose a reason for hiding this comment

Uh oh!

arsenm Sep 25, 2025

Choose a reason for hiding this comment

Uh oh!

arsenm Sep 25, 2025

Choose a reason for hiding this comment

Uh oh!

rampitec commented Sep 25, 2025

Uh oh!

rampitec commented Sep 25, 2025

Uh oh!

rampitec commented Sep 25, 2025

Uh oh!

jayfoad commented Sep 25, 2025

Uh oh!

kosarev commented Sep 25, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants

jwanggit86 commented Sep 24, 2025 •

edited by arsenm

Loading

jwanggit86 commented Sep 24, 2025 •

edited

Loading