Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 1 commit into
base: main
Choose a base branch
from
Open

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
16 changes: 14 additions & 2 deletions llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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");
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

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

}

ParseStatus AMDGPUAsmParser::parseImm(OperandVector &Operands,
bool HasSP3AbsModifier, bool HasLit,
bool HasLit64) {
Expand All @@ -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;
Expand All @@ -3210,6 +3218,10 @@ ParseStatus AMDGPUAsmParser::parseImm(OperandVector &Operands,
// optional sign.

StringRef Num = getTokenStr();
if (Num.str() == "INV2PI")
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

Num = "0.15915494";
else if (Num.str() == "INV2PI64")
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

Num = "0.15915494309189532";
lex();

APFloat RealVal(APFloat::IEEEdouble());
Expand Down
7 changes: 7 additions & 0 deletions llvm/test/MC/AMDGPU/bf16_imm-fake16.s
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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]

Expand Down
3 changes: 3 additions & 0 deletions llvm/test/MC/AMDGPU/bf16_imm.s
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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]

Expand Down
7 changes: 7 additions & 0 deletions llvm/test/MC/AMDGPU/gfx1250_asm_valu_lit64.s
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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]

Expand Down
6 changes: 6 additions & 0 deletions llvm/test/MC/AMDGPU/gfx1250_asm_vop3.s
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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]

Expand Down Expand Up @@ -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]

Expand Down
10 changes: 10 additions & 0 deletions llvm/test/MC/AMDGPU/gfx7_err_pos.s
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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

Expand Down
3 changes: 3 additions & 0 deletions llvm/test/MC/AMDGPU/gfx950_dlops.s
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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]

Expand Down
6 changes: 6 additions & 0 deletions llvm/test/MC/AMDGPU/literal16.s
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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]

Expand Down
36 changes: 36 additions & 0 deletions llvm/test/MC/AMDGPU/literals.s
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
//---------------------------------------------------------------------------//
Expand Down
6 changes: 6 additions & 0 deletions llvm/test/MC/AMDGPU/literalv216-err.s
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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

Expand Down
4 changes: 4 additions & 0 deletions llvm/test/MC/AMDGPU/literalv216.s
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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]
Expand Down
6 changes: 6 additions & 0 deletions llvm/test/MC/AMDGPU/mai-gfx950-err.s
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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

Expand All @@ -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
4 changes: 4 additions & 0 deletions llvm/test/MC/AMDGPU/mubuf.s
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down