1 //===- SPIRVInstructionSelector.cpp ------------------------------*- C++ -*-==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // This file implements the targeting of the InstructionSelector class for
10 // SPIRV.
11 // TODO: This should be generated by TableGen.
12 //
13 //===----------------------------------------------------------------------===//
14
15 #include "MCTargetDesc/SPIRVBaseInfo.h"
16 #include "MCTargetDesc/SPIRVMCTargetDesc.h"
17 #include "SPIRV.h"
18 #include "SPIRVGlobalRegistry.h"
19 #include "SPIRVInstrInfo.h"
20 #include "SPIRVRegisterInfo.h"
21 #include "SPIRVTargetMachine.h"
22 #include "SPIRVUtils.h"
23 #include "llvm/ADT/APFloat.h"
24 #include "llvm/ADT/StringExtras.h"
25 #include "llvm/CodeGen/GlobalISel/GIMatchTableExecutorImpl.h"
26 #include "llvm/CodeGen/GlobalISel/GenericMachineInstrs.h"
27 #include "llvm/CodeGen/GlobalISel/InstructionSelector.h"
28 #include "llvm/CodeGen/MachineInstrBuilder.h"
29 #include "llvm/CodeGen/MachineRegisterInfo.h"
30 #include "llvm/CodeGen/Register.h"
31 #include "llvm/CodeGen/TargetOpcodes.h"
32 #include "llvm/IR/IntrinsicsSPIRV.h"
33 #include "llvm/Support/Debug.h"
34 #include "llvm/Support/ErrorHandling.h"
35
36 #define DEBUG_TYPE "spirv-isel"
37
38 using namespace llvm;
39 namespace CL = SPIRV::OpenCLExtInst;
40 namespace GL = SPIRV::GLSLExtInst;
41
42 using ExtInstList =
43 std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;
44
45 namespace {
46
47 llvm::SPIRV::SelectionControl::SelectionControl
getSelectionOperandForImm(int Imm)48 getSelectionOperandForImm(int Imm) {
49 if (Imm == 2)
50 return SPIRV::SelectionControl::Flatten;
51 if (Imm == 1)
52 return SPIRV::SelectionControl::DontFlatten;
53 if (Imm == 0)
54 return SPIRV::SelectionControl::None;
55 llvm_unreachable("Invalid immediate");
56 }
57
58 #define GET_GLOBALISEL_PREDICATE_BITSET
59 #include "SPIRVGenGlobalISel.inc"
60 #undef GET_GLOBALISEL_PREDICATE_BITSET
61
62 class SPIRVInstructionSelector : public InstructionSelector {
63 const SPIRVSubtarget &STI;
64 const SPIRVInstrInfo &TII;
65 const SPIRVRegisterInfo &TRI;
66 const RegisterBankInfo &RBI;
67 SPIRVGlobalRegistry &GR;
68 MachineRegisterInfo *MRI;
69 MachineFunction *HasVRegsReset = nullptr;
70
71 /// We need to keep track of the number we give to anonymous global values to
72 /// generate the same name every time when this is needed.
73 mutable DenseMap<const GlobalValue *, unsigned> UnnamedGlobalIDs;
74 SmallPtrSet<MachineInstr *, 8> DeadMIs;
75
76 public:
77 SPIRVInstructionSelector(const SPIRVTargetMachine &TM,
78 const SPIRVSubtarget &ST,
79 const RegisterBankInfo &RBI);
80 void setupMF(MachineFunction &MF, GISelValueTracking *VT,
81 CodeGenCoverage *CoverageInfo, ProfileSummaryInfo *PSI,
82 BlockFrequencyInfo *BFI) override;
83 // Common selection code. Instruction-specific selection occurs in spvSelect.
84 bool select(MachineInstr &I) override;
getName()85 static const char *getName() { return DEBUG_TYPE; }
86
87 #define GET_GLOBALISEL_PREDICATES_DECL
88 #include "SPIRVGenGlobalISel.inc"
89 #undef GET_GLOBALISEL_PREDICATES_DECL
90
91 #define GET_GLOBALISEL_TEMPORARIES_DECL
92 #include "SPIRVGenGlobalISel.inc"
93 #undef GET_GLOBALISEL_TEMPORARIES_DECL
94
95 private:
96 void resetVRegsType(MachineFunction &MF);
97
98 // tblgen-erated 'select' implementation, used as the initial selector for
99 // the patterns that don't require complex C++.
100 bool selectImpl(MachineInstr &I, CodeGenCoverage &CoverageInfo) const;
101
102 // All instruction-specific selection that didn't happen in "select()".
103 // Is basically a large Switch/Case delegating to all other select method.
104 bool spvSelect(Register ResVReg, const SPIRVType *ResType,
105 MachineInstr &I) const;
106
107 bool selectFirstBitHigh(Register ResVReg, const SPIRVType *ResType,
108 MachineInstr &I, bool IsSigned) const;
109
110 bool selectFirstBitLow(Register ResVReg, const SPIRVType *ResType,
111 MachineInstr &I) const;
112
113 bool selectFirstBitSet16(Register ResVReg, const SPIRVType *ResType,
114 MachineInstr &I, unsigned ExtendOpcode,
115 unsigned BitSetOpcode) const;
116
117 bool selectFirstBitSet32(Register ResVReg, const SPIRVType *ResType,
118 MachineInstr &I, Register SrcReg,
119 unsigned BitSetOpcode) const;
120
121 bool selectFirstBitSet64(Register ResVReg, const SPIRVType *ResType,
122 MachineInstr &I, Register SrcReg,
123 unsigned BitSetOpcode, bool SwapPrimarySide) const;
124
125 bool selectFirstBitSet64Overflow(Register ResVReg, const SPIRVType *ResType,
126 MachineInstr &I, Register SrcReg,
127 unsigned BitSetOpcode,
128 bool SwapPrimarySide) const;
129
130 bool selectGlobalValue(Register ResVReg, MachineInstr &I,
131 const MachineInstr *Init = nullptr) const;
132
133 bool selectOpWithSrcs(Register ResVReg, const SPIRVType *ResType,
134 MachineInstr &I, std::vector<Register> SrcRegs,
135 unsigned Opcode) const;
136
137 bool selectUnOp(Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
138 unsigned Opcode) const;
139
140 bool selectBitcast(Register ResVReg, const SPIRVType *ResType,
141 MachineInstr &I) const;
142
143 bool selectLoad(Register ResVReg, const SPIRVType *ResType,
144 MachineInstr &I) const;
145 bool selectStore(MachineInstr &I) const;
146
147 bool selectStackSave(Register ResVReg, const SPIRVType *ResType,
148 MachineInstr &I) const;
149 bool selectStackRestore(MachineInstr &I) const;
150
151 bool selectMemOperation(Register ResVReg, MachineInstr &I) const;
152
153 bool selectAtomicRMW(Register ResVReg, const SPIRVType *ResType,
154 MachineInstr &I, unsigned NewOpcode,
155 unsigned NegateOpcode = 0) const;
156
157 bool selectAtomicCmpXchg(Register ResVReg, const SPIRVType *ResType,
158 MachineInstr &I) const;
159
160 bool selectFence(MachineInstr &I) const;
161
162 bool selectAddrSpaceCast(Register ResVReg, const SPIRVType *ResType,
163 MachineInstr &I) const;
164
165 bool selectAnyOrAll(Register ResVReg, const SPIRVType *ResType,
166 MachineInstr &I, unsigned OpType) const;
167
168 bool selectAll(Register ResVReg, const SPIRVType *ResType,
169 MachineInstr &I) const;
170
171 bool selectAny(Register ResVReg, const SPIRVType *ResType,
172 MachineInstr &I) const;
173
174 bool selectBitreverse(Register ResVReg, const SPIRVType *ResType,
175 MachineInstr &I) const;
176
177 bool selectBuildVector(Register ResVReg, const SPIRVType *ResType,
178 MachineInstr &I) const;
179 bool selectSplatVector(Register ResVReg, const SPIRVType *ResType,
180 MachineInstr &I) const;
181
182 bool selectCmp(Register ResVReg, const SPIRVType *ResType,
183 unsigned comparisonOpcode, MachineInstr &I) const;
184 bool selectDiscard(Register ResVReg, const SPIRVType *ResType,
185 MachineInstr &I) const;
186
187 bool selectICmp(Register ResVReg, const SPIRVType *ResType,
188 MachineInstr &I) const;
189 bool selectFCmp(Register ResVReg, const SPIRVType *ResType,
190 MachineInstr &I) const;
191
192 bool selectSign(Register ResVReg, const SPIRVType *ResType,
193 MachineInstr &I) const;
194
195 bool selectFloatDot(Register ResVReg, const SPIRVType *ResType,
196 MachineInstr &I) const;
197
198 bool selectOverflowArith(Register ResVReg, const SPIRVType *ResType,
199 MachineInstr &I, unsigned Opcode) const;
200
201 bool selectIntegerDot(Register ResVReg, const SPIRVType *ResType,
202 MachineInstr &I, bool Signed) const;
203
204 bool selectIntegerDotExpansion(Register ResVReg, const SPIRVType *ResType,
205 MachineInstr &I) const;
206
207 template <bool Signed>
208 bool selectDot4AddPacked(Register ResVReg, const SPIRVType *ResType,
209 MachineInstr &I) const;
210 template <bool Signed>
211 bool selectDot4AddPackedExpansion(Register ResVReg, const SPIRVType *ResType,
212 MachineInstr &I) const;
213
214 bool selectWaveReduceMax(Register ResVReg, const SPIRVType *ResType,
215 MachineInstr &I, bool IsUnsigned) const;
216
217 bool selectWaveReduceSum(Register ResVReg, const SPIRVType *ResType,
218 MachineInstr &I) const;
219
220 bool selectConst(Register ResVReg, const SPIRVType *ResType,
221 MachineInstr &I) const;
222
223 bool selectSelect(Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
224 bool IsSigned) const;
225 bool selectIToF(Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
226 bool IsSigned, unsigned Opcode) const;
227 bool selectExt(Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
228 bool IsSigned) const;
229
230 bool selectTrunc(Register ResVReg, const SPIRVType *ResType,
231 MachineInstr &I) const;
232
233 bool selectSUCmp(Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
234 bool IsSigned) const;
235
236 bool selectIntToBool(Register IntReg, Register ResVReg, MachineInstr &I,
237 const SPIRVType *intTy, const SPIRVType *boolTy) const;
238
239 bool selectOpUndef(Register ResVReg, const SPIRVType *ResType,
240 MachineInstr &I) const;
241 bool selectFreeze(Register ResVReg, const SPIRVType *ResType,
242 MachineInstr &I) const;
243 bool selectIntrinsic(Register ResVReg, const SPIRVType *ResType,
244 MachineInstr &I) const;
245 bool selectExtractVal(Register ResVReg, const SPIRVType *ResType,
246 MachineInstr &I) const;
247 bool selectInsertVal(Register ResVReg, const SPIRVType *ResType,
248 MachineInstr &I) const;
249 bool selectExtractElt(Register ResVReg, const SPIRVType *ResType,
250 MachineInstr &I) const;
251 bool selectInsertElt(Register ResVReg, const SPIRVType *ResType,
252 MachineInstr &I) const;
253 bool selectGEP(Register ResVReg, const SPIRVType *ResType,
254 MachineInstr &I) const;
255
256 bool selectFrameIndex(Register ResVReg, const SPIRVType *ResType,
257 MachineInstr &I) const;
258 bool selectAllocaArray(Register ResVReg, const SPIRVType *ResType,
259 MachineInstr &I) const;
260
261 bool selectBranch(MachineInstr &I) const;
262 bool selectBranchCond(MachineInstr &I) const;
263
264 bool selectPhi(Register ResVReg, const SPIRVType *ResType,
265 MachineInstr &I) const;
266
267 bool selectExtInst(Register ResVReg, const SPIRVType *RestType,
268 MachineInstr &I, GL::GLSLExtInst GLInst) const;
269 bool selectExtInst(Register ResVReg, const SPIRVType *ResType,
270 MachineInstr &I, CL::OpenCLExtInst CLInst) const;
271 bool selectExtInst(Register ResVReg, const SPIRVType *ResType,
272 MachineInstr &I, CL::OpenCLExtInst CLInst,
273 GL::GLSLExtInst GLInst) const;
274 bool selectExtInst(Register ResVReg, const SPIRVType *ResType,
275 MachineInstr &I, const ExtInstList &ExtInsts) const;
276
277 bool selectLog10(Register ResVReg, const SPIRVType *ResType,
278 MachineInstr &I) const;
279
280 bool selectSaturate(Register ResVReg, const SPIRVType *ResType,
281 MachineInstr &I) const;
282
283 bool selectWaveOpInst(Register ResVReg, const SPIRVType *ResType,
284 MachineInstr &I, unsigned Opcode) const;
285
286 bool selectWaveActiveCountBits(Register ResVReg, const SPIRVType *ResType,
287 MachineInstr &I) const;
288
289 bool selectUnmergeValues(MachineInstr &I) const;
290
291 bool selectHandleFromBinding(Register &ResVReg, const SPIRVType *ResType,
292 MachineInstr &I) const;
293
294 bool selectReadImageIntrinsic(Register &ResVReg, const SPIRVType *ResType,
295 MachineInstr &I) const;
296 bool selectImageWriteIntrinsic(MachineInstr &I) const;
297 bool selectResourceGetPointer(Register &ResVReg, const SPIRVType *ResType,
298 MachineInstr &I) const;
299
300 // Utilities
301 std::pair<Register, bool>
302 buildI32Constant(uint32_t Val, MachineInstr &I,
303 const SPIRVType *ResType = nullptr) const;
304
305 Register buildZerosVal(const SPIRVType *ResType, MachineInstr &I) const;
306 Register buildZerosValF(const SPIRVType *ResType, MachineInstr &I) const;
307 Register buildOnesVal(bool AllOnes, const SPIRVType *ResType,
308 MachineInstr &I) const;
309 Register buildOnesValF(const SPIRVType *ResType, MachineInstr &I) const;
310
311 bool wrapIntoSpecConstantOp(MachineInstr &I,
312 SmallVector<Register> &CompositeArgs) const;
313
314 Register getUcharPtrTypeReg(MachineInstr &I,
315 SPIRV::StorageClass::StorageClass SC) const;
316 MachineInstrBuilder buildSpecConstantOp(MachineInstr &I, Register Dest,
317 Register Src, Register DestType,
318 uint32_t Opcode) const;
319 MachineInstrBuilder buildConstGenericPtr(MachineInstr &I, Register SrcPtr,
320 SPIRVType *SrcPtrTy) const;
321 Register buildPointerToResource(const SPIRVType *ResType,
322 SPIRV::StorageClass::StorageClass SC,
323 uint32_t Set, uint32_t Binding,
324 uint32_t ArraySize, Register IndexReg,
325 bool IsNonUniform, StringRef Name,
326 MachineIRBuilder MIRBuilder) const;
327 SPIRVType *widenTypeToVec4(const SPIRVType *Type, MachineInstr &I) const;
328 bool extractSubvector(Register &ResVReg, const SPIRVType *ResType,
329 Register &ReadReg, MachineInstr &InsertionPoint) const;
330 bool generateImageRead(Register &ResVReg, const SPIRVType *ResType,
331 Register ImageReg, Register IdxReg, DebugLoc Loc,
332 MachineInstr &Pos) const;
333 bool BuildCOPY(Register DestReg, Register SrcReg, MachineInstr &I) const;
334 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
335 Register ResVReg, const SPIRVType *ResType,
336 MachineInstr &I) const;
337 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
338 Register ResVReg, const SPIRVType *ResType,
339 MachineInstr &I) const;
340 bool loadHandleBeforePosition(Register &HandleReg, const SPIRVType *ResType,
341 GIntrinsic &HandleDef, MachineInstr &Pos) const;
342 };
343
sampledTypeIsSignedInteger(const llvm::Type * HandleType)344 bool sampledTypeIsSignedInteger(const llvm::Type *HandleType) {
345 const TargetExtType *TET = cast<TargetExtType>(HandleType);
346 if (TET->getTargetExtName() == "spirv.Image") {
347 return false;
348 }
349 assert(TET->getTargetExtName() == "spirv.SignedImage");
350 return TET->getTypeParameter(0)->isIntegerTy();
351 }
352 } // end anonymous namespace
353
354 #define GET_GLOBALISEL_IMPL
355 #include "SPIRVGenGlobalISel.inc"
356 #undef GET_GLOBALISEL_IMPL
357
SPIRVInstructionSelector(const SPIRVTargetMachine & TM,const SPIRVSubtarget & ST,const RegisterBankInfo & RBI)358 SPIRVInstructionSelector::SPIRVInstructionSelector(const SPIRVTargetMachine &TM,
359 const SPIRVSubtarget &ST,
360 const RegisterBankInfo &RBI)
361 : InstructionSelector(), STI(ST), TII(*ST.getInstrInfo()),
362 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
363 #define GET_GLOBALISEL_PREDICATES_INIT
364 #include "SPIRVGenGlobalISel.inc"
365 #undef GET_GLOBALISEL_PREDICATES_INIT
366 #define GET_GLOBALISEL_TEMPORARIES_INIT
367 #include "SPIRVGenGlobalISel.inc"
368 #undef GET_GLOBALISEL_TEMPORARIES_INIT
369 {
370 }
371
setupMF(MachineFunction & MF,GISelValueTracking * VT,CodeGenCoverage * CoverageInfo,ProfileSummaryInfo * PSI,BlockFrequencyInfo * BFI)372 void SPIRVInstructionSelector::setupMF(MachineFunction &MF,
373 GISelValueTracking *VT,
374 CodeGenCoverage *CoverageInfo,
375 ProfileSummaryInfo *PSI,
376 BlockFrequencyInfo *BFI) {
377 MRI = &MF.getRegInfo();
378 GR.setCurrentFunc(MF);
379 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
380 }
381
382 // Ensure that register classes correspond to pattern matching rules.
resetVRegsType(MachineFunction & MF)383 void SPIRVInstructionSelector::resetVRegsType(MachineFunction &MF) {
384 if (HasVRegsReset == &MF)
385 return;
386 HasVRegsReset = &MF;
387
388 MachineRegisterInfo &MRI = MF.getRegInfo();
389 for (unsigned I = 0, E = MRI.getNumVirtRegs(); I != E; ++I) {
390 Register Reg = Register::index2VirtReg(I);
391 LLT RegType = MRI.getType(Reg);
392 if (RegType.isScalar())
393 MRI.setType(Reg, LLT::scalar(64));
394 else if (RegType.isPointer())
395 MRI.setType(Reg, LLT::pointer(0, 64));
396 else if (RegType.isVector())
397 MRI.setType(Reg, LLT::fixed_vector(2, LLT::scalar(64)));
398 }
399 for (const auto &MBB : MF) {
400 for (const auto &MI : MBB) {
401 if (isPreISelGenericOpcode(MI.getOpcode()))
402 GR.erase(&MI);
403 if (MI.getOpcode() != SPIRV::ASSIGN_TYPE)
404 continue;
405
406 Register DstReg = MI.getOperand(0).getReg();
407 LLT DstType = MRI.getType(DstReg);
408 Register SrcReg = MI.getOperand(1).getReg();
409 LLT SrcType = MRI.getType(SrcReg);
410 if (DstType != SrcType)
411 MRI.setType(DstReg, MRI.getType(SrcReg));
412
413 const TargetRegisterClass *DstRC = MRI.getRegClassOrNull(DstReg);
414 const TargetRegisterClass *SrcRC = MRI.getRegClassOrNull(SrcReg);
415 if (DstRC != SrcRC && SrcRC)
416 MRI.setRegClass(DstReg, SrcRC);
417 }
418 }
419 }
420
421 // Return true if the type represents a constant register
isConstReg(MachineRegisterInfo * MRI,MachineInstr * OpDef,SmallPtrSet<SPIRVType *,4> & Visited)422 static bool isConstReg(MachineRegisterInfo *MRI, MachineInstr *OpDef,
423 SmallPtrSet<SPIRVType *, 4> &Visited) {
424 OpDef = passCopy(OpDef, MRI);
425
426 if (Visited.contains(OpDef))
427 return true;
428 Visited.insert(OpDef);
429
430 unsigned Opcode = OpDef->getOpcode();
431 switch (Opcode) {
432 case TargetOpcode::G_CONSTANT:
433 case TargetOpcode::G_FCONSTANT:
434 return true;
435 case TargetOpcode::G_INTRINSIC:
436 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
437 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
438 return cast<GIntrinsic>(*OpDef).getIntrinsicID() ==
439 Intrinsic::spv_const_composite;
440 case TargetOpcode::G_BUILD_VECTOR:
441 case TargetOpcode::G_SPLAT_VECTOR: {
442 for (unsigned i = OpDef->getNumExplicitDefs(); i < OpDef->getNumOperands();
443 i++) {
444 MachineInstr *OpNestedDef =
445 OpDef->getOperand(i).isReg()
446 ? MRI->getVRegDef(OpDef->getOperand(i).getReg())
447 : nullptr;
448 if (OpNestedDef && !isConstReg(MRI, OpNestedDef, Visited))
449 return false;
450 }
451 return true;
452 case SPIRV::OpConstantTrue:
453 case SPIRV::OpConstantFalse:
454 case SPIRV::OpConstantI:
455 case SPIRV::OpConstantF:
456 case SPIRV::OpConstantComposite:
457 case SPIRV::OpConstantCompositeContinuedINTEL:
458 case SPIRV::OpConstantSampler:
459 case SPIRV::OpConstantNull:
460 case SPIRV::OpUndef:
461 case SPIRV::OpConstantFunctionPointerINTEL:
462 return true;
463 }
464 }
465 return false;
466 }
467
468 // Return true if the virtual register represents a constant
isConstReg(MachineRegisterInfo * MRI,Register OpReg)469 static bool isConstReg(MachineRegisterInfo *MRI, Register OpReg) {
470 SmallPtrSet<SPIRVType *, 4> Visited;
471 if (MachineInstr *OpDef = MRI->getVRegDef(OpReg))
472 return isConstReg(MRI, OpDef, Visited);
473 return false;
474 }
475
isDead(const MachineInstr & MI,const MachineRegisterInfo & MRI)476 bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI) {
477 for (const auto &MO : MI.all_defs()) {
478 Register Reg = MO.getReg();
479 if (Reg.isPhysical() || !MRI.use_nodbg_empty(Reg))
480 return false;
481 }
482 if (MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE || MI.isFakeUse() ||
483 MI.isLifetimeMarker())
484 return false;
485 if (MI.isPHI())
486 return true;
487 if (MI.mayStore() || MI.isCall() ||
488 (MI.mayLoad() && MI.hasOrderedMemoryRef()) || MI.isPosition() ||
489 MI.isDebugInstr() || MI.isTerminator() || MI.isJumpTableDebugInfo())
490 return false;
491 return true;
492 }
493
select(MachineInstr & I)494 bool SPIRVInstructionSelector::select(MachineInstr &I) {
495 resetVRegsType(*I.getParent()->getParent());
496
497 assert(I.getParent() && "Instruction should be in a basic block!");
498 assert(I.getParent()->getParent() && "Instruction should be in a function!");
499
500 Register Opcode = I.getOpcode();
501 // If it's not a GMIR instruction, we've selected it already.
502 if (!isPreISelGenericOpcode(Opcode)) {
503 if (Opcode == SPIRV::ASSIGN_TYPE) { // These pseudos aren't needed any more.
504 Register DstReg = I.getOperand(0).getReg();
505 Register SrcReg = I.getOperand(1).getReg();
506 auto *Def = MRI->getVRegDef(SrcReg);
507 if (isTypeFoldingSupported(Def->getOpcode()) &&
508 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
509 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
510 bool Res = selectImpl(I, *CoverageInfo);
511 LLVM_DEBUG({
512 if (!Res && Def->getOpcode() != TargetOpcode::G_CONSTANT) {
513 dbgs() << "Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
514 I.print(dbgs());
515 }
516 });
517 assert(Res || Def->getOpcode() == TargetOpcode::G_CONSTANT);
518 if (Res) {
519 if (!isTriviallyDead(*Def, *MRI) && isDead(*Def, *MRI))
520 DeadMIs.insert(Def);
521 return Res;
522 }
523 }
524 MRI->setRegClass(SrcReg, MRI->getRegClass(DstReg));
525 MRI->replaceRegWith(SrcReg, DstReg);
526 GR.invalidateMachineInstr(&I);
527 I.removeFromParent();
528 return true;
529 } else if (I.getNumDefs() == 1) {
530 // Make all vregs 64 bits (for SPIR-V IDs).
531 MRI->setType(I.getOperand(0).getReg(), LLT::scalar(64));
532 }
533 return constrainSelectedInstRegOperands(I, TII, TRI, RBI);
534 }
535
536 if (DeadMIs.contains(&I)) {
537 // if the instruction has been already made dead by folding it away
538 // erase it
539 LLVM_DEBUG(dbgs() << "Instruction is folded and dead.\n");
540 salvageDebugInfo(*MRI, I);
541 GR.invalidateMachineInstr(&I);
542 I.eraseFromParent();
543 return true;
544 }
545
546 if (I.getNumOperands() != I.getNumExplicitOperands()) {
547 LLVM_DEBUG(errs() << "Generic instr has unexpected implicit operands\n");
548 return false;
549 }
550
551 // Common code for getting return reg+type, and removing selected instr
552 // from parent occurs here. Instr-specific selection happens in spvSelect().
553 bool HasDefs = I.getNumDefs() > 0;
554 Register ResVReg = HasDefs ? I.getOperand(0).getReg() : Register(0);
555 SPIRVType *ResType = HasDefs ? GR.getSPIRVTypeForVReg(ResVReg) : nullptr;
556 assert(!HasDefs || ResType || I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
557 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
558 if (spvSelect(ResVReg, ResType, I)) {
559 if (HasDefs) // Make all vregs 64 bits (for SPIR-V IDs).
560 for (unsigned i = 0; i < I.getNumDefs(); ++i)
561 MRI->setType(I.getOperand(i).getReg(), LLT::scalar(64));
562 GR.invalidateMachineInstr(&I);
563 I.removeFromParent();
564 return true;
565 }
566 return false;
567 }
568
mayApplyGenericSelection(unsigned Opcode)569 static bool mayApplyGenericSelection(unsigned Opcode) {
570 switch (Opcode) {
571 case TargetOpcode::G_CONSTANT:
572 case TargetOpcode::G_FCONSTANT:
573 return false;
574 case TargetOpcode::G_SADDO:
575 case TargetOpcode::G_SSUBO:
576 return true;
577 }
578 return isTypeFoldingSupported(Opcode);
579 }
580
BuildCOPY(Register DestReg,Register SrcReg,MachineInstr & I) const581 bool SPIRVInstructionSelector::BuildCOPY(Register DestReg, Register SrcReg,
582 MachineInstr &I) const {
583 const TargetRegisterClass *DstRC = MRI->getRegClassOrNull(DestReg);
584 const TargetRegisterClass *SrcRC = MRI->getRegClassOrNull(SrcReg);
585 if (DstRC != SrcRC && SrcRC)
586 MRI->setRegClass(DestReg, SrcRC);
587 return BuildMI(*I.getParent(), I, I.getDebugLoc(),
588 TII.get(TargetOpcode::COPY))
589 .addDef(DestReg)
590 .addUse(SrcReg)
591 .constrainAllUses(TII, TRI, RBI);
592 }
593
spvSelect(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const594 bool SPIRVInstructionSelector::spvSelect(Register ResVReg,
595 const SPIRVType *ResType,
596 MachineInstr &I) const {
597 const unsigned Opcode = I.getOpcode();
598 if (mayApplyGenericSelection(Opcode))
599 return selectImpl(I, *CoverageInfo);
600 switch (Opcode) {
601 case TargetOpcode::G_CONSTANT:
602 case TargetOpcode::G_FCONSTANT:
603 return selectConst(ResVReg, ResType, I);
604 case TargetOpcode::G_GLOBAL_VALUE:
605 return selectGlobalValue(ResVReg, I);
606 case TargetOpcode::G_IMPLICIT_DEF:
607 return selectOpUndef(ResVReg, ResType, I);
608 case TargetOpcode::G_FREEZE:
609 return selectFreeze(ResVReg, ResType, I);
610
611 case TargetOpcode::G_INTRINSIC:
612 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
613 case TargetOpcode::G_INTRINSIC_CONVERGENT:
614 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
615 return selectIntrinsic(ResVReg, ResType, I);
616 case TargetOpcode::G_BITREVERSE:
617 return selectBitreverse(ResVReg, ResType, I);
618
619 case TargetOpcode::G_BUILD_VECTOR:
620 return selectBuildVector(ResVReg, ResType, I);
621 case TargetOpcode::G_SPLAT_VECTOR:
622 return selectSplatVector(ResVReg, ResType, I);
623
624 case TargetOpcode::G_SHUFFLE_VECTOR: {
625 MachineBasicBlock &BB = *I.getParent();
626 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpVectorShuffle))
627 .addDef(ResVReg)
628 .addUse(GR.getSPIRVTypeID(ResType))
629 .addUse(I.getOperand(1).getReg())
630 .addUse(I.getOperand(2).getReg());
631 for (auto V : I.getOperand(3).getShuffleMask())
632 MIB.addImm(V);
633 return MIB.constrainAllUses(TII, TRI, RBI);
634 }
635 case TargetOpcode::G_MEMMOVE:
636 case TargetOpcode::G_MEMCPY:
637 case TargetOpcode::G_MEMSET:
638 return selectMemOperation(ResVReg, I);
639
640 case TargetOpcode::G_ICMP:
641 return selectICmp(ResVReg, ResType, I);
642 case TargetOpcode::G_FCMP:
643 return selectFCmp(ResVReg, ResType, I);
644
645 case TargetOpcode::G_FRAME_INDEX:
646 return selectFrameIndex(ResVReg, ResType, I);
647
648 case TargetOpcode::G_LOAD:
649 return selectLoad(ResVReg, ResType, I);
650 case TargetOpcode::G_STORE:
651 return selectStore(I);
652
653 case TargetOpcode::G_BR:
654 return selectBranch(I);
655 case TargetOpcode::G_BRCOND:
656 return selectBranchCond(I);
657
658 case TargetOpcode::G_PHI:
659 return selectPhi(ResVReg, ResType, I);
660
661 case TargetOpcode::G_FPTOSI:
662 return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertFToS);
663 case TargetOpcode::G_FPTOUI:
664 return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertFToU);
665
666 case TargetOpcode::G_SITOFP:
667 return selectIToF(ResVReg, ResType, I, true, SPIRV::OpConvertSToF);
668 case TargetOpcode::G_UITOFP:
669 return selectIToF(ResVReg, ResType, I, false, SPIRV::OpConvertUToF);
670
671 case TargetOpcode::G_CTPOP:
672 return selectUnOp(ResVReg, ResType, I, SPIRV::OpBitCount);
673 case TargetOpcode::G_SMIN:
674 return selectExtInst(ResVReg, ResType, I, CL::s_min, GL::SMin);
675 case TargetOpcode::G_UMIN:
676 return selectExtInst(ResVReg, ResType, I, CL::u_min, GL::UMin);
677
678 case TargetOpcode::G_SMAX:
679 return selectExtInst(ResVReg, ResType, I, CL::s_max, GL::SMax);
680 case TargetOpcode::G_UMAX:
681 return selectExtInst(ResVReg, ResType, I, CL::u_max, GL::UMax);
682
683 case TargetOpcode::G_SCMP:
684 return selectSUCmp(ResVReg, ResType, I, true);
685 case TargetOpcode::G_UCMP:
686 return selectSUCmp(ResVReg, ResType, I, false);
687
688 case TargetOpcode::G_STRICT_FMA:
689 case TargetOpcode::G_FMA:
690 return selectExtInst(ResVReg, ResType, I, CL::fma, GL::Fma);
691
692 case TargetOpcode::G_STRICT_FLDEXP:
693 return selectExtInst(ResVReg, ResType, I, CL::ldexp);
694
695 case TargetOpcode::G_FPOW:
696 return selectExtInst(ResVReg, ResType, I, CL::pow, GL::Pow);
697 case TargetOpcode::G_FPOWI:
698 return selectExtInst(ResVReg, ResType, I, CL::pown);
699
700 case TargetOpcode::G_FEXP:
701 return selectExtInst(ResVReg, ResType, I, CL::exp, GL::Exp);
702 case TargetOpcode::G_FEXP2:
703 return selectExtInst(ResVReg, ResType, I, CL::exp2, GL::Exp2);
704
705 case TargetOpcode::G_FLOG:
706 return selectExtInst(ResVReg, ResType, I, CL::log, GL::Log);
707 case TargetOpcode::G_FLOG2:
708 return selectExtInst(ResVReg, ResType, I, CL::log2, GL::Log2);
709 case TargetOpcode::G_FLOG10:
710 return selectLog10(ResVReg, ResType, I);
711
712 case TargetOpcode::G_FABS:
713 return selectExtInst(ResVReg, ResType, I, CL::fabs, GL::FAbs);
714 case TargetOpcode::G_ABS:
715 return selectExtInst(ResVReg, ResType, I, CL::s_abs, GL::SAbs);
716
717 case TargetOpcode::G_FMINNUM:
718 case TargetOpcode::G_FMINIMUM:
719 return selectExtInst(ResVReg, ResType, I, CL::fmin, GL::NMin);
720 case TargetOpcode::G_FMAXNUM:
721 case TargetOpcode::G_FMAXIMUM:
722 return selectExtInst(ResVReg, ResType, I, CL::fmax, GL::NMax);
723
724 case TargetOpcode::G_FCOPYSIGN:
725 return selectExtInst(ResVReg, ResType, I, CL::copysign);
726
727 case TargetOpcode::G_FCEIL:
728 return selectExtInst(ResVReg, ResType, I, CL::ceil, GL::Ceil);
729 case TargetOpcode::G_FFLOOR:
730 return selectExtInst(ResVReg, ResType, I, CL::floor, GL::Floor);
731
732 case TargetOpcode::G_FCOS:
733 return selectExtInst(ResVReg, ResType, I, CL::cos, GL::Cos);
734 case TargetOpcode::G_FSIN:
735 return selectExtInst(ResVReg, ResType, I, CL::sin, GL::Sin);
736 case TargetOpcode::G_FTAN:
737 return selectExtInst(ResVReg, ResType, I, CL::tan, GL::Tan);
738 case TargetOpcode::G_FACOS:
739 return selectExtInst(ResVReg, ResType, I, CL::acos, GL::Acos);
740 case TargetOpcode::G_FASIN:
741 return selectExtInst(ResVReg, ResType, I, CL::asin, GL::Asin);
742 case TargetOpcode::G_FATAN:
743 return selectExtInst(ResVReg, ResType, I, CL::atan, GL::Atan);
744 case TargetOpcode::G_FATAN2:
745 return selectExtInst(ResVReg, ResType, I, CL::atan2, GL::Atan2);
746 case TargetOpcode::G_FCOSH:
747 return selectExtInst(ResVReg, ResType, I, CL::cosh, GL::Cosh);
748 case TargetOpcode::G_FSINH:
749 return selectExtInst(ResVReg, ResType, I, CL::sinh, GL::Sinh);
750 case TargetOpcode::G_FTANH:
751 return selectExtInst(ResVReg, ResType, I, CL::tanh, GL::Tanh);
752
753 case TargetOpcode::G_STRICT_FSQRT:
754 case TargetOpcode::G_FSQRT:
755 return selectExtInst(ResVReg, ResType, I, CL::sqrt, GL::Sqrt);
756
757 case TargetOpcode::G_CTTZ:
758 case TargetOpcode::G_CTTZ_ZERO_UNDEF:
759 return selectExtInst(ResVReg, ResType, I, CL::ctz);
760 case TargetOpcode::G_CTLZ:
761 case TargetOpcode::G_CTLZ_ZERO_UNDEF:
762 return selectExtInst(ResVReg, ResType, I, CL::clz);
763
764 case TargetOpcode::G_INTRINSIC_ROUND:
765 return selectExtInst(ResVReg, ResType, I, CL::round, GL::Round);
766 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
767 return selectExtInst(ResVReg, ResType, I, CL::rint, GL::RoundEven);
768 case TargetOpcode::G_INTRINSIC_TRUNC:
769 return selectExtInst(ResVReg, ResType, I, CL::trunc, GL::Trunc);
770 case TargetOpcode::G_FRINT:
771 case TargetOpcode::G_FNEARBYINT:
772 return selectExtInst(ResVReg, ResType, I, CL::rint, GL::RoundEven);
773
774 case TargetOpcode::G_SMULH:
775 return selectExtInst(ResVReg, ResType, I, CL::s_mul_hi);
776 case TargetOpcode::G_UMULH:
777 return selectExtInst(ResVReg, ResType, I, CL::u_mul_hi);
778
779 case TargetOpcode::G_SADDSAT:
780 return selectExtInst(ResVReg, ResType, I, CL::s_add_sat);
781 case TargetOpcode::G_UADDSAT:
782 return selectExtInst(ResVReg, ResType, I, CL::u_add_sat);
783 case TargetOpcode::G_SSUBSAT:
784 return selectExtInst(ResVReg, ResType, I, CL::s_sub_sat);
785 case TargetOpcode::G_USUBSAT:
786 return selectExtInst(ResVReg, ResType, I, CL::u_sub_sat);
787
788 case TargetOpcode::G_UADDO:
789 return selectOverflowArith(ResVReg, ResType, I,
790 ResType->getOpcode() == SPIRV::OpTypeVector
791 ? SPIRV::OpIAddCarryV
792 : SPIRV::OpIAddCarryS);
793 case TargetOpcode::G_USUBO:
794 return selectOverflowArith(ResVReg, ResType, I,
795 ResType->getOpcode() == SPIRV::OpTypeVector
796 ? SPIRV::OpISubBorrowV
797 : SPIRV::OpISubBorrowS);
798 case TargetOpcode::G_UMULO:
799 return selectOverflowArith(ResVReg, ResType, I, SPIRV::OpUMulExtended);
800 case TargetOpcode::G_SMULO:
801 return selectOverflowArith(ResVReg, ResType, I, SPIRV::OpSMulExtended);
802
803 case TargetOpcode::G_SEXT:
804 return selectExt(ResVReg, ResType, I, true);
805 case TargetOpcode::G_ANYEXT:
806 case TargetOpcode::G_ZEXT:
807 return selectExt(ResVReg, ResType, I, false);
808 case TargetOpcode::G_TRUNC:
809 return selectTrunc(ResVReg, ResType, I);
810 case TargetOpcode::G_FPTRUNC:
811 case TargetOpcode::G_FPEXT:
812 return selectUnOp(ResVReg, ResType, I, SPIRV::OpFConvert);
813
814 case TargetOpcode::G_PTRTOINT:
815 return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertPtrToU);
816 case TargetOpcode::G_INTTOPTR:
817 return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertUToPtr);
818 case TargetOpcode::G_BITCAST:
819 return selectBitcast(ResVReg, ResType, I);
820 case TargetOpcode::G_ADDRSPACE_CAST:
821 return selectAddrSpaceCast(ResVReg, ResType, I);
822 case TargetOpcode::G_PTR_ADD: {
823 // Currently, we get G_PTR_ADD only applied to global variables.
824 assert(I.getOperand(1).isReg() && I.getOperand(2).isReg());
825 Register GV = I.getOperand(1).getReg();
826 MachineRegisterInfo::def_instr_iterator II = MRI->def_instr_begin(GV);
827 (void)II;
828 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
829 (*II).getOpcode() == TargetOpcode::COPY ||
830 (*II).getOpcode() == SPIRV::OpVariable) &&
831 getImm(I.getOperand(2), MRI));
832 // It may be the initialization of a global variable.
833 bool IsGVInit = false;
834 for (MachineRegisterInfo::use_instr_iterator
835 UseIt = MRI->use_instr_begin(I.getOperand(0).getReg()),
836 UseEnd = MRI->use_instr_end();
837 UseIt != UseEnd; UseIt = std::next(UseIt)) {
838 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
839 (*UseIt).getOpcode() == SPIRV::OpVariable) {
840 IsGVInit = true;
841 break;
842 }
843 }
844 MachineBasicBlock &BB = *I.getParent();
845 if (!IsGVInit) {
846 SPIRVType *GVType = GR.getSPIRVTypeForVReg(GV);
847 SPIRVType *GVPointeeType = GR.getPointeeType(GVType);
848 SPIRVType *ResPointeeType = GR.getPointeeType(ResType);
849 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
850 // Build a new virtual register that is associated with the required
851 // data type.
852 Register NewVReg = MRI->createGenericVirtualRegister(MRI->getType(GV));
853 MRI->setRegClass(NewVReg, MRI->getRegClass(GV));
854 // Having a correctly typed base we are ready to build the actually
855 // required GEP. It may not be a constant though, because all Operands
856 // of OpSpecConstantOp is to originate from other const instructions,
857 // and only the AccessChain named opcodes accept a global OpVariable
858 // instruction. We can't use an AccessChain opcode because of the type
859 // mismatch between result and base types.
860 if (!GR.isBitcastCompatible(ResType, GVType))
861 report_fatal_error(
862 "incompatible result and operand types in a bitcast");
863 Register ResTypeReg = GR.getSPIRVTypeID(ResType);
864 MachineInstrBuilder MIB =
865 BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpBitcast))
866 .addDef(NewVReg)
867 .addUse(ResTypeReg)
868 .addUse(GV);
869 return MIB.constrainAllUses(TII, TRI, RBI) &&
870 BuildMI(BB, I, I.getDebugLoc(),
871 TII.get(STI.isLogicalSPIRV()
872 ? SPIRV::OpInBoundsAccessChain
873 : SPIRV::OpInBoundsPtrAccessChain))
874 .addDef(ResVReg)
875 .addUse(ResTypeReg)
876 .addUse(NewVReg)
877 .addUse(I.getOperand(2).getReg())
878 .constrainAllUses(TII, TRI, RBI);
879 } else {
880 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSpecConstantOp))
881 .addDef(ResVReg)
882 .addUse(GR.getSPIRVTypeID(ResType))
883 .addImm(
884 static_cast<uint32_t>(SPIRV::Opcode::InBoundsPtrAccessChain))
885 .addUse(GV)
886 .addUse(I.getOperand(2).getReg())
887 .constrainAllUses(TII, TRI, RBI);
888 }
889 }
890 // It's possible to translate G_PTR_ADD to OpSpecConstantOp: either to
891 // initialize a global variable with a constant expression (e.g., the test
892 // case opencl/basic/progvar_prog_scope_init.ll), or for another use case
893 Register Idx = buildZerosVal(GR.getOrCreateSPIRVIntegerType(32, I, TII), I);
894 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSpecConstantOp))
895 .addDef(ResVReg)
896 .addUse(GR.getSPIRVTypeID(ResType))
897 .addImm(static_cast<uint32_t>(
898 SPIRV::Opcode::InBoundsPtrAccessChain))
899 .addUse(GV)
900 .addUse(Idx)
901 .addUse(I.getOperand(2).getReg());
902 return MIB.constrainAllUses(TII, TRI, RBI);
903 }
904
905 case TargetOpcode::G_ATOMICRMW_OR:
906 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicOr);
907 case TargetOpcode::G_ATOMICRMW_ADD:
908 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicIAdd);
909 case TargetOpcode::G_ATOMICRMW_AND:
910 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicAnd);
911 case TargetOpcode::G_ATOMICRMW_MAX:
912 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicSMax);
913 case TargetOpcode::G_ATOMICRMW_MIN:
914 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicSMin);
915 case TargetOpcode::G_ATOMICRMW_SUB:
916 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicISub);
917 case TargetOpcode::G_ATOMICRMW_XOR:
918 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicXor);
919 case TargetOpcode::G_ATOMICRMW_UMAX:
920 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicUMax);
921 case TargetOpcode::G_ATOMICRMW_UMIN:
922 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicUMin);
923 case TargetOpcode::G_ATOMICRMW_XCHG:
924 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicExchange);
925 case TargetOpcode::G_ATOMIC_CMPXCHG:
926 return selectAtomicCmpXchg(ResVReg, ResType, I);
927
928 case TargetOpcode::G_ATOMICRMW_FADD:
929 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicFAddEXT);
930 case TargetOpcode::G_ATOMICRMW_FSUB:
931 // Translate G_ATOMICRMW_FSUB to OpAtomicFAddEXT with negative value operand
932 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicFAddEXT,
933 SPIRV::OpFNegate);
934 case TargetOpcode::G_ATOMICRMW_FMIN:
935 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicFMinEXT);
936 case TargetOpcode::G_ATOMICRMW_FMAX:
937 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicFMaxEXT);
938
939 case TargetOpcode::G_FENCE:
940 return selectFence(I);
941
942 case TargetOpcode::G_STACKSAVE:
943 return selectStackSave(ResVReg, ResType, I);
944 case TargetOpcode::G_STACKRESTORE:
945 return selectStackRestore(I);
946
947 case TargetOpcode::G_UNMERGE_VALUES:
948 return selectUnmergeValues(I);
949
950 // Discard gen opcodes for intrinsics which we do not expect to actually
951 // represent code after lowering or intrinsics which are not implemented but
952 // should not crash when found in a customer's LLVM IR input.
953 case TargetOpcode::G_TRAP:
954 case TargetOpcode::G_DEBUGTRAP:
955 case TargetOpcode::G_UBSANTRAP:
956 case TargetOpcode::DBG_LABEL:
957 return true;
958
959 default:
960 return false;
961 }
962 }
963
selectExtInst(Register ResVReg,const SPIRVType * ResType,MachineInstr & I,GL::GLSLExtInst GLInst) const964 bool SPIRVInstructionSelector::selectExtInst(Register ResVReg,
965 const SPIRVType *ResType,
966 MachineInstr &I,
967 GL::GLSLExtInst GLInst) const {
968 if (!STI.canUseExtInstSet(
969 SPIRV::InstructionSet::InstructionSet::GLSL_std_450)) {
970 std::string DiagMsg;
971 raw_string_ostream OS(DiagMsg);
972 I.print(OS, true, false, false, false);
973 DiagMsg += " is only supported with the GLSL extended instruction set.\n";
974 report_fatal_error(DiagMsg.c_str(), false);
975 }
976 return selectExtInst(ResVReg, ResType, I,
977 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}});
978 }
979
selectExtInst(Register ResVReg,const SPIRVType * ResType,MachineInstr & I,CL::OpenCLExtInst CLInst) const980 bool SPIRVInstructionSelector::selectExtInst(Register ResVReg,
981 const SPIRVType *ResType,
982 MachineInstr &I,
983 CL::OpenCLExtInst CLInst) const {
984 return selectExtInst(ResVReg, ResType, I,
985 {{SPIRV::InstructionSet::OpenCL_std, CLInst}});
986 }
987
selectExtInst(Register ResVReg,const SPIRVType * ResType,MachineInstr & I,CL::OpenCLExtInst CLInst,GL::GLSLExtInst GLInst) const988 bool SPIRVInstructionSelector::selectExtInst(Register ResVReg,
989 const SPIRVType *ResType,
990 MachineInstr &I,
991 CL::OpenCLExtInst CLInst,
992 GL::GLSLExtInst GLInst) const {
993 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
994 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
995 return selectExtInst(ResVReg, ResType, I, ExtInsts);
996 }
997
selectExtInst(Register ResVReg,const SPIRVType * ResType,MachineInstr & I,const ExtInstList & Insts) const998 bool SPIRVInstructionSelector::selectExtInst(Register ResVReg,
999 const SPIRVType *ResType,
1000 MachineInstr &I,
1001 const ExtInstList &Insts) const {
1002
1003 for (const auto &Ex : Insts) {
1004 SPIRV::InstructionSet::InstructionSet Set = Ex.first;
1005 uint32_t Opcode = Ex.second;
1006 if (STI.canUseExtInstSet(Set)) {
1007 MachineBasicBlock &BB = *I.getParent();
1008 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
1009 .addDef(ResVReg)
1010 .addUse(GR.getSPIRVTypeID(ResType))
1011 .addImm(static_cast<uint32_t>(Set))
1012 .addImm(Opcode);
1013 const unsigned NumOps = I.getNumOperands();
1014 unsigned Index = 1;
1015 if (Index < NumOps &&
1016 I.getOperand(Index).getType() ==
1017 MachineOperand::MachineOperandType::MO_IntrinsicID)
1018 Index = 2;
1019 for (; Index < NumOps; ++Index)
1020 MIB.add(I.getOperand(Index));
1021 return MIB.constrainAllUses(TII, TRI, RBI);
1022 }
1023 }
1024 return false;
1025 }
1026
selectOpWithSrcs(Register ResVReg,const SPIRVType * ResType,MachineInstr & I,std::vector<Register> Srcs,unsigned Opcode) const1027 bool SPIRVInstructionSelector::selectOpWithSrcs(Register ResVReg,
1028 const SPIRVType *ResType,
1029 MachineInstr &I,
1030 std::vector<Register> Srcs,
1031 unsigned Opcode) const {
1032 auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(Opcode))
1033 .addDef(ResVReg)
1034 .addUse(GR.getSPIRVTypeID(ResType));
1035 for (Register SReg : Srcs) {
1036 MIB.addUse(SReg);
1037 }
1038 return MIB.constrainAllUses(TII, TRI, RBI);
1039 }
1040
selectUnOp(Register ResVReg,const SPIRVType * ResType,MachineInstr & I,unsigned Opcode) const1041 bool SPIRVInstructionSelector::selectUnOp(Register ResVReg,
1042 const SPIRVType *ResType,
1043 MachineInstr &I,
1044 unsigned Opcode) const {
1045 if (STI.isPhysicalSPIRV() && I.getOperand(1).isReg()) {
1046 Register SrcReg = I.getOperand(1).getReg();
1047 bool IsGV = false;
1048 for (MachineRegisterInfo::def_instr_iterator DefIt =
1049 MRI->def_instr_begin(SrcReg);
1050 DefIt != MRI->def_instr_end(); DefIt = std::next(DefIt)) {
1051 if ((*DefIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1052 (*DefIt).getOpcode() == SPIRV::OpVariable) {
1053 IsGV = true;
1054 break;
1055 }
1056 }
1057 if (IsGV) {
1058 uint32_t SpecOpcode = 0;
1059 switch (Opcode) {
1060 case SPIRV::OpConvertPtrToU:
1061 SpecOpcode = static_cast<uint32_t>(SPIRV::Opcode::ConvertPtrToU);
1062 break;
1063 case SPIRV::OpConvertUToPtr:
1064 SpecOpcode = static_cast<uint32_t>(SPIRV::Opcode::ConvertUToPtr);
1065 break;
1066 }
1067 if (SpecOpcode)
1068 return BuildMI(*I.getParent(), I, I.getDebugLoc(),
1069 TII.get(SPIRV::OpSpecConstantOp))
1070 .addDef(ResVReg)
1071 .addUse(GR.getSPIRVTypeID(ResType))
1072 .addImm(SpecOpcode)
1073 .addUse(SrcReg)
1074 .constrainAllUses(TII, TRI, RBI);
1075 }
1076 }
1077 return selectOpWithSrcs(ResVReg, ResType, I, {I.getOperand(1).getReg()},
1078 Opcode);
1079 }
1080
selectBitcast(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const1081 bool SPIRVInstructionSelector::selectBitcast(Register ResVReg,
1082 const SPIRVType *ResType,
1083 MachineInstr &I) const {
1084 Register OpReg = I.getOperand(1).getReg();
1085 SPIRVType *OpType = OpReg.isValid() ? GR.getSPIRVTypeForVReg(OpReg) : nullptr;
1086 if (!GR.isBitcastCompatible(ResType, OpType))
1087 report_fatal_error("incompatible result and operand types in a bitcast");
1088 return selectUnOp(ResVReg, ResType, I, SPIRV::OpBitcast);
1089 }
1090
addMemoryOperands(MachineMemOperand * MemOp,MachineInstrBuilder & MIB,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry & GR)1091 static void addMemoryOperands(MachineMemOperand *MemOp,
1092 MachineInstrBuilder &MIB,
1093 MachineIRBuilder &MIRBuilder,
1094 SPIRVGlobalRegistry &GR) {
1095 uint32_t SpvMemOp = static_cast<uint32_t>(SPIRV::MemoryOperand::None);
1096 if (MemOp->isVolatile())
1097 SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1098 if (MemOp->isNonTemporal())
1099 SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1100 if (MemOp->getAlign().value())
1101 SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1102
1103 [[maybe_unused]] MachineInstr *AliasList = nullptr;
1104 [[maybe_unused]] MachineInstr *NoAliasList = nullptr;
1105 const SPIRVSubtarget *ST =
1106 static_cast<const SPIRVSubtarget *>(&MIRBuilder.getMF().getSubtarget());
1107 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1108 if (auto *MD = MemOp->getAAInfo().Scope) {
1109 AliasList = GR.getOrAddMemAliasingINTELInst(MIRBuilder, MD);
1110 if (AliasList)
1111 SpvMemOp |=
1112 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1113 }
1114 if (auto *MD = MemOp->getAAInfo().NoAlias) {
1115 NoAliasList = GR.getOrAddMemAliasingINTELInst(MIRBuilder, MD);
1116 if (NoAliasList)
1117 SpvMemOp |=
1118 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1119 }
1120 }
1121
1122 if (SpvMemOp != static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1123 MIB.addImm(SpvMemOp);
1124 if (SpvMemOp & static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1125 MIB.addImm(MemOp->getAlign().value());
1126 if (AliasList)
1127 MIB.addUse(AliasList->getOperand(0).getReg());
1128 if (NoAliasList)
1129 MIB.addUse(NoAliasList->getOperand(0).getReg());
1130 }
1131 }
1132
addMemoryOperands(uint64_t Flags,MachineInstrBuilder & MIB)1133 static void addMemoryOperands(uint64_t Flags, MachineInstrBuilder &MIB) {
1134 uint32_t SpvMemOp = static_cast<uint32_t>(SPIRV::MemoryOperand::None);
1135 if (Flags & MachineMemOperand::Flags::MOVolatile)
1136 SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1137 if (Flags & MachineMemOperand::Flags::MONonTemporal)
1138 SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1139
1140 if (SpvMemOp != static_cast<uint32_t>(SPIRV::MemoryOperand::None))
1141 MIB.addImm(SpvMemOp);
1142 }
1143
selectLoad(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const1144 bool SPIRVInstructionSelector::selectLoad(Register ResVReg,
1145 const SPIRVType *ResType,
1146 MachineInstr &I) const {
1147 unsigned OpOffset = isa<GIntrinsic>(I) ? 1 : 0;
1148 Register Ptr = I.getOperand(1 + OpOffset).getReg();
1149
1150 auto *PtrDef = getVRegDef(*MRI, Ptr);
1151 auto *IntPtrDef = dyn_cast<GIntrinsic>(PtrDef);
1152 if (IntPtrDef &&
1153 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1154 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1155 SPIRVType *HandleType = GR.getSPIRVTypeForVReg(HandleReg);
1156 if (HandleType->getOpcode() == SPIRV::OpTypeImage) {
1157 Register NewHandleReg =
1158 MRI->createVirtualRegister(MRI->getRegClass(HandleReg));
1159 auto *HandleDef = cast<GIntrinsic>(getVRegDef(*MRI, HandleReg));
1160 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef, I)) {
1161 return false;
1162 }
1163
1164 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1165 return generateImageRead(ResVReg, ResType, NewHandleReg, IdxReg,
1166 I.getDebugLoc(), I);
1167 }
1168 }
1169
1170 auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpLoad))
1171 .addDef(ResVReg)
1172 .addUse(GR.getSPIRVTypeID(ResType))
1173 .addUse(Ptr);
1174 if (!I.getNumMemOperands()) {
1175 assert(I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1176 I.getOpcode() ==
1177 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1178 addMemoryOperands(I.getOperand(2 + OpOffset).getImm(), MIB);
1179 } else {
1180 MachineIRBuilder MIRBuilder(I);
1181 addMemoryOperands(*I.memoperands_begin(), MIB, MIRBuilder, GR);
1182 }
1183 return MIB.constrainAllUses(TII, TRI, RBI);
1184 }
1185
selectStore(MachineInstr & I) const1186 bool SPIRVInstructionSelector::selectStore(MachineInstr &I) const {
1187 unsigned OpOffset = isa<GIntrinsic>(I) ? 1 : 0;
1188 Register StoreVal = I.getOperand(0 + OpOffset).getReg();
1189 Register Ptr = I.getOperand(1 + OpOffset).getReg();
1190
1191 auto *PtrDef = getVRegDef(*MRI, Ptr);
1192 auto *IntPtrDef = dyn_cast<GIntrinsic>(PtrDef);
1193 if (IntPtrDef &&
1194 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1195 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1196 Register NewHandleReg =
1197 MRI->createVirtualRegister(MRI->getRegClass(HandleReg));
1198 auto *HandleDef = cast<GIntrinsic>(getVRegDef(*MRI, HandleReg));
1199 SPIRVType *HandleType = GR.getSPIRVTypeForVReg(HandleReg);
1200 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef, I)) {
1201 return false;
1202 }
1203
1204 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1205 if (HandleType->getOpcode() == SPIRV::OpTypeImage) {
1206 auto BMI = BuildMI(*I.getParent(), I, I.getDebugLoc(),
1207 TII.get(SPIRV::OpImageWrite))
1208 .addUse(NewHandleReg)
1209 .addUse(IdxReg)
1210 .addUse(StoreVal);
1211
1212 const llvm::Type *LLVMHandleType = GR.getTypeForSPIRVType(HandleType);
1213 if (sampledTypeIsSignedInteger(LLVMHandleType))
1214 BMI.addImm(0x1000); // SignExtend
1215
1216 return BMI.constrainAllUses(TII, TRI, RBI);
1217 }
1218 }
1219
1220 MachineBasicBlock &BB = *I.getParent();
1221 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpStore))
1222 .addUse(Ptr)
1223 .addUse(StoreVal);
1224 if (!I.getNumMemOperands()) {
1225 assert(I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1226 I.getOpcode() ==
1227 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1228 addMemoryOperands(I.getOperand(2 + OpOffset).getImm(), MIB);
1229 } else {
1230 MachineIRBuilder MIRBuilder(I);
1231 addMemoryOperands(*I.memoperands_begin(), MIB, MIRBuilder, GR);
1232 }
1233 return MIB.constrainAllUses(TII, TRI, RBI);
1234 }
1235
selectStackSave(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const1236 bool SPIRVInstructionSelector::selectStackSave(Register ResVReg,
1237 const SPIRVType *ResType,
1238 MachineInstr &I) const {
1239 if (!STI.canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
1240 report_fatal_error(
1241 "llvm.stacksave intrinsic: this instruction requires the following "
1242 "SPIR-V extension: SPV_INTEL_variable_length_array",
1243 false);
1244 MachineBasicBlock &BB = *I.getParent();
1245 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSaveMemoryINTEL))
1246 .addDef(ResVReg)
1247 .addUse(GR.getSPIRVTypeID(ResType))
1248 .constrainAllUses(TII, TRI, RBI);
1249 }
1250
selectStackRestore(MachineInstr & I) const1251 bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &I) const {
1252 if (!STI.canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
1253 report_fatal_error(
1254 "llvm.stackrestore intrinsic: this instruction requires the following "
1255 "SPIR-V extension: SPV_INTEL_variable_length_array",
1256 false);
1257 if (!I.getOperand(0).isReg())
1258 return false;
1259 MachineBasicBlock &BB = *I.getParent();
1260 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpRestoreMemoryINTEL))
1261 .addUse(I.getOperand(0).getReg())
1262 .constrainAllUses(TII, TRI, RBI);
1263 }
1264
selectMemOperation(Register ResVReg,MachineInstr & I) const1265 bool SPIRVInstructionSelector::selectMemOperation(Register ResVReg,
1266 MachineInstr &I) const {
1267 MachineBasicBlock &BB = *I.getParent();
1268 Register SrcReg = I.getOperand(1).getReg();
1269 bool Result = true;
1270 if (I.getOpcode() == TargetOpcode::G_MEMSET) {
1271 MachineIRBuilder MIRBuilder(I);
1272 assert(I.getOperand(1).isReg() && I.getOperand(2).isReg());
1273 unsigned Val = getIConstVal(I.getOperand(1).getReg(), MRI);
1274 unsigned Num = getIConstVal(I.getOperand(2).getReg(), MRI);
1275 Type *ValTy = Type::getInt8Ty(I.getMF()->getFunction().getContext());
1276 Type *ArrTy = ArrayType::get(ValTy, Num);
1277 SPIRVType *VarTy = GR.getOrCreateSPIRVPointerType(
1278 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
1279
1280 SPIRVType *SpvArrTy = GR.getOrCreateSPIRVType(
1281 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None, false);
1282 Register Const = GR.getOrCreateConstIntArray(Val, Num, I, SpvArrTy, TII);
1283 // TODO: check if we have such GV, add init, use buildGlobalVariable.
1284 Function &CurFunction = GR.CurMF->getFunction();
1285 Type *LLVMArrTy =
1286 ArrayType::get(IntegerType::get(CurFunction.getContext(), 8), Num);
1287 // Module takes ownership of the global var.
1288 GlobalVariable *GV = new GlobalVariable(*CurFunction.getParent(), LLVMArrTy,
1289 true, GlobalValue::InternalLinkage,
1290 Constant::getNullValue(LLVMArrTy));
1291 Register VarReg = MRI->createGenericVirtualRegister(LLT::scalar(64));
1292 auto MIBVar =
1293 BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpVariable))
1294 .addDef(VarReg)
1295 .addUse(GR.getSPIRVTypeID(VarTy))
1296 .addImm(SPIRV::StorageClass::UniformConstant)
1297 .addUse(Const);
1298 Result &= MIBVar.constrainAllUses(TII, TRI, RBI);
1299
1300 GR.add(GV, MIBVar);
1301 GR.addGlobalObject(GV, GR.CurMF, VarReg);
1302
1303 buildOpDecorate(VarReg, I, TII, SPIRV::Decoration::Constant, {});
1304 SPIRVType *SourceTy = GR.getOrCreateSPIRVPointerType(
1305 ValTy, I, SPIRV::StorageClass::UniformConstant);
1306 SrcReg = MRI->createGenericVirtualRegister(LLT::scalar(64));
1307 selectOpWithSrcs(SrcReg, SourceTy, I, {VarReg}, SPIRV::OpBitcast);
1308 }
1309 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCopyMemorySized))
1310 .addUse(I.getOperand(0).getReg())
1311 .addUse(SrcReg)
1312 .addUse(I.getOperand(2).getReg());
1313 if (I.getNumMemOperands()) {
1314 MachineIRBuilder MIRBuilder(I);
1315 addMemoryOperands(*I.memoperands_begin(), MIB, MIRBuilder, GR);
1316 }
1317 Result &= MIB.constrainAllUses(TII, TRI, RBI);
1318 if (ResVReg.isValid() && ResVReg != MIB->getOperand(0).getReg())
1319 Result &= BuildCOPY(ResVReg, MIB->getOperand(0).getReg(), I);
1320 return Result;
1321 }
1322
selectAtomicRMW(Register ResVReg,const SPIRVType * ResType,MachineInstr & I,unsigned NewOpcode,unsigned NegateOpcode) const1323 bool SPIRVInstructionSelector::selectAtomicRMW(Register ResVReg,
1324 const SPIRVType *ResType,
1325 MachineInstr &I,
1326 unsigned NewOpcode,
1327 unsigned NegateOpcode) const {
1328 bool Result = true;
1329 assert(I.hasOneMemOperand());
1330 const MachineMemOperand *MemOp = *I.memoperands_begin();
1331 uint32_t Scope = static_cast<uint32_t>(getMemScope(
1332 GR.CurMF->getFunction().getContext(), MemOp->getSyncScopeID()));
1333 auto ScopeConstant = buildI32Constant(Scope, I);
1334 Register ScopeReg = ScopeConstant.first;
1335 Result &= ScopeConstant.second;
1336
1337 Register Ptr = I.getOperand(1).getReg();
1338 // TODO: Changed as it's implemented in the translator. See test/atomicrmw.ll
1339 // auto ScSem =
1340 // getMemSemanticsForStorageClass(GR.getPointerStorageClass(Ptr));
1341 AtomicOrdering AO = MemOp->getSuccessOrdering();
1342 uint32_t MemSem = static_cast<uint32_t>(getMemSemantics(AO));
1343 auto MemSemConstant = buildI32Constant(MemSem /*| ScSem*/, I);
1344 Register MemSemReg = MemSemConstant.first;
1345 Result &= MemSemConstant.second;
1346
1347 Register ValueReg = I.getOperand(2).getReg();
1348 if (NegateOpcode != 0) {
1349 // Translation with negative value operand is requested
1350 Register TmpReg = createVirtualRegister(ResType, &GR, MRI, MRI->getMF());
1351 Result &= selectOpWithSrcs(TmpReg, ResType, I, {ValueReg}, NegateOpcode);
1352 ValueReg = TmpReg;
1353 }
1354
1355 return Result &&
1356 BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(NewOpcode))
1357 .addDef(ResVReg)
1358 .addUse(GR.getSPIRVTypeID(ResType))
1359 .addUse(Ptr)
1360 .addUse(ScopeReg)
1361 .addUse(MemSemReg)
1362 .addUse(ValueReg)
1363 .constrainAllUses(TII, TRI, RBI);
1364 }
1365
selectUnmergeValues(MachineInstr & I) const1366 bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &I) const {
1367 unsigned ArgI = I.getNumOperands() - 1;
1368 Register SrcReg =
1369 I.getOperand(ArgI).isReg() ? I.getOperand(ArgI).getReg() : Register(0);
1370 SPIRVType *DefType =
1371 SrcReg.isValid() ? GR.getSPIRVTypeForVReg(SrcReg) : nullptr;
1372 if (!DefType || DefType->getOpcode() != SPIRV::OpTypeVector)
1373 report_fatal_error(
1374 "cannot select G_UNMERGE_VALUES with a non-vector argument");
1375
1376 SPIRVType *ScalarType =
1377 GR.getSPIRVTypeForVReg(DefType->getOperand(1).getReg());
1378 MachineBasicBlock &BB = *I.getParent();
1379 bool Res = false;
1380 for (unsigned i = 0; i < I.getNumDefs(); ++i) {
1381 Register ResVReg = I.getOperand(i).getReg();
1382 SPIRVType *ResType = GR.getSPIRVTypeForVReg(ResVReg);
1383 if (!ResType) {
1384 // There was no "assign type" actions, let's fix this now
1385 ResType = ScalarType;
1386 MRI->setRegClass(ResVReg, GR.getRegClass(ResType));
1387 MRI->setType(ResVReg, LLT::scalar(GR.getScalarOrVectorBitWidth(ResType)));
1388 GR.assignSPIRVTypeToVReg(ResType, ResVReg, *GR.CurMF);
1389 }
1390 auto MIB =
1391 BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
1392 .addDef(ResVReg)
1393 .addUse(GR.getSPIRVTypeID(ResType))
1394 .addUse(SrcReg)
1395 .addImm(static_cast<int64_t>(i));
1396 Res |= MIB.constrainAllUses(TII, TRI, RBI);
1397 }
1398 return Res;
1399 }
1400
selectFence(MachineInstr & I) const1401 bool SPIRVInstructionSelector::selectFence(MachineInstr &I) const {
1402 AtomicOrdering AO = AtomicOrdering(I.getOperand(0).getImm());
1403 uint32_t MemSem = static_cast<uint32_t>(getMemSemantics(AO));
1404 auto MemSemConstant = buildI32Constant(MemSem, I);
1405 Register MemSemReg = MemSemConstant.first;
1406 bool Result = MemSemConstant.second;
1407 SyncScope::ID Ord = SyncScope::ID(I.getOperand(1).getImm());
1408 uint32_t Scope = static_cast<uint32_t>(
1409 getMemScope(GR.CurMF->getFunction().getContext(), Ord));
1410 auto ScopeConstant = buildI32Constant(Scope, I);
1411 Register ScopeReg = ScopeConstant.first;
1412 Result &= ScopeConstant.second;
1413 MachineBasicBlock &BB = *I.getParent();
1414 return Result &&
1415 BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpMemoryBarrier))
1416 .addUse(ScopeReg)
1417 .addUse(MemSemReg)
1418 .constrainAllUses(TII, TRI, RBI);
1419 }
1420
selectOverflowArith(Register ResVReg,const SPIRVType * ResType,MachineInstr & I,unsigned Opcode) const1421 bool SPIRVInstructionSelector::selectOverflowArith(Register ResVReg,
1422 const SPIRVType *ResType,
1423 MachineInstr &I,
1424 unsigned Opcode) const {
1425 Type *ResTy = nullptr;
1426 StringRef ResName;
1427 if (!GR.findValueAttrs(&I, ResTy, ResName))
1428 report_fatal_error(
1429 "Not enough info to select the arithmetic with overflow instruction");
1430 if (!ResTy || !ResTy->isStructTy())
1431 report_fatal_error("Expect struct type result for the arithmetic "
1432 "with overflow instruction");
1433 // "Result Type must be from OpTypeStruct. The struct must have two members,
1434 // and the two members must be the same type."
1435 Type *ResElemTy = cast<StructType>(ResTy)->getElementType(0);
1436 ResTy = StructType::get(ResElemTy, ResElemTy);
1437 // Build SPIR-V types and constant(s) if needed.
1438 MachineIRBuilder MIRBuilder(I);
1439 SPIRVType *StructType = GR.getOrCreateSPIRVType(
1440 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite, false);
1441 assert(I.getNumDefs() > 1 && "Not enought operands");
1442 SPIRVType *BoolType = GR.getOrCreateSPIRVBoolType(I, TII);
1443 unsigned N = GR.getScalarOrVectorComponentCount(ResType);
1444 if (N > 1)
1445 BoolType = GR.getOrCreateSPIRVVectorType(BoolType, N, I, TII);
1446 Register BoolTypeReg = GR.getSPIRVTypeID(BoolType);
1447 Register ZeroReg = buildZerosVal(ResType, I);
1448 // A new virtual register to store the result struct.
1449 Register StructVReg = MRI->createGenericVirtualRegister(LLT::scalar(64));
1450 MRI->setRegClass(StructVReg, &SPIRV::IDRegClass);
1451 // Build the result name if needed.
1452 if (ResName.size() > 0)
1453 buildOpName(StructVReg, ResName, MIRBuilder);
1454 // Build the arithmetic with overflow instruction.
1455 MachineBasicBlock &BB = *I.getParent();
1456 auto MIB =
1457 BuildMI(BB, MIRBuilder.getInsertPt(), I.getDebugLoc(), TII.get(Opcode))
1458 .addDef(StructVReg)
1459 .addUse(GR.getSPIRVTypeID(StructType));
1460 for (unsigned i = I.getNumDefs(); i < I.getNumOperands(); ++i)
1461 MIB.addUse(I.getOperand(i).getReg());
1462 bool Result = MIB.constrainAllUses(TII, TRI, RBI);
1463 // Build instructions to extract fields of the instruction's result.
1464 // A new virtual register to store the higher part of the result struct.
1465 Register HigherVReg = MRI->createGenericVirtualRegister(LLT::scalar(64));
1466 MRI->setRegClass(HigherVReg, &SPIRV::iIDRegClass);
1467 for (unsigned i = 0; i < I.getNumDefs(); ++i) {
1468 auto MIB =
1469 BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
1470 .addDef(i == 1 ? HigherVReg : I.getOperand(i).getReg())
1471 .addUse(GR.getSPIRVTypeID(ResType))
1472 .addUse(StructVReg)
1473 .addImm(i);
1474 Result &= MIB.constrainAllUses(TII, TRI, RBI);
1475 }
1476 // Build boolean value from the higher part.
1477 return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpINotEqual))
1478 .addDef(I.getOperand(1).getReg())
1479 .addUse(BoolTypeReg)
1480 .addUse(HigherVReg)
1481 .addUse(ZeroReg)
1482 .constrainAllUses(TII, TRI, RBI);
1483 }
1484
selectAtomicCmpXchg(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const1485 bool SPIRVInstructionSelector::selectAtomicCmpXchg(Register ResVReg,
1486 const SPIRVType *ResType,
1487 MachineInstr &I) const {
1488 bool Result = true;
1489 Register ScopeReg;
1490 Register MemSemEqReg;
1491 Register MemSemNeqReg;
1492 Register Ptr = I.getOperand(2).getReg();
1493 if (!isa<GIntrinsic>(I)) {
1494 assert(I.hasOneMemOperand());
1495 const MachineMemOperand *MemOp = *I.memoperands_begin();
1496 unsigned Scope = static_cast<uint32_t>(getMemScope(
1497 GR.CurMF->getFunction().getContext(), MemOp->getSyncScopeID()));
1498 auto ScopeConstant = buildI32Constant(Scope, I);
1499 ScopeReg = ScopeConstant.first;
1500 Result &= ScopeConstant.second;
1501
1502 unsigned ScSem = static_cast<uint32_t>(
1503 getMemSemanticsForStorageClass(GR.getPointerStorageClass(Ptr)));
1504 AtomicOrdering AO = MemOp->getSuccessOrdering();
1505 unsigned MemSemEq = static_cast<uint32_t>(getMemSemantics(AO)) | ScSem;
1506 auto MemSemEqConstant = buildI32Constant(MemSemEq, I);
1507 MemSemEqReg = MemSemEqConstant.first;
1508 Result &= MemSemEqConstant.second;
1509 AtomicOrdering FO = MemOp->getFailureOrdering();
1510 unsigned MemSemNeq = static_cast<uint32_t>(getMemSemantics(FO)) | ScSem;
1511 if (MemSemEq == MemSemNeq)
1512 MemSemNeqReg = MemSemEqReg;
1513 else {
1514 auto MemSemNeqConstant = buildI32Constant(MemSemEq, I);
1515 MemSemNeqReg = MemSemNeqConstant.first;
1516 Result &= MemSemNeqConstant.second;
1517 }
1518 } else {
1519 ScopeReg = I.getOperand(5).getReg();
1520 MemSemEqReg = I.getOperand(6).getReg();
1521 MemSemNeqReg = I.getOperand(7).getReg();
1522 }
1523
1524 Register Cmp = I.getOperand(3).getReg();
1525 Register Val = I.getOperand(4).getReg();
1526 SPIRVType *SpvValTy = GR.getSPIRVTypeForVReg(Val);
1527 Register ACmpRes = createVirtualRegister(SpvValTy, &GR, MRI, *I.getMF());
1528 const DebugLoc &DL = I.getDebugLoc();
1529 Result &=
1530 BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpAtomicCompareExchange))
1531 .addDef(ACmpRes)
1532 .addUse(GR.getSPIRVTypeID(SpvValTy))
1533 .addUse(Ptr)
1534 .addUse(ScopeReg)
1535 .addUse(MemSemEqReg)
1536 .addUse(MemSemNeqReg)
1537 .addUse(Val)
1538 .addUse(Cmp)
1539 .constrainAllUses(TII, TRI, RBI);
1540 SPIRVType *BoolTy = GR.getOrCreateSPIRVBoolType(I, TII);
1541 Register CmpSuccReg = createVirtualRegister(BoolTy, &GR, MRI, *I.getMF());
1542 Result &= BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpIEqual))
1543 .addDef(CmpSuccReg)
1544 .addUse(GR.getSPIRVTypeID(BoolTy))
1545 .addUse(ACmpRes)
1546 .addUse(Cmp)
1547 .constrainAllUses(TII, TRI, RBI);
1548 Register TmpReg = createVirtualRegister(ResType, &GR, MRI, *I.getMF());
1549 Result &= BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpCompositeInsert))
1550 .addDef(TmpReg)
1551 .addUse(GR.getSPIRVTypeID(ResType))
1552 .addUse(ACmpRes)
1553 .addUse(GR.getOrCreateUndef(I, ResType, TII))
1554 .addImm(0)
1555 .constrainAllUses(TII, TRI, RBI);
1556 return Result &&
1557 BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpCompositeInsert))
1558 .addDef(ResVReg)
1559 .addUse(GR.getSPIRVTypeID(ResType))
1560 .addUse(CmpSuccReg)
1561 .addUse(TmpReg)
1562 .addImm(1)
1563 .constrainAllUses(TII, TRI, RBI);
1564 }
1565
isUSMStorageClass(SPIRV::StorageClass::StorageClass SC)1566 static bool isUSMStorageClass(SPIRV::StorageClass::StorageClass SC) {
1567 switch (SC) {
1568 case SPIRV::StorageClass::DeviceOnlyINTEL:
1569 case SPIRV::StorageClass::HostOnlyINTEL:
1570 return true;
1571 default:
1572 return false;
1573 }
1574 }
1575
1576 // Returns true ResVReg is referred only from global vars and OpName's.
isASCastInGVar(MachineRegisterInfo * MRI,Register ResVReg)1577 static bool isASCastInGVar(MachineRegisterInfo *MRI, Register ResVReg) {
1578 bool IsGRef = false;
1579 bool IsAllowedRefs =
1580 llvm::all_of(MRI->use_instructions(ResVReg), [&IsGRef](auto const &It) {
1581 unsigned Opcode = It.getOpcode();
1582 if (Opcode == SPIRV::OpConstantComposite ||
1583 Opcode == SPIRV::OpVariable ||
1584 isSpvIntrinsic(It, Intrinsic::spv_init_global))
1585 return IsGRef = true;
1586 return Opcode == SPIRV::OpName;
1587 });
1588 return IsAllowedRefs && IsGRef;
1589 }
1590
getUcharPtrTypeReg(MachineInstr & I,SPIRV::StorageClass::StorageClass SC) const1591 Register SPIRVInstructionSelector::getUcharPtrTypeReg(
1592 MachineInstr &I, SPIRV::StorageClass::StorageClass SC) const {
1593 return GR.getSPIRVTypeID(GR.getOrCreateSPIRVPointerType(
1594 Type::getInt8Ty(I.getMF()->getFunction().getContext()), I, SC));
1595 }
1596
1597 MachineInstrBuilder
buildSpecConstantOp(MachineInstr & I,Register Dest,Register Src,Register DestType,uint32_t Opcode) const1598 SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &I, Register Dest,
1599 Register Src, Register DestType,
1600 uint32_t Opcode) const {
1601 return BuildMI(*I.getParent(), I, I.getDebugLoc(),
1602 TII.get(SPIRV::OpSpecConstantOp))
1603 .addDef(Dest)
1604 .addUse(DestType)
1605 .addImm(Opcode)
1606 .addUse(Src);
1607 }
1608
1609 MachineInstrBuilder
buildConstGenericPtr(MachineInstr & I,Register SrcPtr,SPIRVType * SrcPtrTy) const1610 SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &I, Register SrcPtr,
1611 SPIRVType *SrcPtrTy) const {
1612 SPIRVType *GenericPtrTy =
1613 GR.changePointerStorageClass(SrcPtrTy, SPIRV::StorageClass::Generic, I);
1614 Register Tmp = MRI->createVirtualRegister(&SPIRV::pIDRegClass);
1615 MRI->setType(Tmp, LLT::pointer(storageClassToAddressSpace(
1616 SPIRV::StorageClass::Generic),
1617 GR.getPointerSize()));
1618 MachineFunction *MF = I.getParent()->getParent();
1619 GR.assignSPIRVTypeToVReg(GenericPtrTy, Tmp, *MF);
1620 MachineInstrBuilder MIB = buildSpecConstantOp(
1621 I, Tmp, SrcPtr, GR.getSPIRVTypeID(GenericPtrTy),
1622 static_cast<uint32_t>(SPIRV::Opcode::PtrCastToGeneric));
1623 GR.add(MIB.getInstr(), MIB);
1624 return MIB;
1625 }
1626
1627 // In SPIR-V address space casting can only happen to and from the Generic
1628 // storage class. We can also only cast Workgroup, CrossWorkgroup, or Function
1629 // pointers to and from Generic pointers. As such, we can convert e.g. from
1630 // Workgroup to Function by going via a Generic pointer as an intermediary. All
1631 // other combinations can only be done by a bitcast, and are probably not safe.
selectAddrSpaceCast(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const1632 bool SPIRVInstructionSelector::selectAddrSpaceCast(Register ResVReg,
1633 const SPIRVType *ResType,
1634 MachineInstr &I) const {
1635 MachineBasicBlock &BB = *I.getParent();
1636 const DebugLoc &DL = I.getDebugLoc();
1637
1638 Register SrcPtr = I.getOperand(1).getReg();
1639 SPIRVType *SrcPtrTy = GR.getSPIRVTypeForVReg(SrcPtr);
1640
1641 // don't generate a cast for a null that may be represented by OpTypeInt
1642 if (SrcPtrTy->getOpcode() != SPIRV::OpTypePointer ||
1643 ResType->getOpcode() != SPIRV::OpTypePointer)
1644 return BuildCOPY(ResVReg, SrcPtr, I);
1645
1646 SPIRV::StorageClass::StorageClass SrcSC = GR.getPointerStorageClass(SrcPtrTy);
1647 SPIRV::StorageClass::StorageClass DstSC = GR.getPointerStorageClass(ResType);
1648
1649 if (isASCastInGVar(MRI, ResVReg)) {
1650 // AddrSpaceCast uses within OpVariable and OpConstantComposite instructions
1651 // are expressed by OpSpecConstantOp with an Opcode.
1652 // TODO: maybe insert a check whether the Kernel capability was declared and
1653 // so PtrCastToGeneric/GenericCastToPtr are available.
1654 unsigned SpecOpcode =
1655 DstSC == SPIRV::StorageClass::Generic && isGenericCastablePtr(SrcSC)
1656 ? static_cast<uint32_t>(SPIRV::Opcode::PtrCastToGeneric)
1657 : (SrcSC == SPIRV::StorageClass::Generic &&
1658 isGenericCastablePtr(DstSC)
1659 ? static_cast<uint32_t>(SPIRV::Opcode::GenericCastToPtr)
1660 : 0);
1661 // TODO: OpConstantComposite expects i8*, so we are forced to forget a
1662 // correct value of ResType and use general i8* instead. Maybe this should
1663 // be addressed in the emit-intrinsic step to infer a correct
1664 // OpConstantComposite type.
1665 if (SpecOpcode) {
1666 return buildSpecConstantOp(I, ResVReg, SrcPtr,
1667 getUcharPtrTypeReg(I, DstSC), SpecOpcode)
1668 .constrainAllUses(TII, TRI, RBI);
1669 } else if (isGenericCastablePtr(SrcSC) && isGenericCastablePtr(DstSC)) {
1670 MachineInstrBuilder MIB = buildConstGenericPtr(I, SrcPtr, SrcPtrTy);
1671 return MIB.constrainAllUses(TII, TRI, RBI) &&
1672 buildSpecConstantOp(
1673 I, ResVReg, MIB->getOperand(0).getReg(),
1674 getUcharPtrTypeReg(I, DstSC),
1675 static_cast<uint32_t>(SPIRV::Opcode::GenericCastToPtr))
1676 .constrainAllUses(TII, TRI, RBI);
1677 }
1678 }
1679
1680 // don't generate a cast between identical storage classes
1681 if (SrcSC == DstSC)
1682 return BuildCOPY(ResVReg, SrcPtr, I);
1683
1684 if ((SrcSC == SPIRV::StorageClass::Function &&
1685 DstSC == SPIRV::StorageClass::Private) ||
1686 (DstSC == SPIRV::StorageClass::Function &&
1687 SrcSC == SPIRV::StorageClass::Private))
1688 return BuildCOPY(ResVReg, SrcPtr, I);
1689
1690 // Casting from an eligible pointer to Generic.
1691 if (DstSC == SPIRV::StorageClass::Generic && isGenericCastablePtr(SrcSC))
1692 return selectUnOp(ResVReg, ResType, I, SPIRV::OpPtrCastToGeneric);
1693 // Casting from Generic to an eligible pointer.
1694 if (SrcSC == SPIRV::StorageClass::Generic && isGenericCastablePtr(DstSC))
1695 return selectUnOp(ResVReg, ResType, I, SPIRV::OpGenericCastToPtr);
1696 // Casting between 2 eligible pointers using Generic as an intermediary.
1697 if (isGenericCastablePtr(SrcSC) && isGenericCastablePtr(DstSC)) {
1698 SPIRVType *GenericPtrTy =
1699 GR.changePointerStorageClass(SrcPtrTy, SPIRV::StorageClass::Generic, I);
1700 Register Tmp = createVirtualRegister(GenericPtrTy, &GR, MRI, MRI->getMF());
1701 bool Result = BuildMI(BB, I, DL, TII.get(SPIRV::OpPtrCastToGeneric))
1702 .addDef(Tmp)
1703 .addUse(GR.getSPIRVTypeID(GenericPtrTy))
1704 .addUse(SrcPtr)
1705 .constrainAllUses(TII, TRI, RBI);
1706 return Result && BuildMI(BB, I, DL, TII.get(SPIRV::OpGenericCastToPtr))
1707 .addDef(ResVReg)
1708 .addUse(GR.getSPIRVTypeID(ResType))
1709 .addUse(Tmp)
1710 .constrainAllUses(TII, TRI, RBI);
1711 }
1712
1713 // Check if instructions from the SPV_INTEL_usm_storage_classes extension may
1714 // be applied
1715 if (isUSMStorageClass(SrcSC) && DstSC == SPIRV::StorageClass::CrossWorkgroup)
1716 return selectUnOp(ResVReg, ResType, I,
1717 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
1718 if (SrcSC == SPIRV::StorageClass::CrossWorkgroup && isUSMStorageClass(DstSC))
1719 return selectUnOp(ResVReg, ResType, I,
1720 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
1721 if (isUSMStorageClass(SrcSC) && DstSC == SPIRV::StorageClass::Generic)
1722 return selectUnOp(ResVReg, ResType, I, SPIRV::OpPtrCastToGeneric);
1723 if (SrcSC == SPIRV::StorageClass::Generic && isUSMStorageClass(DstSC))
1724 return selectUnOp(ResVReg, ResType, I, SPIRV::OpGenericCastToPtr);
1725
1726 // Bitcast for pointers requires that the address spaces must match
1727 return false;
1728 }
1729
getFCmpOpcode(unsigned PredNum)1730 static unsigned getFCmpOpcode(unsigned PredNum) {
1731 auto Pred = static_cast<CmpInst::Predicate>(PredNum);
1732 switch (Pred) {
1733 case CmpInst::FCMP_OEQ:
1734 return SPIRV::OpFOrdEqual;
1735 case CmpInst::FCMP_OGE:
1736 return SPIRV::OpFOrdGreaterThanEqual;
1737 case CmpInst::FCMP_OGT:
1738 return SPIRV::OpFOrdGreaterThan;
1739 case CmpInst::FCMP_OLE:
1740 return SPIRV::OpFOrdLessThanEqual;
1741 case CmpInst::FCMP_OLT:
1742 return SPIRV::OpFOrdLessThan;
1743 case CmpInst::FCMP_ONE:
1744 return SPIRV::OpFOrdNotEqual;
1745 case CmpInst::FCMP_ORD:
1746 return SPIRV::OpOrdered;
1747 case CmpInst::FCMP_UEQ:
1748 return SPIRV::OpFUnordEqual;
1749 case CmpInst::FCMP_UGE:
1750 return SPIRV::OpFUnordGreaterThanEqual;
1751 case CmpInst::FCMP_UGT:
1752 return SPIRV::OpFUnordGreaterThan;
1753 case CmpInst::FCMP_ULE:
1754 return SPIRV::OpFUnordLessThanEqual;
1755 case CmpInst::FCMP_ULT:
1756 return SPIRV::OpFUnordLessThan;
1757 case CmpInst::FCMP_UNE:
1758 return SPIRV::OpFUnordNotEqual;
1759 case CmpInst::FCMP_UNO:
1760 return SPIRV::OpUnordered;
1761 default:
1762 llvm_unreachable("Unknown predicate type for FCmp");
1763 }
1764 }
1765
getICmpOpcode(unsigned PredNum)1766 static unsigned getICmpOpcode(unsigned PredNum) {
1767 auto Pred = static_cast<CmpInst::Predicate>(PredNum);
1768 switch (Pred) {
1769 case CmpInst::ICMP_EQ:
1770 return SPIRV::OpIEqual;
1771 case CmpInst::ICMP_NE:
1772 return SPIRV::OpINotEqual;
1773 case CmpInst::ICMP_SGE:
1774 return SPIRV::OpSGreaterThanEqual;
1775 case CmpInst::ICMP_SGT:
1776 return SPIRV::OpSGreaterThan;
1777 case CmpInst::ICMP_SLE:
1778 return SPIRV::OpSLessThanEqual;
1779 case CmpInst::ICMP_SLT:
1780 return SPIRV::OpSLessThan;
1781 case CmpInst::ICMP_UGE:
1782 return SPIRV::OpUGreaterThanEqual;
1783 case CmpInst::ICMP_UGT:
1784 return SPIRV::OpUGreaterThan;
1785 case CmpInst::ICMP_ULE:
1786 return SPIRV::OpULessThanEqual;
1787 case CmpInst::ICMP_ULT:
1788 return SPIRV::OpULessThan;
1789 default:
1790 llvm_unreachable("Unknown predicate type for ICmp");
1791 }
1792 }
1793
getPtrCmpOpcode(unsigned Pred)1794 static unsigned getPtrCmpOpcode(unsigned Pred) {
1795 switch (static_cast<CmpInst::Predicate>(Pred)) {
1796 case CmpInst::ICMP_EQ:
1797 return SPIRV::OpPtrEqual;
1798 case CmpInst::ICMP_NE:
1799 return SPIRV::OpPtrNotEqual;
1800 default:
1801 llvm_unreachable("Unknown predicate type for pointer comparison");
1802 }
1803 }
1804
1805 // Return the logical operation, or abort if none exists.
getBoolCmpOpcode(unsigned PredNum)1806 static unsigned getBoolCmpOpcode(unsigned PredNum) {
1807 auto Pred = static_cast<CmpInst::Predicate>(PredNum);
1808 switch (Pred) {
1809 case CmpInst::ICMP_EQ:
1810 return SPIRV::OpLogicalEqual;
1811 case CmpInst::ICMP_NE:
1812 return SPIRV::OpLogicalNotEqual;
1813 default:
1814 llvm_unreachable("Unknown predicate type for Bool comparison");
1815 }
1816 }
1817
getZeroFP(const Type * LLVMFloatTy)1818 static APFloat getZeroFP(const Type *LLVMFloatTy) {
1819 if (!LLVMFloatTy)
1820 return APFloat::getZero(APFloat::IEEEsingle());
1821 switch (LLVMFloatTy->getScalarType()->getTypeID()) {
1822 case Type::HalfTyID:
1823 return APFloat::getZero(APFloat::IEEEhalf());
1824 default:
1825 case Type::FloatTyID:
1826 return APFloat::getZero(APFloat::IEEEsingle());
1827 case Type::DoubleTyID:
1828 return APFloat::getZero(APFloat::IEEEdouble());
1829 }
1830 }
1831
getOneFP(const Type * LLVMFloatTy)1832 static APFloat getOneFP(const Type *LLVMFloatTy) {
1833 if (!LLVMFloatTy)
1834 return APFloat::getOne(APFloat::IEEEsingle());
1835 switch (LLVMFloatTy->getScalarType()->getTypeID()) {
1836 case Type::HalfTyID:
1837 return APFloat::getOne(APFloat::IEEEhalf());
1838 default:
1839 case Type::FloatTyID:
1840 return APFloat::getOne(APFloat::IEEEsingle());
1841 case Type::DoubleTyID:
1842 return APFloat::getOne(APFloat::IEEEdouble());
1843 }
1844 }
1845
selectAnyOrAll(Register ResVReg,const SPIRVType * ResType,MachineInstr & I,unsigned OpAnyOrAll) const1846 bool SPIRVInstructionSelector::selectAnyOrAll(Register ResVReg,
1847 const SPIRVType *ResType,
1848 MachineInstr &I,
1849 unsigned OpAnyOrAll) const {
1850 assert(I.getNumOperands() == 3);
1851 assert(I.getOperand(2).isReg());
1852 MachineBasicBlock &BB = *I.getParent();
1853 Register InputRegister = I.getOperand(2).getReg();
1854 SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);
1855
1856 if (!InputType)
1857 report_fatal_error("Input Type could not be determined.");
1858
1859 bool IsBoolTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeBool);
1860 bool IsVectorTy = InputType->getOpcode() == SPIRV::OpTypeVector;
1861 if (IsBoolTy && !IsVectorTy) {
1862 assert(ResVReg == I.getOperand(0).getReg());
1863 return BuildCOPY(ResVReg, InputRegister, I);
1864 }
1865
1866 bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);
1867 unsigned SpirvNotEqualId =
1868 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
1869 SPIRVType *SpvBoolScalarTy = GR.getOrCreateSPIRVBoolType(I, TII);
1870 SPIRVType *SpvBoolTy = SpvBoolScalarTy;
1871 Register NotEqualReg = ResVReg;
1872
1873 if (IsVectorTy) {
1874 NotEqualReg =
1875 IsBoolTy ? InputRegister
1876 : createVirtualRegister(SpvBoolTy, &GR, MRI, MRI->getMF());
1877 const unsigned NumElts = InputType->getOperand(2).getImm();
1878 SpvBoolTy = GR.getOrCreateSPIRVVectorType(SpvBoolTy, NumElts, I, TII);
1879 }
1880
1881 bool Result = true;
1882 if (!IsBoolTy) {
1883 Register ConstZeroReg =
1884 IsFloatTy ? buildZerosValF(InputType, I) : buildZerosVal(InputType, I);
1885
1886 Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SpirvNotEqualId))
1887 .addDef(NotEqualReg)
1888 .addUse(GR.getSPIRVTypeID(SpvBoolTy))
1889 .addUse(InputRegister)
1890 .addUse(ConstZeroReg)
1891 .constrainAllUses(TII, TRI, RBI);
1892 }
1893
1894 if (!IsVectorTy)
1895 return Result;
1896
1897 return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(OpAnyOrAll))
1898 .addDef(ResVReg)
1899 .addUse(GR.getSPIRVTypeID(SpvBoolScalarTy))
1900 .addUse(NotEqualReg)
1901 .constrainAllUses(TII, TRI, RBI);
1902 }
1903
selectAll(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const1904 bool SPIRVInstructionSelector::selectAll(Register ResVReg,
1905 const SPIRVType *ResType,
1906 MachineInstr &I) const {
1907 return selectAnyOrAll(ResVReg, ResType, I, SPIRV::OpAll);
1908 }
1909
selectAny(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const1910 bool SPIRVInstructionSelector::selectAny(Register ResVReg,
1911 const SPIRVType *ResType,
1912 MachineInstr &I) const {
1913 return selectAnyOrAll(ResVReg, ResType, I, SPIRV::OpAny);
1914 }
1915
1916 // Select the OpDot instruction for the given float dot
selectFloatDot(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const1917 bool SPIRVInstructionSelector::selectFloatDot(Register ResVReg,
1918 const SPIRVType *ResType,
1919 MachineInstr &I) const {
1920 assert(I.getNumOperands() == 4);
1921 assert(I.getOperand(2).isReg());
1922 assert(I.getOperand(3).isReg());
1923
1924 [[maybe_unused]] SPIRVType *VecType =
1925 GR.getSPIRVTypeForVReg(I.getOperand(2).getReg());
1926
1927 assert(VecType->getOpcode() == SPIRV::OpTypeVector &&
1928 GR.getScalarOrVectorComponentCount(VecType) > 1 &&
1929 "dot product requires a vector of at least 2 components");
1930
1931 [[maybe_unused]] SPIRVType *EltType =
1932 GR.getSPIRVTypeForVReg(VecType->getOperand(1).getReg());
1933
1934 assert(EltType->getOpcode() == SPIRV::OpTypeFloat);
1935
1936 MachineBasicBlock &BB = *I.getParent();
1937 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpDot))
1938 .addDef(ResVReg)
1939 .addUse(GR.getSPIRVTypeID(ResType))
1940 .addUse(I.getOperand(2).getReg())
1941 .addUse(I.getOperand(3).getReg())
1942 .constrainAllUses(TII, TRI, RBI);
1943 }
1944
selectIntegerDot(Register ResVReg,const SPIRVType * ResType,MachineInstr & I,bool Signed) const1945 bool SPIRVInstructionSelector::selectIntegerDot(Register ResVReg,
1946 const SPIRVType *ResType,
1947 MachineInstr &I,
1948 bool Signed) const {
1949 assert(I.getNumOperands() == 4);
1950 assert(I.getOperand(2).isReg());
1951 assert(I.getOperand(3).isReg());
1952 MachineBasicBlock &BB = *I.getParent();
1953
1954 auto DotOp = Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
1955 return BuildMI(BB, I, I.getDebugLoc(), TII.get(DotOp))
1956 .addDef(ResVReg)
1957 .addUse(GR.getSPIRVTypeID(ResType))
1958 .addUse(I.getOperand(2).getReg())
1959 .addUse(I.getOperand(3).getReg())
1960 .constrainAllUses(TII, TRI, RBI);
1961 }
1962
1963 // Since pre-1.6 SPIRV has no integer dot implementation,
1964 // expand by piecewise multiplying and adding the results
selectIntegerDotExpansion(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const1965 bool SPIRVInstructionSelector::selectIntegerDotExpansion(
1966 Register ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
1967 assert(I.getNumOperands() == 4);
1968 assert(I.getOperand(2).isReg());
1969 assert(I.getOperand(3).isReg());
1970 MachineBasicBlock &BB = *I.getParent();
1971
1972 // Multiply the vectors, then sum the results
1973 Register Vec0 = I.getOperand(2).getReg();
1974 Register Vec1 = I.getOperand(3).getReg();
1975 Register TmpVec = MRI->createVirtualRegister(GR.getRegClass(ResType));
1976 SPIRVType *VecType = GR.getSPIRVTypeForVReg(Vec0);
1977
1978 bool Result = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIMulV))
1979 .addDef(TmpVec)
1980 .addUse(GR.getSPIRVTypeID(VecType))
1981 .addUse(Vec0)
1982 .addUse(Vec1)
1983 .constrainAllUses(TII, TRI, RBI);
1984
1985 assert(VecType->getOpcode() == SPIRV::OpTypeVector &&
1986 GR.getScalarOrVectorComponentCount(VecType) > 1 &&
1987 "dot product requires a vector of at least 2 components");
1988
1989 Register Res = MRI->createVirtualRegister(GR.getRegClass(ResType));
1990 Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
1991 .addDef(Res)
1992 .addUse(GR.getSPIRVTypeID(ResType))
1993 .addUse(TmpVec)
1994 .addImm(0)
1995 .constrainAllUses(TII, TRI, RBI);
1996
1997 for (unsigned i = 1; i < GR.getScalarOrVectorComponentCount(VecType); i++) {
1998 Register Elt = MRI->createVirtualRegister(GR.getRegClass(ResType));
1999
2000 Result &=
2001 BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
2002 .addDef(Elt)
2003 .addUse(GR.getSPIRVTypeID(ResType))
2004 .addUse(TmpVec)
2005 .addImm(i)
2006 .constrainAllUses(TII, TRI, RBI);
2007
2008 Register Sum = i < GR.getScalarOrVectorComponentCount(VecType) - 1
2009 ? MRI->createVirtualRegister(GR.getRegClass(ResType))
2010 : ResVReg;
2011
2012 Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIAddS))
2013 .addDef(Sum)
2014 .addUse(GR.getSPIRVTypeID(ResType))
2015 .addUse(Res)
2016 .addUse(Elt)
2017 .constrainAllUses(TII, TRI, RBI);
2018 Res = Sum;
2019 }
2020
2021 return Result;
2022 }
2023
2024 template <bool Signed>
selectDot4AddPacked(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const2025 bool SPIRVInstructionSelector::selectDot4AddPacked(Register ResVReg,
2026 const SPIRVType *ResType,
2027 MachineInstr &I) const {
2028 assert(I.getNumOperands() == 5);
2029 assert(I.getOperand(2).isReg());
2030 assert(I.getOperand(3).isReg());
2031 assert(I.getOperand(4).isReg());
2032 MachineBasicBlock &BB = *I.getParent();
2033
2034 Register Acc = I.getOperand(2).getReg();
2035 Register X = I.getOperand(3).getReg();
2036 Register Y = I.getOperand(4).getReg();
2037
2038 auto DotOp = Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2039 Register Dot = MRI->createVirtualRegister(GR.getRegClass(ResType));
2040 bool Result = BuildMI(BB, I, I.getDebugLoc(), TII.get(DotOp))
2041 .addDef(Dot)
2042 .addUse(GR.getSPIRVTypeID(ResType))
2043 .addUse(X)
2044 .addUse(Y)
2045 .constrainAllUses(TII, TRI, RBI);
2046
2047 return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIAddS))
2048 .addDef(ResVReg)
2049 .addUse(GR.getSPIRVTypeID(ResType))
2050 .addUse(Dot)
2051 .addUse(Acc)
2052 .constrainAllUses(TII, TRI, RBI);
2053 }
2054
2055 // Since pre-1.6 SPIRV has no DotProductInput4x8BitPacked implementation,
2056 // extract the elements of the packed inputs, multiply them and add the result
2057 // to the accumulator.
2058 template <bool Signed>
selectDot4AddPackedExpansion(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const2059 bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
2060 Register ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
2061 assert(I.getNumOperands() == 5);
2062 assert(I.getOperand(2).isReg());
2063 assert(I.getOperand(3).isReg());
2064 assert(I.getOperand(4).isReg());
2065 MachineBasicBlock &BB = *I.getParent();
2066
2067 bool Result = true;
2068
2069 Register Acc = I.getOperand(2).getReg();
2070 Register X = I.getOperand(3).getReg();
2071 Register Y = I.getOperand(4).getReg();
2072
2073 SPIRVType *EltType = GR.getOrCreateSPIRVIntegerType(8, I, TII);
2074 auto ExtractOp =
2075 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
2076
2077 bool ZeroAsNull = !STI.isShader();
2078 // Extract the i8 element, multiply and add it to the accumulator
2079 for (unsigned i = 0; i < 4; i++) {
2080 // A[i]
2081 Register AElt = MRI->createVirtualRegister(&SPIRV::IDRegClass);
2082 Result &=
2083 BuildMI(BB, I, I.getDebugLoc(), TII.get(ExtractOp))
2084 .addDef(AElt)
2085 .addUse(GR.getSPIRVTypeID(ResType))
2086 .addUse(X)
2087 .addUse(GR.getOrCreateConstInt(i * 8, I, EltType, TII, ZeroAsNull))
2088 .addUse(GR.getOrCreateConstInt(8, I, EltType, TII, ZeroAsNull))
2089 .constrainAllUses(TII, TRI, RBI);
2090
2091 // B[i]
2092 Register BElt = MRI->createVirtualRegister(&SPIRV::IDRegClass);
2093 Result &=
2094 BuildMI(BB, I, I.getDebugLoc(), TII.get(ExtractOp))
2095 .addDef(BElt)
2096 .addUse(GR.getSPIRVTypeID(ResType))
2097 .addUse(Y)
2098 .addUse(GR.getOrCreateConstInt(i * 8, I, EltType, TII, ZeroAsNull))
2099 .addUse(GR.getOrCreateConstInt(8, I, EltType, TII, ZeroAsNull))
2100 .constrainAllUses(TII, TRI, RBI);
2101
2102 // A[i] * B[i]
2103 Register Mul = MRI->createVirtualRegister(&SPIRV::IDRegClass);
2104 Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIMulS))
2105 .addDef(Mul)
2106 .addUse(GR.getSPIRVTypeID(ResType))
2107 .addUse(AElt)
2108 .addUse(BElt)
2109 .constrainAllUses(TII, TRI, RBI);
2110
2111 // Discard 24 highest-bits so that stored i32 register is i8 equivalent
2112 Register MaskMul = MRI->createVirtualRegister(&SPIRV::IDRegClass);
2113 Result &=
2114 BuildMI(BB, I, I.getDebugLoc(), TII.get(ExtractOp))
2115 .addDef(MaskMul)
2116 .addUse(GR.getSPIRVTypeID(ResType))
2117 .addUse(Mul)
2118 .addUse(GR.getOrCreateConstInt(0, I, EltType, TII, ZeroAsNull))
2119 .addUse(GR.getOrCreateConstInt(8, I, EltType, TII, ZeroAsNull))
2120 .constrainAllUses(TII, TRI, RBI);
2121
2122 // Acc = Acc + A[i] * B[i]
2123 Register Sum =
2124 i < 3 ? MRI->createVirtualRegister(&SPIRV::IDRegClass) : ResVReg;
2125 Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIAddS))
2126 .addDef(Sum)
2127 .addUse(GR.getSPIRVTypeID(ResType))
2128 .addUse(Acc)
2129 .addUse(MaskMul)
2130 .constrainAllUses(TII, TRI, RBI);
2131
2132 Acc = Sum;
2133 }
2134
2135 return Result;
2136 }
2137
2138 /// Transform saturate(x) to clamp(x, 0.0f, 1.0f) as SPIRV
2139 /// does not have a saturate builtin.
selectSaturate(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const2140 bool SPIRVInstructionSelector::selectSaturate(Register ResVReg,
2141 const SPIRVType *ResType,
2142 MachineInstr &I) const {
2143 assert(I.getNumOperands() == 3);
2144 assert(I.getOperand(2).isReg());
2145 MachineBasicBlock &BB = *I.getParent();
2146 Register VZero = buildZerosValF(ResType, I);
2147 Register VOne = buildOnesValF(ResType, I);
2148
2149 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
2150 .addDef(ResVReg)
2151 .addUse(GR.getSPIRVTypeID(ResType))
2152 .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::GLSL_std_450))
2153 .addImm(GL::FClamp)
2154 .addUse(I.getOperand(2).getReg())
2155 .addUse(VZero)
2156 .addUse(VOne)
2157 .constrainAllUses(TII, TRI, RBI);
2158 }
2159
selectSign(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const2160 bool SPIRVInstructionSelector::selectSign(Register ResVReg,
2161 const SPIRVType *ResType,
2162 MachineInstr &I) const {
2163 assert(I.getNumOperands() == 3);
2164 assert(I.getOperand(2).isReg());
2165 MachineBasicBlock &BB = *I.getParent();
2166 Register InputRegister = I.getOperand(2).getReg();
2167 SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);
2168 auto &DL = I.getDebugLoc();
2169
2170 if (!InputType)
2171 report_fatal_error("Input Type could not be determined.");
2172
2173 bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);
2174
2175 unsigned SignBitWidth = GR.getScalarOrVectorBitWidth(InputType);
2176 unsigned ResBitWidth = GR.getScalarOrVectorBitWidth(ResType);
2177
2178 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
2179
2180 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
2181 Register SignReg = NeedsConversion
2182 ? MRI->createVirtualRegister(&SPIRV::IDRegClass)
2183 : ResVReg;
2184
2185 bool Result =
2186 BuildMI(BB, I, DL, TII.get(SPIRV::OpExtInst))
2187 .addDef(SignReg)
2188 .addUse(GR.getSPIRVTypeID(InputType))
2189 .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::GLSL_std_450))
2190 .addImm(SignOpcode)
2191 .addUse(InputRegister)
2192 .constrainAllUses(TII, TRI, RBI);
2193
2194 if (NeedsConversion) {
2195 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
2196 Result &= BuildMI(*I.getParent(), I, DL, TII.get(ConvertOpcode))
2197 .addDef(ResVReg)
2198 .addUse(GR.getSPIRVTypeID(ResType))
2199 .addUse(SignReg)
2200 .constrainAllUses(TII, TRI, RBI);
2201 }
2202
2203 return Result;
2204 }
2205
selectWaveOpInst(Register ResVReg,const SPIRVType * ResType,MachineInstr & I,unsigned Opcode) const2206 bool SPIRVInstructionSelector::selectWaveOpInst(Register ResVReg,
2207 const SPIRVType *ResType,
2208 MachineInstr &I,
2209 unsigned Opcode) const {
2210 MachineBasicBlock &BB = *I.getParent();
2211 SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);
2212
2213 auto BMI = BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
2214 .addDef(ResVReg)
2215 .addUse(GR.getSPIRVTypeID(ResType))
2216 .addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I,
2217 IntTy, TII, !STI.isShader()));
2218
2219 for (unsigned J = 2; J < I.getNumOperands(); J++) {
2220 BMI.addUse(I.getOperand(J).getReg());
2221 }
2222
2223 return BMI.constrainAllUses(TII, TRI, RBI);
2224 }
2225
selectWaveActiveCountBits(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const2226 bool SPIRVInstructionSelector::selectWaveActiveCountBits(
2227 Register ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
2228
2229 SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);
2230 SPIRVType *BallotType = GR.getOrCreateSPIRVVectorType(IntTy, 4, I, TII);
2231 Register BallotReg = MRI->createVirtualRegister(GR.getRegClass(BallotType));
2232 bool Result = selectWaveOpInst(BallotReg, BallotType, I,
2233 SPIRV::OpGroupNonUniformBallot);
2234
2235 MachineBasicBlock &BB = *I.getParent();
2236 Result &= BuildMI(BB, I, I.getDebugLoc(),
2237 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
2238 .addDef(ResVReg)
2239 .addUse(GR.getSPIRVTypeID(ResType))
2240 .addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I, IntTy,
2241 TII, !STI.isShader()))
2242 .addImm(SPIRV::GroupOperation::Reduce)
2243 .addUse(BallotReg)
2244 .constrainAllUses(TII, TRI, RBI);
2245
2246 return Result;
2247 }
2248
selectWaveReduceMax(Register ResVReg,const SPIRVType * ResType,MachineInstr & I,bool IsUnsigned) const2249 bool SPIRVInstructionSelector::selectWaveReduceMax(Register ResVReg,
2250 const SPIRVType *ResType,
2251 MachineInstr &I,
2252 bool IsUnsigned) const {
2253 assert(I.getNumOperands() == 3);
2254 assert(I.getOperand(2).isReg());
2255 MachineBasicBlock &BB = *I.getParent();
2256 Register InputRegister = I.getOperand(2).getReg();
2257 SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);
2258
2259 if (!InputType)
2260 report_fatal_error("Input Type could not be determined.");
2261
2262 SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);
2263 // Retreive the operation to use based on input type
2264 bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);
2265 auto IntegerOpcodeType =
2266 IsUnsigned ? SPIRV::OpGroupNonUniformUMax : SPIRV::OpGroupNonUniformSMax;
2267 auto Opcode = IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntegerOpcodeType;
2268 return BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
2269 .addDef(ResVReg)
2270 .addUse(GR.getSPIRVTypeID(ResType))
2271 .addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I, IntTy, TII,
2272 !STI.isShader()))
2273 .addImm(SPIRV::GroupOperation::Reduce)
2274 .addUse(I.getOperand(2).getReg())
2275 .constrainAllUses(TII, TRI, RBI);
2276 }
2277
selectWaveReduceSum(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const2278 bool SPIRVInstructionSelector::selectWaveReduceSum(Register ResVReg,
2279 const SPIRVType *ResType,
2280 MachineInstr &I) const {
2281 assert(I.getNumOperands() == 3);
2282 assert(I.getOperand(2).isReg());
2283 MachineBasicBlock &BB = *I.getParent();
2284 Register InputRegister = I.getOperand(2).getReg();
2285 SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);
2286
2287 if (!InputType)
2288 report_fatal_error("Input Type could not be determined.");
2289
2290 SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);
2291 // Retreive the operation to use based on input type
2292 bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);
2293 auto Opcode =
2294 IsFloatTy ? SPIRV::OpGroupNonUniformFAdd : SPIRV::OpGroupNonUniformIAdd;
2295 return BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
2296 .addDef(ResVReg)
2297 .addUse(GR.getSPIRVTypeID(ResType))
2298 .addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I, IntTy, TII,
2299 !STI.isShader()))
2300 .addImm(SPIRV::GroupOperation::Reduce)
2301 .addUse(I.getOperand(2).getReg());
2302 }
2303
selectBitreverse(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const2304 bool SPIRVInstructionSelector::selectBitreverse(Register ResVReg,
2305 const SPIRVType *ResType,
2306 MachineInstr &I) const {
2307 MachineBasicBlock &BB = *I.getParent();
2308 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpBitReverse))
2309 .addDef(ResVReg)
2310 .addUse(GR.getSPIRVTypeID(ResType))
2311 .addUse(I.getOperand(1).getReg())
2312 .constrainAllUses(TII, TRI, RBI);
2313 }
2314
selectFreeze(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const2315 bool SPIRVInstructionSelector::selectFreeze(Register ResVReg,
2316 const SPIRVType *ResType,
2317 MachineInstr &I) const {
2318 // There is no way to implement `freeze` correctly without support on SPIR-V
2319 // standard side, but we may at least address a simple (static) case when
2320 // undef/poison value presence is obvious. The main benefit of even
2321 // incomplete `freeze` support is preventing of translation from crashing due
2322 // to lack of support on legalization and instruction selection steps.
2323 if (!I.getOperand(0).isReg() || !I.getOperand(1).isReg())
2324 return false;
2325 Register OpReg = I.getOperand(1).getReg();
2326 if (MachineInstr *Def = MRI->getVRegDef(OpReg)) {
2327 if (Def->getOpcode() == TargetOpcode::COPY)
2328 Def = MRI->getVRegDef(Def->getOperand(1).getReg());
2329 Register Reg;
2330 switch (Def->getOpcode()) {
2331 case SPIRV::ASSIGN_TYPE:
2332 if (MachineInstr *AssignToDef =
2333 MRI->getVRegDef(Def->getOperand(1).getReg())) {
2334 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
2335 Reg = Def->getOperand(2).getReg();
2336 }
2337 break;
2338 case SPIRV::OpUndef:
2339 Reg = Def->getOperand(1).getReg();
2340 break;
2341 }
2342 unsigned DestOpCode;
2343 if (Reg.isValid()) {
2344 DestOpCode = SPIRV::OpConstantNull;
2345 } else {
2346 DestOpCode = TargetOpcode::COPY;
2347 Reg = OpReg;
2348 }
2349 return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(DestOpCode))
2350 .addDef(I.getOperand(0).getReg())
2351 .addUse(Reg)
2352 .constrainAllUses(TII, TRI, RBI);
2353 }
2354 return false;
2355 }
2356
selectBuildVector(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const2357 bool SPIRVInstructionSelector::selectBuildVector(Register ResVReg,
2358 const SPIRVType *ResType,
2359 MachineInstr &I) const {
2360 unsigned N = 0;
2361 if (ResType->getOpcode() == SPIRV::OpTypeVector)
2362 N = GR.getScalarOrVectorComponentCount(ResType);
2363 else if (ResType->getOpcode() == SPIRV::OpTypeArray)
2364 N = getArrayComponentCount(MRI, ResType);
2365 else
2366 report_fatal_error("Cannot select G_BUILD_VECTOR with a non-vector result");
2367 if (I.getNumExplicitOperands() - I.getNumExplicitDefs() != N)
2368 report_fatal_error("G_BUILD_VECTOR and the result type are inconsistent");
2369
2370 // check if we may construct a constant vector
2371 bool IsConst = true;
2372 for (unsigned i = I.getNumExplicitDefs();
2373 i < I.getNumExplicitOperands() && IsConst; ++i)
2374 if (!isConstReg(MRI, I.getOperand(i).getReg()))
2375 IsConst = false;
2376
2377 if (!IsConst && N < 2)
2378 report_fatal_error(
2379 "There must be at least two constituent operands in a vector");
2380
2381 MRI->setRegClass(ResVReg, GR.getRegClass(ResType));
2382 auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
2383 TII.get(IsConst ? SPIRV::OpConstantComposite
2384 : SPIRV::OpCompositeConstruct))
2385 .addDef(ResVReg)
2386 .addUse(GR.getSPIRVTypeID(ResType));
2387 for (unsigned i = I.getNumExplicitDefs(); i < I.getNumExplicitOperands(); ++i)
2388 MIB.addUse(I.getOperand(i).getReg());
2389 return MIB.constrainAllUses(TII, TRI, RBI);
2390 }
2391
selectSplatVector(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const2392 bool SPIRVInstructionSelector::selectSplatVector(Register ResVReg,
2393 const SPIRVType *ResType,
2394 MachineInstr &I) const {
2395 unsigned N = 0;
2396 if (ResType->getOpcode() == SPIRV::OpTypeVector)
2397 N = GR.getScalarOrVectorComponentCount(ResType);
2398 else if (ResType->getOpcode() == SPIRV::OpTypeArray)
2399 N = getArrayComponentCount(MRI, ResType);
2400 else
2401 report_fatal_error("Cannot select G_SPLAT_VECTOR with a non-vector result");
2402
2403 unsigned OpIdx = I.getNumExplicitDefs();
2404 if (!I.getOperand(OpIdx).isReg())
2405 report_fatal_error("Unexpected argument in G_SPLAT_VECTOR");
2406
2407 // check if we may construct a constant vector
2408 Register OpReg = I.getOperand(OpIdx).getReg();
2409 bool IsConst = isConstReg(MRI, OpReg);
2410
2411 if (!IsConst && N < 2)
2412 report_fatal_error(
2413 "There must be at least two constituent operands in a vector");
2414
2415 MRI->setRegClass(ResVReg, GR.getRegClass(ResType));
2416 auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
2417 TII.get(IsConst ? SPIRV::OpConstantComposite
2418 : SPIRV::OpCompositeConstruct))
2419 .addDef(ResVReg)
2420 .addUse(GR.getSPIRVTypeID(ResType));
2421 for (unsigned i = 0; i < N; ++i)
2422 MIB.addUse(OpReg);
2423 return MIB.constrainAllUses(TII, TRI, RBI);
2424 }
2425
selectDiscard(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const2426 bool SPIRVInstructionSelector::selectDiscard(Register ResVReg,
2427 const SPIRVType *ResType,
2428 MachineInstr &I) const {
2429
2430 unsigned Opcode;
2431
2432 if (STI.canUseExtension(
2433 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
2434 STI.isAtLeastSPIRVVer(llvm::VersionTuple(1, 6))) {
2435 Opcode = SPIRV::OpDemoteToHelperInvocation;
2436 } else {
2437 Opcode = SPIRV::OpKill;
2438 // OpKill must be the last operation of any basic block.
2439 if (MachineInstr *NextI = I.getNextNode()) {
2440 GR.invalidateMachineInstr(NextI);
2441 NextI->removeFromParent();
2442 }
2443 }
2444
2445 MachineBasicBlock &BB = *I.getParent();
2446 return BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
2447 .constrainAllUses(TII, TRI, RBI);
2448 }
2449
selectCmp(Register ResVReg,const SPIRVType * ResType,unsigned CmpOpc,MachineInstr & I) const2450 bool SPIRVInstructionSelector::selectCmp(Register ResVReg,
2451 const SPIRVType *ResType,
2452 unsigned CmpOpc,
2453 MachineInstr &I) const {
2454 Register Cmp0 = I.getOperand(2).getReg();
2455 Register Cmp1 = I.getOperand(3).getReg();
2456 assert(GR.getSPIRVTypeForVReg(Cmp0)->getOpcode() ==
2457 GR.getSPIRVTypeForVReg(Cmp1)->getOpcode() &&
2458 "CMP operands should have the same type");
2459 return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(CmpOpc))
2460 .addDef(ResVReg)
2461 .addUse(GR.getSPIRVTypeID(ResType))
2462 .addUse(Cmp0)
2463 .addUse(Cmp1)
2464 .constrainAllUses(TII, TRI, RBI);
2465 }
2466
selectICmp(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const2467 bool SPIRVInstructionSelector::selectICmp(Register ResVReg,
2468 const SPIRVType *ResType,
2469 MachineInstr &I) const {
2470 auto Pred = I.getOperand(1).getPredicate();
2471 unsigned CmpOpc;
2472
2473 Register CmpOperand = I.getOperand(2).getReg();
2474 if (GR.isScalarOfType(CmpOperand, SPIRV::OpTypePointer))
2475 CmpOpc = getPtrCmpOpcode(Pred);
2476 else if (GR.isScalarOrVectorOfType(CmpOperand, SPIRV::OpTypeBool))
2477 CmpOpc = getBoolCmpOpcode(Pred);
2478 else
2479 CmpOpc = getICmpOpcode(Pred);
2480 return selectCmp(ResVReg, ResType, CmpOpc, I);
2481 }
2482
2483 std::pair<Register, bool>
buildI32Constant(uint32_t Val,MachineInstr & I,const SPIRVType * ResType) const2484 SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &I,
2485 const SPIRVType *ResType) const {
2486 Type *LLVMTy = IntegerType::get(GR.CurMF->getFunction().getContext(), 32);
2487 const SPIRVType *SpvI32Ty =
2488 ResType ? ResType : GR.getOrCreateSPIRVIntegerType(32, I, TII);
2489 // Find a constant in DT or build a new one.
2490 auto ConstInt = ConstantInt::get(LLVMTy, Val);
2491 Register NewReg = GR.find(ConstInt, GR.CurMF);
2492 bool Result = true;
2493 if (!NewReg.isValid()) {
2494 NewReg = MRI->createGenericVirtualRegister(LLT::scalar(64));
2495 MachineBasicBlock &BB = *I.getParent();
2496 MachineInstr *MI =
2497 Val == 0
2498 ? BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantNull))
2499 .addDef(NewReg)
2500 .addUse(GR.getSPIRVTypeID(SpvI32Ty))
2501 : BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantI))
2502 .addDef(NewReg)
2503 .addUse(GR.getSPIRVTypeID(SpvI32Ty))
2504 .addImm(APInt(32, Val).getZExtValue());
2505 Result &= constrainSelectedInstRegOperands(*MI, TII, TRI, RBI);
2506 GR.add(ConstInt, MI);
2507 }
2508 return {NewReg, Result};
2509 }
2510
selectFCmp(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const2511 bool SPIRVInstructionSelector::selectFCmp(Register ResVReg,
2512 const SPIRVType *ResType,
2513 MachineInstr &I) const {
2514 unsigned CmpOp = getFCmpOpcode(I.getOperand(1).getPredicate());
2515 return selectCmp(ResVReg, ResType, CmpOp, I);
2516 }
2517
buildZerosVal(const SPIRVType * ResType,MachineInstr & I) const2518 Register SPIRVInstructionSelector::buildZerosVal(const SPIRVType *ResType,
2519 MachineInstr &I) const {
2520 // OpenCL uses nulls for Zero. In HLSL we don't use null constants.
2521 bool ZeroAsNull = !STI.isShader();
2522 if (ResType->getOpcode() == SPIRV::OpTypeVector)
2523 return GR.getOrCreateConstVector(0UL, I, ResType, TII, ZeroAsNull);
2524 return GR.getOrCreateConstInt(0, I, ResType, TII, ZeroAsNull);
2525 }
2526
buildZerosValF(const SPIRVType * ResType,MachineInstr & I) const2527 Register SPIRVInstructionSelector::buildZerosValF(const SPIRVType *ResType,
2528 MachineInstr &I) const {
2529 // OpenCL uses nulls for Zero. In HLSL we don't use null constants.
2530 bool ZeroAsNull = !STI.isShader();
2531 APFloat VZero = getZeroFP(GR.getTypeForSPIRVType(ResType));
2532 if (ResType->getOpcode() == SPIRV::OpTypeVector)
2533 return GR.getOrCreateConstVector(VZero, I, ResType, TII, ZeroAsNull);
2534 return GR.getOrCreateConstFP(VZero, I, ResType, TII, ZeroAsNull);
2535 }
2536
buildOnesValF(const SPIRVType * ResType,MachineInstr & I) const2537 Register SPIRVInstructionSelector::buildOnesValF(const SPIRVType *ResType,
2538 MachineInstr &I) const {
2539 // OpenCL uses nulls for Zero. In HLSL we don't use null constants.
2540 bool ZeroAsNull = !STI.isShader();
2541 APFloat VOne = getOneFP(GR.getTypeForSPIRVType(ResType));
2542 if (ResType->getOpcode() == SPIRV::OpTypeVector)
2543 return GR.getOrCreateConstVector(VOne, I, ResType, TII, ZeroAsNull);
2544 return GR.getOrCreateConstFP(VOne, I, ResType, TII, ZeroAsNull);
2545 }
2546
buildOnesVal(bool AllOnes,const SPIRVType * ResType,MachineInstr & I) const2547 Register SPIRVInstructionSelector::buildOnesVal(bool AllOnes,
2548 const SPIRVType *ResType,
2549 MachineInstr &I) const {
2550 unsigned BitWidth = GR.getScalarOrVectorBitWidth(ResType);
2551 APInt One =
2552 AllOnes ? APInt::getAllOnes(BitWidth) : APInt::getOneBitSet(BitWidth, 0);
2553 if (ResType->getOpcode() == SPIRV::OpTypeVector)
2554 return GR.getOrCreateConstVector(One.getZExtValue(), I, ResType, TII);
2555 return GR.getOrCreateConstInt(One.getZExtValue(), I, ResType, TII);
2556 }
2557
selectSelect(Register ResVReg,const SPIRVType * ResType,MachineInstr & I,bool IsSigned) const2558 bool SPIRVInstructionSelector::selectSelect(Register ResVReg,
2559 const SPIRVType *ResType,
2560 MachineInstr &I,
2561 bool IsSigned) const {
2562 // To extend a bool, we need to use OpSelect between constants.
2563 Register ZeroReg = buildZerosVal(ResType, I);
2564 Register OneReg = buildOnesVal(IsSigned, ResType, I);
2565 bool IsScalarBool =
2566 GR.isScalarOfType(I.getOperand(1).getReg(), SPIRV::OpTypeBool);
2567 unsigned Opcode =
2568 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
2569 return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(Opcode))
2570 .addDef(ResVReg)
2571 .addUse(GR.getSPIRVTypeID(ResType))
2572 .addUse(I.getOperand(1).getReg())
2573 .addUse(OneReg)
2574 .addUse(ZeroReg)
2575 .constrainAllUses(TII, TRI, RBI);
2576 }
2577
selectIToF(Register ResVReg,const SPIRVType * ResType,MachineInstr & I,bool IsSigned,unsigned Opcode) const2578 bool SPIRVInstructionSelector::selectIToF(Register ResVReg,
2579 const SPIRVType *ResType,
2580 MachineInstr &I, bool IsSigned,
2581 unsigned Opcode) const {
2582 Register SrcReg = I.getOperand(1).getReg();
2583 // We can convert bool value directly to float type without OpConvert*ToF,
2584 // however the translator generates OpSelect+OpConvert*ToF, so we do the same.
2585 if (GR.isScalarOrVectorOfType(I.getOperand(1).getReg(), SPIRV::OpTypeBool)) {
2586 unsigned BitWidth = GR.getScalarOrVectorBitWidth(ResType);
2587 SPIRVType *TmpType = GR.getOrCreateSPIRVIntegerType(BitWidth, I, TII);
2588 if (ResType->getOpcode() == SPIRV::OpTypeVector) {
2589 const unsigned NumElts = ResType->getOperand(2).getImm();
2590 TmpType = GR.getOrCreateSPIRVVectorType(TmpType, NumElts, I, TII);
2591 }
2592 SrcReg = createVirtualRegister(TmpType, &GR, MRI, MRI->getMF());
2593 selectSelect(SrcReg, TmpType, I, false);
2594 }
2595 return selectOpWithSrcs(ResVReg, ResType, I, {SrcReg}, Opcode);
2596 }
2597
selectExt(Register ResVReg,const SPIRVType * ResType,MachineInstr & I,bool IsSigned) const2598 bool SPIRVInstructionSelector::selectExt(Register ResVReg,
2599 const SPIRVType *ResType,
2600 MachineInstr &I, bool IsSigned) const {
2601 Register SrcReg = I.getOperand(1).getReg();
2602 if (GR.isScalarOrVectorOfType(SrcReg, SPIRV::OpTypeBool))
2603 return selectSelect(ResVReg, ResType, I, IsSigned);
2604
2605 SPIRVType *SrcType = GR.getSPIRVTypeForVReg(SrcReg);
2606 if (SrcType == ResType)
2607 return BuildCOPY(ResVReg, SrcReg, I);
2608
2609 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
2610 return selectUnOp(ResVReg, ResType, I, Opcode);
2611 }
2612
selectSUCmp(Register ResVReg,const SPIRVType * ResType,MachineInstr & I,bool IsSigned) const2613 bool SPIRVInstructionSelector::selectSUCmp(Register ResVReg,
2614 const SPIRVType *ResType,
2615 MachineInstr &I,
2616 bool IsSigned) const {
2617 MachineIRBuilder MIRBuilder(I);
2618 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
2619 MachineBasicBlock &BB = *I.getParent();
2620 // Ensure we have bool.
2621 SPIRVType *BoolType = GR.getOrCreateSPIRVBoolType(I, TII);
2622 unsigned N = GR.getScalarOrVectorComponentCount(ResType);
2623 if (N > 1)
2624 BoolType = GR.getOrCreateSPIRVVectorType(BoolType, N, I, TII);
2625 Register BoolTypeReg = GR.getSPIRVTypeID(BoolType);
2626 // Build less-than-equal and less-than.
2627 // TODO: replace with one-liner createVirtualRegister() from
2628 // llvm/lib/Target/SPIRV/SPIRVUtils.cpp when PR #116609 is merged.
2629 Register IsLessEqReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
2630 MRI->setType(IsLessEqReg, LLT::scalar(64));
2631 GR.assignSPIRVTypeToVReg(ResType, IsLessEqReg, MIRBuilder.getMF());
2632 bool Result = BuildMI(BB, I, I.getDebugLoc(),
2633 TII.get(IsSigned ? SPIRV::OpSLessThanEqual
2634 : SPIRV::OpULessThanEqual))
2635 .addDef(IsLessEqReg)
2636 .addUse(BoolTypeReg)
2637 .addUse(I.getOperand(1).getReg())
2638 .addUse(I.getOperand(2).getReg())
2639 .constrainAllUses(TII, TRI, RBI);
2640 Register IsLessReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
2641 MRI->setType(IsLessReg, LLT::scalar(64));
2642 GR.assignSPIRVTypeToVReg(ResType, IsLessReg, MIRBuilder.getMF());
2643 Result &= BuildMI(BB, I, I.getDebugLoc(),
2644 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
2645 .addDef(IsLessReg)
2646 .addUse(BoolTypeReg)
2647 .addUse(I.getOperand(1).getReg())
2648 .addUse(I.getOperand(2).getReg())
2649 .constrainAllUses(TII, TRI, RBI);
2650 // Build selects.
2651 Register ResTypeReg = GR.getSPIRVTypeID(ResType);
2652 Register NegOneOrZeroReg =
2653 MRI->createVirtualRegister(GR.getRegClass(ResType));
2654 MRI->setType(NegOneOrZeroReg, LLT::scalar(64));
2655 GR.assignSPIRVTypeToVReg(ResType, NegOneOrZeroReg, MIRBuilder.getMF());
2656 unsigned SelectOpcode =
2657 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
2658 Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SelectOpcode))
2659 .addDef(NegOneOrZeroReg)
2660 .addUse(ResTypeReg)
2661 .addUse(IsLessReg)
2662 .addUse(buildOnesVal(true, ResType, I)) // -1
2663 .addUse(buildZerosVal(ResType, I))
2664 .constrainAllUses(TII, TRI, RBI);
2665 return Result & BuildMI(BB, I, I.getDebugLoc(), TII.get(SelectOpcode))
2666 .addDef(ResVReg)
2667 .addUse(ResTypeReg)
2668 .addUse(IsLessEqReg)
2669 .addUse(NegOneOrZeroReg) // -1 or 0
2670 .addUse(buildOnesVal(false, ResType, I))
2671 .constrainAllUses(TII, TRI, RBI);
2672 }
2673
selectIntToBool(Register IntReg,Register ResVReg,MachineInstr & I,const SPIRVType * IntTy,const SPIRVType * BoolTy) const2674 bool SPIRVInstructionSelector::selectIntToBool(Register IntReg,
2675 Register ResVReg,
2676 MachineInstr &I,
2677 const SPIRVType *IntTy,
2678 const SPIRVType *BoolTy) const {
2679 // To truncate to a bool, we use OpBitwiseAnd 1 and OpINotEqual to zero.
2680 Register BitIntReg = createVirtualRegister(IntTy, &GR, MRI, MRI->getMF());
2681 bool IsVectorTy = IntTy->getOpcode() == SPIRV::OpTypeVector;
2682 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
2683 Register Zero = buildZerosVal(IntTy, I);
2684 Register One = buildOnesVal(false, IntTy, I);
2685 MachineBasicBlock &BB = *I.getParent();
2686 bool Result = BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
2687 .addDef(BitIntReg)
2688 .addUse(GR.getSPIRVTypeID(IntTy))
2689 .addUse(IntReg)
2690 .addUse(One)
2691 .constrainAllUses(TII, TRI, RBI);
2692 return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpINotEqual))
2693 .addDef(ResVReg)
2694 .addUse(GR.getSPIRVTypeID(BoolTy))
2695 .addUse(BitIntReg)
2696 .addUse(Zero)
2697 .constrainAllUses(TII, TRI, RBI);
2698 }
2699
selectTrunc(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const2700 bool SPIRVInstructionSelector::selectTrunc(Register ResVReg,
2701 const SPIRVType *ResType,
2702 MachineInstr &I) const {
2703 Register IntReg = I.getOperand(1).getReg();
2704 const SPIRVType *ArgType = GR.getSPIRVTypeForVReg(IntReg);
2705 if (GR.isScalarOrVectorOfType(ResVReg, SPIRV::OpTypeBool))
2706 return selectIntToBool(IntReg, ResVReg, I, ArgType, ResType);
2707 if (ArgType == ResType)
2708 return BuildCOPY(ResVReg, IntReg, I);
2709 bool IsSigned = GR.isScalarOrVectorSigned(ResType);
2710 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
2711 return selectUnOp(ResVReg, ResType, I, Opcode);
2712 }
2713
selectConst(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const2714 bool SPIRVInstructionSelector::selectConst(Register ResVReg,
2715 const SPIRVType *ResType,
2716 MachineInstr &I) const {
2717 unsigned Opcode = I.getOpcode();
2718 unsigned TpOpcode = ResType->getOpcode();
2719 Register Reg;
2720 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
2721 assert(Opcode == TargetOpcode::G_CONSTANT &&
2722 I.getOperand(1).getCImm()->isZero());
2723 MachineBasicBlock &DepMBB = I.getMF()->front();
2724 MachineIRBuilder MIRBuilder(DepMBB, DepMBB.getFirstNonPHI());
2725 Reg = GR.getOrCreateConstNullPtr(MIRBuilder, ResType);
2726 } else if (Opcode == TargetOpcode::G_FCONSTANT) {
2727 Reg = GR.getOrCreateConstFP(I.getOperand(1).getFPImm()->getValue(), I,
2728 ResType, TII, !STI.isShader());
2729 } else {
2730 Reg = GR.getOrCreateConstInt(I.getOperand(1).getCImm()->getZExtValue(), I,
2731 ResType, TII, !STI.isShader());
2732 }
2733 return Reg == ResVReg ? true : BuildCOPY(ResVReg, Reg, I);
2734 }
2735
selectOpUndef(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const2736 bool SPIRVInstructionSelector::selectOpUndef(Register ResVReg,
2737 const SPIRVType *ResType,
2738 MachineInstr &I) const {
2739 return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpUndef))
2740 .addDef(ResVReg)
2741 .addUse(GR.getSPIRVTypeID(ResType))
2742 .constrainAllUses(TII, TRI, RBI);
2743 }
2744
selectInsertVal(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const2745 bool SPIRVInstructionSelector::selectInsertVal(Register ResVReg,
2746 const SPIRVType *ResType,
2747 MachineInstr &I) const {
2748 MachineBasicBlock &BB = *I.getParent();
2749 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeInsert))
2750 .addDef(ResVReg)
2751 .addUse(GR.getSPIRVTypeID(ResType))
2752 // object to insert
2753 .addUse(I.getOperand(3).getReg())
2754 // composite to insert into
2755 .addUse(I.getOperand(2).getReg());
2756 for (unsigned i = 4; i < I.getNumOperands(); i++)
2757 MIB.addImm(foldImm(I.getOperand(i), MRI));
2758 return MIB.constrainAllUses(TII, TRI, RBI);
2759 }
2760
selectExtractVal(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const2761 bool SPIRVInstructionSelector::selectExtractVal(Register ResVReg,
2762 const SPIRVType *ResType,
2763 MachineInstr &I) const {
2764 MachineBasicBlock &BB = *I.getParent();
2765 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
2766 .addDef(ResVReg)
2767 .addUse(GR.getSPIRVTypeID(ResType))
2768 .addUse(I.getOperand(2).getReg());
2769 for (unsigned i = 3; i < I.getNumOperands(); i++)
2770 MIB.addImm(foldImm(I.getOperand(i), MRI));
2771 return MIB.constrainAllUses(TII, TRI, RBI);
2772 }
2773
selectInsertElt(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const2774 bool SPIRVInstructionSelector::selectInsertElt(Register ResVReg,
2775 const SPIRVType *ResType,
2776 MachineInstr &I) const {
2777 if (getImm(I.getOperand(4), MRI))
2778 return selectInsertVal(ResVReg, ResType, I);
2779 MachineBasicBlock &BB = *I.getParent();
2780 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpVectorInsertDynamic))
2781 .addDef(ResVReg)
2782 .addUse(GR.getSPIRVTypeID(ResType))
2783 .addUse(I.getOperand(2).getReg())
2784 .addUse(I.getOperand(3).getReg())
2785 .addUse(I.getOperand(4).getReg())
2786 .constrainAllUses(TII, TRI, RBI);
2787 }
2788
selectExtractElt(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const2789 bool SPIRVInstructionSelector::selectExtractElt(Register ResVReg,
2790 const SPIRVType *ResType,
2791 MachineInstr &I) const {
2792 if (getImm(I.getOperand(3), MRI))
2793 return selectExtractVal(ResVReg, ResType, I);
2794 MachineBasicBlock &BB = *I.getParent();
2795 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpVectorExtractDynamic))
2796 .addDef(ResVReg)
2797 .addUse(GR.getSPIRVTypeID(ResType))
2798 .addUse(I.getOperand(2).getReg())
2799 .addUse(I.getOperand(3).getReg())
2800 .constrainAllUses(TII, TRI, RBI);
2801 }
2802
selectGEP(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const2803 bool SPIRVInstructionSelector::selectGEP(Register ResVReg,
2804 const SPIRVType *ResType,
2805 MachineInstr &I) const {
2806 const bool IsGEPInBounds = I.getOperand(2).getImm();
2807
2808 // OpAccessChain could be used for OpenCL, but the SPIRV-LLVM Translator only
2809 // relies on PtrAccessChain, so we'll try not to deviate. For Vulkan however,
2810 // we have to use Op[InBounds]AccessChain.
2811 const unsigned Opcode = STI.isLogicalSPIRV()
2812 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
2813 : SPIRV::OpAccessChain)
2814 : (IsGEPInBounds ? SPIRV::OpInBoundsPtrAccessChain
2815 : SPIRV::OpPtrAccessChain);
2816
2817 auto Res = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(Opcode))
2818 .addDef(ResVReg)
2819 .addUse(GR.getSPIRVTypeID(ResType))
2820 // Object to get a pointer to.
2821 .addUse(I.getOperand(3).getReg());
2822 // Adding indices.
2823 const unsigned StartingIndex =
2824 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
2825 ? 5
2826 : 4;
2827 for (unsigned i = StartingIndex; i < I.getNumExplicitOperands(); ++i)
2828 Res.addUse(I.getOperand(i).getReg());
2829 return Res.constrainAllUses(TII, TRI, RBI);
2830 }
2831
2832 // Maybe wrap a value into OpSpecConstantOp
wrapIntoSpecConstantOp(MachineInstr & I,SmallVector<Register> & CompositeArgs) const2833 bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
2834 MachineInstr &I, SmallVector<Register> &CompositeArgs) const {
2835 bool Result = true;
2836 unsigned Lim = I.getNumExplicitOperands();
2837 for (unsigned i = I.getNumExplicitDefs() + 1; i < Lim; ++i) {
2838 Register OpReg = I.getOperand(i).getReg();
2839 MachineInstr *OpDefine = MRI->getVRegDef(OpReg);
2840 SPIRVType *OpType = GR.getSPIRVTypeForVReg(OpReg);
2841 SmallPtrSet<SPIRVType *, 4> Visited;
2842 if (!OpDefine || !OpType || isConstReg(MRI, OpDefine, Visited) ||
2843 OpDefine->getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
2844 GR.isAggregateType(OpType)) {
2845 // The case of G_ADDRSPACE_CAST inside spv_const_composite() is processed
2846 // by selectAddrSpaceCast()
2847 CompositeArgs.push_back(OpReg);
2848 continue;
2849 }
2850 MachineFunction *MF = I.getMF();
2851 Register WrapReg = GR.find(OpDefine, MF);
2852 if (WrapReg.isValid()) {
2853 CompositeArgs.push_back(WrapReg);
2854 continue;
2855 }
2856 // Create a new register for the wrapper
2857 WrapReg = MRI->createVirtualRegister(GR.getRegClass(OpType));
2858 CompositeArgs.push_back(WrapReg);
2859 // Decorate the wrapper register and generate a new instruction
2860 MRI->setType(WrapReg, LLT::pointer(0, 64));
2861 GR.assignSPIRVTypeToVReg(OpType, WrapReg, *MF);
2862 auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
2863 TII.get(SPIRV::OpSpecConstantOp))
2864 .addDef(WrapReg)
2865 .addUse(GR.getSPIRVTypeID(OpType))
2866 .addImm(static_cast<uint32_t>(SPIRV::Opcode::Bitcast))
2867 .addUse(OpReg);
2868 GR.add(OpDefine, MIB);
2869 Result = MIB.constrainAllUses(TII, TRI, RBI);
2870 if (!Result)
2871 break;
2872 }
2873 return Result;
2874 }
2875
selectIntrinsic(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const2876 bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
2877 const SPIRVType *ResType,
2878 MachineInstr &I) const {
2879 MachineBasicBlock &BB = *I.getParent();
2880 Intrinsic::ID IID = cast<GIntrinsic>(I).getIntrinsicID();
2881 switch (IID) {
2882 case Intrinsic::spv_load:
2883 return selectLoad(ResVReg, ResType, I);
2884 case Intrinsic::spv_store:
2885 return selectStore(I);
2886 case Intrinsic::spv_extractv:
2887 return selectExtractVal(ResVReg, ResType, I);
2888 case Intrinsic::spv_insertv:
2889 return selectInsertVal(ResVReg, ResType, I);
2890 case Intrinsic::spv_extractelt:
2891 return selectExtractElt(ResVReg, ResType, I);
2892 case Intrinsic::spv_insertelt:
2893 return selectInsertElt(ResVReg, ResType, I);
2894 case Intrinsic::spv_gep:
2895 return selectGEP(ResVReg, ResType, I);
2896 case Intrinsic::spv_unref_global:
2897 case Intrinsic::spv_init_global: {
2898 MachineInstr *MI = MRI->getVRegDef(I.getOperand(1).getReg());
2899 MachineInstr *Init = I.getNumExplicitOperands() > 2
2900 ? MRI->getVRegDef(I.getOperand(2).getReg())
2901 : nullptr;
2902 assert(MI);
2903 Register GVarVReg = MI->getOperand(0).getReg();
2904 bool Res = selectGlobalValue(GVarVReg, *MI, Init);
2905 // We violate SSA form by inserting OpVariable and still having a gMIR
2906 // instruction %vreg = G_GLOBAL_VALUE @gvar. We need to fix this by erasing
2907 // the duplicated definition.
2908 if (MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
2909 GR.invalidateMachineInstr(MI);
2910 MI->removeFromParent();
2911 }
2912 return Res;
2913 }
2914 case Intrinsic::spv_undef: {
2915 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpUndef))
2916 .addDef(ResVReg)
2917 .addUse(GR.getSPIRVTypeID(ResType));
2918 return MIB.constrainAllUses(TII, TRI, RBI);
2919 }
2920 case Intrinsic::spv_const_composite: {
2921 // If no values are attached, the composite is null constant.
2922 bool IsNull = I.getNumExplicitDefs() + 1 == I.getNumExplicitOperands();
2923 SmallVector<Register> CompositeArgs;
2924 MRI->setRegClass(ResVReg, GR.getRegClass(ResType));
2925
2926 // skip type MD node we already used when generated assign.type for this
2927 if (!IsNull) {
2928 if (!wrapIntoSpecConstantOp(I, CompositeArgs))
2929 return false;
2930 MachineIRBuilder MIR(I);
2931 SmallVector<MachineInstr *, 4> Instructions = createContinuedInstructions(
2932 MIR, SPIRV::OpConstantComposite, 3,
2933 SPIRV::OpConstantCompositeContinuedINTEL, CompositeArgs, ResVReg,
2934 GR.getSPIRVTypeID(ResType));
2935 for (auto *Instr : Instructions) {
2936 Instr->setDebugLoc(I.getDebugLoc());
2937 if (!constrainSelectedInstRegOperands(*Instr, TII, TRI, RBI))
2938 return false;
2939 }
2940 return true;
2941 } else {
2942 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantNull))
2943 .addDef(ResVReg)
2944 .addUse(GR.getSPIRVTypeID(ResType));
2945 return MIB.constrainAllUses(TII, TRI, RBI);
2946 }
2947 }
2948 case Intrinsic::spv_assign_name: {
2949 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpName));
2950 MIB.addUse(I.getOperand(I.getNumExplicitDefs() + 1).getReg());
2951 for (unsigned i = I.getNumExplicitDefs() + 2;
2952 i < I.getNumExplicitOperands(); ++i) {
2953 MIB.addImm(I.getOperand(i).getImm());
2954 }
2955 return MIB.constrainAllUses(TII, TRI, RBI);
2956 }
2957 case Intrinsic::spv_switch: {
2958 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSwitch));
2959 for (unsigned i = 1; i < I.getNumExplicitOperands(); ++i) {
2960 if (I.getOperand(i).isReg())
2961 MIB.addReg(I.getOperand(i).getReg());
2962 else if (I.getOperand(i).isCImm())
2963 addNumImm(I.getOperand(i).getCImm()->getValue(), MIB);
2964 else if (I.getOperand(i).isMBB())
2965 MIB.addMBB(I.getOperand(i).getMBB());
2966 else
2967 llvm_unreachable("Unexpected OpSwitch operand");
2968 }
2969 return MIB.constrainAllUses(TII, TRI, RBI);
2970 }
2971 case Intrinsic::spv_loop_merge: {
2972 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpLoopMerge));
2973 for (unsigned i = 1; i < I.getNumExplicitOperands(); ++i) {
2974 if (I.getOperand(i).isMBB())
2975 MIB.addMBB(I.getOperand(i).getMBB());
2976 else
2977 MIB.addImm(foldImm(I.getOperand(i), MRI));
2978 }
2979 return MIB.constrainAllUses(TII, TRI, RBI);
2980 }
2981 case Intrinsic::spv_selection_merge: {
2982 auto MIB =
2983 BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSelectionMerge));
2984 assert(I.getOperand(1).isMBB() &&
2985 "operand 1 to spv_selection_merge must be a basic block");
2986 MIB.addMBB(I.getOperand(1).getMBB());
2987 MIB.addImm(getSelectionOperandForImm(I.getOperand(2).getImm()));
2988 return MIB.constrainAllUses(TII, TRI, RBI);
2989 }
2990 case Intrinsic::spv_cmpxchg:
2991 return selectAtomicCmpXchg(ResVReg, ResType, I);
2992 case Intrinsic::spv_unreachable:
2993 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpUnreachable))
2994 .constrainAllUses(TII, TRI, RBI);
2995 case Intrinsic::spv_alloca:
2996 return selectFrameIndex(ResVReg, ResType, I);
2997 case Intrinsic::spv_alloca_array:
2998 return selectAllocaArray(ResVReg, ResType, I);
2999 case Intrinsic::spv_assume:
3000 if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_expect_assume))
3001 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpAssumeTrueKHR))
3002 .addUse(I.getOperand(1).getReg())
3003 .constrainAllUses(TII, TRI, RBI);
3004 break;
3005 case Intrinsic::spv_expect:
3006 if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_expect_assume))
3007 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExpectKHR))
3008 .addDef(ResVReg)
3009 .addUse(GR.getSPIRVTypeID(ResType))
3010 .addUse(I.getOperand(2).getReg())
3011 .addUse(I.getOperand(3).getReg())
3012 .constrainAllUses(TII, TRI, RBI);
3013 break;
3014 case Intrinsic::arithmetic_fence:
3015 if (STI.canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence))
3016 return BuildMI(BB, I, I.getDebugLoc(),
3017 TII.get(SPIRV::OpArithmeticFenceEXT))
3018 .addDef(ResVReg)
3019 .addUse(GR.getSPIRVTypeID(ResType))
3020 .addUse(I.getOperand(2).getReg())
3021 .constrainAllUses(TII, TRI, RBI);
3022 else
3023 return BuildCOPY(ResVReg, I.getOperand(2).getReg(), I);
3024 break;
3025 case Intrinsic::spv_thread_id:
3026 // The HLSL SV_DispatchThreadID semantic is lowered to llvm.spv.thread.id
3027 // intrinsic in LLVM IR for SPIR-V backend.
3028 //
3029 // In SPIR-V backend, llvm.spv.thread.id is now correctly translated to a
3030 // `GlobalInvocationId` builtin variable
3031 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
3032 ResType, I);
3033 case Intrinsic::spv_thread_id_in_group:
3034 // The HLSL SV_GroupThreadId semantic is lowered to
3035 // llvm.spv.thread.id.in.group intrinsic in LLVM IR for SPIR-V backend.
3036 //
3037 // In SPIR-V backend, llvm.spv.thread.id.in.group is now correctly
3038 // translated to a `LocalInvocationId` builtin variable
3039 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
3040 ResType, I);
3041 case Intrinsic::spv_group_id:
3042 // The HLSL SV_GroupId semantic is lowered to
3043 // llvm.spv.group.id intrinsic in LLVM IR for SPIR-V backend.
3044 //
3045 // In SPIR-V backend, llvm.spv.group.id is now translated to a `WorkgroupId`
3046 // builtin variable
3047 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
3048 I);
3049 case Intrinsic::spv_flattened_thread_id_in_group:
3050 // The HLSL SV_GroupIndex semantic is lowered to
3051 // llvm.spv.flattened.thread.id.in.group() intrinsic in LLVM IR for SPIR-V
3052 // backend.
3053 //
3054 // In SPIR-V backend, llvm.spv.flattened.thread.id.in.group is translated to
3055 // a `LocalInvocationIndex` builtin variable
3056 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
3057 ResType, I);
3058 case Intrinsic::spv_workgroup_size:
3059 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
3060 ResType, I);
3061 case Intrinsic::spv_global_size:
3062 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
3063 I);
3064 case Intrinsic::spv_global_offset:
3065 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
3066 ResType, I);
3067 case Intrinsic::spv_num_workgroups:
3068 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
3069 ResType, I);
3070 case Intrinsic::spv_subgroup_size:
3071 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
3072 I);
3073 case Intrinsic::spv_num_subgroups:
3074 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
3075 I);
3076 case Intrinsic::spv_subgroup_id:
3077 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType, I);
3078 case Intrinsic::spv_subgroup_local_invocation_id:
3079 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
3080 ResVReg, ResType, I);
3081 case Intrinsic::spv_subgroup_max_size:
3082 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
3083 I);
3084 case Intrinsic::spv_fdot:
3085 return selectFloatDot(ResVReg, ResType, I);
3086 case Intrinsic::spv_udot:
3087 case Intrinsic::spv_sdot:
3088 if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
3089 STI.isAtLeastSPIRVVer(VersionTuple(1, 6)))
3090 return selectIntegerDot(ResVReg, ResType, I,
3091 /*Signed=*/IID == Intrinsic::spv_sdot);
3092 return selectIntegerDotExpansion(ResVReg, ResType, I);
3093 case Intrinsic::spv_dot4add_i8packed:
3094 if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
3095 STI.isAtLeastSPIRVVer(VersionTuple(1, 6)))
3096 return selectDot4AddPacked<true>(ResVReg, ResType, I);
3097 return selectDot4AddPackedExpansion<true>(ResVReg, ResType, I);
3098 case Intrinsic::spv_dot4add_u8packed:
3099 if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
3100 STI.isAtLeastSPIRVVer(VersionTuple(1, 6)))
3101 return selectDot4AddPacked<false>(ResVReg, ResType, I);
3102 return selectDot4AddPackedExpansion<false>(ResVReg, ResType, I);
3103 case Intrinsic::spv_all:
3104 return selectAll(ResVReg, ResType, I);
3105 case Intrinsic::spv_any:
3106 return selectAny(ResVReg, ResType, I);
3107 case Intrinsic::spv_cross:
3108 return selectExtInst(ResVReg, ResType, I, CL::cross, GL::Cross);
3109 case Intrinsic::spv_distance:
3110 return selectExtInst(ResVReg, ResType, I, CL::distance, GL::Distance);
3111 case Intrinsic::spv_lerp:
3112 return selectExtInst(ResVReg, ResType, I, CL::mix, GL::FMix);
3113 case Intrinsic::spv_length:
3114 return selectExtInst(ResVReg, ResType, I, CL::length, GL::Length);
3115 case Intrinsic::spv_degrees:
3116 return selectExtInst(ResVReg, ResType, I, CL::degrees, GL::Degrees);
3117 case Intrinsic::spv_faceforward:
3118 return selectExtInst(ResVReg, ResType, I, GL::FaceForward);
3119 case Intrinsic::spv_frac:
3120 return selectExtInst(ResVReg, ResType, I, CL::fract, GL::Fract);
3121 case Intrinsic::spv_normalize:
3122 return selectExtInst(ResVReg, ResType, I, CL::normalize, GL::Normalize);
3123 case Intrinsic::spv_reflect:
3124 return selectExtInst(ResVReg, ResType, I, GL::Reflect);
3125 case Intrinsic::spv_rsqrt:
3126 return selectExtInst(ResVReg, ResType, I, CL::rsqrt, GL::InverseSqrt);
3127 case Intrinsic::spv_sign:
3128 return selectSign(ResVReg, ResType, I);
3129 case Intrinsic::spv_smoothstep:
3130 return selectExtInst(ResVReg, ResType, I, CL::smoothstep, GL::SmoothStep);
3131 case Intrinsic::spv_firstbituhigh: // There is no CL equivalent of FindUMsb
3132 return selectFirstBitHigh(ResVReg, ResType, I, /*IsSigned=*/false);
3133 case Intrinsic::spv_firstbitshigh: // There is no CL equivalent of FindSMsb
3134 return selectFirstBitHigh(ResVReg, ResType, I, /*IsSigned=*/true);
3135 case Intrinsic::spv_firstbitlow: // There is no CL equivlent of FindILsb
3136 return selectFirstBitLow(ResVReg, ResType, I);
3137 case Intrinsic::spv_group_memory_barrier_with_group_sync: {
3138 bool Result = true;
3139 auto MemSemConstant =
3140 buildI32Constant(SPIRV::MemorySemantics::SequentiallyConsistent, I);
3141 Register MemSemReg = MemSemConstant.first;
3142 Result &= MemSemConstant.second;
3143 auto ScopeConstant = buildI32Constant(SPIRV::Scope::Workgroup, I);
3144 Register ScopeReg = ScopeConstant.first;
3145 Result &= ScopeConstant.second;
3146 MachineBasicBlock &BB = *I.getParent();
3147 return Result &&
3148 BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpControlBarrier))
3149 .addUse(ScopeReg)
3150 .addUse(ScopeReg)
3151 .addUse(MemSemReg)
3152 .constrainAllUses(TII, TRI, RBI);
3153 }
3154 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
3155 Register PtrReg = I.getOperand(I.getNumExplicitDefs() + 1).getReg();
3156 SPIRV::StorageClass::StorageClass ResSC =
3157 GR.getPointerStorageClass(ResType);
3158 if (!isGenericCastablePtr(ResSC))
3159 report_fatal_error("The target storage class is not castable from the "
3160 "Generic storage class");
3161 return BuildMI(BB, I, I.getDebugLoc(),
3162 TII.get(SPIRV::OpGenericCastToPtrExplicit))
3163 .addDef(ResVReg)
3164 .addUse(GR.getSPIRVTypeID(ResType))
3165 .addUse(PtrReg)
3166 .addImm(ResSC)
3167 .constrainAllUses(TII, TRI, RBI);
3168 }
3169 case Intrinsic::spv_lifetime_start:
3170 case Intrinsic::spv_lifetime_end: {
3171 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
3172 : SPIRV::OpLifetimeStop;
3173 int64_t Size = I.getOperand(I.getNumExplicitDefs() + 1).getImm();
3174 Register PtrReg = I.getOperand(I.getNumExplicitDefs() + 2).getReg();
3175 if (Size == -1)
3176 Size = 0;
3177 return BuildMI(BB, I, I.getDebugLoc(), TII.get(Op))
3178 .addUse(PtrReg)
3179 .addImm(Size)
3180 .constrainAllUses(TII, TRI, RBI);
3181 }
3182 case Intrinsic::spv_saturate:
3183 return selectSaturate(ResVReg, ResType, I);
3184 case Intrinsic::spv_nclamp:
3185 return selectExtInst(ResVReg, ResType, I, CL::fclamp, GL::NClamp);
3186 case Intrinsic::spv_uclamp:
3187 return selectExtInst(ResVReg, ResType, I, CL::u_clamp, GL::UClamp);
3188 case Intrinsic::spv_sclamp:
3189 return selectExtInst(ResVReg, ResType, I, CL::s_clamp, GL::SClamp);
3190 case Intrinsic::spv_wave_active_countbits:
3191 return selectWaveActiveCountBits(ResVReg, ResType, I);
3192 case Intrinsic::spv_wave_all:
3193 return selectWaveOpInst(ResVReg, ResType, I, SPIRV::OpGroupNonUniformAll);
3194 case Intrinsic::spv_wave_any:
3195 return selectWaveOpInst(ResVReg, ResType, I, SPIRV::OpGroupNonUniformAny);
3196 case Intrinsic::spv_wave_is_first_lane:
3197 return selectWaveOpInst(ResVReg, ResType, I, SPIRV::OpGroupNonUniformElect);
3198 case Intrinsic::spv_wave_reduce_umax:
3199 return selectWaveReduceMax(ResVReg, ResType, I, /*IsUnsigned*/ true);
3200 case Intrinsic::spv_wave_reduce_max:
3201 return selectWaveReduceMax(ResVReg, ResType, I, /*IsUnsigned*/ false);
3202 case Intrinsic::spv_wave_reduce_sum:
3203 return selectWaveReduceSum(ResVReg, ResType, I);
3204 case Intrinsic::spv_wave_readlane:
3205 return selectWaveOpInst(ResVReg, ResType, I,
3206 SPIRV::OpGroupNonUniformShuffle);
3207 case Intrinsic::spv_step:
3208 return selectExtInst(ResVReg, ResType, I, CL::step, GL::Step);
3209 case Intrinsic::spv_radians:
3210 return selectExtInst(ResVReg, ResType, I, CL::radians, GL::Radians);
3211 // Discard intrinsics which we do not expect to actually represent code after
3212 // lowering or intrinsics which are not implemented but should not crash when
3213 // found in a customer's LLVM IR input.
3214 case Intrinsic::instrprof_increment:
3215 case Intrinsic::instrprof_increment_step:
3216 case Intrinsic::instrprof_value_profile:
3217 break;
3218 // Discard internal intrinsics.
3219 case Intrinsic::spv_value_md:
3220 break;
3221 case Intrinsic::spv_resource_handlefrombinding: {
3222 return selectHandleFromBinding(ResVReg, ResType, I);
3223 }
3224 case Intrinsic::spv_resource_store_typedbuffer: {
3225 return selectImageWriteIntrinsic(I);
3226 }
3227 case Intrinsic::spv_resource_load_typedbuffer: {
3228 return selectReadImageIntrinsic(ResVReg, ResType, I);
3229 }
3230 case Intrinsic::spv_resource_getpointer: {
3231 return selectResourceGetPointer(ResVReg, ResType, I);
3232 }
3233 case Intrinsic::spv_discard: {
3234 return selectDiscard(ResVReg, ResType, I);
3235 }
3236 default: {
3237 std::string DiagMsg;
3238 raw_string_ostream OS(DiagMsg);
3239 I.print(OS);
3240 DiagMsg = "Intrinsic selection not implemented: " + DiagMsg;
3241 report_fatal_error(DiagMsg.c_str(), false);
3242 }
3243 }
3244 return true;
3245 }
3246
selectHandleFromBinding(Register & ResVReg,const SPIRVType * ResType,MachineInstr & I) const3247 bool SPIRVInstructionSelector::selectHandleFromBinding(Register &ResVReg,
3248 const SPIRVType *ResType,
3249 MachineInstr &I) const {
3250 // The images need to be loaded in the same basic block as their use. We defer
3251 // loading the image to the intrinsic that uses it.
3252 if (ResType->getOpcode() == SPIRV::OpTypeImage)
3253 return true;
3254
3255 return loadHandleBeforePosition(ResVReg, GR.getSPIRVTypeForVReg(ResVReg),
3256 *cast<GIntrinsic>(&I), I);
3257 }
3258
selectReadImageIntrinsic(Register & ResVReg,const SPIRVType * ResType,MachineInstr & I) const3259 bool SPIRVInstructionSelector::selectReadImageIntrinsic(
3260 Register &ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
3261
3262 // If the load of the image is in a different basic block, then
3263 // this will generate invalid code. A proper solution is to move
3264 // the OpLoad from selectHandleFromBinding here. However, to do
3265 // that we will need to change the return type of the intrinsic.
3266 // We will do that when we can, but for now trying to move forward with other
3267 // issues.
3268 Register ImageReg = I.getOperand(2).getReg();
3269 auto *ImageDef = cast<GIntrinsic>(getVRegDef(*MRI, ImageReg));
3270 Register NewImageReg = MRI->createVirtualRegister(MRI->getRegClass(ImageReg));
3271 if (!loadHandleBeforePosition(NewImageReg, GR.getSPIRVTypeForVReg(ImageReg),
3272 *ImageDef, I)) {
3273 return false;
3274 }
3275
3276 Register IdxReg = I.getOperand(3).getReg();
3277 DebugLoc Loc = I.getDebugLoc();
3278 MachineInstr &Pos = I;
3279
3280 return generateImageRead(ResVReg, ResType, NewImageReg, IdxReg, Loc, Pos);
3281 }
3282
generateImageRead(Register & ResVReg,const SPIRVType * ResType,Register ImageReg,Register IdxReg,DebugLoc Loc,MachineInstr & Pos) const3283 bool SPIRVInstructionSelector::generateImageRead(Register &ResVReg,
3284 const SPIRVType *ResType,
3285 Register ImageReg,
3286 Register IdxReg, DebugLoc Loc,
3287 MachineInstr &Pos) const {
3288 SPIRVType *ImageType = GR.getSPIRVTypeForVReg(ImageReg);
3289 assert(ImageType && ImageType->getOpcode() == SPIRV::OpTypeImage &&
3290 "ImageReg is not an image type.");
3291 bool IsSignedInteger =
3292 sampledTypeIsSignedInteger(GR.getTypeForSPIRVType(ImageType));
3293
3294 uint64_t ResultSize = GR.getScalarOrVectorComponentCount(ResType);
3295 if (ResultSize == 4) {
3296 auto BMI = BuildMI(*Pos.getParent(), Pos, Loc, TII.get(SPIRV::OpImageRead))
3297 .addDef(ResVReg)
3298 .addUse(GR.getSPIRVTypeID(ResType))
3299 .addUse(ImageReg)
3300 .addUse(IdxReg);
3301
3302 if (IsSignedInteger)
3303 BMI.addImm(0x1000); // SignExtend
3304 return BMI.constrainAllUses(TII, TRI, RBI);
3305 }
3306
3307 SPIRVType *ReadType = widenTypeToVec4(ResType, Pos);
3308 Register ReadReg = MRI->createVirtualRegister(GR.getRegClass(ReadType));
3309 auto BMI = BuildMI(*Pos.getParent(), Pos, Loc, TII.get(SPIRV::OpImageRead))
3310 .addDef(ReadReg)
3311 .addUse(GR.getSPIRVTypeID(ReadType))
3312 .addUse(ImageReg)
3313 .addUse(IdxReg);
3314 if (IsSignedInteger)
3315 BMI.addImm(0x1000); // SignExtend
3316 bool Succeed = BMI.constrainAllUses(TII, TRI, RBI);
3317 if (!Succeed)
3318 return false;
3319
3320 if (ResultSize == 1) {
3321 return BuildMI(*Pos.getParent(), Pos, Loc,
3322 TII.get(SPIRV::OpCompositeExtract))
3323 .addDef(ResVReg)
3324 .addUse(GR.getSPIRVTypeID(ResType))
3325 .addUse(ReadReg)
3326 .addImm(0)
3327 .constrainAllUses(TII, TRI, RBI);
3328 }
3329 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
3330 }
3331
selectResourceGetPointer(Register & ResVReg,const SPIRVType * ResType,MachineInstr & I) const3332 bool SPIRVInstructionSelector::selectResourceGetPointer(
3333 Register &ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
3334 Register ResourcePtr = I.getOperand(2).getReg();
3335 SPIRVType *RegType = GR.getSPIRVTypeForVReg(ResourcePtr, I.getMF());
3336 if (RegType->getOpcode() == SPIRV::OpTypeImage) {
3337 // For texel buffers, the index into the image is part of the OpImageRead or
3338 // OpImageWrite instructions. So we will do nothing in this case. This
3339 // intrinsic will be combined with the load or store when selecting the load
3340 // or store.
3341 return true;
3342 }
3343
3344 assert(ResType->getOpcode() == SPIRV::OpTypePointer);
3345 MachineIRBuilder MIRBuilder(I);
3346
3347 Register IndexReg = I.getOperand(3).getReg();
3348 Register ZeroReg =
3349 buildZerosVal(GR.getOrCreateSPIRVIntegerType(32, I, TII), I);
3350 return BuildMI(*I.getParent(), I, I.getDebugLoc(),
3351 TII.get(SPIRV::OpAccessChain))
3352 .addDef(ResVReg)
3353 .addUse(GR.getSPIRVTypeID(ResType))
3354 .addUse(ResourcePtr)
3355 .addUse(ZeroReg)
3356 .addUse(IndexReg)
3357 .constrainAllUses(TII, TRI, RBI);
3358 }
3359
extractSubvector(Register & ResVReg,const SPIRVType * ResType,Register & ReadReg,MachineInstr & InsertionPoint) const3360 bool SPIRVInstructionSelector::extractSubvector(
3361 Register &ResVReg, const SPIRVType *ResType, Register &ReadReg,
3362 MachineInstr &InsertionPoint) const {
3363 SPIRVType *InputType = GR.getResultType(ReadReg);
3364 [[maybe_unused]] uint64_t InputSize =
3365 GR.getScalarOrVectorComponentCount(InputType);
3366 uint64_t ResultSize = GR.getScalarOrVectorComponentCount(ResType);
3367 assert(InputSize > 1 && "The input must be a vector.");
3368 assert(ResultSize > 1 && "The result must be a vector.");
3369 assert(ResultSize < InputSize &&
3370 "Cannot extract more element than there are in the input.");
3371 SmallVector<Register> ComponentRegisters;
3372 SPIRVType *ScalarType = GR.getScalarOrVectorComponentType(ResType);
3373 const TargetRegisterClass *ScalarRegClass = GR.getRegClass(ScalarType);
3374 for (uint64_t I = 0; I < ResultSize; I++) {
3375 Register ComponentReg = MRI->createVirtualRegister(ScalarRegClass);
3376 bool Succeed = BuildMI(*InsertionPoint.getParent(), InsertionPoint,
3377 InsertionPoint.getDebugLoc(),
3378 TII.get(SPIRV::OpCompositeExtract))
3379 .addDef(ComponentReg)
3380 .addUse(ScalarType->getOperand(0).getReg())
3381 .addUse(ReadReg)
3382 .addImm(I)
3383 .constrainAllUses(TII, TRI, RBI);
3384 if (!Succeed)
3385 return false;
3386 ComponentRegisters.emplace_back(ComponentReg);
3387 }
3388
3389 MachineInstrBuilder MIB = BuildMI(*InsertionPoint.getParent(), InsertionPoint,
3390 InsertionPoint.getDebugLoc(),
3391 TII.get(SPIRV::OpCompositeConstruct))
3392 .addDef(ResVReg)
3393 .addUse(GR.getSPIRVTypeID(ResType));
3394
3395 for (Register ComponentReg : ComponentRegisters)
3396 MIB.addUse(ComponentReg);
3397 return MIB.constrainAllUses(TII, TRI, RBI);
3398 }
3399
selectImageWriteIntrinsic(MachineInstr & I) const3400 bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
3401 MachineInstr &I) const {
3402 // If the load of the image is in a different basic block, then
3403 // this will generate invalid code. A proper solution is to move
3404 // the OpLoad from selectHandleFromBinding here. However, to do
3405 // that we will need to change the return type of the intrinsic.
3406 // We will do that when we can, but for now trying to move forward with other
3407 // issues.
3408 Register ImageReg = I.getOperand(1).getReg();
3409 auto *ImageDef = cast<GIntrinsic>(getVRegDef(*MRI, ImageReg));
3410 Register NewImageReg = MRI->createVirtualRegister(MRI->getRegClass(ImageReg));
3411 if (!loadHandleBeforePosition(NewImageReg, GR.getSPIRVTypeForVReg(ImageReg),
3412 *ImageDef, I)) {
3413 return false;
3414 }
3415
3416 Register CoordinateReg = I.getOperand(2).getReg();
3417 Register DataReg = I.getOperand(3).getReg();
3418 assert(GR.getResultType(DataReg)->getOpcode() == SPIRV::OpTypeVector);
3419 assert(GR.getScalarOrVectorComponentCount(GR.getResultType(DataReg)) == 4);
3420 return BuildMI(*I.getParent(), I, I.getDebugLoc(),
3421 TII.get(SPIRV::OpImageWrite))
3422 .addUse(NewImageReg)
3423 .addUse(CoordinateReg)
3424 .addUse(DataReg)
3425 .constrainAllUses(TII, TRI, RBI);
3426 }
3427
buildPointerToResource(const SPIRVType * SpirvResType,SPIRV::StorageClass::StorageClass SC,uint32_t Set,uint32_t Binding,uint32_t ArraySize,Register IndexReg,bool IsNonUniform,StringRef Name,MachineIRBuilder MIRBuilder) const3428 Register SPIRVInstructionSelector::buildPointerToResource(
3429 const SPIRVType *SpirvResType, SPIRV::StorageClass::StorageClass SC,
3430 uint32_t Set, uint32_t Binding, uint32_t ArraySize, Register IndexReg,
3431 bool IsNonUniform, StringRef Name, MachineIRBuilder MIRBuilder) const {
3432 const Type *ResType = GR.getTypeForSPIRVType(SpirvResType);
3433 if (ArraySize == 1) {
3434 SPIRVType *PtrType =
3435 GR.getOrCreateSPIRVPointerType(ResType, MIRBuilder, SC);
3436 assert(GR.getPointeeType(PtrType) == SpirvResType &&
3437 "SpirvResType did not have an explicit layout.");
3438 return GR.getOrCreateGlobalVariableWithBinding(PtrType, Set, Binding, Name,
3439 MIRBuilder);
3440 }
3441
3442 const Type *VarType = ArrayType::get(const_cast<Type *>(ResType), ArraySize);
3443 SPIRVType *VarPointerType =
3444 GR.getOrCreateSPIRVPointerType(VarType, MIRBuilder, SC);
3445 Register VarReg = GR.getOrCreateGlobalVariableWithBinding(
3446 VarPointerType, Set, Binding, Name, MIRBuilder);
3447
3448 SPIRVType *ResPointerType =
3449 GR.getOrCreateSPIRVPointerType(ResType, MIRBuilder, SC);
3450
3451 Register AcReg = MRI->createVirtualRegister(GR.getRegClass(ResPointerType));
3452 if (IsNonUniform) {
3453 // It is unclear which value needs to be marked an non-uniform, so both
3454 // the index and the access changed are decorated as non-uniform.
3455 buildOpDecorate(IndexReg, MIRBuilder, SPIRV::Decoration::NonUniformEXT, {});
3456 buildOpDecorate(AcReg, MIRBuilder, SPIRV::Decoration::NonUniformEXT, {});
3457 }
3458
3459 MIRBuilder.buildInstr(SPIRV::OpAccessChain)
3460 .addDef(AcReg)
3461 .addUse(GR.getSPIRVTypeID(ResPointerType))
3462 .addUse(VarReg)
3463 .addUse(IndexReg);
3464
3465 return AcReg;
3466 }
3467
selectFirstBitSet16(Register ResVReg,const SPIRVType * ResType,MachineInstr & I,unsigned ExtendOpcode,unsigned BitSetOpcode) const3468 bool SPIRVInstructionSelector::selectFirstBitSet16(
3469 Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
3470 unsigned ExtendOpcode, unsigned BitSetOpcode) const {
3471 Register ExtReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
3472 bool Result = selectOpWithSrcs(ExtReg, ResType, I, {I.getOperand(2).getReg()},
3473 ExtendOpcode);
3474
3475 return Result &&
3476 selectFirstBitSet32(ResVReg, ResType, I, ExtReg, BitSetOpcode);
3477 }
3478
selectFirstBitSet32(Register ResVReg,const SPIRVType * ResType,MachineInstr & I,Register SrcReg,unsigned BitSetOpcode) const3479 bool SPIRVInstructionSelector::selectFirstBitSet32(
3480 Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
3481 Register SrcReg, unsigned BitSetOpcode) const {
3482 return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
3483 .addDef(ResVReg)
3484 .addUse(GR.getSPIRVTypeID(ResType))
3485 .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::GLSL_std_450))
3486 .addImm(BitSetOpcode)
3487 .addUse(SrcReg)
3488 .constrainAllUses(TII, TRI, RBI);
3489 }
3490
selectFirstBitSet64Overflow(Register ResVReg,const SPIRVType * ResType,MachineInstr & I,Register SrcReg,unsigned BitSetOpcode,bool SwapPrimarySide) const3491 bool SPIRVInstructionSelector::selectFirstBitSet64Overflow(
3492 Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
3493 Register SrcReg, unsigned BitSetOpcode, bool SwapPrimarySide) const {
3494
3495 // SPIR-V allow vectors of size 2,3,4 only. Calling with a larger vectors
3496 // requires creating a param register and return register with an invalid
3497 // vector size. If that is resolved, then this function can be used for
3498 // vectors of any component size.
3499 unsigned ComponentCount = GR.getScalarOrVectorComponentCount(ResType);
3500 assert(ComponentCount < 5 && "Vec 5+ will generate invalid SPIR-V ops");
3501
3502 MachineIRBuilder MIRBuilder(I);
3503 SPIRVType *BaseType = GR.retrieveScalarOrVectorIntType(ResType);
3504 SPIRVType *I64Type = GR.getOrCreateSPIRVIntegerType(64, MIRBuilder);
3505 SPIRVType *I64x2Type =
3506 GR.getOrCreateSPIRVVectorType(I64Type, 2, MIRBuilder, false);
3507 SPIRVType *Vec2ResType =
3508 GR.getOrCreateSPIRVVectorType(BaseType, 2, MIRBuilder, false);
3509
3510 std::vector<Register> PartialRegs;
3511
3512 // Loops 0, 2, 4, ... but stops one loop early when ComponentCount is odd
3513 unsigned CurrentComponent = 0;
3514 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
3515 // This register holds the firstbitX result for each of the i64x2 vectors
3516 // extracted from SrcReg
3517 Register BitSetResult =
3518 MRI->createVirtualRegister(GR.getRegClass(I64x2Type));
3519
3520 auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
3521 TII.get(SPIRV::OpVectorShuffle))
3522 .addDef(BitSetResult)
3523 .addUse(GR.getSPIRVTypeID(I64x2Type))
3524 .addUse(SrcReg)
3525 .addUse(SrcReg)
3526 .addImm(CurrentComponent)
3527 .addImm(CurrentComponent + 1);
3528
3529 if (!MIB.constrainAllUses(TII, TRI, RBI))
3530 return false;
3531
3532 Register SubVecBitSetReg =
3533 MRI->createVirtualRegister(GR.getRegClass(Vec2ResType));
3534
3535 if (!selectFirstBitSet64(SubVecBitSetReg, Vec2ResType, I, BitSetResult,
3536 BitSetOpcode, SwapPrimarySide))
3537 return false;
3538
3539 PartialRegs.push_back(SubVecBitSetReg);
3540 }
3541
3542 // On odd component counts we need to handle one more component
3543 if (CurrentComponent != ComponentCount) {
3544 bool ZeroAsNull = !STI.isShader();
3545 Register FinalElemReg = MRI->createVirtualRegister(GR.getRegClass(I64Type));
3546 Register ConstIntLastIdx = GR.getOrCreateConstInt(
3547 ComponentCount - 1, I, BaseType, TII, ZeroAsNull);
3548
3549 if (!selectOpWithSrcs(FinalElemReg, I64Type, I, {SrcReg, ConstIntLastIdx},
3550 SPIRV::OpVectorExtractDynamic))
3551 return false;
3552
3553 Register FinalElemBitSetReg =
3554 MRI->createVirtualRegister(GR.getRegClass(BaseType));
3555
3556 if (!selectFirstBitSet64(FinalElemBitSetReg, BaseType, I, FinalElemReg,
3557 BitSetOpcode, SwapPrimarySide))
3558 return false;
3559
3560 PartialRegs.push_back(FinalElemBitSetReg);
3561 }
3562
3563 // Join all the resulting registers back into the return type in order
3564 // (ie i32x2, i32x2, i32x1 -> i32x5)
3565 return selectOpWithSrcs(ResVReg, ResType, I, PartialRegs,
3566 SPIRV::OpCompositeConstruct);
3567 }
3568
selectFirstBitSet64(Register ResVReg,const SPIRVType * ResType,MachineInstr & I,Register SrcReg,unsigned BitSetOpcode,bool SwapPrimarySide) const3569 bool SPIRVInstructionSelector::selectFirstBitSet64(
3570 Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
3571 Register SrcReg, unsigned BitSetOpcode, bool SwapPrimarySide) const {
3572 unsigned ComponentCount = GR.getScalarOrVectorComponentCount(ResType);
3573 SPIRVType *BaseType = GR.retrieveScalarOrVectorIntType(ResType);
3574 bool ZeroAsNull = !STI.isShader();
3575 Register ConstIntZero =
3576 GR.getOrCreateConstInt(0, I, BaseType, TII, ZeroAsNull);
3577 Register ConstIntOne =
3578 GR.getOrCreateConstInt(1, I, BaseType, TII, ZeroAsNull);
3579
3580 // SPIRV doesn't support vectors with more than 4 components. Since the
3581 // algoritm below converts i64 -> i32x2 and i64x4 -> i32x8 it can only
3582 // operate on vectors with 2 or less components. When largers vectors are
3583 // seen. Split them, recurse, then recombine them.
3584 if (ComponentCount > 2) {
3585 return selectFirstBitSet64Overflow(ResVReg, ResType, I, SrcReg,
3586 BitSetOpcode, SwapPrimarySide);
3587 }
3588
3589 // 1. Split int64 into 2 pieces using a bitcast
3590 MachineIRBuilder MIRBuilder(I);
3591 SPIRVType *PostCastType = GR.getOrCreateSPIRVVectorType(
3592 BaseType, 2 * ComponentCount, MIRBuilder, false);
3593 Register BitcastReg =
3594 MRI->createVirtualRegister(GR.getRegClass(PostCastType));
3595
3596 if (!selectOpWithSrcs(BitcastReg, PostCastType, I, {SrcReg},
3597 SPIRV::OpBitcast))
3598 return false;
3599
3600 // 2. Find the first set bit from the primary side for all the pieces in #1
3601 Register FBSReg = MRI->createVirtualRegister(GR.getRegClass(PostCastType));
3602 if (!selectFirstBitSet32(FBSReg, PostCastType, I, BitcastReg, BitSetOpcode))
3603 return false;
3604
3605 // 3. Split result vector into high bits and low bits
3606 Register HighReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
3607 Register LowReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
3608
3609 bool IsScalarRes = ResType->getOpcode() != SPIRV::OpTypeVector;
3610 if (IsScalarRes) {
3611 // if scalar do a vector extract
3612 if (!selectOpWithSrcs(HighReg, ResType, I, {FBSReg, ConstIntZero},
3613 SPIRV::OpVectorExtractDynamic))
3614 return false;
3615 if (!selectOpWithSrcs(LowReg, ResType, I, {FBSReg, ConstIntOne},
3616 SPIRV::OpVectorExtractDynamic))
3617 return false;
3618 } else {
3619 // if vector do a shufflevector
3620 auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
3621 TII.get(SPIRV::OpVectorShuffle))
3622 .addDef(HighReg)
3623 .addUse(GR.getSPIRVTypeID(ResType))
3624 .addUse(FBSReg)
3625 // Per the spec, repeat the vector if only one vec is needed
3626 .addUse(FBSReg);
3627
3628 // high bits are stored in even indexes. Extract them from FBSReg
3629 for (unsigned J = 0; J < ComponentCount * 2; J += 2) {
3630 MIB.addImm(J);
3631 }
3632
3633 if (!MIB.constrainAllUses(TII, TRI, RBI))
3634 return false;
3635
3636 MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
3637 TII.get(SPIRV::OpVectorShuffle))
3638 .addDef(LowReg)
3639 .addUse(GR.getSPIRVTypeID(ResType))
3640 .addUse(FBSReg)
3641 // Per the spec, repeat the vector if only one vec is needed
3642 .addUse(FBSReg);
3643
3644 // low bits are stored in odd indexes. Extract them from FBSReg
3645 for (unsigned J = 1; J < ComponentCount * 2; J += 2) {
3646 MIB.addImm(J);
3647 }
3648 if (!MIB.constrainAllUses(TII, TRI, RBI))
3649 return false;
3650 }
3651
3652 // 4. Check the result. When primary bits == -1 use secondary, otherwise use
3653 // primary
3654 SPIRVType *BoolType = GR.getOrCreateSPIRVBoolType(I, TII);
3655 Register NegOneReg;
3656 Register Reg0;
3657 Register Reg32;
3658 unsigned SelectOp;
3659 unsigned AddOp;
3660
3661 if (IsScalarRes) {
3662 NegOneReg =
3663 GR.getOrCreateConstInt((unsigned)-1, I, ResType, TII, ZeroAsNull);
3664 Reg0 = GR.getOrCreateConstInt(0, I, ResType, TII, ZeroAsNull);
3665 Reg32 = GR.getOrCreateConstInt(32, I, ResType, TII, ZeroAsNull);
3666 SelectOp = SPIRV::OpSelectSISCond;
3667 AddOp = SPIRV::OpIAddS;
3668 } else {
3669 BoolType = GR.getOrCreateSPIRVVectorType(BoolType, ComponentCount,
3670 MIRBuilder, false);
3671 NegOneReg =
3672 GR.getOrCreateConstVector((unsigned)-1, I, ResType, TII, ZeroAsNull);
3673 Reg0 = GR.getOrCreateConstVector(0, I, ResType, TII, ZeroAsNull);
3674 Reg32 = GR.getOrCreateConstVector(32, I, ResType, TII, ZeroAsNull);
3675 SelectOp = SPIRV::OpSelectVIVCond;
3676 AddOp = SPIRV::OpIAddV;
3677 }
3678
3679 Register PrimaryReg = HighReg;
3680 Register SecondaryReg = LowReg;
3681 Register PrimaryShiftReg = Reg32;
3682 Register SecondaryShiftReg = Reg0;
3683
3684 // By default the emitted opcodes check for the set bit from the MSB side.
3685 // Setting SwapPrimarySide checks the set bit from the LSB side
3686 if (SwapPrimarySide) {
3687 PrimaryReg = LowReg;
3688 SecondaryReg = HighReg;
3689 PrimaryShiftReg = Reg0;
3690 SecondaryShiftReg = Reg32;
3691 }
3692
3693 // Check if the primary bits are == -1
3694 Register BReg = MRI->createVirtualRegister(GR.getRegClass(BoolType));
3695 if (!selectOpWithSrcs(BReg, BoolType, I, {PrimaryReg, NegOneReg},
3696 SPIRV::OpIEqual))
3697 return false;
3698
3699 // Select secondary bits if true in BReg, otherwise primary bits
3700 Register TmpReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
3701 if (!selectOpWithSrcs(TmpReg, ResType, I, {BReg, SecondaryReg, PrimaryReg},
3702 SelectOp))
3703 return false;
3704
3705 // 5. Add 32 when high bits are used, otherwise 0 for low bits
3706 Register ValReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
3707 if (!selectOpWithSrcs(ValReg, ResType, I,
3708 {BReg, SecondaryShiftReg, PrimaryShiftReg}, SelectOp))
3709 return false;
3710
3711 return selectOpWithSrcs(ResVReg, ResType, I, {ValReg, TmpReg}, AddOp);
3712 }
3713
selectFirstBitHigh(Register ResVReg,const SPIRVType * ResType,MachineInstr & I,bool IsSigned) const3714 bool SPIRVInstructionSelector::selectFirstBitHigh(Register ResVReg,
3715 const SPIRVType *ResType,
3716 MachineInstr &I,
3717 bool IsSigned) const {
3718 // FindUMsb and FindSMsb intrinsics only support 32 bit integers
3719 Register OpReg = I.getOperand(2).getReg();
3720 SPIRVType *OpType = GR.getSPIRVTypeForVReg(OpReg);
3721 // zero or sign extend
3722 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
3723 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
3724
3725 switch (GR.getScalarOrVectorBitWidth(OpType)) {
3726 case 16:
3727 return selectFirstBitSet16(ResVReg, ResType, I, ExtendOpcode, BitSetOpcode);
3728 case 32:
3729 return selectFirstBitSet32(ResVReg, ResType, I, OpReg, BitSetOpcode);
3730 case 64:
3731 return selectFirstBitSet64(ResVReg, ResType, I, OpReg, BitSetOpcode,
3732 /*SwapPrimarySide=*/false);
3733 default:
3734 report_fatal_error(
3735 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
3736 }
3737 }
3738
selectFirstBitLow(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const3739 bool SPIRVInstructionSelector::selectFirstBitLow(Register ResVReg,
3740 const SPIRVType *ResType,
3741 MachineInstr &I) const {
3742 // FindILsb intrinsic only supports 32 bit integers
3743 Register OpReg = I.getOperand(2).getReg();
3744 SPIRVType *OpType = GR.getSPIRVTypeForVReg(OpReg);
3745 // OpUConvert treats the operand bits as an unsigned i16 and zero extends it
3746 // to an unsigned i32. As this leaves all the least significant bits unchanged
3747 // so the first set bit from the LSB side doesn't change.
3748 unsigned ExtendOpcode = SPIRV::OpUConvert;
3749 unsigned BitSetOpcode = GL::FindILsb;
3750
3751 switch (GR.getScalarOrVectorBitWidth(OpType)) {
3752 case 16:
3753 return selectFirstBitSet16(ResVReg, ResType, I, ExtendOpcode, BitSetOpcode);
3754 case 32:
3755 return selectFirstBitSet32(ResVReg, ResType, I, OpReg, BitSetOpcode);
3756 case 64:
3757 return selectFirstBitSet64(ResVReg, ResType, I, OpReg, BitSetOpcode,
3758 /*SwapPrimarySide=*/true);
3759 default:
3760 report_fatal_error("spv_firstbitlow only supports 16,32,64 bits.");
3761 }
3762 }
3763
selectAllocaArray(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const3764 bool SPIRVInstructionSelector::selectAllocaArray(Register ResVReg,
3765 const SPIRVType *ResType,
3766 MachineInstr &I) const {
3767 // there was an allocation size parameter to the allocation instruction
3768 // that is not 1
3769 MachineBasicBlock &BB = *I.getParent();
3770 bool Res = BuildMI(BB, I, I.getDebugLoc(),
3771 TII.get(SPIRV::OpVariableLengthArrayINTEL))
3772 .addDef(ResVReg)
3773 .addUse(GR.getSPIRVTypeID(ResType))
3774 .addUse(I.getOperand(2).getReg())
3775 .constrainAllUses(TII, TRI, RBI);
3776 if (!STI.isShader()) {
3777 unsigned Alignment = I.getOperand(3).getImm();
3778 buildOpDecorate(ResVReg, I, TII, SPIRV::Decoration::Alignment, {Alignment});
3779 }
3780 return Res;
3781 }
3782
selectFrameIndex(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const3783 bool SPIRVInstructionSelector::selectFrameIndex(Register ResVReg,
3784 const SPIRVType *ResType,
3785 MachineInstr &I) const {
3786 // Change order of instructions if needed: all OpVariable instructions in a
3787 // function must be the first instructions in the first block
3788 auto It = getOpVariableMBBIt(I);
3789 bool Res = BuildMI(*It->getParent(), It, It->getDebugLoc(),
3790 TII.get(SPIRV::OpVariable))
3791 .addDef(ResVReg)
3792 .addUse(GR.getSPIRVTypeID(ResType))
3793 .addImm(static_cast<uint32_t>(SPIRV::StorageClass::Function))
3794 .constrainAllUses(TII, TRI, RBI);
3795 if (!STI.isShader()) {
3796 unsigned Alignment = I.getOperand(2).getImm();
3797 buildOpDecorate(ResVReg, *It, TII, SPIRV::Decoration::Alignment,
3798 {Alignment});
3799 }
3800 return Res;
3801 }
3802
selectBranch(MachineInstr & I) const3803 bool SPIRVInstructionSelector::selectBranch(MachineInstr &I) const {
3804 // InstructionSelector walks backwards through the instructions. We can use
3805 // both a G_BR and a G_BRCOND to create an OpBranchConditional. We hit G_BR
3806 // first, so can generate an OpBranchConditional here. If there is no
3807 // G_BRCOND, we just use OpBranch for a regular unconditional branch.
3808 const MachineInstr *PrevI = I.getPrevNode();
3809 MachineBasicBlock &MBB = *I.getParent();
3810 if (PrevI != nullptr && PrevI->getOpcode() == TargetOpcode::G_BRCOND) {
3811 return BuildMI(MBB, I, I.getDebugLoc(), TII.get(SPIRV::OpBranchConditional))
3812 .addUse(PrevI->getOperand(0).getReg())
3813 .addMBB(PrevI->getOperand(1).getMBB())
3814 .addMBB(I.getOperand(0).getMBB())
3815 .constrainAllUses(TII, TRI, RBI);
3816 }
3817 return BuildMI(MBB, I, I.getDebugLoc(), TII.get(SPIRV::OpBranch))
3818 .addMBB(I.getOperand(0).getMBB())
3819 .constrainAllUses(TII, TRI, RBI);
3820 }
3821
selectBranchCond(MachineInstr & I) const3822 bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &I) const {
3823 // InstructionSelector walks backwards through the instructions. For an
3824 // explicit conditional branch with no fallthrough, we use both a G_BR and a
3825 // G_BRCOND to create an OpBranchConditional. We should hit G_BR first, and
3826 // generate the OpBranchConditional in selectBranch above.
3827 //
3828 // If an OpBranchConditional has been generated, we simply return, as the work
3829 // is alread done. If there is no OpBranchConditional, LLVM must be relying on
3830 // implicit fallthrough to the next basic block, so we need to create an
3831 // OpBranchConditional with an explicit "false" argument pointing to the next
3832 // basic block that LLVM would fall through to.
3833 const MachineInstr *NextI = I.getNextNode();
3834 // Check if this has already been successfully selected.
3835 if (NextI != nullptr && NextI->getOpcode() == SPIRV::OpBranchConditional)
3836 return true;
3837 // Must be relying on implicit block fallthrough, so generate an
3838 // OpBranchConditional with the "next" basic block as the "false" target.
3839 MachineBasicBlock &MBB = *I.getParent();
3840 unsigned NextMBBNum = MBB.getNextNode()->getNumber();
3841 MachineBasicBlock *NextMBB = I.getMF()->getBlockNumbered(NextMBBNum);
3842 return BuildMI(MBB, I, I.getDebugLoc(), TII.get(SPIRV::OpBranchConditional))
3843 .addUse(I.getOperand(0).getReg())
3844 .addMBB(I.getOperand(1).getMBB())
3845 .addMBB(NextMBB)
3846 .constrainAllUses(TII, TRI, RBI);
3847 }
3848
selectPhi(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const3849 bool SPIRVInstructionSelector::selectPhi(Register ResVReg,
3850 const SPIRVType *ResType,
3851 MachineInstr &I) const {
3852 auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpPhi))
3853 .addDef(ResVReg)
3854 .addUse(GR.getSPIRVTypeID(ResType));
3855 const unsigned NumOps = I.getNumOperands();
3856 for (unsigned i = 1; i < NumOps; i += 2) {
3857 MIB.addUse(I.getOperand(i + 0).getReg());
3858 MIB.addMBB(I.getOperand(i + 1).getMBB());
3859 }
3860 bool Res = MIB.constrainAllUses(TII, TRI, RBI);
3861 MIB->setDesc(TII.get(TargetOpcode::PHI));
3862 MIB->removeOperand(1);
3863 return Res;
3864 }
3865
selectGlobalValue(Register ResVReg,MachineInstr & I,const MachineInstr * Init) const3866 bool SPIRVInstructionSelector::selectGlobalValue(
3867 Register ResVReg, MachineInstr &I, const MachineInstr *Init) const {
3868 // FIXME: don't use MachineIRBuilder here, replace it with BuildMI.
3869 MachineIRBuilder MIRBuilder(I);
3870 const GlobalValue *GV = I.getOperand(1).getGlobal();
3871 Type *GVType = toTypedPointer(GR.getDeducedGlobalValueType(GV));
3872
3873 std::string GlobalIdent;
3874 if (!GV->hasName()) {
3875 unsigned &ID = UnnamedGlobalIDs[GV];
3876 if (ID == 0)
3877 ID = UnnamedGlobalIDs.size();
3878 GlobalIdent = "__unnamed_" + Twine(ID).str();
3879 } else {
3880 GlobalIdent = GV->getName();
3881 }
3882
3883 // Behaviour of functions as operands depends on availability of the
3884 // corresponding extension (SPV_INTEL_function_pointers):
3885 // - If there is an extension to operate with functions as operands:
3886 // We create a proper constant operand and evaluate a correct type for a
3887 // function pointer.
3888 // - Without the required extension:
3889 // We have functions as operands in tests with blocks of instruction e.g. in
3890 // transcoding/global_block.ll. These operands are not used and should be
3891 // substituted by zero constants. Their type is expected to be always
3892 // OpTypePointer Function %uchar.
3893 if (isa<Function>(GV)) {
3894 const Constant *ConstVal = GV;
3895 MachineBasicBlock &BB = *I.getParent();
3896 Register NewReg = GR.find(ConstVal, GR.CurMF);
3897 if (!NewReg.isValid()) {
3898 Register NewReg = ResVReg;
3899 const Function *GVFun =
3900 STI.canUseExtension(SPIRV::Extension::SPV_INTEL_function_pointers)
3901 ? dyn_cast<Function>(GV)
3902 : nullptr;
3903 SPIRVType *ResType = GR.getOrCreateSPIRVPointerType(
3904 GVType, I,
3905 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
3906 : addressSpaceToStorageClass(GV->getAddressSpace(), STI));
3907 if (GVFun) {
3908 // References to a function via function pointers generate virtual
3909 // registers without a definition. We will resolve it later, during
3910 // module analysis stage.
3911 Register ResTypeReg = GR.getSPIRVTypeID(ResType);
3912 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
3913 Register FuncVReg =
3914 MRI->createGenericVirtualRegister(GR.getRegType(ResType));
3915 MRI->setRegClass(FuncVReg, &SPIRV::pIDRegClass);
3916 MachineInstrBuilder MIB1 =
3917 BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpUndef))
3918 .addDef(FuncVReg)
3919 .addUse(ResTypeReg);
3920 MachineInstrBuilder MIB2 =
3921 BuildMI(BB, I, I.getDebugLoc(),
3922 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
3923 .addDef(NewReg)
3924 .addUse(ResTypeReg)
3925 .addUse(FuncVReg);
3926 GR.add(ConstVal, MIB2);
3927 // mapping the function pointer to the used Function
3928 GR.recordFunctionPointer(&MIB2.getInstr()->getOperand(2), GVFun);
3929 return MIB1.constrainAllUses(TII, TRI, RBI) &&
3930 MIB2.constrainAllUses(TII, TRI, RBI);
3931 }
3932 MachineInstrBuilder MIB3 =
3933 BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantNull))
3934 .addDef(NewReg)
3935 .addUse(GR.getSPIRVTypeID(ResType));
3936 GR.add(ConstVal, MIB3);
3937 return MIB3.constrainAllUses(TII, TRI, RBI);
3938 }
3939 assert(NewReg != ResVReg);
3940 return BuildCOPY(ResVReg, NewReg, I);
3941 }
3942 auto GlobalVar = cast<GlobalVariable>(GV);
3943 assert(GlobalVar->getName() != "llvm.global.annotations");
3944
3945 // Skip empty declaration for GVs with initializers till we get the decl with
3946 // passed initializer.
3947 if (hasInitializer(GlobalVar) && !Init)
3948 return true;
3949
3950 bool HasLnkTy = !GV->hasInternalLinkage() && !GV->hasPrivateLinkage() &&
3951 !GV->hasHiddenVisibility();
3952 SPIRV::LinkageType::LinkageType LnkType =
3953 GV->isDeclarationForLinker()
3954 ? SPIRV::LinkageType::Import
3955 : (GV->hasLinkOnceODRLinkage() &&
3956 STI.canUseExtension(SPIRV::Extension::SPV_KHR_linkonce_odr)
3957 ? SPIRV::LinkageType::LinkOnceODR
3958 : SPIRV::LinkageType::Export);
3959
3960 const unsigned AddrSpace = GV->getAddressSpace();
3961 SPIRV::StorageClass::StorageClass StorageClass =
3962 addressSpaceToStorageClass(AddrSpace, STI);
3963 SPIRVType *ResType = GR.getOrCreateSPIRVPointerType(GVType, I, StorageClass);
3964 Register Reg = GR.buildGlobalVariable(
3965 ResVReg, ResType, GlobalIdent, GV, StorageClass, Init,
3966 GlobalVar->isConstant(), HasLnkTy, LnkType, MIRBuilder, true);
3967 return Reg.isValid();
3968 }
3969
selectLog10(Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const3970 bool SPIRVInstructionSelector::selectLog10(Register ResVReg,
3971 const SPIRVType *ResType,
3972 MachineInstr &I) const {
3973 if (STI.canUseExtInstSet(SPIRV::InstructionSet::OpenCL_std)) {
3974 return selectExtInst(ResVReg, ResType, I, CL::log10);
3975 }
3976
3977 // There is no log10 instruction in the GLSL Extended Instruction set, so it
3978 // is implemented as:
3979 // log10(x) = log2(x) * (1 / log2(10))
3980 // = log2(x) * 0.30103
3981
3982 MachineIRBuilder MIRBuilder(I);
3983 MachineBasicBlock &BB = *I.getParent();
3984
3985 // Build log2(x).
3986 Register VarReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
3987 bool Result =
3988 BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
3989 .addDef(VarReg)
3990 .addUse(GR.getSPIRVTypeID(ResType))
3991 .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::GLSL_std_450))
3992 .addImm(GL::Log2)
3993 .add(I.getOperand(1))
3994 .constrainAllUses(TII, TRI, RBI);
3995
3996 // Build 0.30103.
3997 assert(ResType->getOpcode() == SPIRV::OpTypeVector ||
3998 ResType->getOpcode() == SPIRV::OpTypeFloat);
3999 // TODO: Add matrix implementation once supported by the HLSL frontend.
4000 const SPIRVType *SpirvScalarType =
4001 ResType->getOpcode() == SPIRV::OpTypeVector
4002 ? GR.getSPIRVTypeForVReg(ResType->getOperand(1).getReg())
4003 : ResType;
4004 Register ScaleReg =
4005 GR.buildConstantFP(APFloat(0.30103f), MIRBuilder, SpirvScalarType);
4006
4007 // Multiply log2(x) by 0.30103 to get log10(x) result.
4008 auto Opcode = ResType->getOpcode() == SPIRV::OpTypeVector
4009 ? SPIRV::OpVectorTimesScalar
4010 : SPIRV::OpFMulS;
4011 return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
4012 .addDef(ResVReg)
4013 .addUse(GR.getSPIRVTypeID(ResType))
4014 .addUse(VarReg)
4015 .addUse(ScaleReg)
4016 .constrainAllUses(TII, TRI, RBI);
4017 }
4018
4019 // Generate the instructions to load 3-element vector builtin input
4020 // IDs/Indices.
4021 // Like: GlobalInvocationId, LocalInvocationId, etc....
4022
loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const4023 bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
4024 SPIRV::BuiltIn::BuiltIn BuiltInValue, Register ResVReg,
4025 const SPIRVType *ResType, MachineInstr &I) const {
4026 MachineIRBuilder MIRBuilder(I);
4027 const SPIRVType *Vec3Ty =
4028 GR.getOrCreateSPIRVVectorType(ResType, 3, MIRBuilder, false);
4029 const SPIRVType *PtrType = GR.getOrCreateSPIRVPointerType(
4030 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
4031
4032 // Create new register for the input ID builtin variable.
4033 Register NewRegister =
4034 MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::iIDRegClass);
4035 MIRBuilder.getMRI()->setType(NewRegister, LLT::pointer(0, 64));
4036 GR.assignSPIRVTypeToVReg(PtrType, NewRegister, MIRBuilder.getMF());
4037
4038 // Build global variable with the necessary decorations for the input ID
4039 // builtin variable.
4040 Register Variable = GR.buildGlobalVariable(
4041 NewRegister, PtrType, getLinkStringForBuiltIn(BuiltInValue), nullptr,
4042 SPIRV::StorageClass::Input, nullptr, true, false,
4043 SPIRV::LinkageType::Import, MIRBuilder, false);
4044
4045 // Create new register for loading value.
4046 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
4047 Register LoadedRegister = MRI->createVirtualRegister(&SPIRV::iIDRegClass);
4048 MIRBuilder.getMRI()->setType(LoadedRegister, LLT::pointer(0, 64));
4049 GR.assignSPIRVTypeToVReg(Vec3Ty, LoadedRegister, MIRBuilder.getMF());
4050
4051 // Load v3uint value from the global variable.
4052 bool Result =
4053 BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpLoad))
4054 .addDef(LoadedRegister)
4055 .addUse(GR.getSPIRVTypeID(Vec3Ty))
4056 .addUse(Variable);
4057
4058 // Get the input ID index. Expecting operand is a constant immediate value,
4059 // wrapped in a type assignment.
4060 assert(I.getOperand(2).isReg());
4061 const uint32_t ThreadId = foldImm(I.getOperand(2), MRI);
4062
4063 // Extract the input ID from the loaded vector value.
4064 MachineBasicBlock &BB = *I.getParent();
4065 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
4066 .addDef(ResVReg)
4067 .addUse(GR.getSPIRVTypeID(ResType))
4068 .addUse(LoadedRegister)
4069 .addImm(ThreadId);
4070 return Result && MIB.constrainAllUses(TII, TRI, RBI);
4071 }
4072
4073 // Generate the instructions to load 32-bit integer builtin input IDs/Indices.
4074 // Like LocalInvocationIndex
loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,Register ResVReg,const SPIRVType * ResType,MachineInstr & I) const4075 bool SPIRVInstructionSelector::loadBuiltinInputID(
4076 SPIRV::BuiltIn::BuiltIn BuiltInValue, Register ResVReg,
4077 const SPIRVType *ResType, MachineInstr &I) const {
4078 MachineIRBuilder MIRBuilder(I);
4079 const SPIRVType *PtrType = GR.getOrCreateSPIRVPointerType(
4080 ResType, MIRBuilder, SPIRV::StorageClass::Input);
4081
4082 // Create new register for the input ID builtin variable.
4083 Register NewRegister =
4084 MIRBuilder.getMRI()->createVirtualRegister(GR.getRegClass(PtrType));
4085 MIRBuilder.getMRI()->setType(
4086 NewRegister,
4087 LLT::pointer(storageClassToAddressSpace(SPIRV::StorageClass::Input),
4088 GR.getPointerSize()));
4089 GR.assignSPIRVTypeToVReg(PtrType, NewRegister, MIRBuilder.getMF());
4090
4091 // Build global variable with the necessary decorations for the input ID
4092 // builtin variable.
4093 Register Variable = GR.buildGlobalVariable(
4094 NewRegister, PtrType, getLinkStringForBuiltIn(BuiltInValue), nullptr,
4095 SPIRV::StorageClass::Input, nullptr, true, false,
4096 SPIRV::LinkageType::Import, MIRBuilder, false);
4097
4098 // Load uint value from the global variable.
4099 auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpLoad))
4100 .addDef(ResVReg)
4101 .addUse(GR.getSPIRVTypeID(ResType))
4102 .addUse(Variable);
4103
4104 return MIB.constrainAllUses(TII, TRI, RBI);
4105 }
4106
widenTypeToVec4(const SPIRVType * Type,MachineInstr & I) const4107 SPIRVType *SPIRVInstructionSelector::widenTypeToVec4(const SPIRVType *Type,
4108 MachineInstr &I) const {
4109 MachineIRBuilder MIRBuilder(I);
4110 if (Type->getOpcode() != SPIRV::OpTypeVector)
4111 return GR.getOrCreateSPIRVVectorType(Type, 4, MIRBuilder, false);
4112
4113 uint64_t VectorSize = Type->getOperand(2).getImm();
4114 if (VectorSize == 4)
4115 return Type;
4116
4117 Register ScalarTypeReg = Type->getOperand(1).getReg();
4118 const SPIRVType *ScalarType = GR.getSPIRVTypeForVReg(ScalarTypeReg);
4119 return GR.getOrCreateSPIRVVectorType(ScalarType, 4, MIRBuilder, false);
4120 }
4121
loadHandleBeforePosition(Register & HandleReg,const SPIRVType * ResType,GIntrinsic & HandleDef,MachineInstr & Pos) const4122 bool SPIRVInstructionSelector::loadHandleBeforePosition(
4123 Register &HandleReg, const SPIRVType *ResType, GIntrinsic &HandleDef,
4124 MachineInstr &Pos) const {
4125
4126 assert(HandleDef.getIntrinsicID() ==
4127 Intrinsic::spv_resource_handlefrombinding);
4128 uint32_t Set = foldImm(HandleDef.getOperand(2), MRI);
4129 uint32_t Binding = foldImm(HandleDef.getOperand(3), MRI);
4130 uint32_t ArraySize = foldImm(HandleDef.getOperand(4), MRI);
4131 Register IndexReg = HandleDef.getOperand(5).getReg();
4132 bool IsNonUniform = ArraySize > 1 && foldImm(HandleDef.getOperand(6), MRI);
4133 std::string Name =
4134 getStringValueFromReg(HandleDef.getOperand(7).getReg(), *MRI);
4135
4136 bool IsStructuredBuffer = ResType->getOpcode() == SPIRV::OpTypePointer;
4137 MachineIRBuilder MIRBuilder(HandleDef);
4138 SPIRVType *VarType = ResType;
4139 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
4140
4141 if (IsStructuredBuffer) {
4142 VarType = GR.getPointeeType(ResType);
4143 SC = GR.getPointerStorageClass(ResType);
4144 }
4145
4146 Register VarReg =
4147 buildPointerToResource(VarType, SC, Set, Binding, ArraySize, IndexReg,
4148 IsNonUniform, Name, MIRBuilder);
4149
4150 if (IsNonUniform)
4151 buildOpDecorate(HandleReg, HandleDef, TII, SPIRV::Decoration::NonUniformEXT,
4152 {});
4153
4154 // The handle for the buffer is the pointer to the resource. For an image, the
4155 // handle is the image object. So images get an extra load.
4156 uint32_t LoadOpcode =
4157 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
4158 GR.assignSPIRVTypeToVReg(ResType, HandleReg, *Pos.getMF());
4159 return BuildMI(*Pos.getParent(), Pos, HandleDef.getDebugLoc(),
4160 TII.get(LoadOpcode))
4161 .addDef(HandleReg)
4162 .addUse(GR.getSPIRVTypeID(ResType))
4163 .addUse(VarReg)
4164 .constrainAllUses(TII, TRI, RBI);
4165 }
4166
4167 namespace llvm {
4168 InstructionSelector *
createSPIRVInstructionSelector(const SPIRVTargetMachine & TM,const SPIRVSubtarget & Subtarget,const RegisterBankInfo & RBI)4169 createSPIRVInstructionSelector(const SPIRVTargetMachine &TM,
4170 const SPIRVSubtarget &Subtarget,
4171 const RegisterBankInfo &RBI) {
4172 return new SPIRVInstructionSelector(TM, Subtarget, RBI);
4173 }
4174 } // namespace llvm
4175