bool CC_X86_32_VectorCall(unsigned &ValNo, MVT &ValVT, MVT &LocVT,
CCValAssign::LocInfo &LocInfo,
ISD::ArgFlagsTy &ArgFlags, CCState &State) {
// On the second pass, go through the HVAs only.
if (ArgFlags.isSecArgPass()) {
if (ArgFlags.isHva())
return CC_X86_VectorCallAssignRegister(ValNo, ValVT, LocVT, LocInfo,
ArgFlags, State);
return true;
}
// Process only vector types as defined by vectorcall spec:
// "A vector type is either a floating point type, for example,
// a float or double, or an SIMD vector type, for example, __m128 or __m256".
if (!(ValVT.isFloatingPoint() ||
(ValVT.isVector() && ValVT.getSizeInBits() >= 128))) {
return false;
}
if (ArgFlags.isHva())
return true; // If this is an HVA - Stop the search.
// Assign XMM register.
if (unsigned Reg = State.AllocateReg(CC_X86_VectorCallGetSSEs(ValVT))) {
State.addLoc(CCValAssign::getReg(ValNo, ValVT, Reg, LocVT, LocInfo));
return true;
}
// In case we did not find an available XMM register for a vector -
// pass it indirectly.
// It is similar to CCPassIndirect, with the addition of inreg.
if (!ValVT.isFloatingPoint()) {
LocVT = MVT::i32;
LocInfo = CCValAssign::Indirect;
ArgFlags.setInReg();
}
return false; // No register was assigned - Continue the search.
}
bool CC_X86_64_VectorCall(unsigned &ValNo, MVT &ValVT, MVT &LocVT,
CCValAssign::LocInfo &LocInfo,
ISD::ArgFlagsTy &ArgFlags, CCState &State) {
// On the second pass, go through the HVAs only.
if (ArgFlags.isSecArgPass()) {
if (ArgFlags.isHva())
return CC_X86_VectorCallAssignRegister(ValNo, ValVT, LocVT, LocInfo,
ArgFlags, State);
return true;
}
// Process only vector types as defined by vectorcall spec:
// "A vector type is either a floating-point type, for example,
// a float or double, or an SIMD vector type, for example, __m128 or __m256".
if (!(ValVT.isFloatingPoint() ||
(ValVT.isVector() && ValVT.getSizeInBits() >= 128))) {
// If R9 was already assigned it means that we are after the fourth element
// and because this is not an HVA / Vector type, we need to allocate
// shadow XMM register.
if (State.isAllocated(X86::R9)) {
// Assign shadow XMM register.
(void)State.AllocateReg(CC_X86_VectorCallGetSSEs(ValVT));
}
return false;
}
if (!ArgFlags.isHva() || ArgFlags.isHvaStart()) {
// Assign shadow GPR register.
(void)State.AllocateReg(CC_X86_64_VectorCallGetGPRs());
// Assign XMM register - (shadow for HVA and non-shadow for non HVA).
if (unsigned Reg = State.AllocateReg(CC_X86_VectorCallGetSSEs(ValVT))) {
// In Vectorcall Calling convention, additional shadow stack can be
// created on top of the basic 32 bytes of win64.
// It can happen if the fifth or sixth argument is vector type or HVA.
// At that case for each argument a shadow stack of 8 bytes is allocated.
if (Reg == X86::XMM4 || Reg == X86::XMM5)
State.AllocateStack(8, 8);
if (!ArgFlags.isHva()) {
State.addLoc(CCValAssign::getReg(ValNo, ValVT, Reg, LocVT, LocInfo));
return true; // Allocated a register - Stop the search.
}
}
}
// If this is an HVA - Stop the search,
// otherwise continue the search.
return ArgFlags.isHva();
}
SDNode *NVPTXDAGToDAGISel::SelectStore(SDNode *N) {
DebugLoc dl = N->getDebugLoc();
StoreSDNode *ST = cast<StoreSDNode>(N);
EVT StoreVT = ST->getMemoryVT();
SDNode *NVPTXST = NULL;
// do not support pre/post inc/dec
if (ST->isIndexed())
return NULL;
if (!StoreVT.isSimple())
return NULL;
// Address Space Setting
unsigned int codeAddrSpace = getCodeAddrSpace(ST, Subtarget);
// Volatile Setting
// - .volatile is only availalble for .global and .shared
bool isVolatile = ST->isVolatile();
if (codeAddrSpace != NVPTX::PTXLdStInstCode::GLOBAL &&
codeAddrSpace != NVPTX::PTXLdStInstCode::SHARED &&
codeAddrSpace != NVPTX::PTXLdStInstCode::GENERIC)
isVolatile = false;
// Vector Setting
MVT SimpleVT = StoreVT.getSimpleVT();
unsigned vecType = NVPTX::PTXLdStInstCode::Scalar;
if (SimpleVT.isVector()) {
unsigned num = SimpleVT.getVectorNumElements();
if (num == 2)
vecType = NVPTX::PTXLdStInstCode::V2;
else if (num == 4)
vecType = NVPTX::PTXLdStInstCode::V4;
else
return NULL;
}
// Type Setting: toType + toTypeWidth
// - for integer type, always use 'u'
//
MVT ScalarVT = SimpleVT.getScalarType();
unsigned toTypeWidth = ScalarVT.getSizeInBits();
unsigned int toType;
if (ScalarVT.isFloatingPoint())
toType = NVPTX::PTXLdStInstCode::Float;
else
toType = NVPTX::PTXLdStInstCode::Unsigned;
// Create the machine instruction DAG
SDValue Chain = N->getOperand(0);
SDValue N1 = N->getOperand(1);
SDValue N2 = N->getOperand(2);
SDValue Addr;
SDValue Offset, Base;
unsigned Opcode;
MVT::SimpleValueType SourceVT =
N1.getNode()->getValueType(0).getSimpleVT().SimpleTy;
if (SelectDirectAddr(N2, Addr)) {
switch (SourceVT) {
case MVT::i8:
Opcode = NVPTX::ST_i8_avar;
break;
case MVT::i16:
Opcode = NVPTX::ST_i16_avar;
break;
case MVT::i32:
Opcode = NVPTX::ST_i32_avar;
break;
case MVT::i64:
Opcode = NVPTX::ST_i64_avar;
break;
case MVT::f32:
Opcode = NVPTX::ST_f32_avar;
break;
case MVT::f64:
Opcode = NVPTX::ST_f64_avar;
break;
default:
return NULL;
}
SDValue Ops[] = { N1, getI32Imm(isVolatile), getI32Imm(codeAddrSpace),
getI32Imm(vecType), getI32Imm(toType),
getI32Imm(toTypeWidth), Addr, Chain };
NVPTXST = CurDAG->getMachineNode(Opcode, dl, MVT::Other, Ops, 8);
} else if (Subtarget.is64Bit()
? SelectADDRsi64(N2.getNode(), N2, Base, Offset)
: SelectADDRsi(N2.getNode(), N2, Base, Offset)) {
switch (SourceVT) {
case MVT::i8:
Opcode = NVPTX::ST_i8_asi;
break;
case MVT::i16:
Opcode = NVPTX::ST_i16_asi;
break;
case MVT::i32:
Opcode = NVPTX::ST_i32_asi;
break;
case MVT::i64:
Opcode = NVPTX::ST_i64_asi;
break;
case MVT::f32:
Opcode = NVPTX::ST_f32_asi;
break;
case MVT::f64:
Opcode = NVPTX::ST_f64_asi;
break;
default:
return NULL;
}
SDValue Ops[] = { N1, getI32Imm(isVolatile), getI32Imm(codeAddrSpace),
getI32Imm(vecType), getI32Imm(toType),
getI32Imm(toTypeWidth), Base, Offset, Chain };
NVPTXST = CurDAG->getMachineNode(Opcode, dl, MVT::Other, Ops, 9);
} else if (Subtarget.is64Bit()
? SelectADDRri64(N2.getNode(), N2, Base, Offset)
: SelectADDRri(N2.getNode(), N2, Base, Offset)) {
if (Subtarget.is64Bit()) {
switch (SourceVT) {
case MVT::i8:
Opcode = NVPTX::ST_i8_ari_64;
break;
case MVT::i16:
Opcode = NVPTX::ST_i16_ari_64;
break;
case MVT::i32:
Opcode = NVPTX::ST_i32_ari_64;
break;
case MVT::i64:
Opcode = NVPTX::ST_i64_ari_64;
break;
case MVT::f32:
Opcode = NVPTX::ST_f32_ari_64;
break;
case MVT::f64:
Opcode = NVPTX::ST_f64_ari_64;
break;
default:
return NULL;
}
} else {
switch (SourceVT) {
case MVT::i8:
Opcode = NVPTX::ST_i8_ari;
break;
case MVT::i16:
Opcode = NVPTX::ST_i16_ari;
break;
case MVT::i32:
Opcode = NVPTX::ST_i32_ari;
break;
case MVT::i64:
Opcode = NVPTX::ST_i64_ari;
break;
case MVT::f32:
Opcode = NVPTX::ST_f32_ari;
break;
case MVT::f64:
Opcode = NVPTX::ST_f64_ari;
break;
default:
return NULL;
}
}
SDValue Ops[] = { N1, getI32Imm(isVolatile), getI32Imm(codeAddrSpace),
getI32Imm(vecType), getI32Imm(toType),
getI32Imm(toTypeWidth), Base, Offset, Chain };
NVPTXST = CurDAG->getMachineNode(Opcode, dl, MVT::Other, Ops, 9);
} else {
if (Subtarget.is64Bit()) {
switch (SourceVT) {
case MVT::i8:
Opcode = NVPTX::ST_i8_areg_64;
break;
case MVT::i16:
Opcode = NVPTX::ST_i16_areg_64;
break;
case MVT::i32:
Opcode = NVPTX::ST_i32_areg_64;
break;
case MVT::i64:
Opcode = NVPTX::ST_i64_areg_64;
break;
case MVT::f32:
Opcode = NVPTX::ST_f32_areg_64;
break;
case MVT::f64:
Opcode = NVPTX::ST_f64_areg_64;
break;
default:
return NULL;
}
} else {
switch (SourceVT) {
case MVT::i8:
Opcode = NVPTX::ST_i8_areg;
break;
case MVT::i16:
Opcode = NVPTX::ST_i16_areg;
break;
case MVT::i32:
Opcode = NVPTX::ST_i32_areg;
break;
case MVT::i64:
Opcode = NVPTX::ST_i64_areg;
break;
case MVT::f32:
Opcode = NVPTX::ST_f32_areg;
break;
case MVT::f64:
Opcode = NVPTX::ST_f64_areg;
break;
default:
return NULL;
}
}
SDValue Ops[] = { N1, getI32Imm(isVolatile), getI32Imm(codeAddrSpace),
getI32Imm(vecType), getI32Imm(toType),
getI32Imm(toTypeWidth), N2, Chain };
NVPTXST = CurDAG->getMachineNode(Opcode, dl, MVT::Other, Ops, 8);
}
if (NVPTXST != NULL) {
MachineSDNode::mmo_iterator MemRefs0 = MF->allocateMemRefsArray(1);
MemRefs0[0] = cast<MemSDNode>(N)->getMemOperand();
cast<MachineSDNode>(NVPTXST)->setMemRefs(MemRefs0, MemRefs0 + 1);
}
return NVPTXST;
}
SDNode *NVPTXDAGToDAGISel::SelectLoadVector(SDNode *N) {
SDValue Chain = N->getOperand(0);
SDValue Op1 = N->getOperand(1);
SDValue Addr, Offset, Base;
unsigned Opcode;
DebugLoc DL = N->getDebugLoc();
SDNode *LD;
MemSDNode *MemSD = cast<MemSDNode>(N);
EVT LoadedVT = MemSD->getMemoryVT();
if (!LoadedVT.isSimple())
return NULL;
// Address Space Setting
unsigned int CodeAddrSpace = getCodeAddrSpace(MemSD, Subtarget);
// Volatile Setting
// - .volatile is only availalble for .global and .shared
bool IsVolatile = MemSD->isVolatile();
if (CodeAddrSpace != NVPTX::PTXLdStInstCode::GLOBAL &&
CodeAddrSpace != NVPTX::PTXLdStInstCode::SHARED &&
CodeAddrSpace != NVPTX::PTXLdStInstCode::GENERIC)
IsVolatile = false;
// Vector Setting
MVT SimpleVT = LoadedVT.getSimpleVT();
// Type Setting: fromType + fromTypeWidth
//
// Sign : ISD::SEXTLOAD
// Unsign : ISD::ZEXTLOAD, ISD::NON_EXTLOAD or ISD::EXTLOAD and the
// type is integer
// Float : ISD::NON_EXTLOAD or ISD::EXTLOAD and the type is float
MVT ScalarVT = SimpleVT.getScalarType();
unsigned FromTypeWidth = ScalarVT.getSizeInBits();
unsigned int FromType;
// The last operand holds the original LoadSDNode::getExtensionType() value
unsigned ExtensionType = cast<ConstantSDNode>(
N->getOperand(N->getNumOperands() - 1))->getZExtValue();
if (ExtensionType == ISD::SEXTLOAD)
FromType = NVPTX::PTXLdStInstCode::Signed;
else if (ScalarVT.isFloatingPoint())
FromType = NVPTX::PTXLdStInstCode::Float;
else
FromType = NVPTX::PTXLdStInstCode::Unsigned;
unsigned VecType;
switch (N->getOpcode()) {
case NVPTXISD::LoadV2:
VecType = NVPTX::PTXLdStInstCode::V2;
break;
case NVPTXISD::LoadV4:
VecType = NVPTX::PTXLdStInstCode::V4;
break;
default:
return NULL;
}
EVT EltVT = N->getValueType(0);
if (SelectDirectAddr(Op1, Addr)) {
switch (N->getOpcode()) {
default:
return NULL;
case NVPTXISD::LoadV2:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::LDV_i8_v2_avar;
break;
case MVT::i16:
Opcode = NVPTX::LDV_i16_v2_avar;
break;
case MVT::i32:
Opcode = NVPTX::LDV_i32_v2_avar;
break;
case MVT::i64:
Opcode = NVPTX::LDV_i64_v2_avar;
break;
case MVT::f32:
Opcode = NVPTX::LDV_f32_v2_avar;
break;
case MVT::f64:
Opcode = NVPTX::LDV_f64_v2_avar;
break;
}
break;
case NVPTXISD::LoadV4:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::LDV_i8_v4_avar;
break;
case MVT::i16:
Opcode = NVPTX::LDV_i16_v4_avar;
break;
case MVT::i32:
Opcode = NVPTX::LDV_i32_v4_avar;
break;
case MVT::f32:
Opcode = NVPTX::LDV_f32_v4_avar;
break;
}
break;
}
SDValue Ops[] = { getI32Imm(IsVolatile), getI32Imm(CodeAddrSpace),
getI32Imm(VecType), getI32Imm(FromType),
getI32Imm(FromTypeWidth), Addr, Chain };
LD = CurDAG->getMachineNode(Opcode, DL, N->getVTList(), Ops, 7);
} else if (Subtarget.is64Bit()
? SelectADDRsi64(Op1.getNode(), Op1, Base, Offset)
: SelectADDRsi(Op1.getNode(), Op1, Base, Offset)) {
switch (N->getOpcode()) {
default:
return NULL;
case NVPTXISD::LoadV2:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::LDV_i8_v2_asi;
break;
case MVT::i16:
Opcode = NVPTX::LDV_i16_v2_asi;
break;
case MVT::i32:
Opcode = NVPTX::LDV_i32_v2_asi;
break;
case MVT::i64:
Opcode = NVPTX::LDV_i64_v2_asi;
break;
case MVT::f32:
Opcode = NVPTX::LDV_f32_v2_asi;
break;
case MVT::f64:
Opcode = NVPTX::LDV_f64_v2_asi;
break;
}
break;
case NVPTXISD::LoadV4:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::LDV_i8_v4_asi;
break;
case MVT::i16:
Opcode = NVPTX::LDV_i16_v4_asi;
break;
case MVT::i32:
Opcode = NVPTX::LDV_i32_v4_asi;
break;
case MVT::f32:
Opcode = NVPTX::LDV_f32_v4_asi;
break;
}
break;
}
SDValue Ops[] = { getI32Imm(IsVolatile), getI32Imm(CodeAddrSpace),
getI32Imm(VecType), getI32Imm(FromType),
getI32Imm(FromTypeWidth), Base, Offset, Chain };
LD = CurDAG->getMachineNode(Opcode, DL, N->getVTList(), Ops, 8);
} else if (Subtarget.is64Bit()
? SelectADDRri64(Op1.getNode(), Op1, Base, Offset)
: SelectADDRri(Op1.getNode(), Op1, Base, Offset)) {
if (Subtarget.is64Bit()) {
switch (N->getOpcode()) {
default:
return NULL;
case NVPTXISD::LoadV2:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::LDV_i8_v2_ari_64;
break;
case MVT::i16:
Opcode = NVPTX::LDV_i16_v2_ari_64;
break;
case MVT::i32:
Opcode = NVPTX::LDV_i32_v2_ari_64;
break;
case MVT::i64:
Opcode = NVPTX::LDV_i64_v2_ari_64;
break;
case MVT::f32:
Opcode = NVPTX::LDV_f32_v2_ari_64;
break;
case MVT::f64:
Opcode = NVPTX::LDV_f64_v2_ari_64;
break;
}
break;
case NVPTXISD::LoadV4:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::LDV_i8_v4_ari_64;
break;
case MVT::i16:
Opcode = NVPTX::LDV_i16_v4_ari_64;
break;
case MVT::i32:
Opcode = NVPTX::LDV_i32_v4_ari_64;
break;
case MVT::f32:
Opcode = NVPTX::LDV_f32_v4_ari_64;
break;
}
break;
}
} else {
switch (N->getOpcode()) {
default:
return NULL;
case NVPTXISD::LoadV2:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::LDV_i8_v2_ari;
break;
case MVT::i16:
Opcode = NVPTX::LDV_i16_v2_ari;
break;
case MVT::i32:
Opcode = NVPTX::LDV_i32_v2_ari;
break;
case MVT::i64:
Opcode = NVPTX::LDV_i64_v2_ari;
break;
case MVT::f32:
Opcode = NVPTX::LDV_f32_v2_ari;
break;
case MVT::f64:
Opcode = NVPTX::LDV_f64_v2_ari;
break;
}
break;
case NVPTXISD::LoadV4:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::LDV_i8_v4_ari;
break;
case MVT::i16:
Opcode = NVPTX::LDV_i16_v4_ari;
break;
case MVT::i32:
Opcode = NVPTX::LDV_i32_v4_ari;
break;
case MVT::f32:
Opcode = NVPTX::LDV_f32_v4_ari;
break;
}
break;
}
}
SDValue Ops[] = { getI32Imm(IsVolatile), getI32Imm(CodeAddrSpace),
getI32Imm(VecType), getI32Imm(FromType),
getI32Imm(FromTypeWidth), Base, Offset, Chain };
LD = CurDAG->getMachineNode(Opcode, DL, N->getVTList(), Ops, 8);
} else {
if (Subtarget.is64Bit()) {
switch (N->getOpcode()) {
default:
return NULL;
case NVPTXISD::LoadV2:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::LDV_i8_v2_areg_64;
break;
case MVT::i16:
Opcode = NVPTX::LDV_i16_v2_areg_64;
break;
case MVT::i32:
Opcode = NVPTX::LDV_i32_v2_areg_64;
break;
case MVT::i64:
Opcode = NVPTX::LDV_i64_v2_areg_64;
break;
case MVT::f32:
Opcode = NVPTX::LDV_f32_v2_areg_64;
break;
case MVT::f64:
Opcode = NVPTX::LDV_f64_v2_areg_64;
break;
}
break;
case NVPTXISD::LoadV4:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::LDV_i8_v4_areg_64;
break;
case MVT::i16:
Opcode = NVPTX::LDV_i16_v4_areg_64;
break;
case MVT::i32:
Opcode = NVPTX::LDV_i32_v4_areg_64;
break;
case MVT::f32:
Opcode = NVPTX::LDV_f32_v4_areg_64;
break;
}
break;
}
} else {
switch (N->getOpcode()) {
default:
return NULL;
case NVPTXISD::LoadV2:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::LDV_i8_v2_areg;
break;
case MVT::i16:
Opcode = NVPTX::LDV_i16_v2_areg;
break;
case MVT::i32:
Opcode = NVPTX::LDV_i32_v2_areg;
break;
case MVT::i64:
Opcode = NVPTX::LDV_i64_v2_areg;
break;
case MVT::f32:
Opcode = NVPTX::LDV_f32_v2_areg;
break;
case MVT::f64:
Opcode = NVPTX::LDV_f64_v2_areg;
break;
}
break;
case NVPTXISD::LoadV4:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::LDV_i8_v4_areg;
break;
case MVT::i16:
Opcode = NVPTX::LDV_i16_v4_areg;
break;
case MVT::i32:
Opcode = NVPTX::LDV_i32_v4_areg;
break;
case MVT::f32:
Opcode = NVPTX::LDV_f32_v4_areg;
break;
}
break;
}
}
SDValue Ops[] = { getI32Imm(IsVolatile), getI32Imm(CodeAddrSpace),
getI32Imm(VecType), getI32Imm(FromType),
getI32Imm(FromTypeWidth), Op1, Chain };
LD = CurDAG->getMachineNode(Opcode, DL, N->getVTList(), Ops, 7);
}
MachineSDNode::mmo_iterator MemRefs0 = MF->allocateMemRefsArray(1);
MemRefs0[0] = cast<MemSDNode>(N)->getMemOperand();
cast<MachineSDNode>(LD)->setMemRefs(MemRefs0, MemRefs0 + 1);
return LD;
}
SDNode *NVPTXDAGToDAGISel::SelectLoad(SDNode *N) {
DebugLoc dl = N->getDebugLoc();
LoadSDNode *LD = cast<LoadSDNode>(N);
EVT LoadedVT = LD->getMemoryVT();
SDNode *NVPTXLD = NULL;
// do not support pre/post inc/dec
if (LD->isIndexed())
return NULL;
if (!LoadedVT.isSimple())
return NULL;
// Address Space Setting
unsigned int codeAddrSpace = getCodeAddrSpace(LD, Subtarget);
// Volatile Setting
// - .volatile is only availalble for .global and .shared
bool isVolatile = LD->isVolatile();
if (codeAddrSpace != NVPTX::PTXLdStInstCode::GLOBAL &&
codeAddrSpace != NVPTX::PTXLdStInstCode::SHARED &&
codeAddrSpace != NVPTX::PTXLdStInstCode::GENERIC)
isVolatile = false;
// Vector Setting
MVT SimpleVT = LoadedVT.getSimpleVT();
unsigned vecType = NVPTX::PTXLdStInstCode::Scalar;
if (SimpleVT.isVector()) {
unsigned num = SimpleVT.getVectorNumElements();
if (num == 2)
vecType = NVPTX::PTXLdStInstCode::V2;
else if (num == 4)
vecType = NVPTX::PTXLdStInstCode::V4;
else
return NULL;
}
// Type Setting: fromType + fromTypeWidth
//
// Sign : ISD::SEXTLOAD
// Unsign : ISD::ZEXTLOAD, ISD::NON_EXTLOAD or ISD::EXTLOAD and the
// type is integer
// Float : ISD::NON_EXTLOAD or ISD::EXTLOAD and the type is float
MVT ScalarVT = SimpleVT.getScalarType();
unsigned fromTypeWidth = ScalarVT.getSizeInBits();
unsigned int fromType;
if ((LD->getExtensionType() == ISD::SEXTLOAD))
fromType = NVPTX::PTXLdStInstCode::Signed;
else if (ScalarVT.isFloatingPoint())
fromType = NVPTX::PTXLdStInstCode::Float;
else
fromType = NVPTX::PTXLdStInstCode::Unsigned;
// Create the machine instruction DAG
SDValue Chain = N->getOperand(0);
SDValue N1 = N->getOperand(1);
SDValue Addr;
SDValue Offset, Base;
unsigned Opcode;
MVT::SimpleValueType TargetVT = LD->getValueType(0).getSimpleVT().SimpleTy;
if (SelectDirectAddr(N1, Addr)) {
switch (TargetVT) {
case MVT::i8:
Opcode = NVPTX::LD_i8_avar;
break;
case MVT::i16:
Opcode = NVPTX::LD_i16_avar;
break;
case MVT::i32:
Opcode = NVPTX::LD_i32_avar;
break;
case MVT::i64:
Opcode = NVPTX::LD_i64_avar;
break;
case MVT::f32:
Opcode = NVPTX::LD_f32_avar;
break;
case MVT::f64:
Opcode = NVPTX::LD_f64_avar;
break;
default:
return NULL;
}
SDValue Ops[] = { getI32Imm(isVolatile), getI32Imm(codeAddrSpace),
getI32Imm(vecType), getI32Imm(fromType),
getI32Imm(fromTypeWidth), Addr, Chain };
NVPTXLD = CurDAG->getMachineNode(Opcode, dl, TargetVT, MVT::Other, Ops, 7);
} else if (Subtarget.is64Bit()
? SelectADDRsi64(N1.getNode(), N1, Base, Offset)
: SelectADDRsi(N1.getNode(), N1, Base, Offset)) {
switch (TargetVT) {
case MVT::i8:
Opcode = NVPTX::LD_i8_asi;
break;
case MVT::i16:
Opcode = NVPTX::LD_i16_asi;
break;
case MVT::i32:
Opcode = NVPTX::LD_i32_asi;
break;
case MVT::i64:
Opcode = NVPTX::LD_i64_asi;
break;
case MVT::f32:
Opcode = NVPTX::LD_f32_asi;
break;
case MVT::f64:
Opcode = NVPTX::LD_f64_asi;
break;
default:
return NULL;
}
SDValue Ops[] = { getI32Imm(isVolatile), getI32Imm(codeAddrSpace),
getI32Imm(vecType), getI32Imm(fromType),
getI32Imm(fromTypeWidth), Base, Offset, Chain };
NVPTXLD = CurDAG->getMachineNode(Opcode, dl, TargetVT, MVT::Other, Ops, 8);
} else if (Subtarget.is64Bit()
? SelectADDRri64(N1.getNode(), N1, Base, Offset)
: SelectADDRri(N1.getNode(), N1, Base, Offset)) {
if (Subtarget.is64Bit()) {
switch (TargetVT) {
case MVT::i8:
Opcode = NVPTX::LD_i8_ari_64;
break;
case MVT::i16:
Opcode = NVPTX::LD_i16_ari_64;
break;
case MVT::i32:
Opcode = NVPTX::LD_i32_ari_64;
break;
case MVT::i64:
Opcode = NVPTX::LD_i64_ari_64;
break;
case MVT::f32:
Opcode = NVPTX::LD_f32_ari_64;
break;
case MVT::f64:
Opcode = NVPTX::LD_f64_ari_64;
break;
default:
return NULL;
}
} else {
switch (TargetVT) {
case MVT::i8:
Opcode = NVPTX::LD_i8_ari;
break;
case MVT::i16:
Opcode = NVPTX::LD_i16_ari;
break;
case MVT::i32:
Opcode = NVPTX::LD_i32_ari;
break;
case MVT::i64:
Opcode = NVPTX::LD_i64_ari;
break;
case MVT::f32:
Opcode = NVPTX::LD_f32_ari;
break;
case MVT::f64:
Opcode = NVPTX::LD_f64_ari;
break;
default:
return NULL;
}
}
SDValue Ops[] = { getI32Imm(isVolatile), getI32Imm(codeAddrSpace),
getI32Imm(vecType), getI32Imm(fromType),
getI32Imm(fromTypeWidth), Base, Offset, Chain };
NVPTXLD = CurDAG->getMachineNode(Opcode, dl, TargetVT, MVT::Other, Ops, 8);
} else {
if (Subtarget.is64Bit()) {
switch (TargetVT) {
case MVT::i8:
Opcode = NVPTX::LD_i8_areg_64;
break;
case MVT::i16:
Opcode = NVPTX::LD_i16_areg_64;
break;
case MVT::i32:
Opcode = NVPTX::LD_i32_areg_64;
break;
case MVT::i64:
Opcode = NVPTX::LD_i64_areg_64;
break;
case MVT::f32:
Opcode = NVPTX::LD_f32_areg_64;
break;
case MVT::f64:
Opcode = NVPTX::LD_f64_areg_64;
break;
default:
return NULL;
}
} else {
switch (TargetVT) {
case MVT::i8:
Opcode = NVPTX::LD_i8_areg;
break;
case MVT::i16:
Opcode = NVPTX::LD_i16_areg;
break;
case MVT::i32:
Opcode = NVPTX::LD_i32_areg;
break;
case MVT::i64:
Opcode = NVPTX::LD_i64_areg;
break;
case MVT::f32:
Opcode = NVPTX::LD_f32_areg;
break;
case MVT::f64:
Opcode = NVPTX::LD_f64_areg;
break;
default:
return NULL;
}
}
SDValue Ops[] = { getI32Imm(isVolatile), getI32Imm(codeAddrSpace),
getI32Imm(vecType), getI32Imm(fromType),
getI32Imm(fromTypeWidth), N1, Chain };
NVPTXLD = CurDAG->getMachineNode(Opcode, dl, TargetVT, MVT::Other, Ops, 7);
}
if (NVPTXLD != NULL) {
MachineSDNode::mmo_iterator MemRefs0 = MF->allocateMemRefsArray(1);
MemRefs0[0] = cast<MemSDNode>(N)->getMemOperand();
cast<MachineSDNode>(NVPTXLD)->setMemRefs(MemRefs0, MemRefs0 + 1);
}
return NVPTXLD;
}
SDNode *NVPTXDAGToDAGISel::SelectStoreVector(SDNode *N) {
SDValue Chain = N->getOperand(0);
SDValue Op1 = N->getOperand(1);
SDValue Addr, Offset, Base;
unsigned Opcode;
DebugLoc DL = N->getDebugLoc();
SDNode *ST;
EVT EltVT = Op1.getValueType();
MemSDNode *MemSD = cast<MemSDNode>(N);
EVT StoreVT = MemSD->getMemoryVT();
// Address Space Setting
unsigned CodeAddrSpace = getCodeAddrSpace(MemSD, Subtarget);
if (CodeAddrSpace == NVPTX::PTXLdStInstCode::CONSTANT) {
report_fatal_error("Cannot store to pointer that points to constant "
"memory space");
}
// Volatile Setting
// - .volatile is only availalble for .global and .shared
bool IsVolatile = MemSD->isVolatile();
if (CodeAddrSpace != NVPTX::PTXLdStInstCode::GLOBAL &&
CodeAddrSpace != NVPTX::PTXLdStInstCode::SHARED &&
CodeAddrSpace != NVPTX::PTXLdStInstCode::GENERIC)
IsVolatile = false;
// Type Setting: toType + toTypeWidth
// - for integer type, always use 'u'
assert(StoreVT.isSimple() && "Store value is not simple");
MVT ScalarVT = StoreVT.getSimpleVT().getScalarType();
unsigned ToTypeWidth = ScalarVT.getSizeInBits();
unsigned ToType;
if (ScalarVT.isFloatingPoint())
ToType = NVPTX::PTXLdStInstCode::Float;
else
ToType = NVPTX::PTXLdStInstCode::Unsigned;
SmallVector<SDValue, 12> StOps;
SDValue N2;
unsigned VecType;
switch (N->getOpcode()) {
case NVPTXISD::StoreV2:
VecType = NVPTX::PTXLdStInstCode::V2;
StOps.push_back(N->getOperand(1));
StOps.push_back(N->getOperand(2));
N2 = N->getOperand(3);
break;
case NVPTXISD::StoreV4:
VecType = NVPTX::PTXLdStInstCode::V4;
StOps.push_back(N->getOperand(1));
StOps.push_back(N->getOperand(2));
StOps.push_back(N->getOperand(3));
StOps.push_back(N->getOperand(4));
N2 = N->getOperand(5);
break;
default:
return NULL;
}
StOps.push_back(getI32Imm(IsVolatile));
StOps.push_back(getI32Imm(CodeAddrSpace));
StOps.push_back(getI32Imm(VecType));
StOps.push_back(getI32Imm(ToType));
StOps.push_back(getI32Imm(ToTypeWidth));
if (SelectDirectAddr(N2, Addr)) {
switch (N->getOpcode()) {
default:
return NULL;
case NVPTXISD::StoreV2:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::STV_i8_v2_avar;
break;
case MVT::i16:
Opcode = NVPTX::STV_i16_v2_avar;
break;
case MVT::i32:
Opcode = NVPTX::STV_i32_v2_avar;
break;
case MVT::i64:
Opcode = NVPTX::STV_i64_v2_avar;
break;
case MVT::f32:
Opcode = NVPTX::STV_f32_v2_avar;
break;
case MVT::f64:
Opcode = NVPTX::STV_f64_v2_avar;
break;
}
break;
case NVPTXISD::StoreV4:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::STV_i8_v4_avar;
break;
case MVT::i16:
Opcode = NVPTX::STV_i16_v4_avar;
break;
case MVT::i32:
Opcode = NVPTX::STV_i32_v4_avar;
break;
case MVT::f32:
Opcode = NVPTX::STV_f32_v4_avar;
break;
}
break;
}
StOps.push_back(Addr);
} else if (Subtarget.is64Bit()
? SelectADDRsi64(N2.getNode(), N2, Base, Offset)
: SelectADDRsi(N2.getNode(), N2, Base, Offset)) {
switch (N->getOpcode()) {
default:
return NULL;
case NVPTXISD::StoreV2:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::STV_i8_v2_asi;
break;
case MVT::i16:
Opcode = NVPTX::STV_i16_v2_asi;
break;
case MVT::i32:
Opcode = NVPTX::STV_i32_v2_asi;
break;
case MVT::i64:
Opcode = NVPTX::STV_i64_v2_asi;
break;
case MVT::f32:
Opcode = NVPTX::STV_f32_v2_asi;
break;
case MVT::f64:
Opcode = NVPTX::STV_f64_v2_asi;
break;
}
break;
case NVPTXISD::StoreV4:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::STV_i8_v4_asi;
break;
case MVT::i16:
Opcode = NVPTX::STV_i16_v4_asi;
break;
case MVT::i32:
Opcode = NVPTX::STV_i32_v4_asi;
break;
case MVT::f32:
Opcode = NVPTX::STV_f32_v4_asi;
break;
}
break;
}
StOps.push_back(Base);
StOps.push_back(Offset);
} else if (Subtarget.is64Bit()
? SelectADDRri64(N2.getNode(), N2, Base, Offset)
: SelectADDRri(N2.getNode(), N2, Base, Offset)) {
if (Subtarget.is64Bit()) {
switch (N->getOpcode()) {
default:
return NULL;
case NVPTXISD::StoreV2:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::STV_i8_v2_ari_64;
break;
case MVT::i16:
Opcode = NVPTX::STV_i16_v2_ari_64;
break;
case MVT::i32:
Opcode = NVPTX::STV_i32_v2_ari_64;
break;
case MVT::i64:
Opcode = NVPTX::STV_i64_v2_ari_64;
break;
case MVT::f32:
Opcode = NVPTX::STV_f32_v2_ari_64;
break;
case MVT::f64:
Opcode = NVPTX::STV_f64_v2_ari_64;
break;
}
break;
case NVPTXISD::StoreV4:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::STV_i8_v4_ari_64;
break;
case MVT::i16:
Opcode = NVPTX::STV_i16_v4_ari_64;
break;
case MVT::i32:
Opcode = NVPTX::STV_i32_v4_ari_64;
break;
case MVT::f32:
Opcode = NVPTX::STV_f32_v4_ari_64;
break;
}
break;
}
} else {
switch (N->getOpcode()) {
default:
return NULL;
case NVPTXISD::StoreV2:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::STV_i8_v2_ari;
break;
case MVT::i16:
Opcode = NVPTX::STV_i16_v2_ari;
break;
case MVT::i32:
Opcode = NVPTX::STV_i32_v2_ari;
break;
case MVT::i64:
Opcode = NVPTX::STV_i64_v2_ari;
break;
case MVT::f32:
Opcode = NVPTX::STV_f32_v2_ari;
break;
case MVT::f64:
Opcode = NVPTX::STV_f64_v2_ari;
break;
}
break;
case NVPTXISD::StoreV4:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::STV_i8_v4_ari;
break;
case MVT::i16:
Opcode = NVPTX::STV_i16_v4_ari;
break;
case MVT::i32:
Opcode = NVPTX::STV_i32_v4_ari;
break;
case MVT::f32:
Opcode = NVPTX::STV_f32_v4_ari;
break;
}
break;
}
}
StOps.push_back(Base);
StOps.push_back(Offset);
} else {
if (Subtarget.is64Bit()) {
switch (N->getOpcode()) {
default:
return NULL;
case NVPTXISD::StoreV2:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::STV_i8_v2_areg_64;
break;
case MVT::i16:
Opcode = NVPTX::STV_i16_v2_areg_64;
break;
case MVT::i32:
Opcode = NVPTX::STV_i32_v2_areg_64;
break;
case MVT::i64:
Opcode = NVPTX::STV_i64_v2_areg_64;
break;
case MVT::f32:
Opcode = NVPTX::STV_f32_v2_areg_64;
break;
case MVT::f64:
Opcode = NVPTX::STV_f64_v2_areg_64;
break;
}
break;
case NVPTXISD::StoreV4:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::STV_i8_v4_areg_64;
break;
case MVT::i16:
Opcode = NVPTX::STV_i16_v4_areg_64;
break;
case MVT::i32:
Opcode = NVPTX::STV_i32_v4_areg_64;
break;
case MVT::f32:
Opcode = NVPTX::STV_f32_v4_areg_64;
break;
}
break;
}
} else {
switch (N->getOpcode()) {
default:
return NULL;
case NVPTXISD::StoreV2:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::STV_i8_v2_areg;
break;
case MVT::i16:
Opcode = NVPTX::STV_i16_v2_areg;
break;
case MVT::i32:
Opcode = NVPTX::STV_i32_v2_areg;
break;
case MVT::i64:
Opcode = NVPTX::STV_i64_v2_areg;
break;
case MVT::f32:
Opcode = NVPTX::STV_f32_v2_areg;
break;
case MVT::f64:
Opcode = NVPTX::STV_f64_v2_areg;
break;
}
break;
case NVPTXISD::StoreV4:
switch (EltVT.getSimpleVT().SimpleTy) {
default:
return NULL;
case MVT::i8:
Opcode = NVPTX::STV_i8_v4_areg;
break;
case MVT::i16:
Opcode = NVPTX::STV_i16_v4_areg;
break;
case MVT::i32:
Opcode = NVPTX::STV_i32_v4_areg;
break;
case MVT::f32:
Opcode = NVPTX::STV_f32_v4_areg;
break;
}
break;
}
}
StOps.push_back(N2);
}
StOps.push_back(Chain);
ST = CurDAG->getMachineNode(Opcode, DL, MVT::Other, &StOps[0], StOps.size());
MachineSDNode::mmo_iterator MemRefs0 = MF->allocateMemRefsArray(1);
MemRefs0[0] = cast<MemSDNode>(N)->getMemOperand();
cast<MachineSDNode>(ST)->setMemRefs(MemRefs0, MemRefs0 + 1);
return ST;
}
