Skip to content

Commit 1488b4e

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 dd1897c commit 1488b4e

File tree

17 files changed

+379
-56
lines changed

17 files changed

+379
-56
lines changed

‎clang/lib/CodeGen/CGBuiltin.cpp‎

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -5912,8 +5912,6 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
59125912
}
59135913
}
59145914

5915-
assert(ArgValue->getType()->canLosslesslyBitCastTo(PTy) &&
5916-
"Must be able to losslessly bit cast to param");
59175915
// Cast vector type (e.g., v256i32) to x86_amx, this only happen
59185916
// in amx intrinsics.
59195917
if (PTy->isX86_AMXTy())
@@ -5943,8 +5941,6 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
59435941
}
59445942
}
59455943

5946-
assert(V->getType()->canLosslesslyBitCastTo(RetTy) &&
5947-
"Must be able to losslessly bit cast result type");
59485944
// Cast x86_amx to vector type (e.g., v256i32), this only happen
59495945
// in amx intrinsics.
59505946
if (V->getType()->isX86_AMXTy())

‎clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-gfx11.cl‎

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -11,7 +11,10 @@ typedef unsigned short __attribute__((ext_vector_type(2))) ushort2;
1111
// CHECK: call float @llvm.amdgcn.fdot2(<2 x half> %v2hA, <2 x half> %v2hB, float %fC, i1 false)
1212
// CHECK: call float @llvm.amdgcn.fdot2(<2 x half> %v2hA, <2 x half> %v2hB, float %fC, i1 true)
1313
// CHECK: call half @llvm.amdgcn.fdot2.f16.f16(<2 x half> %v2hA, <2 x half> %v2hB, half %hC)
14-
// CHECK: call i16 @llvm.amdgcn.fdot2.bf16.bf16(<2 x i16> %v2ssA, <2 x i16> %v2ssB, i16 %sC)
14+
// CHECK: [[s1:%[0-9]+]] = bitcast <2 x i16> %v2ssA to <2 x bfloat>
15+
// CHECK-NEXT: [[s2:%[0-9]+]] = bitcast <2 x i16> %v2ssB to <2 x bfloat>
16+
// CHECK-NEXT: [[s3:%[0-9]+]] = bitcast i16 %sC to bfloat
17+
// CHECK-NEXT: [[d:%[0-9]+]] = tail call bfloat @llvm.amdgcn.fdot2.bf16.bf16(<2 x bfloat> [[s1]], <2 x bfloat> [[s2]], bfloat [[s3]])
1518
// CHECK: call float @llvm.amdgcn.fdot2.f32.bf16(<2 x i16> %v2ssA, <2 x i16> %v2ssB, float %fC, i1 false)
1619
// CHECK: call float @llvm.amdgcn.fdot2.f32.bf16(<2 x i16> %v2ssA, <2 x i16> %v2ssB, float %fC, i1 true)
1720
// CHECK: call i32 @llvm.amdgcn.udot4(i32 %uiA, i32 %uiB, i32 %uiC, i1 false)

‎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/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp‎

Lines changed: 92 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:
@@ -2295,6 +2355,22 @@ void AMDGPUOperand::addLiteralImmOperand(MCInst &Inst, int64_t Val, bool ApplyMo
22952355
setImmKindLiteral();
22962356
return;
22972357

2358+
case AMDGPU::OPERAND_REG_IMM_BF16:
2359+
case AMDGPU::OPERAND_REG_IMM_BF16_DEFERRED:
2360+
case AMDGPU::OPERAND_REG_INLINE_C_BF16:
2361+
case AMDGPU::OPERAND_REG_INLINE_AC_BF16:
2362+
if (isSafeTruncation(Val, 16) &&
2363+
AMDGPU::isInlinableLiteralBF16(static_cast<int16_t>(Val),
2364+
AsmParser->hasInv2PiInlineImm())) {
2365+
Inst.addOperand(MCOperand::createImm(Val));
2366+
setImmKindConst();
2367+
return;
2368+
}
2369+
2370+
Inst.addOperand(MCOperand::createImm(Val & 0xffff));
2371+
setImmKindLiteral();
2372+
return;
2373+
22982374
case AMDGPU::OPERAND_REG_INLINE_C_V2INT16:
22992375
case AMDGPU::OPERAND_REG_INLINE_C_V2FP16:
23002376
case AMDGPU::OPERAND_REG_INLINE_AC_V2INT16:
@@ -2306,6 +2382,17 @@ void AMDGPUOperand::addLiteralImmOperand(MCInst &Inst, int64_t Val, bool ApplyMo
23062382
Inst.addOperand(MCOperand::createImm(Val));
23072383
return;
23082384
}
2385+
2386+
case AMDGPU::OPERAND_REG_INLINE_C_V2BF16:
2387+
case AMDGPU::OPERAND_REG_INLINE_AC_V2BF16: {
2388+
assert(isSafeTruncation(Val, 16));
2389+
assert(AMDGPU::isInlinableLiteralBF16(static_cast<int16_t>(Val),
2390+
AsmParser->hasInv2PiInlineImm()));
2391+
2392+
Inst.addOperand(MCOperand::createImm(Val));
2393+
return;
2394+
}
2395+
23092396
case AMDGPU::OPERAND_KIMM32:
23102397
Inst.addOperand(MCOperand::createImm(Literal.getLoBits(32).getZExtValue()));
23112398
setImmKindMandatoryLiteral();
@@ -3429,6 +3516,11 @@ bool AMDGPUAsmParser::isInlineConstant(const MCInst &Inst,
34293516
OperandType == AMDGPU::OPERAND_REG_IMM_V2FP16)
34303517
return AMDGPU::isInlinableLiteralV2F16(Val);
34313518

3519+
if (OperandType == AMDGPU::OPERAND_REG_INLINE_C_V2BF16 ||
3520+
OperandType == AMDGPU::OPERAND_REG_INLINE_AC_V2BF16 ||
3521+
OperandType == AMDGPU::OPERAND_REG_IMM_V2BF16)
3522+
return AMDGPU::isInlinableLiteralV2BF16(Val);
3523+
34323524
return AMDGPU::isInlinableLiteral16(Val, hasInv2PiInlineImm());
34333525
}
34343526
default:

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

