Skip to content

Commit d14668f

Browse files
shiltianshiltian
committed
[RFC][WIP][AMDGPU] Use bf16 instead of i16 for bfloat
Currently it looks like we generally use `i16` to represent `bf16` in those tablegen files. I'm not sure of the reason behind it. My wild guess is the type `bf16` was not available when we enabled the support. This patch is trying to use `bf16` directly in those tablegen files, aiming at fixing #79369. Of course for #79369 a workaround can be to treat all `INT16` variants as `BFloat` in `getOpFltSemantics`, but it doesn't look good IMHO. Since I'm fairly new to AMDGPU backend, I'd appreciate it if you can point out where I don't understand correctly.
1 parent bdde5f9 commit d14668f

File tree

14 files changed

+191
-56
lines changed

14 files changed

+191
-56
lines changed

‎clang/lib/CodeGen/CGBuiltin.cpp

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -5908,8 +5908,6 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
59085908
}
59095909
}
59105910

5911-
assert(ArgValue->getType()->canLosslesslyBitCastTo(PTy) &&
5912-
"Must be able to losslessly bit cast to param");
59135911
// Cast vector type (e.g., v256i32) to x86_amx, this only happen
59145912
// in amx intrinsics.
59155913
if (PTy->isX86_AMXTy())
@@ -5939,8 +5937,6 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
59395937
}
59405938
}
59415939

5942-
assert(V->getType()->canLosslesslyBitCastTo(RetTy) &&
5943-
"Must be able to losslessly bit cast result type");
59445940
// Cast x86_amx to vector type (e.g., v256i32), this only happen
59455941
// in amx intrinsics.
59465942
if (V->getType()->isX86_AMXTy())

‎llvm/include/llvm/IR/IntrinsicsAMDGPU.td

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -2819,11 +2819,11 @@ def int_amdgcn_fdot2_f16_f16 :
28192819
def int_amdgcn_fdot2_bf16_bf16 :
28202820
ClangBuiltin<"__builtin_amdgcn_fdot2_bf16_bf16">,
28212821
DefaultAttrsIntrinsic<
2822-
[llvm_i16_ty], // %r
2822+
[llvm_bfloat_ty], // %r
28232823
[
2824-
llvm_v2i16_ty, // %a
2825-
llvm_v2i16_ty, // %b
2826-
llvm_i16_ty // %c
2824+
llvm_v2bf16_ty, // %a
2825+
llvm_v2bf16_ty, // %b
2826+
llvm_bfloat_ty // %c
28272827
],
28282828
[IntrNoMem, IntrSpeculatable]
28292829
>;

