[AMDGPU] Implement 'llvm.get.fpenv' and 'llvm.set.fpenv' #83906
Merged
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
jhuber6
requested review from
arsenm,
jayfoad,
JonChesterfield,
kzhuravl,
rampitec and
shiltian
llvmbot
added
clang
|
@llvm/pr-subscribers-clang @llvm/pr-subscribers-backend-amdgpu Author: Joseph Huber (jhuber6) ChangesSummary: Patch is 20.08 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/83906.diff 13 Files Affected:
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index 213311b96df74f..6628e8f265fe48 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -325,6 +325,9 @@ BUILTIN(__builtin_amdgcn_read_exec_hi, "Ui", "nc")
BUILTIN(__builtin_amdgcn_endpgm, "v", "nr")
+BUILTIN(__builtin_amdgcn_get_fpenv, "WUi", "n")
+BUILTIN(__builtin_amdgcn_set_fpenv, "vWUi", "n")
+
//===----------------------------------------------------------------------===//
// R600-NI only builtins.
//===----------------------------------------------------------------------===//
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 9ee51ca7142c77..63a0ac9a0afe4a 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -18406,6 +18406,17 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
CGM.getIntrinsic(Intrinsic::amdgcn_global_load_tr, {ArgTy});
return Builder.CreateCall(F, {Addr});
}
+ case AMDGPU::BI__builtin_amdgcn_get_fpenv: {
+ Function *F = CGM.getIntrinsic(Intrinsic::get_fpenv,
+ {llvm::Type::getInt64Ty(getLLVMContext())});
+ return Builder.CreateCall(F);
+ }
+ case AMDGPU::BI__builtin_amdgcn_set_fpenv: {
+ Function *F = CGM.getIntrinsic(Intrinsic::set_fpenv,
+ {llvm::Type::getInt64Ty(getLLVMContext())});
+ llvm::Value *Env = EmitScalarExpr(E->getArg(0));
+ return Builder.CreateCall(F, {Env});
+ }
case AMDGPU::BI__builtin_amdgcn_read_exec:
return EmitAMDGCNBallotForExec(*this, E, Int64Ty, Int64Ty, false);
case AMDGPU::BI__builtin_amdgcn_read_exec_lo:
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
index 7d9010ee9067d6..8a4533633706b2 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -839,6 +839,18 @@ unsigned test_wavefrontsize() {
return __builtin_amdgcn_wavefrontsize();
}
+// CHECK-LABEL test_get_fpenv(
+unsigned long test_get_fpenv() {
+ // CHECK: call i64 @llvm.get.fpenv.i64()
+ return __builtin_amdgcn_get_fpenv();
+}
+
+// CHECK-LABEL test_set_fpenv(
+void test_set_fpenv(unsigned long env) {
+ // CHECK: call void @llvm.set.fpenv.i64(i64 %[[ENV:.+]])
+ __builtin_amdgcn_set_fpenv(env);
+}
+
// CHECK-DAG: [[$WI_RANGE]] = !{i32 0, i32 1024}
// CHECK-DAG: [[$WS_RANGE]] = !{i16 1, i16 1025}
// CHECK-DAG: attributes #[[$NOUNWIND_READONLY]] = { convergent mustprogress nocallback nofree nounwind willreturn memory(none) }
diff --git a/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst
index 0c588c84958624..ed41be4b08dee8 100644
--- a/llvm/docs/AMDGPUUsage.rst
+++ b/llvm/docs/AMDGPUUsage.rst
@@ -1151,6 +1151,13 @@ The AMDGPU backend implements the following LLVM IR intrinsics.
register do not exactly match the FLT_ROUNDS values,
so a conversion is performed.
+ :ref:`llvm.get.fpenv<int_get_fpenv>` Returns the current value of the AMDGPU floating point environment.
+ This stores information related to the current rounding mode,
+ denormalization mode, enabled traps, and floating point exceptions.
+ The format is a 64-bit concatenation of the MODE and TRAPSTS registers.
+
+ :ref:`llvm.set.fpenv<int_set_fpenv>` Sets the floating point environment to the specifies state.
+
llvm.amdgcn.wave.reduce.umin Performs an arithmetic unsigned min reduction on the unsigned values
provided by each lane in the wavefront.
Intrinsic takes a hint for reduction strategy using second operand
diff --git a/llvm/docs/LangRef.rst b/llvm/docs/LangRef.rst
index f56d4ed28f2855..3926cdedd34c5f 100644
--- a/llvm/docs/LangRef.rst
+++ b/llvm/docs/LangRef.rst
@@ -26427,6 +26427,7 @@ similar to C library function 'fesetround', however this intrinsic does not
return any value and uses platform-independent representation of IEEE rounding
modes.
+.. _int_get_fpenv:
'``llvm.get.fpenv``' Intrinsic
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -26450,6 +26451,7 @@ Semantics:
The '``llvm.get.fpenv``' intrinsic reads the current floating-point environment
and returns it as an integer value.
+.. _int_set_fpenv:
'``llvm.set.fpenv``' Intrinsic
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
diff --git a/llvm/docs/ReleaseNotes.rst b/llvm/docs/ReleaseNotes.rst
index 8ce6ee5cebb266..56abe7fe40270f 100644
--- a/llvm/docs/ReleaseNotes.rst
+++ b/llvm/docs/ReleaseNotes.rst
@@ -70,6 +70,8 @@ Changes to the AArch64 Backend
Changes to the AMDGPU Backend
-----------------------------
+* Implemented the ``llvm.get.fpenv`` and ``llvm.set.fpenv`` intrinsics.
+
Changes to the ARM Backend
--------------------------
diff --git a/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp
index 4c3b983f2960df..deef6611c18438 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp
@@ -905,6 +905,11 @@ AMDGPULegalizerInfo::AMDGPULegalizerInfo(const GCNSubtarget &ST_,
getActionDefinitionsBuilder(G_STACKRESTORE)
.legalFor({PrivatePtr});
+ getActionDefinitionsBuilder(G_GET_FPENV)
+ .customFor({S64});
+ getActionDefinitionsBuilder(G_SET_FPENV)
+ .customFor({S64});
+
getActionDefinitionsBuilder(G_GLOBAL_VALUE)
.customIf(typeIsNot(0, PrivatePtr));
@@ -2128,6 +2133,10 @@ bool AMDGPULegalizerInfo::legalizeCustom(
return legalizeFPTruncRound(MI, B);
case TargetOpcode::G_STACKSAVE:
return legalizeStackSave(MI, B);
+ case TargetOpcode::G_GET_FPENV:
+ return legalizeGetFPEnv(MI, B);
+ case TargetOpcode::G_SET_FPENV:
+ return legalizeSetFPEnv(MI, B);
default:
return false;
}
@@ -6940,6 +6949,42 @@ bool AMDGPULegalizerInfo::legalizeWaveID(MachineInstr &MI,
return true;
}
+static constexpr unsigned FPEnvModeBitField =
+ AMDGPU::Hwreg::HwregEncoding::encode(AMDGPU::Hwreg::ID_MODE, 0, 23);
+
+static constexpr unsigned FPEnvTrapBitField =
+ AMDGPU::Hwreg::HwregEncoding::encode(AMDGPU::Hwreg::ID_TRAPSTS, 0, 5);
+
+bool AMDGPULegalizerInfo::legalizeGetFPEnv(MachineInstr &MI,
+ MachineIRBuilder &B) const {
+ auto ModeReg =
+ B.buildIntrinsic(Intrinsic::amdgcn_s_getreg, {S32},
+ /*HasSideEffects=*/true, /*isConvergent=*/false)
+ .addImm(FPEnvModeBitField);
+ auto TrapReg =
+ B.buildIntrinsic(Intrinsic::amdgcn_s_getreg, {S32},
+ /*HasSideEffects=*/true, /*isConvergent=*/false)
+ .addImm(FPEnvTrapBitField);
+ B.buildMergeLikeInstr(MI.getOperand(0).getReg(), {ModeReg, TrapReg});
+ MI.eraseFromParent();
+ return true;
+}
+
+bool AMDGPULegalizerInfo::legalizeSetFPEnv(MachineInstr &MI,
+ MachineIRBuilder &B) const {
+ auto Unmerge = B.buildUnmerge({S32, S32}, MI.getOperand(0));
+ B.buildIntrinsic(Intrinsic::amdgcn_s_setreg, ArrayRef<DstOp>(),
+ /*HasSideEffects=*/true, /*isConvergent=*/false)
+ .addImm(FPEnvModeBitField)
+ .addReg(Unmerge.getReg(0));
+ B.buildIntrinsic(Intrinsic::amdgcn_s_setreg, ArrayRef<DstOp>(),
+ /*HasSideEffects=*/true, /*isConvergent=*/false)
+ .addImm(FPEnvTrapBitField)
+ .addReg(Unmerge.getReg(1));
+ MI.eraseFromParent();
+ return true;
+}
+
bool AMDGPULegalizerInfo::legalizeIntrinsic(LegalizerHelper &Helper,
MachineInstr &MI) const {
MachineIRBuilder &B = Helper.MIRBuilder;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.h b/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.h
index ecbe42681c6690..62b6fa12eb8751 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.h
@@ -214,6 +214,9 @@ class AMDGPULegalizerInfo final : public LegalizerInfo {
bool legalizeStackSave(MachineInstr &MI, MachineIRBuilder &B) const;
bool legalizeWaveID(MachineInstr &MI, MachineIRBuilder &B) const;
+ bool legalizeGetFPEnv(MachineInstr &MI, MachineIRBuilder &B) const;
+ bool legalizeSetFPEnv(MachineInstr &MI, MachineIRBuilder &B) const;
+
bool legalizeImageIntrinsic(
MachineInstr &MI, MachineIRBuilder &B,
GISelChangeObserver &Observer,
diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
index 30a65bb332652e..c45a65f451cd06 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -876,6 +876,8 @@ SITargetLowering::SITargetLowering(const TargetMachine &TM,
setOperationAction(ISD::STACKSAVE, MVT::Other, Custom);
setOperationAction(ISD::GET_ROUNDING, MVT::i32, Custom);
+ setOperationAction(ISD::GET_FPENV, MVT::i64, Custom);
+ setOperationAction(ISD::SET_FPENV, MVT::i64, Custom);
// TODO: Could move this to custom lowering, could benefit from combines on
// extract of relevant bits.
@@ -4079,6 +4081,70 @@ SDValue SITargetLowering::lowerFP_EXTEND(SDValue Op, SelectionDAG &DAG) const {
return DAG.getNode(ISD::BF16_TO_FP, SL, DstVT, BitCast);
}
+SDValue SITargetLowering::lowerGET_FPENV(SDValue Op, SelectionDAG &DAG) const {
+ SDLoc SL(Op);
+ assert(Op.getValueType() == MVT::i64);
+
+ uint32_t ModeHwReg =
+ AMDGPU::Hwreg::HwregEncoding::encode(AMDGPU::Hwreg::ID_MODE, 0, 23);
+ SDValue ModeHwRegImm = DAG.getTargetConstant(ModeHwReg, SL, MVT::i32);
+ uint32_t TrapHwReg =
+ AMDGPU::Hwreg::HwregEncoding::encode(AMDGPU::Hwreg::ID_TRAPSTS, 0, 5);
+ SDValue TrapHwRegImm = DAG.getTargetConstant(TrapHwReg, SL, MVT::i32);
+
+ SDVTList VTList = DAG.getVTList(MVT::i32, MVT::Other);
+ SDValue IntrinID =
+ DAG.getTargetConstant(Intrinsic::amdgcn_s_getreg, SL, MVT::i32);
+ SDValue GetModeReg = DAG.getNode(ISD::INTRINSIC_W_CHAIN, SL, VTList,
+ Op.getOperand(0), IntrinID, ModeHwRegImm);
+ SDValue GetTrapReg = DAG.getNode(ISD::INTRINSIC_W_CHAIN, SL, VTList,
+ Op.getOperand(0), IntrinID, TrapHwRegImm);
+ SDValue MergedReg =
+ DAG.getMergeValues({GetModeReg.getValue(1), GetTrapReg.getValue(1)}, SL);
+
+ SDValue CvtPtr =
+ DAG.getNode(ISD::BUILD_VECTOR, SL, MVT::v2i32, GetModeReg, GetTrapReg);
+ SDValue Result = DAG.getNode(ISD::BITCAST, SL, MVT::i64, CvtPtr);
+
+ return DAG.getMergeValues({Result, MergedReg}, SL);
+}
+
+SDValue SITargetLowering::lowerSET_FPENV(SDValue Op, SelectionDAG &DAG) const {
+ SDLoc SL(Op);
+ assert(Op.getOperand(1).getValueType() == MVT::i64);
+
+ SDValue Input = DAG.getNode(ISD::BITCAST, SL, MVT::v2i32, Op.getOperand(1));
+ SDValue NewModeReg = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, SL, MVT::i32, Input,
+ DAG.getConstant(0, SL, MVT::i32));
+ SDValue NewTrapReg = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, SL, MVT::i32, Input,
+ DAG.getConstant(1, SL, MVT::i32));
+
+ SDValue ReadFirstLaneID =
+ DAG.getTargetConstant(Intrinsic::amdgcn_readfirstlane, SL, MVT::i32);
+ NewModeReg = DAG.getNode(ISD::INTRINSIC_WO_CHAIN, SL, MVT::i32,
+ ReadFirstLaneID, NewModeReg);
+ NewTrapReg = DAG.getNode(ISD::INTRINSIC_WO_CHAIN, SL, MVT::i32,
+ ReadFirstLaneID, NewTrapReg);
+
+ unsigned ModeHwReg =
+ AMDGPU::Hwreg::HwregEncoding::encode(AMDGPU::Hwreg::ID_MODE, 0, 23);
+ SDValue ModeHwRegImm = DAG.getTargetConstant(ModeHwReg, SL, MVT::i32);
+ unsigned TrapHwReg =
+ AMDGPU::Hwreg::HwregEncoding::encode(AMDGPU::Hwreg::ID_TRAPSTS, 0, 5);
+ SDValue TrapHwRegImm = DAG.getTargetConstant(TrapHwReg, SL, MVT::i32);
+
+ SDVTList VTList = DAG.getVTList(MVT::Other);
+ SDValue IntrinID =
+ DAG.getTargetConstant(Intrinsic::amdgcn_s_setreg, SL, MVT::i32);
+ SDValue SetModeReg =
+ DAG.getNode(ISD::INTRINSIC_VOID, SL, VTList, Op.getOperand(0), IntrinID,
+ ModeHwRegImm, NewModeReg);
+ SDValue SetTrapReg =
+ DAG.getNode(ISD::INTRINSIC_VOID, SL, VTList, Op.getOperand(0), IntrinID,
+ TrapHwRegImm, NewTrapReg);
+ return DAG.getNode(ISD::TokenFactor, SL, MVT::Other, SetTrapReg, SetModeReg);
+}
+
Register SITargetLowering::getRegisterByName(const char* RegName, LLT VT,
const MachineFunction &MF) const {
Register Reg = StringSwitch<Register>(RegName)
@@ -5663,6 +5729,10 @@ SDValue SITargetLowering::LowerOperation(SDValue Op, SelectionDAG &DAG) const {
case ISD::FP_EXTEND:
case ISD::STRICT_FP_EXTEND:
return lowerFP_EXTEND(Op, DAG);
+ case ISD::GET_FPENV:
+ return lowerGET_FPENV(Op, DAG);
+ case ISD::SET_FPENV:
+ return lowerSET_FPENV(Op, DAG);
}
return SDValue();
}
diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.h b/llvm/lib/Target/AMDGPU/SIISelLowering.h
index fc90a208fa0b3a..a20442e3737ee1 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.h
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.h
@@ -425,6 +425,8 @@ class SITargetLowering final : public AMDGPUTargetLowering {
SDValue lowerPREFETCH(SDValue Op, SelectionDAG &DAG) const;
SDValue lowerFP_EXTEND(SDValue Op, SelectionDAG &DAG) const;
+ SDValue lowerGET_FPENV(SDValue Op, SelectionDAG &DAG) const;
+ SDValue lowerSET_FPENV(SDValue Op, SelectionDAG &DAG) const;
Register getRegisterByName(const char* RegName, LLT VT,
const MachineFunction &MF) const override;
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.h b/llvm/lib/Target/AMDGPU/SIInstrInfo.h
index dab2cb2946ac97..ce4992e03f3783 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.h
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.h
@@ -864,7 +864,10 @@ class SIInstrInfo final : public AMDGPUGenInstrInfo {
return Opcode == AMDGPU::S_CMPK_EQ_U32 || Opcode == AMDGPU::S_CMPK_LG_U32 ||
Opcode == AMDGPU::S_CMPK_GT_U32 || Opcode == AMDGPU::S_CMPK_GE_U32 ||
Opcode == AMDGPU::S_CMPK_LT_U32 || Opcode == AMDGPU::S_CMPK_LE_U32 ||
- Opcode == AMDGPU::S_GETREG_B32;
+ Opcode == AMDGPU::S_GETREG_B32 || Opcode == AMDGPU::S_SETREG_B32 ||
+ Opcode == AMDGPU::S_SETREG_B32_mode ||
+ Opcode == AMDGPU::S_SETREG_IMM32_B32 ||
+ Opcode == AMDGPU::S_SETREG_IMM32_B32_mode;
}
/// \returns true if this is an s_store_dword* instruction. This is more
diff --git a/llvm/lib/Target/AMDGPU/SOPInstructions.td b/llvm/lib/Target/AMDGPU/SOPInstructions.td
index b5de311f8c58ce..22f35ee7169017 100644
--- a/llvm/lib/Target/AMDGPU/SOPInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SOPInstructions.td
@@ -1122,7 +1122,7 @@ class S_SETREG_B32_Pseudo <list<dag> pattern=[]> : SOPK_Pseudo <
pattern>;
def S_SETREG_B32 : S_SETREG_B32_Pseudo <
- [(int_amdgcn_s_setreg (i32 SIMM16bit:$simm16), i32:$sdst)]> {
+ [(int_amdgcn_s_setreg (i32 timm:$simm16), i32:$sdst)]> {
// Use custom inserter to optimize some cases to
// S_DENORM_MODE/S_ROUND_MODE/S_SETREG_B32_mode.
let usesCustomInserter = 1;
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.fpenv.ll b/llvm/test/CodeGen/AMDGPU/llvm.fpenv.ll
new file mode 100644
index 00000000000000..6625a52ccdb896
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.fpenv.ll
@@ -0,0 +1,114 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 4
+; RUN: llc -mtriple=amdgcn -mcpu=tahiti < %s | FileCheck -check-prefixes=GFX6 %s
+; RUN: llc -mtriple=amdgcn -global-isel -mcpu=tahiti < %s | FileCheck -check-prefixes=GFX6 %s
+; RUN: llc -mtriple=amdgcn -mcpu=fiji < %s | FileCheck -check-prefixes=GFX8 %s
+; RUN: llc -mtriple=amdgcn -global-isel -mcpu=fiji < %s | FileCheck -check-prefixes=GFX8 %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx900 < %s | FileCheck -check-prefixes=GFX9 %s
+; RUN: llc -mtriple=amdgcn -global-isel -mcpu=gfx900 < %s | FileCheck -check-prefixes=GFX9 %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx1030 < %s | FileCheck -check-prefixes=GFX10 %s
+; RUN: llc -mtriple=amdgcn -global-isel -mcpu=gfx1030 < %s | FileCheck -check-prefixes=GFX10 %s
+; RUN: llc -mtriple=amdgcn -amdgpu-enable-delay-alu=0 -mcpu=gfx1100 < %s | FileCheck -check-prefixes=GFX11 %s
+; RUN: llc -mtriple=amdgcn -amdgpu-enable-delay-alu=0 -global-isel -mcpu=gfx1100 < %s | FileCheck -check-prefixes=GFX11 %s
+
+declare i64 @llvm.get.fpenv.i64()
+
+define i64 @get_fpenv() {
+; GFX6-LABEL: get_fpenv:
+; GFX6: ; %bb.0: ; %entry
+; GFX6-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX6-NEXT: s_getreg_b32 s4, hwreg(HW_REG_MODE, 0, 23)
+; GFX6-NEXT: s_getreg_b32 s5, hwreg(HW_REG_TRAPSTS, 0, 5)
+; GFX6-NEXT: v_mov_b32_e32 v0, s4
+; GFX6-NEXT: v_mov_b32_e32 v1, s5
+; GFX6-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX8-LABEL: get_fpenv:
+; GFX8: ; %bb.0: ; %entry
+; GFX8-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX8-NEXT: s_getreg_b32 s4, hwreg(HW_REG_MODE, 0, 23)
+; GFX8-NEXT: s_getreg_b32 s5, hwreg(HW_REG_TRAPSTS, 0, 5)
+; GFX8-NEXT: v_mov_b32_e32 v0, s4
+; GFX8-NEXT: v_mov_b32_e32 v1, s5
+; GFX8-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX9-LABEL: get_fpenv:
+; GFX9: ; %bb.0: ; %entry
+; GFX9-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX9-NEXT: s_getreg_b32 s4, hwreg(HW_REG_MODE, 0, 23)
+; GFX9-NEXT: s_getreg_b32 s5, hwreg(HW_REG_TRAPSTS, 0, 5)
+; GFX9-NEXT: v_mov_b32_e32 v0, s4
+; GFX9-NEXT: v_mov_b32_e32 v1, s5
+; GFX9-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX10-LABEL: get_fpenv:
+; GFX10: ; %bb.0: ; %entry
+; GFX10-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX10-NEXT: s_getreg_b32 s4, hwreg(HW_REG_MODE, 0, 23)
+; GFX10-NEXT: s_getreg_b32 s5, hwreg(HW_REG_TRAPSTS, 0, 5)
+; GFX10-NEXT: v_mov_b32_e32 v0, s4
+; GFX10-NEXT: v_mov_b32_e32 v1, s5
+; GFX10-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX11-LABEL: get_fpenv:
+; GFX11: ; %bb.0: ; %entry
+; GFX11-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX11-NEXT: s_getreg_b32 s0, hwreg(HW_REG_MODE, 0, 23)
+; GFX11-NEXT: s_getreg_b32 s1, hwreg(HW_REG_TRAPSTS, 0, 5)
+; GFX11-NEXT: v_dual_mov_b32 v0, s0 :: v_dual_mov_b32 v1, s1
+; GFX11-NEXT: s_setpc_b64 s[30:31]
+entry:
+ %0 = tail call i64 @llvm.get.fpenv.i64()
+ ret i64 %0
+}
+
+declare void @llvm.set.fpenv.i64(i64)
+
+define void @set_fpenv(i64 %env) {
+; GFX6-LABEL: set_fpenv:
+; GFX6: ; %bb.0: ; %entry
+; GFX6-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX6-NEXT: v_readfirstlane_b32 s4, v0
+; GFX6-NEXT: s_setreg_b32 hwreg(HW_REG_MODE, 0, 23), s4
+; GFX6-NEXT: v_readfirstlane_b32 s4, v1
+; GFX6-NEXT: s_setreg_b32 hwreg(HW_REG_TRAPSTS, 0, 5), s4
+; GFX6-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX8-LABEL: set_fpenv:
+; GFX8: ; %bb.0: ; %entry
+; GFX8-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX8-NEXT: v_readfirstlane_b32 s4, v0
+; GFX8-NEXT: s_setreg_b32 hwreg(HW_REG_MODE, 0, 23), s4
+; GFX8-NEXT: v_readfirstlane_b32 s4, v1
+; GFX8-NEXT: s_setreg_b32 hwreg(HW_REG_TRAPSTS, 0, 5), s4
+; GFX8-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX9-LABEL: set_fpenv:
+; GFX9: ; %bb.0: ; %entry
+; GFX9-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX9-NEXT: v_readfirstlane_b32 s4, v0
+; GFX9-NEXT: s_setreg_b32 hwreg(HW_REG_MODE, 0, 23), s4
+; GFX9-NEXT: v_readfirstlane_b32 s4, v1
+; GFX9-NEXT: s_setreg_b32 hwreg(HW_REG_TRAPSTS, 0, 5), s4
+; GFX9-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX10-LABEL: set_fpenv:
+; GFX10: ; %bb.0: ; %entry
+; GFX10-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX10-NEXT: v_readfirstlane_b32 s4, v0
+; GFX10-NEXT: s_setreg_b32 hwreg(HW_REG_MODE, 0, 23), s4
+; GFX10-NEXT: v_readfirstlane_b32 s4, v1
+; GFX10-NEXT: s_setreg_b32 hwreg(HW_REG_TRAPSTS, 0, 5), s4
+; GFX10-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX11-LABEL: set_fpenv:
+; GFX11: ; %bb.0: ; %entry
+; GFX11-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX11-NEXT: v_readfirstlane_b32 s0, v0
+; GFX11-NEXT: s_setreg_b32 hwreg(HW_REG_MODE, 0, 23), s0
+; GFX11-NEXT: v_readfirstlane_b32 s0, v1
+; GFX11-NEXT: s_setreg_b32 hwreg(HW_REG_TRAPSTS, 0, 5), s0
+; GFX11-NEXT: s_setpc...
[truncated]
|
|
@llvm/pr-subscribers-clang-codegen Author: Joseph Huber (jhuber6) ChangesSummary: Patch is 20.08 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/83906.diff 13 Files Affected:
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index 213311b96df74f..6628e8f265fe48 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -325,6 +325,9 @@ BUILTIN(__builtin_amdgcn_read_exec_hi, "Ui", "nc")
BUILTIN(__builtin_amdgcn_endpgm, "v", "nr")
+BUILTIN(__builtin_amdgcn_get_fpenv, "WUi", "n")
+BUILTIN(__builtin_amdgcn_set_fpenv, "vWUi", "n")
+
//===----------------------------------------------------------------------===//
// R600-NI only builtins.
//===----------------------------------------------------------------------===//
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 9ee51ca7142c77..63a0ac9a0afe4a 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -18406,6 +18406,17 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
CGM.getIntrinsic(Intrinsic::amdgcn_global_load_tr, {ArgTy});
return Builder.CreateCall(F, {Addr});
}
+ case AMDGPU::BI__builtin_amdgcn_get_fpenv: {
+ Function *F = CGM.getIntrinsic(Intrinsic::get_fpenv,
+ {llvm::Type::getInt64Ty(getLLVMContext())});
+ return Builder.CreateCall(F);
+ }
+ case AMDGPU::BI__builtin_amdgcn_set_fpenv: {
+ Function *F = CGM.getIntrinsic(Intrinsic::set_fpenv,
+ {llvm::Type::getInt64Ty(getLLVMContext())});
+ llvm::Value *Env = EmitScalarExpr(E->getArg(0));
+ return Builder.CreateCall(F, {Env});
+ }
case AMDGPU::BI__builtin_amdgcn_read_exec:
return EmitAMDGCNBallotForExec(*this, E, Int64Ty, Int64Ty, false);
case AMDGPU::BI__builtin_amdgcn_read_exec_lo:
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
index 7d9010ee9067d6..8a4533633706b2 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -839,6 +839,18 @@ unsigned test_wavefrontsize() {
return __builtin_amdgcn_wavefrontsize();
}
+// CHECK-LABEL test_get_fpenv(
+unsigned long test_get_fpenv() {
+ // CHECK: call i64 @llvm.get.fpenv.i64()
+ return __builtin_amdgcn_get_fpenv();
+}
+
+// CHECK-LABEL test_set_fpenv(
+void test_set_fpenv(unsigned long env) {
+ // CHECK: call void @llvm.set.fpenv.i64(i64 %[[ENV:.+]])
+ __builtin_amdgcn_set_fpenv(env);
+}
+
// CHECK-DAG: [[$WI_RANGE]] = !{i32 0, i32 1024}
// CHECK-DAG: [[$WS_RANGE]] = !{i16 1, i16 1025}
// CHECK-DAG: attributes #[[$NOUNWIND_READONLY]] = { convergent mustprogress nocallback nofree nounwind willreturn memory(none) }
diff --git a/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst
index 0c588c84958624..ed41be4b08dee8 100644
--- a/llvm/docs/AMDGPUUsage.rst
+++ b/llvm/docs/AMDGPUUsage.rst
@@ -1151,6 +1151,13 @@ The AMDGPU backend implements the following LLVM IR intrinsics.
register do not exactly match the FLT_ROUNDS values,
so a conversion is performed.
+ :ref:`llvm.get.fpenv<int_get_fpenv>` Returns the current value of the AMDGPU floating point environment.
+ This stores information related to the current rounding mode,
+ denormalization mode, enabled traps, and floating point exceptions.
+ The format is a 64-bit concatenation of the MODE and TRAPSTS registers.
+
+ :ref:`llvm.set.fpenv<int_set_fpenv>` Sets the floating point environment to the specifies state.
+
llvm.amdgcn.wave.reduce.umin Performs an arithmetic unsigned min reduction on the unsigned values
provided by each lane in the wavefront.
Intrinsic takes a hint for reduction strategy using second operand
diff --git a/llvm/docs/LangRef.rst b/llvm/docs/LangRef.rst
index f56d4ed28f2855..3926cdedd34c5f 100644
--- a/llvm/docs/LangRef.rst
+++ b/llvm/docs/LangRef.rst
@@ -26427,6 +26427,7 @@ similar to C library function 'fesetround', however this intrinsic does not
return any value and uses platform-independent representation of IEEE rounding
modes.
+.. _int_get_fpenv:
'``llvm.get.fpenv``' Intrinsic
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -26450,6 +26451,7 @@ Semantics:
The '``llvm.get.fpenv``' intrinsic reads the current floating-point environment
and returns it as an integer value.
+.. _int_set_fpenv:
'``llvm.set.fpenv``' Intrinsic
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
diff --git a/llvm/docs/ReleaseNotes.rst b/llvm/docs/ReleaseNotes.rst
index 8ce6ee5cebb266..56abe7fe40270f 100644
--- a/llvm/docs/ReleaseNotes.rst
+++ b/llvm/docs/ReleaseNotes.rst
@@ -70,6 +70,8 @@ Changes to the AArch64 Backend
Changes to the AMDGPU Backend
-----------------------------
+* Implemented the ``llvm.get.fpenv`` and ``llvm.set.fpenv`` intrinsics.
+
Changes to the ARM Backend
--------------------------
diff --git a/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp
index 4c3b983f2960df..deef6611c18438 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp
@@ -905,6 +905,11 @@ AMDGPULegalizerInfo::AMDGPULegalizerInfo(const GCNSubtarget &ST_,
getActionDefinitionsBuilder(G_STACKRESTORE)
.legalFor({PrivatePtr});
+ getActionDefinitionsBuilder(G_GET_FPENV)
+ .customFor({S64});
+ getActionDefinitionsBuilder(G_SET_FPENV)
+ .customFor({S64});
+
getActionDefinitionsBuilder(G_GLOBAL_VALUE)
.customIf(typeIsNot(0, PrivatePtr));
@@ -2128,6 +2133,10 @@ bool AMDGPULegalizerInfo::legalizeCustom(
return legalizeFPTruncRound(MI, B);
case TargetOpcode::G_STACKSAVE:
return legalizeStackSave(MI, B);
+ case TargetOpcode::G_GET_FPENV:
+ return legalizeGetFPEnv(MI, B);
+ case TargetOpcode::G_SET_FPENV:
+ return legalizeSetFPEnv(MI, B);
default:
return false;
}
@@ -6940,6 +6949,42 @@ bool AMDGPULegalizerInfo::legalizeWaveID(MachineInstr &MI,
return true;
}
+static constexpr unsigned FPEnvModeBitField =
+ AMDGPU::Hwreg::HwregEncoding::encode(AMDGPU::Hwreg::ID_MODE, 0, 23);
+
+static constexpr unsigned FPEnvTrapBitField =
+ AMDGPU::Hwreg::HwregEncoding::encode(AMDGPU::Hwreg::ID_TRAPSTS, 0, 5);
+
+bool AMDGPULegalizerInfo::legalizeGetFPEnv(MachineInstr &MI,
+ MachineIRBuilder &B) const {
+ auto ModeReg =
+ B.buildIntrinsic(Intrinsic::amdgcn_s_getreg, {S32},
+ /*HasSideEffects=*/true, /*isConvergent=*/false)
+ .addImm(FPEnvModeBitField);
+ auto TrapReg =
+ B.buildIntrinsic(Intrinsic::amdgcn_s_getreg, {S32},
+ /*HasSideEffects=*/true, /*isConvergent=*/false)
+ .addImm(FPEnvTrapBitField);
+ B.buildMergeLikeInstr(MI.getOperand(0).getReg(), {ModeReg, TrapReg});
+ MI.eraseFromParent();
+ return true;
+}
+
+bool AMDGPULegalizerInfo::legalizeSetFPEnv(MachineInstr &MI,
+ MachineIRBuilder &B) const {
+ auto Unmerge = B.buildUnmerge({S32, S32}, MI.getOperand(0));
+ B.buildIntrinsic(Intrinsic::amdgcn_s_setreg, ArrayRef<DstOp>(),
+ /*HasSideEffects=*/true, /*isConvergent=*/false)
+ .addImm(FPEnvModeBitField)
+ .addReg(Unmerge.getReg(0));
+ B.buildIntrinsic(Intrinsic::amdgcn_s_setreg, ArrayRef<DstOp>(),
+ /*HasSideEffects=*/true, /*isConvergent=*/false)
+ .addImm(FPEnvTrapBitField)
+ .addReg(Unmerge.getReg(1));
+ MI.eraseFromParent();
+ return true;
+}
+
bool AMDGPULegalizerInfo::legalizeIntrinsic(LegalizerHelper &Helper,
MachineInstr &MI) const {
MachineIRBuilder &B = Helper.MIRBuilder;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.h b/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.h
index ecbe42681c6690..62b6fa12eb8751 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.h
@@ -214,6 +214,9 @@ class AMDGPULegalizerInfo final : public LegalizerInfo {
bool legalizeStackSave(MachineInstr &MI, MachineIRBuilder &B) const;
bool legalizeWaveID(MachineInstr &MI, MachineIRBuilder &B) const;
+ bool legalizeGetFPEnv(MachineInstr &MI, MachineIRBuilder &B) const;
+ bool legalizeSetFPEnv(MachineInstr &MI, MachineIRBuilder &B) const;
+
bool legalizeImageIntrinsic(
MachineInstr &MI, MachineIRBuilder &B,
GISelChangeObserver &Observer,
diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
index 30a65bb332652e..c45a65f451cd06 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -876,6 +876,8 @@ SITargetLowering::SITargetLowering(const TargetMachine &TM,
setOperationAction(ISD::STACKSAVE, MVT::Other, Custom);
setOperationAction(ISD::GET_ROUNDING, MVT::i32, Custom);
+ setOperationAction(ISD::GET_FPENV, MVT::i64, Custom);
+ setOperationAction(ISD::SET_FPENV, MVT::i64, Custom);
// TODO: Could move this to custom lowering, could benefit from combines on
// extract of relevant bits.
@@ -4079,6 +4081,70 @@ SDValue SITargetLowering::lowerFP_EXTEND(SDValue Op, SelectionDAG &DAG) const {
return DAG.getNode(ISD::BF16_TO_FP, SL, DstVT, BitCast);
}
+SDValue SITargetLowering::lowerGET_FPENV(SDValue Op, SelectionDAG &DAG) const {
+ SDLoc SL(Op);
+ assert(Op.getValueType() == MVT::i64);
+
+ uint32_t ModeHwReg =
+ AMDGPU::Hwreg::HwregEncoding::encode(AMDGPU::Hwreg::ID_MODE, 0, 23);
+ SDValue ModeHwRegImm = DAG.getTargetConstant(ModeHwReg, SL, MVT::i32);
+ uint32_t TrapHwReg =
+ AMDGPU::Hwreg::HwregEncoding::encode(AMDGPU::Hwreg::ID_TRAPSTS, 0, 5);
+ SDValue TrapHwRegImm = DAG.getTargetConstant(TrapHwReg, SL, MVT::i32);
+
+ SDVTList VTList = DAG.getVTList(MVT::i32, MVT::Other);
+ SDValue IntrinID =
+ DAG.getTargetConstant(Intrinsic::amdgcn_s_getreg, SL, MVT::i32);
+ SDValue GetModeReg = DAG.getNode(ISD::INTRINSIC_W_CHAIN, SL, VTList,
+ Op.getOperand(0), IntrinID, ModeHwRegImm);
+ SDValue GetTrapReg = DAG.getNode(ISD::INTRINSIC_W_CHAIN, SL, VTList,
+ Op.getOperand(0), IntrinID, TrapHwRegImm);
+ SDValue MergedReg =
+ DAG.getMergeValues({GetModeReg.getValue(1), GetTrapReg.getValue(1)}, SL);
+
+ SDValue CvtPtr =
+ DAG.getNode(ISD::BUILD_VECTOR, SL, MVT::v2i32, GetModeReg, GetTrapReg);
+ SDValue Result = DAG.getNode(ISD::BITCAST, SL, MVT::i64, CvtPtr);
+
+ return DAG.getMergeValues({Result, MergedReg}, SL);
+}
+
+SDValue SITargetLowering::lowerSET_FPENV(SDValue Op, SelectionDAG &DAG) const {
+ SDLoc SL(Op);
+ assert(Op.getOperand(1).getValueType() == MVT::i64);
+
+ SDValue Input = DAG.getNode(ISD::BITCAST, SL, MVT::v2i32, Op.getOperand(1));
+ SDValue NewModeReg = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, SL, MVT::i32, Input,
+ DAG.getConstant(0, SL, MVT::i32));
+ SDValue NewTrapReg = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, SL, MVT::i32, Input,
+ DAG.getConstant(1, SL, MVT::i32));
+
+ SDValue ReadFirstLaneID =
+ DAG.getTargetConstant(Intrinsic::amdgcn_readfirstlane, SL, MVT::i32);
+ NewModeReg = DAG.getNode(ISD::INTRINSIC_WO_CHAIN, SL, MVT::i32,
+ ReadFirstLaneID, NewModeReg);
+ NewTrapReg = DAG.getNode(ISD::INTRINSIC_WO_CHAIN, SL, MVT::i32,
+ ReadFirstLaneID, NewTrapReg);
+
+ unsigned ModeHwReg =
+ AMDGPU::Hwreg::HwregEncoding::encode(AMDGPU::Hwreg::ID_MODE, 0, 23);
+ SDValue ModeHwRegImm = DAG.getTargetConstant(ModeHwReg, SL, MVT::i32);
+ unsigned TrapHwReg =
+ AMDGPU::Hwreg::HwregEncoding::encode(AMDGPU::Hwreg::ID_TRAPSTS, 0, 5);
+ SDValue TrapHwRegImm = DAG.getTargetConstant(TrapHwReg, SL, MVT::i32);
+
+ SDVTList VTList = DAG.getVTList(MVT::Other);
+ SDValue IntrinID =
+ DAG.getTargetConstant(Intrinsic::amdgcn_s_setreg, SL, MVT::i32);
+ SDValue SetModeReg =
+ DAG.getNode(ISD::INTRINSIC_VOID, SL, VTList, Op.getOperand(0), IntrinID,
+ ModeHwRegImm, NewModeReg);
+ SDValue SetTrapReg =
+ DAG.getNode(ISD::INTRINSIC_VOID, SL, VTList, Op.getOperand(0), IntrinID,
+ TrapHwRegImm, NewTrapReg);
+ return DAG.getNode(ISD::TokenFactor, SL, MVT::Other, SetTrapReg, SetModeReg);
+}
+
Register SITargetLowering::getRegisterByName(const char* RegName, LLT VT,
const MachineFunction &MF) const {
Register Reg = StringSwitch<Register>(RegName)
@@ -5663,6 +5729,10 @@ SDValue SITargetLowering::LowerOperation(SDValue Op, SelectionDAG &DAG) const {
case ISD::FP_EXTEND:
case ISD::STRICT_FP_EXTEND:
return lowerFP_EXTEND(Op, DAG);
+ case ISD::GET_FPENV:
+ return lowerGET_FPENV(Op, DAG);
+ case ISD::SET_FPENV:
+ return lowerSET_FPENV(Op, DAG);
}
return SDValue();
}
diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.h b/llvm/lib/Target/AMDGPU/SIISelLowering.h
index fc90a208fa0b3a..a20442e3737ee1 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.h
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.h
@@ -425,6 +425,8 @@ class SITargetLowering final : public AMDGPUTargetLowering {
SDValue lowerPREFETCH(SDValue Op, SelectionDAG &DAG) const;
SDValue lowerFP_EXTEND(SDValue Op, SelectionDAG &DAG) const;
+ SDValue lowerGET_FPENV(SDValue Op, SelectionDAG &DAG) const;
+ SDValue lowerSET_FPENV(SDValue Op, SelectionDAG &DAG) const;
Register getRegisterByName(const char* RegName, LLT VT,
const MachineFunction &MF) const override;
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.h b/llvm/lib/Target/AMDGPU/SIInstrInfo.h
index dab2cb2946ac97..ce4992e03f3783 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.h
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.h
@@ -864,7 +864,10 @@ class SIInstrInfo final : public AMDGPUGenInstrInfo {
return Opcode == AMDGPU::S_CMPK_EQ_U32 || Opcode == AMDGPU::S_CMPK_LG_U32 ||
Opcode == AMDGPU::S_CMPK_GT_U32 || Opcode == AMDGPU::S_CMPK_GE_U32 ||
Opcode == AMDGPU::S_CMPK_LT_U32 || Opcode == AMDGPU::S_CMPK_LE_U32 ||
- Opcode == AMDGPU::S_GETREG_B32;
+ Opcode == AMDGPU::S_GETREG_B32 || Opcode == AMDGPU::S_SETREG_B32 ||
+ Opcode == AMDGPU::S_SETREG_B32_mode ||
+ Opcode == AMDGPU::S_SETREG_IMM32_B32 ||
+ Opcode == AMDGPU::S_SETREG_IMM32_B32_mode;
}
/// \returns true if this is an s_store_dword* instruction. This is more
diff --git a/llvm/lib/Target/AMDGPU/SOPInstructions.td b/llvm/lib/Target/AMDGPU/SOPInstructions.td
index b5de311f8c58ce..22f35ee7169017 100644
--- a/llvm/lib/Target/AMDGPU/SOPInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SOPInstructions.td
@@ -1122,7 +1122,7 @@ class S_SETREG_B32_Pseudo <list<dag> pattern=[]> : SOPK_Pseudo <
pattern>;
def S_SETREG_B32 : S_SETREG_B32_Pseudo <
- [(int_amdgcn_s_setreg (i32 SIMM16bit:$simm16), i32:$sdst)]> {
+ [(int_amdgcn_s_setreg (i32 timm:$simm16), i32:$sdst)]> {
// Use custom inserter to optimize some cases to
// S_DENORM_MODE/S_ROUND_MODE/S_SETREG_B32_mode.
let usesCustomInserter = 1;
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.fpenv.ll b/llvm/test/CodeGen/AMDGPU/llvm.fpenv.ll
new file mode 100644
index 00000000000000..6625a52ccdb896
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.fpenv.ll
@@ -0,0 +1,114 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 4
+; RUN: llc -mtriple=amdgcn -mcpu=tahiti < %s | FileCheck -check-prefixes=GFX6 %s
+; RUN: llc -mtriple=amdgcn -global-isel -mcpu=tahiti < %s | FileCheck -check-prefixes=GFX6 %s
+; RUN: llc -mtriple=amdgcn -mcpu=fiji < %s | FileCheck -check-prefixes=GFX8 %s
+; RUN: llc -mtriple=amdgcn -global-isel -mcpu=fiji < %s | FileCheck -check-prefixes=GFX8 %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx900 < %s | FileCheck -check-prefixes=GFX9 %s
+; RUN: llc -mtriple=amdgcn -global-isel -mcpu=gfx900 < %s | FileCheck -check-prefixes=GFX9 %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx1030 < %s | FileCheck -check-prefixes=GFX10 %s
+; RUN: llc -mtriple=amdgcn -global-isel -mcpu=gfx1030 < %s | FileCheck -check-prefixes=GFX10 %s
+; RUN: llc -mtriple=amdgcn -amdgpu-enable-delay-alu=0 -mcpu=gfx1100 < %s | FileCheck -check-prefixes=GFX11 %s
+; RUN: llc -mtriple=amdgcn -amdgpu-enable-delay-alu=0 -global-isel -mcpu=gfx1100 < %s | FileCheck -check-prefixes=GFX11 %s
+
+declare i64 @llvm.get.fpenv.i64()
+
+define i64 @get_fpenv() {
+; GFX6-LABEL: get_fpenv:
+; GFX6: ; %bb.0: ; %entry
+; GFX6-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX6-NEXT: s_getreg_b32 s4, hwreg(HW_REG_MODE, 0, 23)
+; GFX6-NEXT: s_getreg_b32 s5, hwreg(HW_REG_TRAPSTS, 0, 5)
+; GFX6-NEXT: v_mov_b32_e32 v0, s4
+; GFX6-NEXT: v_mov_b32_e32 v1, s5
+; GFX6-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX8-LABEL: get_fpenv:
+; GFX8: ; %bb.0: ; %entry
+; GFX8-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX8-NEXT: s_getreg_b32 s4, hwreg(HW_REG_MODE, 0, 23)
+; GFX8-NEXT: s_getreg_b32 s5, hwreg(HW_REG_TRAPSTS, 0, 5)
+; GFX8-NEXT: v_mov_b32_e32 v0, s4
+; GFX8-NEXT: v_mov_b32_e32 v1, s5
+; GFX8-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX9-LABEL: get_fpenv:
+; GFX9: ; %bb.0: ; %entry
+; GFX9-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX9-NEXT: s_getreg_b32 s4, hwreg(HW_REG_MODE, 0, 23)
+; GFX9-NEXT: s_getreg_b32 s5, hwreg(HW_REG_TRAPSTS, 0, 5)
+; GFX9-NEXT: v_mov_b32_e32 v0, s4
+; GFX9-NEXT: v_mov_b32_e32 v1, s5
+; GFX9-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX10-LABEL: get_fpenv:
+; GFX10: ; %bb.0: ; %entry
+; GFX10-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX10-NEXT: s_getreg_b32 s4, hwreg(HW_REG_MODE, 0, 23)
+; GFX10-NEXT: s_getreg_b32 s5, hwreg(HW_REG_TRAPSTS, 0, 5)
+; GFX10-NEXT: v_mov_b32_e32 v0, s4
+; GFX10-NEXT: v_mov_b32_e32 v1, s5
+; GFX10-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX11-LABEL: get_fpenv:
+; GFX11: ; %bb.0: ; %entry
+; GFX11-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX11-NEXT: s_getreg_b32 s0, hwreg(HW_REG_MODE, 0, 23)
+; GFX11-NEXT: s_getreg_b32 s1, hwreg(HW_REG_TRAPSTS, 0, 5)
+; GFX11-NEXT: v_dual_mov_b32 v0, s0 :: v_dual_mov_b32 v1, s1
+; GFX11-NEXT: s_setpc_b64 s[30:31]
+entry:
+ %0 = tail call i64 @llvm.get.fpenv.i64()
+ ret i64 %0
+}
+
+declare void @llvm.set.fpenv.i64(i64)
+
+define void @set_fpenv(i64 %env) {
+; GFX6-LABEL: set_fpenv:
+; GFX6: ; %bb.0: ; %entry
+; GFX6-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX6-NEXT: v_readfirstlane_b32 s4, v0
+; GFX6-NEXT: s_setreg_b32 hwreg(HW_REG_MODE, 0, 23), s4
+; GFX6-NEXT: v_readfirstlane_b32 s4, v1
+; GFX6-NEXT: s_setreg_b32 hwreg(HW_REG_TRAPSTS, 0, 5), s4
+; GFX6-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX8-LABEL: set_fpenv:
+; GFX8: ; %bb.0: ; %entry
+; GFX8-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX8-NEXT: v_readfirstlane_b32 s4, v0
+; GFX8-NEXT: s_setreg_b32 hwreg(HW_REG_MODE, 0, 23), s4
+; GFX8-NEXT: v_readfirstlane_b32 s4, v1
+; GFX8-NEXT: s_setreg_b32 hwreg(HW_REG_TRAPSTS, 0, 5), s4
+; GFX8-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX9-LABEL: set_fpenv:
+; GFX9: ; %bb.0: ; %entry
+; GFX9-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX9-NEXT: v_readfirstlane_b32 s4, v0
+; GFX9-NEXT: s_setreg_b32 hwreg(HW_REG_MODE, 0, 23), s4
+; GFX9-NEXT: v_readfirstlane_b32 s4, v1
+; GFX9-NEXT: s_setreg_b32 hwreg(HW_REG_TRAPSTS, 0, 5), s4
+; GFX9-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX10-LABEL: set_fpenv:
+; GFX10: ; %bb.0: ; %entry
+; GFX10-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX10-NEXT: v_readfirstlane_b32 s4, v0
+; GFX10-NEXT: s_setreg_b32 hwreg(HW_REG_MODE, 0, 23), s4
+; GFX10-NEXT: v_readfirstlane_b32 s4, v1
+; GFX10-NEXT: s_setreg_b32 hwreg(HW_REG_TRAPSTS, 0, 5), s4
+; GFX10-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX11-LABEL: set_fpenv:
+; GFX11: ; %bb.0: ; %entry
+; GFX11-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX11-NEXT: v_readfirstlane_b32 s0, v0
+; GFX11-NEXT: s_setreg_b32 hwreg(HW_REG_MODE, 0, 23), s0
+; GFX11-NEXT: v_readfirstlane_b32 s0, v1
+; GFX11-NEXT: s_setreg_b32 hwreg(HW_REG_TRAPSTS, 0, 5), s0
+; GFX11-NEXT: s_setpc...
[truncated]
|
|
Note that this patch is not quite ready to land. I encountered issues when working with |
|
✅ With the latest revision this PR passed the C/C++ code formatter. |
rampitec
reviewed
Mar 4, 2024
| @@ -1122,7 +1122,7 @@ class S_SETREG_B32_Pseudo <list<dag> pattern=[]> : SOPK_Pseudo < | |||
| pattern>; | |||
|
|
|||
| def S_SETREG_B32 : S_SETREG_B32_Pseudo < | |||
| [(int_amdgcn_s_setreg (i32 SIMM16bit:$simm16), i32:$sdst)]> { | |||
| [(int_amdgcn_s_setreg (i32 timm:$simm16), i32:$sdst)]> { | |||
There was a problem hiding this comment.
This just reverts my patch #77997 and reintroduces the original problem.
There was a problem hiding this comment.
Okay, so it's expected to be negative even though s_getreg isn't? We can keep it, but the problem I was having was that in global-isel it would error out constantly with the current usage. I couldn't find how I was expected to pass the correct sign to that function. What if we just permitted either way in the legalizer? It seems to work both ways.
There was a problem hiding this comment.
It is not expected to be negative, the original problem was that we used to force users to use negative constants. Now we can accept something like 0xf000 instead of a negative value.
There was a problem hiding this comment.
I hit the SOPK error when trying to emit this in globalIsel for the TRAPSTS register.
> llc -mtriple=amdgcn -mcpu=tahiti ../llvm/test/CodeGen/AMDGPU/llvm.fpenv.ll -global-isel
error: Illegal instruction detected: invalid immediate for SOPK instruction
S_SETREG_B32 killed renamable $sgpr4, 45057, implicit-def $mode, implicit $mode
I tried casting them so they would get sign extended, but I got similar issues.
There was a problem hiding this comment.
If it is sign extended, it should work.
There was a problem hiding this comment.
Okay, I think I got it to work by casting it to int16_t only on the store instruction. My previous mistake seems to have been changing the constant for both set and `get. It passes all the AMDGPU codegen tests now.
arsenm
reviewed
Mar 5, 2024
| @@ -325,6 +325,9 @@ BUILTIN(__builtin_amdgcn_read_exec_hi, "Ui", "nc") | |||
|
|
|||
| BUILTIN(__builtin_amdgcn_endpgm, "v", "nr") | |||
|
|
|||
| BUILTIN(__builtin_amdgcn_get_fpenv, "WUi", "n") | |||
There was a problem hiding this comment.
@spavloff was there supposed to be a generic builtin for these?
There was a problem hiding this comment.
There's no builtin as far as I'm aware. I think there might be some pragmas however. Because the intrinsic takes different kinds I figured it was easiest to just specify my use-case. The alternative would be to make this generic and have a LUT by target triple in CGBuiltin that just contains the one entry for AMDGPU.
arsenm
reviewed
Mar 5, 2024
jhuber6
force-pushed
the
AMDGPUEnv
branch
2 times, most recently
from
7808b8a
to
169f891
Compare
arsenm
reviewed
Mar 5, 2024
arsenm
reviewed
Mar 5, 2024
arsenm
reviewed
Mar 5, 2024
Summary: This patch implements the LLVM floating point environment control intrinsics and also exposes it through clang. We encode the floating point environment as a 64-bit value that simply concatenates the values of the mode registers and the current trap status. We only fetch the bits relevant for floating point instructions. That is, rounding mode, denormalization mode, ieee, dx10 clamp, debug, enabled traps, f16 overflow, and active exceptions.
arsenm
approved these changes
Mar 6, 2024
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Summary:
This patch implements the LLVM floating point environment control
intrinsics and also exposes it through clang. We encode the floating
point environment as a 64-bit value that simply concatenates the values
of the mode registers and the current trap status. We only fetch the
bits relevant for floating point instructions. That is, rounding mode,
denormalization mode, ieee, dx10 clamp, debug, enabled traps, f16
overflow, and active exceptions.