Lines changed: 57 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -488,6 +488,47 @@ static bool printImmediateFloat16(uint32_t Imm, const MCSubtargetInfo &STI,
488488
return true;
489489
}
490490

491+
static bool printImmediateBFloat16(uint32_t Imm, const MCSubtargetInfo &STI,
492+
raw_ostream &O) {
493+
if (Imm == 0x3F80)
494+
O << "1.0";
495+
else if (Imm == 0xBF80)
496+
O << "-1.0";
497+
else if (Imm == 0x3F00)
498+
O << "0.5";
499+
else if (Imm == 0xBF00)
500+
O << "-0.5";
501+
else if (Imm == 0x4000)
502+
O << "2.0";
503+
else if (Imm == 0xC000)
504+
O << "-2.0";
505+
else if (Imm == 0x4080)
506+
O << "4.0";
507+
else if (Imm == 0xC080)
508+
O << "-4.0";
509+
else if (Imm == 0x3E22 && STI.hasFeature(AMDGPU::FeatureInv2PiInlineImm))
510+
O << "0.15915494";
511+
else
512+
return false;
513+
514+
return true;
515+
}
516+
517+
void AMDGPUInstPrinter::printImmediateBF16(uint32_t Imm,
518+
const MCSubtargetInfo &STI,
519+
raw_ostream &O) {
520+
int16_t SImm = static_cast<int16_t>(Imm);
521+
if (isInlinableIntLiteral(SImm)) {
522+
O << SImm;
523+
return;
524+
}
525+
526+
if (printImmediateBFloat16(static_cast<uint16_t>(Imm), STI, O))
527+
return;
528+
529+
O << formatHex(static_cast<uint64_t>(Imm));
530+
}
531+
491532
void AMDGPUInstPrinter::printImmediate16(uint32_t Imm,
492533
const MCSubtargetInfo &STI,
493534
raw_ostream &O) {
@@ -528,6 +569,13 @@ void AMDGPUInstPrinter::printImmediateV216(uint32_t Imm, uint8_t OpType,
528569
printImmediateFloat16(static_cast<uint16_t>(Imm), STI, O))
529570
return;
530571
break;
572+
case AMDGPU::OPERAND_REG_IMM_V2BF16:
573+
case AMDGPU::OPERAND_REG_INLINE_C_V2BF16:
574+
case AMDGPU::OPERAND_REG_INLINE_AC_V2BF16:
575+
if (isUInt<16>(Imm) &&
576+
printImmediateBFloat16(static_cast<uint16_t>(Imm), STI, O))
577+
return;
578+
break;
531579
default:
532580
llvm_unreachable("bad operand type");
533581
}
@@ -799,11 +847,20 @@ void AMDGPUInstPrinter::printRegularOperand(const MCInst *MI, unsigned OpNo,
799847
case AMDGPU::OPERAND_REG_IMM_FP16_DEFERRED:
800848
printImmediate16(Op.getImm(), STI, O);
801849
break;
850+
case AMDGPU::OPERAND_REG_INLINE_C_BF16:
851+
case AMDGPU::OPERAND_REG_INLINE_AC_BF16:
852+
case AMDGPU::OPERAND_REG_IMM_BF16:
853+
case AMDGPU::OPERAND_REG_IMM_BF16_DEFERRED:
854+
printImmediateBF16(Op.getImm(), STI, O);
855+
break;
802856
case AMDGPU::OPERAND_REG_IMM_V2INT16:
857+
case AMDGPU::OPERAND_REG_IMM_V2BF16:
803858
case AMDGPU::OPERAND_REG_IMM_V2FP16:
804859
case AMDGPU::OPERAND_REG_INLINE_C_V2INT16:
805860
case AMDGPU::OPERAND_REG_INLINE_AC_V2INT16:
861+
case AMDGPU::OPERAND_REG_INLINE_C_V2BF16:
806862
case AMDGPU::OPERAND_REG_INLINE_C_V2FP16:
863+
case AMDGPU::OPERAND_REG_INLINE_AC_V2BF16:
807864
case AMDGPU::OPERAND_REG_INLINE_AC_V2FP16:
808865
printImmediateV216(Op.getImm(), OpTy, STI, O);
809866
break;

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

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -88,6 +88,8 @@ class AMDGPUInstPrinter : public MCInstPrinter {
8888
raw_ostream &O);
8989
void printImmediate16(uint32_t Imm, const MCSubtargetInfo &STI,
9090
raw_ostream &O);
91+
void printImmediateBF16(uint32_t Imm, const MCSubtargetInfo &STI,
92+
raw_ostream &O);
9193
void printImmediateV216(uint32_t Imm, uint8_t OpType,
9294
const MCSubtargetInfo &STI, raw_ostream &O);
9395
bool printImmediateFloat32(uint32_t Imm, const MCSubtargetInfo &STI,

0 commit comments

Comments
 (0)