OSchip
/
llvm-project
mirror of https://github.com/THU-DSP-LAB/llvm-project.git
13
0
Fork 6
Code Issues Packages Projects Releases Wiki Activity

Revert "[VENTUS][RISCV][feat] Legalized vector parameters" 

This reverts commit 7bd98c0ff8.
This commit is contained in:
zhoujingya 2023-10-08 17:31:53 +08:00
parent 8bea355dbe
commit 0a45eabde0
3 changed files with 3 additions and 70 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

47
llvm/lib/Target/RISCV/RISCVISelLowering.cpp
View File

@ -7460,46 +7460,6 @@ SDValue RISCVTargetLowering::lowerKernArgParameterPtr(SelectionDAG &DAG,
return DAG.getObjectPtrOffset(SL, BasePtr, TypeSize::Fixed(Offset)); return DAG.getObjectPtrOffset(SL, BasePtr, TypeSize::Fixed(Offset));
} }
SDValue RISCVTargetLowering::getFPExtOrFPRound(SelectionDAG &DAG,
SDValue Op,
const SDLoc &DL,
EVT VT) const {
return Op.getValueType().bitsLE(VT) ?
DAG.getNode(ISD::FP_EXTEND, DL, VT, Op) :
DAG.getNode(ISD::FP_ROUND, DL, VT, Op,
DAG.getTargetConstant(0, DL, MVT::i32));
}
SDValue RISCVTargetLowering::convertArgType(SelectionDAG &DAG, EVT VT, EVT MemVT,
const SDLoc &SL, SDValue Val,
bool Signed,
const ISD::InputArg *Arg) const {
// First, if it is a widened vector, narrow it.
if (VT.isVector() &&
VT.getVectorNumElements() != MemVT.getVectorNumElements()) {
EVT NarrowedVT =
EVT::getVectorVT(*DAG.getContext(), MemVT.getVectorElementType(),
VT.getVectorNumElements());
Val = DAG.getNode(ISD::EXTRACT_SUBVECTOR, SL, NarrowedVT, Val,
DAG.getConstant(0, SL, MVT::i32));
}
// Then convert the vector elements or scalar value.
if (Arg && (Arg->Flags.isSExt() || Arg->Flags.isZExt()) &&
VT.bitsLT(MemVT)) {
unsigned Opc = Arg->Flags.isZExt() ? ISD::AssertZext : ISD::AssertSext;
Val = DAG.getNode(Opc, SL, MemVT, Val, DAG.getValueType(VT));
}
if (MemVT.isFloatingPoint())
Val = getFPExtOrFPRound(DAG, Val, SL, VT);
else if (Signed)
Val = DAG.getSExtOrTrunc(Val, SL, VT);
else
Val = DAG.getZExtOrTrunc(Val, SL, VT);
return Val;
}
SDValue RISCVTargetLowering::lowerKernargMemParameter( SDValue RISCVTargetLowering::lowerKernargMemParameter(
SelectionDAG &DAG, EVT VT, EVT MemVT, const SDLoc &SL, SDValue Chain, SelectionDAG &DAG, EVT VT, EVT MemVT, const SDLoc &SL, SDValue Chain,
@ -7530,7 +7490,7 @@ SDValue RISCVTargetLowering::lowerKernargMemParameter(
SDValue ArgVal = DAG.getNode(ISD::TRUNCATE, SL, IntVT, Extract); SDValue ArgVal = DAG.getNode(ISD::TRUNCATE, SL, IntVT, Extract);
ArgVal = DAG.getNode(ISD::BITCAST, SL, MemVT, ArgVal); ArgVal = DAG.getNode(ISD::BITCAST, SL, MemVT, ArgVal);
// TODO: Support vector and half type. // TODO: Support vector and half type.
ArgVal = convertArgType(DAG, VT, MemVT, SL, ArgVal, Signed, Arg); //ArgVal = convertArgType(DAG, VT, MemVT, SL, ArgVal, Signed, Arg);
return DAG.getMergeValues({ ArgVal, Load.getValue(1) }, SL); return DAG.getMergeValues({ ArgVal, Load.getValue(1) }, SL);
} }
@ -7540,9 +7500,8 @@ SDValue RISCVTargetLowering::lowerKernargMemParameter(
MachineMemOperand::MODereferenceable | MachineMemOperand::MODereferenceable |
MachineMemOperand::MOInvariant); MachineMemOperand::MOInvariant);
SDValue Val = convertArgType(DAG, VT, MemVT, SL, Load, Signed, Arg); // SDValue Val = convertArgType(DAG, VT, MemVT, SL, Load, Signed, Arg);
// return DAG.getMergeValues({ Load, Load.getValue(1) }, SL); return DAG.getMergeValues({ Load, Load.getValue(1) }, SL);
return DAG.getMergeValues({ Val, Load.getValue(1) }, SL);
} }
// Returns the opcode of the target-specific SDNode that implements the 32-bit // Returns the opcode of the target-specific SDNode that implements the 32-bit

7
llvm/lib/Target/RISCV/RISCVISelLowering.h
View File

@ -722,13 +722,6 @@ private:
SDValue expandUnalignedRVVStore(SDValue Op, SelectionDAG &DAG) const; SDValue expandUnalignedRVVStore(SDValue Op, SelectionDAG &DAG) const;
SDValue lowerKernArgParameterPtr(SelectionDAG &DAG, const SDLoc &SL, SDValue lowerKernArgParameterPtr(SelectionDAG &DAG, const SDLoc &SL,
SDValue Chain, uint64_t Offset) const; SDValue Chain, uint64_t Offset) const;
SDValue getFPExtOrFPRound(SelectionDAG &DAG, SDValue Op,
const SDLoc &DL,
EVT VT) const;
SDValue convertArgType(SelectionDAG &DAG, EVT VT, EVT MemVT,
const SDLoc &SL, SDValue Val,
bool Signed,
const ISD::InputArg *Arg) const;
SDValue lowerKernargMemParameter(SelectionDAG &DAG, EVT VT, EVT MemVT, SDValue lowerKernargMemParameter(SelectionDAG &DAG, EVT VT, EVT MemVT,
const SDLoc &SL, SDValue Chain, const SDLoc &SL, SDValue Chain,
uint64_t Offset, Align Alignment, uint64_t Offset, Align Alignment,

19
llvm/test/CodeGen/RISCV/VentusGPGPU/vecargtest.cl
View File

@ -1,19 +0,0 @@
// RUN: clang -target riscv32 -mcpu=ventus-gpgpu < %s \
// RUN: | FileCheck -check-prefix=VENTUS %s
__kernel void test_kernel(
char2 c, uchar2 uc, short2 s, ushort2 us, int2 i, uint2 ui, float2 f,
__global float2 *result)
{
// there is no FileCheck now
result[0] = convert_float2(c);
result[1] = convert_float2(uc);
result[2] = convert_float2(s);
result[3] = convert_float2(us);
result[4] = convert_float2(i);
result[5] = convert_float2(ui);
result[6] = f;
}