‎llvm/lib/CodeGen/GlobalISel/IRTranslator.cpp

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1562,8 +1562,9 @@ bool IRTranslator::translateBitCast(const User &U,
15621562

15631563
bool IRTranslator::translateCast(unsigned Opcode, const User &U,
15641564
MachineIRBuilder &MIRBuilder) {
1565-
if (U.getType()->getScalarType()->isBFloatTy() ||
1566-
U.getOperand(0)->getType()->getScalarType()->isBFloatTy())
1565+
if (Opcode != TargetOpcode::G_BITCAST &&
1566+
(U.getType()->getScalarType()->isBFloatTy() ||
1567+
U.getOperand(0)->getType()->getScalarType()->isBFloatTy()))
15671568
return false;
15681569
Register Op = getOrCreateVReg(*U.getOperand(0));
15691570
Register Res = getOrCreateVReg(U);

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

Lines changed: 66 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -475,6 +475,8 @@ class AMDGPUOperand : public MCParsedAsmOperand {
475475

476476
bool isSSrcF64() const { return isSCSrc_b64() || isLiteralImm(MVT::f64); }
477477

478+
bool isSSrc_bf16() const { return isSCSrcB16() || isLiteralImm(MVT::bf16); }
479+
478480
bool isSSrc_f16() const { return isSCSrcB16() || isLiteralImm(MVT::f16); }
479481

480482
bool isSSrcV2F16() const {
@@ -541,22 +543,40 @@ class AMDGPUOperand : public MCParsedAsmOperand {
541543
return isRegOrInlineNoMods(AMDGPU::VS_64RegClassID, MVT::f64);
542544
}
543545

546+
bool isVCSrcTBF16() const {
547+
return isRegOrInlineNoMods(AMDGPU::VS_16RegClassID, MVT::bf16);
548+
}
549+
544550
bool isVCSrcTF16() const {
545551
return isRegOrInlineNoMods(AMDGPU::VS_16RegClassID, MVT::f16);
546552
}
547553

554+
bool isVCSrcTBF16_Lo128() const {
555+
return isRegOrInlineNoMods(AMDGPU::VS_16_Lo128RegClassID, MVT::bf16);
556+
}
557+
548558
bool isVCSrcTF16_Lo128() const {
549559
return isRegOrInlineNoMods(AMDGPU::VS_16_Lo128RegClassID, MVT::f16);
550560
}
551561

562+
bool isVCSrcFake16BF16_Lo128() const {
563+
return isRegOrInlineNoMods(AMDGPU::VS_32_Lo128RegClassID, MVT::bf16);
564+
}
565+
552566
bool isVCSrcFake16F16_Lo128() const {
553567
return isRegOrInlineNoMods(AMDGPU::VS_32_Lo128RegClassID, MVT::f16);
554568
}
555569

570+
bool isVCSrc_bf16() const {
571+
return isRegOrInlineNoMods(AMDGPU::VS_32RegClassID, MVT::bf16);
572+
}
573+
556574
bool isVCSrc_f16() const {
557575
return isRegOrInlineNoMods(AMDGPU::VS_32RegClassID, MVT::f16);
558576
}
559577

578+
bool isVCSrc_v2bf16() const { return isVCSrc_bf16(); }
579+
560580
bool isVCSrc_v2f16() const { return isVCSrc_f16(); }
561581

562582
bool isVSrc_b32() const {
@@ -597,18 +617,34 @@ class AMDGPUOperand : public MCParsedAsmOperand {
597617

598618
bool isVSrc_f64() const { return isVCSrcF64() || isLiteralImm(MVT::f64); }
599619

620+
bool isVSrcT_bf16() const { return isVCSrcTBF16() || isLiteralImm(MVT::bf16); }
621+
600622
bool isVSrcT_f16() const { return isVCSrcTF16() || isLiteralImm(MVT::f16); }
601623

624+
bool isVSrcT_bf16_Lo128() const {
625+
return isVCSrcTBF16_Lo128() || isLiteralImm(MVT::bf16);
626+
}
627+
602628
bool isVSrcT_f16_Lo128() const {
603629
return isVCSrcTF16_Lo128() || isLiteralImm(MVT::f16);
604630
}
605631

632+
bool isVSrcFake16_bf16_Lo128() const {
633+
return isVCSrcFake16BF16_Lo128() || isLiteralImm(MVT::bf16);
634+
}
635+
606636
bool isVSrcFake16_f16_Lo128() const {
607637
return isVCSrcFake16F16_Lo128() || isLiteralImm(MVT::f16);
608638
}
609639

640+
bool isVSrc_bf16() const { return isVCSrc_bf16() || isLiteralImm(MVT::bf16); }
641+
610642
bool isVSrc_f16() const { return isVCSrc_f16() || isLiteralImm(MVT::f16); }
611643

644+
bool isVSrc_v2bf16() const {
645+
return isVSrc_bf16() || isLiteralImm(MVT::v2bf16);
646+
}
647+
612648
bool isVSrc_v2f16() const { return isVSrc_f16() || isLiteralImm(MVT::v2f16); }
613649

614650
bool isVISrcB32() const {
@@ -635,6 +671,10 @@ class AMDGPUOperand : public MCParsedAsmOperand {
635671
return isVISrcF16() || isVISrcB32();
636672
}
637673

674+
bool isVISrc_64_bf16() const {
675+
return isRegOrInlineNoMods(AMDGPU::VReg_64RegClassID, MVT::bf16);
676+
}
677+
638678
bool isVISrc_64_f16() const {
639679
return isRegOrInlineNoMods(AMDGPU::VReg_64RegClassID, MVT::f16);
640680
}
@@ -803,6 +843,10 @@ class AMDGPUOperand : public MCParsedAsmOperand {
803843
return isAISrc_128F16() || isAISrc_128_b32();
804844
}
805845

846+
bool isVISrc_128_bf16() const {
847+
return isRegOrInlineNoMods(AMDGPU::VReg_128RegClassID, MVT::bf16);
848+
}
849+
806850
bool isVISrc_128_f16() const {
807851
return isRegOrInlineNoMods(AMDGPU::VReg_128RegClassID, MVT::f16);
808852
}
@@ -1890,6 +1934,14 @@ static const fltSemantics *getOpFltSemantics(uint8_t OperandType) {
18901934
case AMDGPU::OPERAND_REG_IMM_V2FP16:
18911935
case AMDGPU::OPERAND_KIMM16:
18921936
return &APFloat::IEEEhalf();
1937+
case AMDGPU::OPERAND_REG_IMM_BF16:
1938+
case AMDGPU::OPERAND_REG_IMM_BF16_DEFERRED:
1939+
case AMDGPU::OPERAND_REG_INLINE_C_BF16:
1940+
case AMDGPU::OPERAND_REG_INLINE_C_V2BF16:
1941+
case AMDGPU::OPERAND_REG_INLINE_AC_BF16:
1942+
case AMDGPU::OPERAND_REG_INLINE_AC_V2BF16:
1943+
case AMDGPU::OPERAND_REG_IMM_V2BF16:
1944+
return &APFloat::BFloat();
18931945
default:
18941946
llvm_unreachable("unsupported fp type");
18951947
}
@@ -2186,17 +2238,24 @@ void AMDGPUOperand::addLiteralImmOperand(MCInst &Inst, int64_t Val, bool ApplyMo
21862238
case AMDGPU::OPERAND_REG_INLINE_AC_INT32:
21872239
case AMDGPU::OPERAND_REG_INLINE_AC_FP32:
21882240
case AMDGPU::OPERAND_REG_IMM_INT16:
2241+
case AMDGPU::OPERAND_REG_IMM_BF16:
21892242
case AMDGPU::OPERAND_REG_IMM_FP16:
2243+
case AMDGPU::OPERAND_REG_IMM_BF16_DEFERRED:
21902244
case AMDGPU::OPERAND_REG_IMM_FP16_DEFERRED:
21912245
case AMDGPU::OPERAND_REG_INLINE_C_INT16:
2246+
case AMDGPU::OPERAND_REG_INLINE_C_BF16:
21922247
case AMDGPU::OPERAND_REG_INLINE_C_FP16:
21932248
case AMDGPU::OPERAND_REG_INLINE_C_V2INT16:
2249+
case AMDGPU::OPERAND_REG_INLINE_C_V2BF16:
21942250
case AMDGPU::OPERAND_REG_INLINE_C_V2FP16:
21952251
case AMDGPU::OPERAND_REG_INLINE_AC_INT16:
2252+
case AMDGPU::OPERAND_REG_INLINE_AC_BF16:
21962253
case AMDGPU::OPERAND_REG_INLINE_AC_FP16:
21972254
case AMDGPU::OPERAND_REG_INLINE_AC_V2INT16:
2255+
case AMDGPU::OPERAND_REG_INLINE_AC_V2BF16:
21982256
case AMDGPU::OPERAND_REG_INLINE_AC_V2FP16:
21992257
case AMDGPU::OPERAND_REG_IMM_V2INT16:
2258+
case AMDGPU::OPERAND_REG_IMM_V2BF16:
22002259
case AMDGPU::OPERAND_REG_IMM_V2FP16:
22012260
case AMDGPU::OPERAND_REG_INLINE_C_V2FP32:
22022261
case AMDGPU::OPERAND_REG_IMM_V2FP32:
@@ -2240,6 +2299,7 @@ void AMDGPUOperand::addLiteralImmOperand(MCInst &Inst, int64_t Val, bool ApplyMo
22402299
case AMDGPU::OPERAND_REG_INLINE_AC_INT32:
22412300
case AMDGPU::OPERAND_REG_INLINE_AC_FP32:
22422301
case AMDGPU::OPERAND_REG_IMM_V2INT16:
2302+
case AMDGPU::OPERAND_REG_IMM_V2BF16:
22432303
case AMDGPU::OPERAND_REG_IMM_V2FP16:
22442304
case AMDGPU::OPERAND_REG_IMM_V2FP32:
22452305
case AMDGPU::OPERAND_REG_INLINE_C_V2FP32:
@@ -2277,11 +2337,15 @@ void AMDGPUOperand::addLiteralImmOperand(MCInst &Inst, int64_t Val, bool ApplyMo
22772337
return;
22782338

22792339
case AMDGPU::OPERAND_REG_IMM_INT16:
2340+
case AMDGPU::OPERAND_REG_IMM_BF16:
22802341
case AMDGPU::OPERAND_REG_IMM_FP16:
2342+
case AMDGPU::OPERAND_REG_IMM_BF16_DEFERRED:
22812343
case AMDGPU::OPERAND_REG_IMM_FP16_DEFERRED:
22822344
case AMDGPU::OPERAND_REG_INLINE_C_INT16:
2345+
case AMDGPU::OPERAND_REG_INLINE_C_BF16:
22832346
case AMDGPU::OPERAND_REG_INLINE_C_FP16:
22842347
case AMDGPU::OPERAND_REG_INLINE_AC_INT16:
2348+
case AMDGPU::OPERAND_REG_INLINE_AC_BF16:
22852349
case AMDGPU::OPERAND_REG_INLINE_AC_FP16:
22862350
if (isSafeTruncation(Val, 16) &&
22872351
AMDGPU::isInlinableLiteral16(static_cast<int16_t>(Val),
@@ -2296,8 +2360,10 @@ void AMDGPUOperand::addLiteralImmOperand(MCInst &Inst, int64_t Val, bool ApplyMo
22962360
return;
22972361

22982362
case AMDGPU::OPERAND_REG_INLINE_C_V2INT16:
2363+
case AMDGPU::OPERAND_REG_INLINE_C_V2BF16:
22992364
case AMDGPU::OPERAND_REG_INLINE_C_V2FP16:
23002365
case AMDGPU::OPERAND_REG_INLINE_AC_V2INT16:
2366+
case AMDGPU::OPERAND_REG_INLINE_AC_V2BF16:
23012367
case AMDGPU::OPERAND_REG_INLINE_AC_V2FP16: {
23022368
assert(isSafeTruncation(Val, 16));
23032369
assert(AMDGPU::isInlinableLiteral16(static_cast<int16_t>(Val),

‎llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUInstPrinter.cpp

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -521,8 +521,11 @@ void AMDGPUInstPrinter::printImmediateV216(uint32_t Imm, uint8_t OpType,
521521
if (printImmediateFloat32(Imm, STI, O))
522522
return;
523523
break;
524+
case AMDGPU::OPERAND_REG_IMM_V2BF16:
524525
case AMDGPU::OPERAND_REG_IMM_V2FP16:
526+
case AMDGPU::OPERAND_REG_INLINE_C_V2BF16:
525527
case AMDGPU::OPERAND_REG_INLINE_C_V2FP16:
528+
case AMDGPU::OPERAND_REG_INLINE_AC_V2BF16:
526529
case AMDGPU::OPERAND_REG_INLINE_AC_V2FP16:
527530
if (isUInt<16>(Imm) &&
528531
printImmediateFloat16(static_cast<uint16_t>(Imm), STI, O))
@@ -792,17 +795,24 @@ void AMDGPUInstPrinter::printRegularOperand(const MCInst *MI, unsigned OpNo,
792795
case AMDGPU::OPERAND_REG_IMM_INT16:
793796
printImmediateInt16(Op.getImm(), STI, O);
794797
break;
798+
case AMDGPU::OPERAND_REG_INLINE_C_BF16:
795799
case AMDGPU::OPERAND_REG_INLINE_C_FP16:
800+
case AMDGPU::OPERAND_REG_INLINE_AC_BF16:
796801
case AMDGPU::OPERAND_REG_INLINE_AC_FP16:
802+
case AMDGPU::OPERAND_REG_IMM_BF16:
797803
case AMDGPU::OPERAND_REG_IMM_FP16:
804+
case AMDGPU::OPERAND_REG_IMM_BF16_DEFERRED:
798805
case AMDGPU::OPERAND_REG_IMM_FP16_DEFERRED:
799806
printImmediate16(Op.getImm(), STI, O);
800807
break;
801808
case AMDGPU::OPERAND_REG_IMM_V2INT16:
809+
case AMDGPU::OPERAND_REG_IMM_V2BF16:
802810
case AMDGPU::OPERAND_REG_IMM_V2FP16:
803811
case AMDGPU::OPERAND_REG_INLINE_C_V2INT16:
804812
case AMDGPU::OPERAND_REG_INLINE_AC_V2INT16:
813+
case AMDGPU::OPERAND_REG_INLINE_C_V2BF16:
805814
case AMDGPU::OPERAND_REG_INLINE_C_V2FP16:
815+
case AMDGPU::OPERAND_REG_INLINE_AC_V2BF16:
806816
case AMDGPU::OPERAND_REG_INLINE_AC_V2FP16:
807817
printImmediateV216(Op.getImm(), OpTy, STI, O);
808818
break;

‎llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUMCCodeEmitter.cpp

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -276,9 +276,13 @@ AMDGPUMCCodeEmitter::getLitEncoding(const MCOperand &MO,
276276
case AMDGPU::OPERAND_REG_INLINE_C_INT16:
277277
case AMDGPU::OPERAND_REG_INLINE_AC_INT16:
278278
return getLit16IntEncoding(static_cast<uint16_t>(Imm), STI);
279+
case AMDGPU::OPERAND_REG_IMM_BF16:
279280
case AMDGPU::OPERAND_REG_IMM_FP16:
281+
case AMDGPU::OPERAND_REG_IMM_BF16_DEFERRED:
280282
case AMDGPU::OPERAND_REG_IMM_FP16_DEFERRED:
283+
case AMDGPU::OPERAND_REG_INLINE_C_BF16:
281284
case AMDGPU::OPERAND_REG_INLINE_C_FP16:
285+
case AMDGPU::OPERAND_REG_INLINE_AC_BF16:
282286
case AMDGPU::OPERAND_REG_INLINE_AC_FP16:
283287
// FIXME Is this correct? What do inline immediates do on SI for f16 src
284288
// which does not have f16 support?
@@ -288,8 +292,11 @@ AMDGPUMCCodeEmitter::getLitEncoding(const MCOperand &MO,
288292
case AMDGPU::OPERAND_REG_INLINE_AC_V2INT16:
289293
return AMDGPU::getInlineEncodingV2I16(static_cast<uint32_t>(Imm))
290294
.value_or(255);
295+
case AMDGPU::OPERAND_REG_IMM_V2BF16:
291296
case AMDGPU::OPERAND_REG_IMM_V2FP16:
297+
case AMDGPU::OPERAND_REG_INLINE_C_V2BF16:
292298
case AMDGPU::OPERAND_REG_INLINE_C_V2FP16:
299+
case AMDGPU::OPERAND_REG_INLINE_AC_V2BF16:
293300
case AMDGPU::OPERAND_REG_INLINE_AC_V2FP16:
294301
return AMDGPU::getInlineEncodingV2F16(static_cast<uint32_t>(Imm))
295302
.value_or(255);

‎llvm/lib/Target/AMDGPU/SIDefines.h

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -196,9 +196,12 @@ enum OperandType : unsigned {
196196
OPERAND_REG_IMM_INT16,
197197
OPERAND_REG_IMM_FP32,
198198
OPERAND_REG_IMM_FP64,
199+
OPERAND_REG_IMM_BF16,
199200
OPERAND_REG_IMM_FP16,
201+
OPERAND_REG_IMM_BF16_DEFERRED,
200202
OPERAND_REG_IMM_FP16_DEFERRED,
201203
OPERAND_REG_IMM_FP32_DEFERRED,
204+
OPERAND_REG_IMM_V2BF16,
202205
OPERAND_REG_IMM_V2FP16,
203206
OPERAND_REG_IMM_V2INT16,
204207
OPERAND_REG_IMM_V2INT32,
@@ -208,10 +211,12 @@ enum OperandType : unsigned {
208211
OPERAND_REG_INLINE_C_INT16,
209212
OPERAND_REG_INLINE_C_INT32,
210213
OPERAND_REG_INLINE_C_INT64,
214+
OPERAND_REG_INLINE_C_BF16,
211215
OPERAND_REG_INLINE_C_FP16,
212216
OPERAND_REG_INLINE_C_FP32,
213217
OPERAND_REG_INLINE_C_FP64,
214218
OPERAND_REG_INLINE_C_V2INT16,
219+
OPERAND_REG_INLINE_C_V2BF16,
215220
OPERAND_REG_INLINE_C_V2FP16,
216221
OPERAND_REG_INLINE_C_V2INT32,
217222
OPERAND_REG_INLINE_C_V2FP32,
@@ -226,10 +231,12 @@ enum OperandType : unsigned {
226231
/// Operands with an AccVGPR register or inline constant
227232
OPERAND_REG_INLINE_AC_INT16,
228233
OPERAND_REG_INLINE_AC_INT32,
234+
OPERAND_REG_INLINE_AC_BF16,
229235
OPERAND_REG_INLINE_AC_FP16,
230236
OPERAND_REG_INLINE_AC_FP32,
231237
OPERAND_REG_INLINE_AC_FP64,
232238
OPERAND_REG_INLINE_AC_V2INT16,
239+
OPERAND_REG_INLINE_AC_V2BF16,
233240
OPERAND_REG_INLINE_AC_V2FP16,
234241
OPERAND_REG_INLINE_AC_V2INT32,
235242
OPERAND_REG_INLINE_AC_V2FP32,

‎llvm/lib/Target/AMDGPU/SIInstrInfo.cpp

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4181,13 +4181,20 @@ bool SIInstrInfo::isInlineConstant(const MachineOperand &MO,
41814181
case AMDGPU::OPERAND_REG_INLINE_C_V2INT16:
41824182
case AMDGPU::OPERAND_REG_INLINE_AC_V2INT16:
41834183
return AMDGPU::isInlinableLiteralV2I16(Imm);
4184+
case AMDGPU::OPERAND_REG_IMM_V2BF16:
41844185
case AMDGPU::OPERAND_REG_IMM_V2FP16:
4186+
case AMDGPU::OPERAND_REG_INLINE_C_V2BF16:
41854187
case AMDGPU::OPERAND_REG_INLINE_C_V2FP16:
4188+
case AMDGPU::OPERAND_REG_INLINE_AC_V2BF16:
41864189
case AMDGPU::OPERAND_REG_INLINE_AC_V2FP16:
41874190
return AMDGPU::isInlinableLiteralV2F16(Imm);
4191+
case AMDGPU::OPERAND_REG_IMM_BF16:
41884192
case AMDGPU::OPERAND_REG_IMM_FP16:
4193+
case AMDGPU::OPERAND_REG_IMM_BF16_DEFERRED:
41894194
case AMDGPU::OPERAND_REG_IMM_FP16_DEFERRED:
4195+
case AMDGPU::OPERAND_REG_INLINE_C_BF16:
41904196
case AMDGPU::OPERAND_REG_INLINE_C_FP16:
4197+
case AMDGPU::OPERAND_REG_INLINE_AC_BF16:
41914198
case AMDGPU::OPERAND_REG_INLINE_AC_FP16: {
41924199
if (isInt<16>(Imm) || isUInt<16>(Imm)) {
41934200
// A few special case instructions have 16-bit operands on subtargets

0 commit comments

Comments
 (0)