LLVM 23.0.0git
SPIRVInstructionSelector.cpp
Go to the documentation of this file.
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
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"
32#include "llvm/IR/IntrinsicsSPIRV.h"
33#include "llvm/Support/Debug.h"
35
36#define DEBUG_TYPE "spirv-isel"
37
38using namespace llvm;
39namespace CL = SPIRV::OpenCLExtInst;
40namespace GL = SPIRV::GLSLExtInst;
41
43 std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;
44
45namespace {
46
47llvm::SPIRV::SelectionControl::SelectionControl
48getSelectionOperandForImm(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
62class SPIRVInstructionSelector : public InstructionSelector {
63 const SPIRVSubtarget &STI;
64 const SPIRVInstrInfo &TII;
66 const RegisterBankInfo &RBI;
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;
75
76public:
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;
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
95private:
96 void resetVRegsType(MachineFunction &MF);
97 void removeDeadInstruction(MachineInstr &MI) const;
98 void removeOpNamesForDeadMI(MachineInstr &MI) const;
99
100 // tblgen-erated 'select' implementation, used as the initial selector for
101 // the patterns that don't require complex C++.
102 bool selectImpl(MachineInstr &I, CodeGenCoverage &CoverageInfo) const;
103
104 // All instruction-specific selection that didn't happen in "select()".
105 // Is basically a large Switch/Case delegating to all other select method.
106 bool spvSelect(Register ResVReg, const SPIRVType *ResType,
107 MachineInstr &I) const;
108
109 bool selectFirstBitHigh(Register ResVReg, const SPIRVType *ResType,
110 MachineInstr &I, bool IsSigned) const;
111
112 bool selectFirstBitLow(Register ResVReg, const SPIRVType *ResType,
113 MachineInstr &I) const;
114
115 bool selectFirstBitSet16(Register ResVReg, const SPIRVType *ResType,
116 MachineInstr &I, unsigned ExtendOpcode,
117 unsigned BitSetOpcode) const;
118
119 bool selectFirstBitSet32(Register ResVReg, const SPIRVType *ResType,
120 MachineInstr &I, Register SrcReg,
121 unsigned BitSetOpcode) const;
122
123 bool selectFirstBitSet64(Register ResVReg, const SPIRVType *ResType,
124 MachineInstr &I, Register SrcReg,
125 unsigned BitSetOpcode, bool SwapPrimarySide) const;
126
127 bool selectFirstBitSet64Overflow(Register ResVReg, const SPIRVType *ResType,
128 MachineInstr &I, Register SrcReg,
129 unsigned BitSetOpcode,
130 bool SwapPrimarySide) const;
131
132 bool selectGlobalValue(Register ResVReg, MachineInstr &I,
133 const MachineInstr *Init = nullptr) const;
134
135 bool selectOpWithSrcs(Register ResVReg, const SPIRVType *ResType,
136 MachineInstr &I, std::vector<Register> SrcRegs,
137 unsigned Opcode) const;
138
139 bool selectUnOp(Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
140 unsigned Opcode) const;
141
142 bool selectBitcast(Register ResVReg, const SPIRVType *ResType,
143 MachineInstr &I) const;
144
145 bool selectLoad(Register ResVReg, const SPIRVType *ResType,
146 MachineInstr &I) const;
147 bool selectStore(MachineInstr &I) const;
148
149 bool selectStackSave(Register ResVReg, const SPIRVType *ResType,
150 MachineInstr &I) const;
151 bool selectStackRestore(MachineInstr &I) const;
152
153 bool selectMemOperation(Register ResVReg, MachineInstr &I) const;
154 Register getOrCreateMemSetGlobal(MachineInstr &I) const;
155 bool selectCopyMemory(MachineInstr &I, Register SrcReg) const;
156 bool selectCopyMemorySized(MachineInstr &I, Register SrcReg) const;
157
158 bool selectAtomicRMW(Register ResVReg, const SPIRVType *ResType,
159 MachineInstr &I, unsigned NewOpcode,
160 unsigned NegateOpcode = 0) const;
161
162 bool selectAtomicCmpXchg(Register ResVReg, const SPIRVType *ResType,
163 MachineInstr &I) const;
164
165 bool selectFence(MachineInstr &I) const;
166
167 bool selectAddrSpaceCast(Register ResVReg, const SPIRVType *ResType,
168 MachineInstr &I) const;
169
170 bool selectAnyOrAll(Register ResVReg, const SPIRVType *ResType,
171 MachineInstr &I, unsigned OpType) const;
172
173 bool selectAll(Register ResVReg, const SPIRVType *ResType,
174 MachineInstr &I) const;
175
176 bool selectAny(Register ResVReg, const SPIRVType *ResType,
177 MachineInstr &I) const;
178
179 bool selectBitreverse(Register ResVReg, const SPIRVType *ResType,
180 MachineInstr &I) const;
181
182 bool selectBuildVector(Register ResVReg, const SPIRVType *ResType,
183 MachineInstr &I) const;
184 bool selectSplatVector(Register ResVReg, const SPIRVType *ResType,
185 MachineInstr &I) const;
186
187 bool selectCmp(Register ResVReg, const SPIRVType *ResType,
188 unsigned comparisonOpcode, MachineInstr &I) const;
189 bool selectDiscard(Register ResVReg, const SPIRVType *ResType,
190 MachineInstr &I) const;
191
192 bool selectICmp(Register ResVReg, const SPIRVType *ResType,
193 MachineInstr &I) const;
194 bool selectFCmp(Register ResVReg, const SPIRVType *ResType,
195 MachineInstr &I) const;
196
197 bool selectSign(Register ResVReg, const SPIRVType *ResType,
198 MachineInstr &I) const;
199
200 bool selectFloatDot(Register ResVReg, const SPIRVType *ResType,
201 MachineInstr &I) const;
202
203 bool selectOverflowArith(Register ResVReg, const SPIRVType *ResType,
204 MachineInstr &I, unsigned Opcode) const;
205 bool selectDebugTrap(Register ResVReg, const SPIRVType *ResType,
206 MachineInstr &I) const;
207
208 bool selectIntegerDot(Register ResVReg, const SPIRVType *ResType,
209 MachineInstr &I, bool Signed) const;
210
211 bool selectIntegerDotExpansion(Register ResVReg, const SPIRVType *ResType,
212 MachineInstr &I) const;
213
214 bool selectOpIsInf(Register ResVReg, const SPIRVType *ResType,
215 MachineInstr &I) const;
216
217 bool selectOpIsNan(Register ResVReg, const SPIRVType *ResType,
218 MachineInstr &I) const;
219
220 template <bool Signed>
221 bool selectDot4AddPacked(Register ResVReg, const SPIRVType *ResType,
222 MachineInstr &I) const;
223 template <bool Signed>
224 bool selectDot4AddPackedExpansion(Register ResVReg, const SPIRVType *ResType,
225 MachineInstr &I) const;
226
227 bool selectWavePrefixBitCount(Register ResVReg, const SPIRVType *ResType,
228 MachineInstr &I) const;
229
230 bool selectWaveReduceMax(Register ResVReg, const SPIRVType *ResType,
231 MachineInstr &I, bool IsUnsigned) const;
232
233 bool selectWaveReduceMin(Register ResVReg, const SPIRVType *ResType,
234 MachineInstr &I, bool IsUnsigned) const;
235
236 bool selectWaveReduceSum(Register ResVReg, const SPIRVType *ResType,
237 MachineInstr &I) const;
238
239 bool selectConst(Register ResVReg, const SPIRVType *ResType,
240 MachineInstr &I) const;
241
242 bool selectSelect(Register ResVReg, const SPIRVType *ResType,
243 MachineInstr &I) const;
244 bool selectSelectDefaultArgs(Register ResVReg, const SPIRVType *ResType,
245 MachineInstr &I, bool IsSigned) const;
246 bool selectIToF(Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
247 bool IsSigned, unsigned Opcode) const;
248 bool selectExt(Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
249 bool IsSigned) const;
250
251 bool selectTrunc(Register ResVReg, const SPIRVType *ResType,
252 MachineInstr &I) const;
253
254 bool selectSUCmp(Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
255 bool IsSigned) const;
256
257 bool selectIntToBool(Register IntReg, Register ResVReg, MachineInstr &I,
258 const SPIRVType *intTy, const SPIRVType *boolTy) const;
259
260 bool selectOpUndef(Register ResVReg, const SPIRVType *ResType,
261 MachineInstr &I) const;
262 bool selectFreeze(Register ResVReg, const SPIRVType *ResType,
263 MachineInstr &I) const;
264 bool selectIntrinsic(Register ResVReg, const SPIRVType *ResType,
265 MachineInstr &I) const;
266 bool selectExtractVal(Register ResVReg, const SPIRVType *ResType,
267 MachineInstr &I) const;
268 bool selectInsertVal(Register ResVReg, const SPIRVType *ResType,
269 MachineInstr &I) const;
270 bool selectExtractElt(Register ResVReg, const SPIRVType *ResType,
271 MachineInstr &I) const;
272 bool selectInsertElt(Register ResVReg, const SPIRVType *ResType,
273 MachineInstr &I) const;
274 bool selectGEP(Register ResVReg, const SPIRVType *ResType,
275 MachineInstr &I) const;
276
277 bool selectFrameIndex(Register ResVReg, const SPIRVType *ResType,
278 MachineInstr &I) const;
279 bool selectAllocaArray(Register ResVReg, const SPIRVType *ResType,
280 MachineInstr &I) const;
281
282 bool selectBranch(MachineInstr &I) const;
283 bool selectBranchCond(MachineInstr &I) const;
284
285 bool selectPhi(Register ResVReg, const SPIRVType *ResType,
286 MachineInstr &I) const;
287
288 bool selectExtInst(Register ResVReg, const SPIRVType *RestType,
289 MachineInstr &I, GL::GLSLExtInst GLInst) const;
290 bool selectExtInst(Register ResVReg, const SPIRVType *ResType,
291 MachineInstr &I, CL::OpenCLExtInst CLInst) const;
292 bool selectExtInst(Register ResVReg, const SPIRVType *ResType,
293 MachineInstr &I, CL::OpenCLExtInst CLInst,
294 GL::GLSLExtInst GLInst) const;
295 bool selectExtInst(Register ResVReg, const SPIRVType *ResType,
296 MachineInstr &I, const ExtInstList &ExtInsts) const;
297 bool selectExtInstForLRound(Register ResVReg, const SPIRVType *ResType,
298 MachineInstr &I, CL::OpenCLExtInst CLInst,
299 GL::GLSLExtInst GLInst) const;
300 bool selectExtInstForLRound(Register ResVReg, const SPIRVType *ResType,
302 const ExtInstList &ExtInsts) const;
303
304 bool selectLog10(Register ResVReg, const SPIRVType *ResType,
305 MachineInstr &I) const;
306
307 bool selectSaturate(Register ResVReg, const SPIRVType *ResType,
308 MachineInstr &I) const;
309
310 bool selectWaveOpInst(Register ResVReg, const SPIRVType *ResType,
311 MachineInstr &I, unsigned Opcode) const;
312
313 bool selectWaveActiveCountBits(Register ResVReg, const SPIRVType *ResType,
314 MachineInstr &I) const;
315
317
318 bool selectHandleFromBinding(Register &ResVReg, const SPIRVType *ResType,
319 MachineInstr &I) const;
320
321 bool selectCounterHandleFromBinding(Register &ResVReg,
322 const SPIRVType *ResType,
323 MachineInstr &I) const;
324
325 bool selectReadImageIntrinsic(Register &ResVReg, const SPIRVType *ResType,
326 MachineInstr &I) const;
327 bool selectSampleIntrinsic(Register &ResVReg, const SPIRVType *ResType,
328 MachineInstr &I) const;
329 bool selectImageWriteIntrinsic(MachineInstr &I) const;
330 bool selectResourceGetPointer(Register &ResVReg, const SPIRVType *ResType,
331 MachineInstr &I) const;
332 bool selectPushConstantGetPointer(Register &ResVReg, const SPIRVType *ResType,
333 MachineInstr &I) const;
334 bool selectResourceNonUniformIndex(Register &ResVReg,
335 const SPIRVType *ResType,
336 MachineInstr &I) const;
337 bool selectModf(Register ResVReg, const SPIRVType *ResType,
338 MachineInstr &I) const;
339 bool selectUpdateCounter(Register &ResVReg, const SPIRVType *ResType,
340 MachineInstr &I) const;
341 bool selectFrexp(Register ResVReg, const SPIRVType *ResType,
342 MachineInstr &I) const;
343 bool selectDerivativeInst(Register ResVReg, const SPIRVType *ResType,
344 MachineInstr &I, const unsigned DPdOpCode) const;
345 // Utilities
346 std::pair<Register, bool>
347 buildI32Constant(uint32_t Val, MachineInstr &I,
348 const SPIRVType *ResType = nullptr) const;
349
350 Register buildZerosVal(const SPIRVType *ResType, MachineInstr &I) const;
351 bool isScalarOrVectorIntConstantZero(Register Reg) const;
352 Register buildZerosValF(const SPIRVType *ResType, MachineInstr &I) const;
353 Register buildOnesVal(bool AllOnes, const SPIRVType *ResType,
354 MachineInstr &I) const;
355 Register buildOnesValF(const SPIRVType *ResType, MachineInstr &I) const;
356
357 bool wrapIntoSpecConstantOp(MachineInstr &I,
358 SmallVector<Register> &CompositeArgs) const;
359
360 Register getUcharPtrTypeReg(MachineInstr &I,
361 SPIRV::StorageClass::StorageClass SC) const;
362 MachineInstrBuilder buildSpecConstantOp(MachineInstr &I, Register Dest,
363 Register Src, Register DestType,
364 uint32_t Opcode) const;
365 MachineInstrBuilder buildConstGenericPtr(MachineInstr &I, Register SrcPtr,
366 SPIRVType *SrcPtrTy) const;
367 Register buildPointerToResource(const SPIRVType *ResType,
368 SPIRV::StorageClass::StorageClass SC,
370 uint32_t ArraySize, Register IndexReg,
371 StringRef Name,
372 MachineIRBuilder MIRBuilder) const;
373 SPIRVType *widenTypeToVec4(const SPIRVType *Type, MachineInstr &I) const;
374 bool extractSubvector(Register &ResVReg, const SPIRVType *ResType,
375 Register &ReadReg, MachineInstr &InsertionPoint) const;
376 bool generateImageReadOrFetch(Register &ResVReg, const SPIRVType *ResType,
377 Register ImageReg, Register IdxReg,
378 DebugLoc Loc, MachineInstr &Pos) const;
379 bool BuildCOPY(Register DestReg, Register SrcReg, MachineInstr &I) const;
380 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
381 Register ResVReg, const SPIRVType *ResType,
382 MachineInstr &I) const;
383 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
384 Register ResVReg, const SPIRVType *ResType,
385 MachineInstr &I) const;
386 bool loadHandleBeforePosition(Register &HandleReg, const SPIRVType *ResType,
387 GIntrinsic &HandleDef, MachineInstr &Pos) const;
388 void decorateUsesAsNonUniform(Register &NonUniformReg) const;
389 void errorIfInstrOutsideShader(MachineInstr &I) const;
390};
391
392bool sampledTypeIsSignedInteger(const llvm::Type *HandleType) {
393 const TargetExtType *TET = cast<TargetExtType>(HandleType);
394 if (TET->getTargetExtName() == "spirv.Image") {
395 return false;
396 }
397 assert(TET->getTargetExtName() == "spirv.SignedImage");
398 return TET->getTypeParameter(0)->isIntegerTy();
399}
400} // end anonymous namespace
401
402#define GET_GLOBALISEL_IMPL
403#include "SPIRVGenGlobalISel.inc"
404#undef GET_GLOBALISEL_IMPL
405
406SPIRVInstructionSelector::SPIRVInstructionSelector(const SPIRVTargetMachine &TM,
407 const SPIRVSubtarget &ST,
408 const RegisterBankInfo &RBI)
409 : InstructionSelector(), STI(ST), TII(*ST.getInstrInfo()),
410 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
411 MRI(nullptr),
413#include "SPIRVGenGlobalISel.inc"
416#include "SPIRVGenGlobalISel.inc"
418{
419}
420
421void SPIRVInstructionSelector::setupMF(MachineFunction &MF,
423 CodeGenCoverage *CoverageInfo,
425 BlockFrequencyInfo *BFI) {
426 MRI = &MF.getRegInfo();
427 GR.setCurrentFunc(MF);
428 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
429}
430
431// Ensure that register classes correspond to pattern matching rules.
432void SPIRVInstructionSelector::resetVRegsType(MachineFunction &MF) {
433 if (HasVRegsReset == &MF)
434 return;
435 HasVRegsReset = &MF;
436
437 MachineRegisterInfo &MRI = MF.getRegInfo();
438 for (unsigned I = 0, E = MRI.getNumVirtRegs(); I != E; ++I) {
439 Register Reg = Register::index2VirtReg(I);
440 LLT RegType = MRI.getType(Reg);
441 if (RegType.isScalar())
442 MRI.setType(Reg, LLT::scalar(64));
443 else if (RegType.isPointer())
444 MRI.setType(Reg, LLT::pointer(0, 64));
445 else if (RegType.isVector())
446 MRI.setType(Reg, LLT::fixed_vector(2, LLT::scalar(64)));
447 }
448 for (const auto &MBB : MF) {
449 for (const auto &MI : MBB) {
450 if (isPreISelGenericOpcode(MI.getOpcode()))
451 GR.erase(&MI);
452 if (MI.getOpcode() != SPIRV::ASSIGN_TYPE)
453 continue;
454
455 Register DstReg = MI.getOperand(0).getReg();
456 LLT DstType = MRI.getType(DstReg);
457 Register SrcReg = MI.getOperand(1).getReg();
458 LLT SrcType = MRI.getType(SrcReg);
459 if (DstType != SrcType)
460 MRI.setType(DstReg, MRI.getType(SrcReg));
461
462 const TargetRegisterClass *DstRC = MRI.getRegClassOrNull(DstReg);
463 const TargetRegisterClass *SrcRC = MRI.getRegClassOrNull(SrcReg);
464 if (DstRC != SrcRC && SrcRC)
465 MRI.setRegClass(DstReg, SrcRC);
466 }
467 }
468}
469
470// Return true if the type represents a constant register
473 OpDef = passCopy(OpDef, MRI);
474
475 if (Visited.contains(OpDef))
476 return true;
477 Visited.insert(OpDef);
478
479 unsigned Opcode = OpDef->getOpcode();
480 switch (Opcode) {
481 case TargetOpcode::G_CONSTANT:
482 case TargetOpcode::G_FCONSTANT:
483 case TargetOpcode::G_IMPLICIT_DEF:
484 return true;
485 case TargetOpcode::G_INTRINSIC:
486 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
487 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
488 return cast<GIntrinsic>(*OpDef).getIntrinsicID() ==
489 Intrinsic::spv_const_composite;
490 case TargetOpcode::G_BUILD_VECTOR:
491 case TargetOpcode::G_SPLAT_VECTOR: {
492 for (unsigned i = OpDef->getNumExplicitDefs(); i < OpDef->getNumOperands();
493 i++) {
494 MachineInstr *OpNestedDef =
495 OpDef->getOperand(i).isReg()
496 ? MRI->getVRegDef(OpDef->getOperand(i).getReg())
497 : nullptr;
498 if (OpNestedDef && !isConstReg(MRI, OpNestedDef, Visited))
499 return false;
500 }
501 return true;
502 case SPIRV::OpConstantTrue:
503 case SPIRV::OpConstantFalse:
504 case SPIRV::OpConstantI:
505 case SPIRV::OpConstantF:
506 case SPIRV::OpConstantComposite:
507 case SPIRV::OpConstantCompositeContinuedINTEL:
508 case SPIRV::OpConstantSampler:
509 case SPIRV::OpConstantNull:
510 case SPIRV::OpUndef:
511 case SPIRV::OpConstantFunctionPointerINTEL:
512 return true;
513 }
514 }
515 return false;
516}
517
518// Return true if the virtual register represents a constant
521 if (MachineInstr *OpDef = MRI->getVRegDef(OpReg))
522 return isConstReg(MRI, OpDef, Visited);
523 return false;
524}
525
526// TODO(168736): We should make this either a flag in tabelgen
527// or reduce our dependence on the global registry, so we can remove this
528// function. It can easily be missed when new intrinsics are added.
529
530// Most SPIR-V intrinsics are considered to have side-effects in their tablegen
531// definition because they are referenced in the global registry. This is a list
532// of intrinsics that have no side effects other than their references in the
533// global registry.
535 switch (ID) {
536 // This is not an exhaustive list and may need to be updated.
537 case Intrinsic::spv_all:
538 case Intrinsic::spv_alloca:
539 case Intrinsic::spv_any:
540 case Intrinsic::spv_bitcast:
541 case Intrinsic::spv_const_composite:
542 case Intrinsic::spv_cross:
543 case Intrinsic::spv_degrees:
544 case Intrinsic::spv_distance:
545 case Intrinsic::spv_extractelt:
546 case Intrinsic::spv_extractv:
547 case Intrinsic::spv_faceforward:
548 case Intrinsic::spv_fdot:
549 case Intrinsic::spv_firstbitlow:
550 case Intrinsic::spv_firstbitshigh:
551 case Intrinsic::spv_firstbituhigh:
552 case Intrinsic::spv_frac:
553 case Intrinsic::spv_gep:
554 case Intrinsic::spv_global_offset:
555 case Intrinsic::spv_global_size:
556 case Intrinsic::spv_group_id:
557 case Intrinsic::spv_insertelt:
558 case Intrinsic::spv_insertv:
559 case Intrinsic::spv_isinf:
560 case Intrinsic::spv_isnan:
561 case Intrinsic::spv_lerp:
562 case Intrinsic::spv_length:
563 case Intrinsic::spv_normalize:
564 case Intrinsic::spv_num_subgroups:
565 case Intrinsic::spv_num_workgroups:
566 case Intrinsic::spv_ptrcast:
567 case Intrinsic::spv_radians:
568 case Intrinsic::spv_reflect:
569 case Intrinsic::spv_refract:
570 case Intrinsic::spv_resource_getpointer:
571 case Intrinsic::spv_resource_handlefrombinding:
572 case Intrinsic::spv_resource_handlefromimplicitbinding:
573 case Intrinsic::spv_resource_nonuniformindex:
574 case Intrinsic::spv_resource_sample:
575 case Intrinsic::spv_rsqrt:
576 case Intrinsic::spv_saturate:
577 case Intrinsic::spv_sdot:
578 case Intrinsic::spv_sign:
579 case Intrinsic::spv_smoothstep:
580 case Intrinsic::spv_step:
581 case Intrinsic::spv_subgroup_id:
582 case Intrinsic::spv_subgroup_local_invocation_id:
583 case Intrinsic::spv_subgroup_max_size:
584 case Intrinsic::spv_subgroup_size:
585 case Intrinsic::spv_thread_id:
586 case Intrinsic::spv_thread_id_in_group:
587 case Intrinsic::spv_udot:
588 case Intrinsic::spv_undef:
589 case Intrinsic::spv_value_md:
590 case Intrinsic::spv_workgroup_size:
591 return false;
592 default:
593 return true;
594 }
595}
596
597// TODO(168736): We should make this either a flag in tabelgen
598// or reduce our dependence on the global registry, so we can remove this
599// function. It can easily be missed when new intrinsics are added.
600static bool isOpcodeWithNoSideEffects(unsigned Opcode) {
601 switch (Opcode) {
602 case SPIRV::OpTypeVoid:
603 case SPIRV::OpTypeBool:
604 case SPIRV::OpTypeInt:
605 case SPIRV::OpTypeFloat:
606 case SPIRV::OpTypeVector:
607 case SPIRV::OpTypeMatrix:
608 case SPIRV::OpTypeImage:
609 case SPIRV::OpTypeSampler:
610 case SPIRV::OpTypeSampledImage:
611 case SPIRV::OpTypeArray:
612 case SPIRV::OpTypeRuntimeArray:
613 case SPIRV::OpTypeStruct:
614 case SPIRV::OpTypeOpaque:
615 case SPIRV::OpTypePointer:
616 case SPIRV::OpTypeFunction:
617 case SPIRV::OpTypeEvent:
618 case SPIRV::OpTypeDeviceEvent:
619 case SPIRV::OpTypeReserveId:
620 case SPIRV::OpTypeQueue:
621 case SPIRV::OpTypePipe:
622 case SPIRV::OpTypeForwardPointer:
623 case SPIRV::OpTypePipeStorage:
624 case SPIRV::OpTypeNamedBarrier:
625 case SPIRV::OpTypeAccelerationStructureNV:
626 case SPIRV::OpTypeCooperativeMatrixNV:
627 case SPIRV::OpTypeCooperativeMatrixKHR:
628 return true;
629 default:
630 return false;
631 }
632}
633
635 // If there are no definitions, then assume there is some other
636 // side-effect that makes this instruction live.
637 if (MI.getNumDefs() == 0)
638 return false;
639
640 for (const auto &MO : MI.all_defs()) {
641 Register Reg = MO.getReg();
642 if (Reg.isPhysical()) {
643 LLVM_DEBUG(dbgs() << "Not dead: def of physical register " << Reg);
644 return false;
645 }
646 for (const auto &UseMI : MRI.use_nodbg_instructions(Reg)) {
647 if (UseMI.getOpcode() != SPIRV::OpName) {
648 LLVM_DEBUG(dbgs() << "Not dead: def " << MO << " has use in " << UseMI);
649 return false;
650 }
651 }
652 }
653
654 if (MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE || MI.isFakeUse() ||
655 MI.isLifetimeMarker()) {
657 dbgs()
658 << "Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
659 return false;
660 }
661 if (MI.isPHI()) {
662 LLVM_DEBUG(dbgs() << "Dead: Phi instruction with no uses.\n");
663 return true;
664 }
665
666 // It is possible that the only side effect is that the instruction is
667 // referenced in the global registry. If that is the only side effect, the
668 // intrinsic is dead.
669 if (MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
670 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
671 const auto &Intr = cast<GIntrinsic>(MI);
672 if (!intrinsicHasSideEffects(Intr.getIntrinsicID())) {
673 LLVM_DEBUG(dbgs() << "Dead: Intrinsic with no real side effects.\n");
674 return true;
675 }
676 }
677
678 if (MI.mayStore() || MI.isCall() ||
679 (MI.mayLoad() && MI.hasOrderedMemoryRef()) || MI.isPosition() ||
680 MI.isDebugInstr() || MI.isTerminator() || MI.isJumpTableDebugInfo()) {
681 LLVM_DEBUG(dbgs() << "Not dead: instruction has side effects.\n");
682 return false;
683 }
684
685 if (isPreISelGenericOpcode(MI.getOpcode())) {
686 // TODO: Is there a generic way to check if the opcode has side effects?
687 LLVM_DEBUG(dbgs() << "Dead: Generic opcode with no uses.\n");
688 return true;
689 }
690
691 if (isOpcodeWithNoSideEffects(MI.getOpcode())) {
692 LLVM_DEBUG(dbgs() << "Dead: known opcode with no side effects\n");
693 return true;
694 }
695
696 return false;
697}
698
699void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &MI) const {
700 // Delete the OpName that uses the result if there is one.
701 for (const auto &MO : MI.all_defs()) {
702 Register Reg = MO.getReg();
703 if (Reg.isPhysical())
704 continue;
705 SmallVector<MachineInstr *, 4> UselessOpNames;
706 for (MachineInstr &UseMI : MRI->use_nodbg_instructions(Reg)) {
707 assert(UseMI.getOpcode() == SPIRV::OpName &&
708 "There is still a use of the dead function.");
709 UselessOpNames.push_back(&UseMI);
710 }
711 for (MachineInstr *OpNameMI : UselessOpNames) {
712 GR.invalidateMachineInstr(OpNameMI);
713 OpNameMI->eraseFromParent();
714 }
715 }
716}
717
718void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &MI) const {
721 removeOpNamesForDeadMI(MI);
722 MI.eraseFromParent();
723}
724
725bool SPIRVInstructionSelector::select(MachineInstr &I) {
726 resetVRegsType(*I.getParent()->getParent());
727
728 assert(I.getParent() && "Instruction should be in a basic block!");
729 assert(I.getParent()->getParent() && "Instruction should be in a function!");
730
731 LLVM_DEBUG(dbgs() << "Checking if instruction is dead: " << I;);
732 if (isDead(I, *MRI)) {
733 LLVM_DEBUG(dbgs() << "Instruction is dead.\n");
734 removeDeadInstruction(I);
735 return true;
736 }
737
738 Register Opcode = I.getOpcode();
739 // If it's not a GMIR instruction, we've selected it already.
740 if (!isPreISelGenericOpcode(Opcode)) {
741 if (Opcode == SPIRV::ASSIGN_TYPE) { // These pseudos aren't needed any more.
742 Register DstReg = I.getOperand(0).getReg();
743 Register SrcReg = I.getOperand(1).getReg();
744 auto *Def = MRI->getVRegDef(SrcReg);
745 if (isTypeFoldingSupported(Def->getOpcode()) &&
746 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
747 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
748 bool Res = false;
749 if (Def->getOpcode() == TargetOpcode::G_SELECT) {
750 Register SelectDstReg = Def->getOperand(0).getReg();
751 Res = selectSelect(SelectDstReg, GR.getSPIRVTypeForVReg(SelectDstReg),
752 *Def);
754 Def->removeFromParent();
755 MRI->replaceRegWith(DstReg, SelectDstReg);
757 I.removeFromParent();
758 } else
759 Res = selectImpl(I, *CoverageInfo);
760 LLVM_DEBUG({
761 if (!Res && Def->getOpcode() != TargetOpcode::G_CONSTANT) {
762 dbgs() << "Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
763 I.print(dbgs());
764 }
765 });
766 assert(Res || Def->getOpcode() == TargetOpcode::G_CONSTANT);
767 if (Res) {
768 if (!isTriviallyDead(*Def, *MRI) && isDead(*Def, *MRI))
769 DeadMIs.insert(Def);
770 return Res;
771 }
772 }
773 MRI->setRegClass(SrcReg, MRI->getRegClass(DstReg));
774 MRI->replaceRegWith(SrcReg, DstReg);
776 I.removeFromParent();
777 return true;
778 } else if (I.getNumDefs() == 1) {
779 // Make all vregs 64 bits (for SPIR-V IDs).
780 MRI->setType(I.getOperand(0).getReg(), LLT::scalar(64));
781 }
783 }
784
785 if (DeadMIs.contains(&I)) {
786 // if the instruction has been already made dead by folding it away
787 // erase it
788 LLVM_DEBUG(dbgs() << "Instruction is folded and dead.\n");
789 removeDeadInstruction(I);
790 return true;
791 }
792
793 if (I.getNumOperands() != I.getNumExplicitOperands()) {
794 LLVM_DEBUG(errs() << "Generic instr has unexpected implicit operands\n");
795 return false;
796 }
797
798 // Common code for getting return reg+type, and removing selected instr
799 // from parent occurs here. Instr-specific selection happens in spvSelect().
800 bool HasDefs = I.getNumDefs() > 0;
801 Register ResVReg = HasDefs ? I.getOperand(0).getReg() : Register(0);
802 SPIRVType *ResType = HasDefs ? GR.getSPIRVTypeForVReg(ResVReg) : nullptr;
803 assert(!HasDefs || ResType || I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
804 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
805 if (spvSelect(ResVReg, ResType, I)) {
806 if (HasDefs) // Make all vregs 64 bits (for SPIR-V IDs).
807 for (unsigned i = 0; i < I.getNumDefs(); ++i)
808 MRI->setType(I.getOperand(i).getReg(), LLT::scalar(64));
810 I.removeFromParent();
811 return true;
812 }
813 return false;
814}
815
816static bool mayApplyGenericSelection(unsigned Opcode) {
817 switch (Opcode) {
818 case TargetOpcode::G_CONSTANT:
819 case TargetOpcode::G_FCONSTANT:
820 return false;
821 case TargetOpcode::G_SADDO:
822 case TargetOpcode::G_SSUBO:
823 return true;
824 }
825 return isTypeFoldingSupported(Opcode);
826}
827
828bool SPIRVInstructionSelector::BuildCOPY(Register DestReg, Register SrcReg,
829 MachineInstr &I) const {
830 const TargetRegisterClass *DstRC = MRI->getRegClassOrNull(DestReg);
831 const TargetRegisterClass *SrcRC = MRI->getRegClassOrNull(SrcReg);
832 if (DstRC != SrcRC && SrcRC)
833 MRI->setRegClass(DestReg, SrcRC);
834 return BuildMI(*I.getParent(), I, I.getDebugLoc(),
835 TII.get(TargetOpcode::COPY))
836 .addDef(DestReg)
837 .addUse(SrcReg)
838 .constrainAllUses(TII, TRI, RBI);
839}
840
841bool SPIRVInstructionSelector::spvSelect(Register ResVReg,
842 const SPIRVType *ResType,
843 MachineInstr &I) const {
844 const unsigned Opcode = I.getOpcode();
845 if (mayApplyGenericSelection(Opcode))
846 return selectImpl(I, *CoverageInfo);
847 switch (Opcode) {
848 case TargetOpcode::G_CONSTANT:
849 case TargetOpcode::G_FCONSTANT:
850 return selectConst(ResVReg, ResType, I);
851 case TargetOpcode::G_GLOBAL_VALUE:
852 return selectGlobalValue(ResVReg, I);
853 case TargetOpcode::G_IMPLICIT_DEF:
854 return selectOpUndef(ResVReg, ResType, I);
855 case TargetOpcode::G_FREEZE:
856 return selectFreeze(ResVReg, ResType, I);
857
858 case TargetOpcode::G_INTRINSIC:
859 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
860 case TargetOpcode::G_INTRINSIC_CONVERGENT:
861 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
862 return selectIntrinsic(ResVReg, ResType, I);
863 case TargetOpcode::G_BITREVERSE:
864 return selectBitreverse(ResVReg, ResType, I);
865
866 case TargetOpcode::G_BUILD_VECTOR:
867 return selectBuildVector(ResVReg, ResType, I);
868 case TargetOpcode::G_SPLAT_VECTOR:
869 return selectSplatVector(ResVReg, ResType, I);
870
871 case TargetOpcode::G_SHUFFLE_VECTOR: {
872 MachineBasicBlock &BB = *I.getParent();
873 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpVectorShuffle))
874 .addDef(ResVReg)
875 .addUse(GR.getSPIRVTypeID(ResType))
876 .addUse(I.getOperand(1).getReg())
877 .addUse(I.getOperand(2).getReg());
878 for (auto V : I.getOperand(3).getShuffleMask())
879 MIB.addImm(V);
880 return MIB.constrainAllUses(TII, TRI, RBI);
881 }
882 case TargetOpcode::G_MEMMOVE:
883 case TargetOpcode::G_MEMCPY:
884 case TargetOpcode::G_MEMSET:
885 return selectMemOperation(ResVReg, I);
886
887 case TargetOpcode::G_ICMP:
888 return selectICmp(ResVReg, ResType, I);
889 case TargetOpcode::G_FCMP:
890 return selectFCmp(ResVReg, ResType, I);
891
892 case TargetOpcode::G_FRAME_INDEX:
893 return selectFrameIndex(ResVReg, ResType, I);
894
895 case TargetOpcode::G_LOAD:
896 return selectLoad(ResVReg, ResType, I);
897 case TargetOpcode::G_STORE:
898 return selectStore(I);
899
900 case TargetOpcode::G_BR:
901 return selectBranch(I);
902 case TargetOpcode::G_BRCOND:
903 return selectBranchCond(I);
904
905 case TargetOpcode::G_PHI:
906 return selectPhi(ResVReg, ResType, I);
907
908 case TargetOpcode::G_FPTOSI:
909 return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertFToS);
910 case TargetOpcode::G_FPTOUI:
911 return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertFToU);
912
913 case TargetOpcode::G_FPTOSI_SAT:
914 return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertFToS);
915 case TargetOpcode::G_FPTOUI_SAT:
916 return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertFToU);
917
918 case TargetOpcode::G_SITOFP:
919 return selectIToF(ResVReg, ResType, I, true, SPIRV::OpConvertSToF);
920 case TargetOpcode::G_UITOFP:
921 return selectIToF(ResVReg, ResType, I, false, SPIRV::OpConvertUToF);
922
923 case TargetOpcode::G_CTPOP:
924 return selectUnOp(ResVReg, ResType, I, SPIRV::OpBitCount);
925 case TargetOpcode::G_SMIN:
926 return selectExtInst(ResVReg, ResType, I, CL::s_min, GL::SMin);
927 case TargetOpcode::G_UMIN:
928 return selectExtInst(ResVReg, ResType, I, CL::u_min, GL::UMin);
929
930 case TargetOpcode::G_SMAX:
931 return selectExtInst(ResVReg, ResType, I, CL::s_max, GL::SMax);
932 case TargetOpcode::G_UMAX:
933 return selectExtInst(ResVReg, ResType, I, CL::u_max, GL::UMax);
934
935 case TargetOpcode::G_SCMP:
936 return selectSUCmp(ResVReg, ResType, I, true);
937 case TargetOpcode::G_UCMP:
938 return selectSUCmp(ResVReg, ResType, I, false);
939 case TargetOpcode::G_LROUND:
940 case TargetOpcode::G_LLROUND: {
941 Register regForLround =
942 MRI->createVirtualRegister(MRI->getRegClass(ResVReg), "lround");
943 MRI->setRegClass(regForLround, &SPIRV::iIDRegClass);
944 GR.assignSPIRVTypeToVReg(GR.getSPIRVTypeForVReg(I.getOperand(1).getReg()),
945 regForLround, *(I.getParent()->getParent()));
946 selectExtInstForLRound(regForLround, GR.getSPIRVTypeForVReg(regForLround),
947 I, CL::round, GL::Round);
948 MachineBasicBlock &BB = *I.getParent();
949 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConvertFToS))
950 .addDef(ResVReg)
951 .addUse(GR.getSPIRVTypeID(ResType))
952 .addUse(regForLround);
953 return MIB.constrainAllUses(TII, TRI, RBI);
954 }
955 case TargetOpcode::G_STRICT_FMA:
956 case TargetOpcode::G_FMA: {
957 if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_fma)) {
958 MachineBasicBlock &BB = *I.getParent();
959 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpFmaKHR))
960 .addDef(ResVReg)
961 .addUse(GR.getSPIRVTypeID(ResType))
962 .addUse(I.getOperand(1).getReg())
963 .addUse(I.getOperand(2).getReg())
964 .addUse(I.getOperand(3).getReg())
965 .setMIFlags(I.getFlags());
966 return MIB.constrainAllUses(TII, TRI, RBI);
967 }
968 return selectExtInst(ResVReg, ResType, I, CL::fma, GL::Fma);
969 }
970
971 case TargetOpcode::G_STRICT_FLDEXP:
972 return selectExtInst(ResVReg, ResType, I, CL::ldexp);
973
974 case TargetOpcode::G_FPOW:
975 return selectExtInst(ResVReg, ResType, I, CL::pow, GL::Pow);
976 case TargetOpcode::G_FPOWI:
977 return selectExtInst(ResVReg, ResType, I, CL::pown);
978
979 case TargetOpcode::G_FEXP:
980 return selectExtInst(ResVReg, ResType, I, CL::exp, GL::Exp);
981 case TargetOpcode::G_FEXP2:
982 return selectExtInst(ResVReg, ResType, I, CL::exp2, GL::Exp2);
983 case TargetOpcode::G_FMODF:
984 return selectModf(ResVReg, ResType, I);
985
986 case TargetOpcode::G_FLOG:
987 return selectExtInst(ResVReg, ResType, I, CL::log, GL::Log);
988 case TargetOpcode::G_FLOG2:
989 return selectExtInst(ResVReg, ResType, I, CL::log2, GL::Log2);
990 case TargetOpcode::G_FLOG10:
991 return selectLog10(ResVReg, ResType, I);
992
993 case TargetOpcode::G_FABS:
994 return selectExtInst(ResVReg, ResType, I, CL::fabs, GL::FAbs);
995 case TargetOpcode::G_ABS:
996 return selectExtInst(ResVReg, ResType, I, CL::s_abs, GL::SAbs);
997
998 case TargetOpcode::G_FMINNUM:
999 case TargetOpcode::G_FMINIMUM:
1000 return selectExtInst(ResVReg, ResType, I, CL::fmin, GL::NMin);
1001 case TargetOpcode::G_FMAXNUM:
1002 case TargetOpcode::G_FMAXIMUM:
1003 return selectExtInst(ResVReg, ResType, I, CL::fmax, GL::NMax);
1004
1005 case TargetOpcode::G_FCOPYSIGN:
1006 return selectExtInst(ResVReg, ResType, I, CL::copysign);
1007
1008 case TargetOpcode::G_FCEIL:
1009 return selectExtInst(ResVReg, ResType, I, CL::ceil, GL::Ceil);
1010 case TargetOpcode::G_FFLOOR:
1011 return selectExtInst(ResVReg, ResType, I, CL::floor, GL::Floor);
1012
1013 case TargetOpcode::G_FCOS:
1014 return selectExtInst(ResVReg, ResType, I, CL::cos, GL::Cos);
1015 case TargetOpcode::G_FSIN:
1016 return selectExtInst(ResVReg, ResType, I, CL::sin, GL::Sin);
1017 case TargetOpcode::G_FTAN:
1018 return selectExtInst(ResVReg, ResType, I, CL::tan, GL::Tan);
1019 case TargetOpcode::G_FACOS:
1020 return selectExtInst(ResVReg, ResType, I, CL::acos, GL::Acos);
1021 case TargetOpcode::G_FASIN:
1022 return selectExtInst(ResVReg, ResType, I, CL::asin, GL::Asin);
1023 case TargetOpcode::G_FATAN:
1024 return selectExtInst(ResVReg, ResType, I, CL::atan, GL::Atan);
1025 case TargetOpcode::G_FATAN2:
1026 return selectExtInst(ResVReg, ResType, I, CL::atan2, GL::Atan2);
1027 case TargetOpcode::G_FCOSH:
1028 return selectExtInst(ResVReg, ResType, I, CL::cosh, GL::Cosh);
1029 case TargetOpcode::G_FSINH:
1030 return selectExtInst(ResVReg, ResType, I, CL::sinh, GL::Sinh);
1031 case TargetOpcode::G_FTANH:
1032 return selectExtInst(ResVReg, ResType, I, CL::tanh, GL::Tanh);
1033
1034 case TargetOpcode::G_STRICT_FSQRT:
1035 case TargetOpcode::G_FSQRT:
1036 return selectExtInst(ResVReg, ResType, I, CL::sqrt, GL::Sqrt);
1037
1038 case TargetOpcode::G_CTTZ:
1039 case TargetOpcode::G_CTTZ_ZERO_UNDEF:
1040 return selectExtInst(ResVReg, ResType, I, CL::ctz);
1041 case TargetOpcode::G_CTLZ:
1042 case TargetOpcode::G_CTLZ_ZERO_UNDEF:
1043 return selectExtInst(ResVReg, ResType, I, CL::clz);
1044
1045 case TargetOpcode::G_INTRINSIC_ROUND:
1046 return selectExtInst(ResVReg, ResType, I, CL::round, GL::Round);
1047 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1048 return selectExtInst(ResVReg, ResType, I, CL::rint, GL::RoundEven);
1049 case TargetOpcode::G_INTRINSIC_TRUNC:
1050 return selectExtInst(ResVReg, ResType, I, CL::trunc, GL::Trunc);
1051 case TargetOpcode::G_FRINT:
1052 case TargetOpcode::G_FNEARBYINT:
1053 return selectExtInst(ResVReg, ResType, I, CL::rint, GL::RoundEven);
1054
1055 case TargetOpcode::G_SMULH:
1056 return selectExtInst(ResVReg, ResType, I, CL::s_mul_hi);
1057 case TargetOpcode::G_UMULH:
1058 return selectExtInst(ResVReg, ResType, I, CL::u_mul_hi);
1059
1060 case TargetOpcode::G_SADDSAT:
1061 return selectExtInst(ResVReg, ResType, I, CL::s_add_sat);
1062 case TargetOpcode::G_UADDSAT:
1063 return selectExtInst(ResVReg, ResType, I, CL::u_add_sat);
1064 case TargetOpcode::G_SSUBSAT:
1065 return selectExtInst(ResVReg, ResType, I, CL::s_sub_sat);
1066 case TargetOpcode::G_USUBSAT:
1067 return selectExtInst(ResVReg, ResType, I, CL::u_sub_sat);
1068
1069 case TargetOpcode::G_FFREXP:
1070 return selectFrexp(ResVReg, ResType, I);
1071
1072 case TargetOpcode::G_UADDO:
1073 return selectOverflowArith(ResVReg, ResType, I,
1074 ResType->getOpcode() == SPIRV::OpTypeVector
1075 ? SPIRV::OpIAddCarryV
1076 : SPIRV::OpIAddCarryS);
1077 case TargetOpcode::G_USUBO:
1078 return selectOverflowArith(ResVReg, ResType, I,
1079 ResType->getOpcode() == SPIRV::OpTypeVector
1080 ? SPIRV::OpISubBorrowV
1081 : SPIRV::OpISubBorrowS);
1082 case TargetOpcode::G_UMULO:
1083 return selectOverflowArith(ResVReg, ResType, I, SPIRV::OpUMulExtended);
1084 case TargetOpcode::G_SMULO:
1085 return selectOverflowArith(ResVReg, ResType, I, SPIRV::OpSMulExtended);
1086
1087 case TargetOpcode::G_SEXT:
1088 return selectExt(ResVReg, ResType, I, true);
1089 case TargetOpcode::G_ANYEXT:
1090 case TargetOpcode::G_ZEXT:
1091 return selectExt(ResVReg, ResType, I, false);
1092 case TargetOpcode::G_TRUNC:
1093 return selectTrunc(ResVReg, ResType, I);
1094 case TargetOpcode::G_FPTRUNC:
1095 case TargetOpcode::G_FPEXT:
1096 return selectUnOp(ResVReg, ResType, I, SPIRV::OpFConvert);
1097
1098 case TargetOpcode::G_PTRTOINT:
1099 return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertPtrToU);
1100 case TargetOpcode::G_INTTOPTR:
1101 return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertUToPtr);
1102 case TargetOpcode::G_BITCAST:
1103 return selectBitcast(ResVReg, ResType, I);
1104 case TargetOpcode::G_ADDRSPACE_CAST:
1105 return selectAddrSpaceCast(ResVReg, ResType, I);
1106 case TargetOpcode::G_PTR_ADD: {
1107 // Currently, we get G_PTR_ADD only applied to global variables.
1108 assert(I.getOperand(1).isReg() && I.getOperand(2).isReg());
1109 Register GV = I.getOperand(1).getReg();
1110 MachineRegisterInfo::def_instr_iterator II = MRI->def_instr_begin(GV);
1111 (void)II;
1112 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1113 (*II).getOpcode() == TargetOpcode::COPY ||
1114 (*II).getOpcode() == SPIRV::OpVariable) &&
1115 getImm(I.getOperand(2), MRI));
1116 // It may be the initialization of a global variable.
1117 bool IsGVInit = false;
1119 UseIt = MRI->use_instr_begin(I.getOperand(0).getReg()),
1120 UseEnd = MRI->use_instr_end();
1121 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1122 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1123 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1124 (*UseIt).getOpcode() == SPIRV::OpVariable) {
1125 IsGVInit = true;
1126 break;
1127 }
1128 }
1129 MachineBasicBlock &BB = *I.getParent();
1130 if (!IsGVInit) {
1131 SPIRVType *GVType = GR.getSPIRVTypeForVReg(GV);
1132 SPIRVType *GVPointeeType = GR.getPointeeType(GVType);
1133 SPIRVType *ResPointeeType = GR.getPointeeType(ResType);
1134 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1135 // Build a new virtual register that is associated with the required
1136 // data type.
1137 Register NewVReg = MRI->createGenericVirtualRegister(MRI->getType(GV));
1138 MRI->setRegClass(NewVReg, MRI->getRegClass(GV));
1139 // Having a correctly typed base we are ready to build the actually
1140 // required GEP. It may not be a constant though, because all Operands
1141 // of OpSpecConstantOp is to originate from other const instructions,
1142 // and only the AccessChain named opcodes accept a global OpVariable
1143 // instruction. We can't use an AccessChain opcode because of the type
1144 // mismatch between result and base types.
1145 if (!GR.isBitcastCompatible(ResType, GVType))
1147 "incompatible result and operand types in a bitcast");
1148 Register ResTypeReg = GR.getSPIRVTypeID(ResType);
1149 MachineInstrBuilder MIB =
1150 BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpBitcast))
1151 .addDef(NewVReg)
1152 .addUse(ResTypeReg)
1153 .addUse(GV);
1154 return MIB.constrainAllUses(TII, TRI, RBI) &&
1155 BuildMI(BB, I, I.getDebugLoc(),
1156 TII.get(STI.isLogicalSPIRV()
1157 ? SPIRV::OpInBoundsAccessChain
1158 : SPIRV::OpInBoundsPtrAccessChain))
1159 .addDef(ResVReg)
1160 .addUse(ResTypeReg)
1161 .addUse(NewVReg)
1162 .addUse(I.getOperand(2).getReg())
1163 .constrainAllUses(TII, TRI, RBI);
1164 } else {
1165 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSpecConstantOp))
1166 .addDef(ResVReg)
1167 .addUse(GR.getSPIRVTypeID(ResType))
1168 .addImm(
1169 static_cast<uint32_t>(SPIRV::Opcode::InBoundsPtrAccessChain))
1170 .addUse(GV)
1171 .addUse(I.getOperand(2).getReg())
1172 .constrainAllUses(TII, TRI, RBI);
1173 }
1174 }
1175 // It's possible to translate G_PTR_ADD to OpSpecConstantOp: either to
1176 // initialize a global variable with a constant expression (e.g., the test
1177 // case opencl/basic/progvar_prog_scope_init.ll), or for another use case
1178 Register Idx = buildZerosVal(GR.getOrCreateSPIRVIntegerType(32, I, TII), I);
1179 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSpecConstantOp))
1180 .addDef(ResVReg)
1181 .addUse(GR.getSPIRVTypeID(ResType))
1182 .addImm(static_cast<uint32_t>(
1183 SPIRV::Opcode::InBoundsPtrAccessChain))
1184 .addUse(GV)
1185 .addUse(Idx)
1186 .addUse(I.getOperand(2).getReg());
1187 return MIB.constrainAllUses(TII, TRI, RBI);
1188 }
1189
1190 case TargetOpcode::G_ATOMICRMW_OR:
1191 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicOr);
1192 case TargetOpcode::G_ATOMICRMW_ADD:
1193 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicIAdd);
1194 case TargetOpcode::G_ATOMICRMW_AND:
1195 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicAnd);
1196 case TargetOpcode::G_ATOMICRMW_MAX:
1197 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicSMax);
1198 case TargetOpcode::G_ATOMICRMW_MIN:
1199 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicSMin);
1200 case TargetOpcode::G_ATOMICRMW_SUB:
1201 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicISub);
1202 case TargetOpcode::G_ATOMICRMW_XOR:
1203 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicXor);
1204 case TargetOpcode::G_ATOMICRMW_UMAX:
1205 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicUMax);
1206 case TargetOpcode::G_ATOMICRMW_UMIN:
1207 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicUMin);
1208 case TargetOpcode::G_ATOMICRMW_XCHG:
1209 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicExchange);
1210 case TargetOpcode::G_ATOMIC_CMPXCHG:
1211 return selectAtomicCmpXchg(ResVReg, ResType, I);
1212
1213 case TargetOpcode::G_ATOMICRMW_FADD:
1214 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicFAddEXT);
1215 case TargetOpcode::G_ATOMICRMW_FSUB:
1216 // Translate G_ATOMICRMW_FSUB to OpAtomicFAddEXT with negative value operand
1217 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicFAddEXT,
1218 ResType->getOpcode() == SPIRV::OpTypeVector
1219 ? SPIRV::OpFNegateV
1220 : SPIRV::OpFNegate);
1221 case TargetOpcode::G_ATOMICRMW_FMIN:
1222 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicFMinEXT);
1223 case TargetOpcode::G_ATOMICRMW_FMAX:
1224 return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicFMaxEXT);
1225
1226 case TargetOpcode::G_FENCE:
1227 return selectFence(I);
1228
1229 case TargetOpcode::G_STACKSAVE:
1230 return selectStackSave(ResVReg, ResType, I);
1231 case TargetOpcode::G_STACKRESTORE:
1232 return selectStackRestore(I);
1233
1234 case TargetOpcode::G_UNMERGE_VALUES:
1235 return selectUnmergeValues(I);
1236
1237 // Discard gen opcodes for intrinsics which we do not expect to actually
1238 // represent code after lowering or intrinsics which are not implemented but
1239 // should not crash when found in a customer's LLVM IR input.
1240 case TargetOpcode::G_TRAP:
1241 case TargetOpcode::G_UBSANTRAP:
1242 case TargetOpcode::DBG_LABEL:
1243 return true;
1244 case TargetOpcode::G_DEBUGTRAP:
1245 return selectDebugTrap(ResVReg, ResType, I);
1246
1247 default:
1248 return false;
1249 }
1250}
1251
1252bool SPIRVInstructionSelector::selectDebugTrap(Register ResVReg,
1253 const SPIRVType *ResType,
1254 MachineInstr &I) const {
1255 unsigned Opcode = SPIRV::OpNop;
1256 MachineBasicBlock &BB = *I.getParent();
1257 return BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
1258 .constrainAllUses(TII, TRI, RBI);
1259}
1260
1261bool SPIRVInstructionSelector::selectExtInst(Register ResVReg,
1262 const SPIRVType *ResType,
1263 MachineInstr &I,
1264 GL::GLSLExtInst GLInst) const {
1265 if (!STI.canUseExtInstSet(
1266 SPIRV::InstructionSet::InstructionSet::GLSL_std_450)) {
1267 std::string DiagMsg;
1268 raw_string_ostream OS(DiagMsg);
1269 I.print(OS, true, false, false, false);
1270 DiagMsg += " is only supported with the GLSL extended instruction set.\n";
1271 report_fatal_error(DiagMsg.c_str(), false);
1272 }
1273 return selectExtInst(ResVReg, ResType, I,
1274 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}});
1275}
1276
1277bool SPIRVInstructionSelector::selectExtInst(Register ResVReg,
1278 const SPIRVType *ResType,
1279 MachineInstr &I,
1280 CL::OpenCLExtInst CLInst) const {
1281 return selectExtInst(ResVReg, ResType, I,
1282 {{SPIRV::InstructionSet::OpenCL_std, CLInst}});
1283}
1284
1285bool SPIRVInstructionSelector::selectExtInst(Register ResVReg,
1286 const SPIRVType *ResType,
1287 MachineInstr &I,
1288 CL::OpenCLExtInst CLInst,
1289 GL::GLSLExtInst GLInst) const {
1290 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1291 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1292 return selectExtInst(ResVReg, ResType, I, ExtInsts);
1293}
1294
1295bool SPIRVInstructionSelector::selectExtInst(Register ResVReg,
1296 const SPIRVType *ResType,
1297 MachineInstr &I,
1298 const ExtInstList &Insts) const {
1299
1300 for (const auto &Ex : Insts) {
1301 SPIRV::InstructionSet::InstructionSet Set = Ex.first;
1302 uint32_t Opcode = Ex.second;
1303 if (STI.canUseExtInstSet(Set)) {
1304 MachineBasicBlock &BB = *I.getParent();
1305 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
1306 .addDef(ResVReg)
1307 .addUse(GR.getSPIRVTypeID(ResType))
1308 .addImm(static_cast<uint32_t>(Set))
1309 .addImm(Opcode)
1310 .setMIFlags(I.getFlags());
1311 const unsigned NumOps = I.getNumOperands();
1312 unsigned Index = 1;
1313 if (Index < NumOps &&
1314 I.getOperand(Index).getType() ==
1315 MachineOperand::MachineOperandType::MO_IntrinsicID)
1316 Index = 2;
1317 for (; Index < NumOps; ++Index)
1318 MIB.add(I.getOperand(Index));
1319 return MIB.constrainAllUses(TII, TRI, RBI);
1320 }
1321 }
1322 return false;
1323}
1324bool SPIRVInstructionSelector::selectExtInstForLRound(
1325 Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
1326 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst) const {
1327 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1328 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1329 return selectExtInstForLRound(ResVReg, ResType, I, ExtInsts);
1330}
1331
1332bool SPIRVInstructionSelector::selectExtInstForLRound(
1333 Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
1334 const ExtInstList &Insts) const {
1335 for (const auto &Ex : Insts) {
1336 SPIRV::InstructionSet::InstructionSet Set = Ex.first;
1337 uint32_t Opcode = Ex.second;
1338 if (STI.canUseExtInstSet(Set)) {
1339 MachineBasicBlock &BB = *I.getParent();
1340 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
1341 .addDef(ResVReg)
1342 .addUse(GR.getSPIRVTypeID(ResType))
1343 .addImm(static_cast<uint32_t>(Set))
1344 .addImm(Opcode);
1345 const unsigned NumOps = I.getNumOperands();
1346 unsigned Index = 1;
1347 if (Index < NumOps &&
1348 I.getOperand(Index).getType() ==
1349 MachineOperand::MachineOperandType::MO_IntrinsicID)
1350 Index = 2;
1351 for (; Index < NumOps; ++Index)
1352 MIB.add(I.getOperand(Index));
1353 MIB.constrainAllUses(TII, TRI, RBI);
1354 return true;
1355 }
1356 }
1357 return false;
1358}
1359
1360bool SPIRVInstructionSelector::selectFrexp(Register ResVReg,
1361 const SPIRVType *ResType,
1362 MachineInstr &I) const {
1363 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1364 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1365 for (const auto &Ex : ExtInsts) {
1366 SPIRV::InstructionSet::InstructionSet Set = Ex.first;
1367 uint32_t Opcode = Ex.second;
1368 if (!STI.canUseExtInstSet(Set))
1369 continue;
1370
1371 MachineIRBuilder MIRBuilder(I);
1372 SPIRVType *PointeeTy = GR.getSPIRVTypeForVReg(I.getOperand(1).getReg());
1374 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1375 Register PointerVReg =
1376 createVirtualRegister(PointerType, &GR, MRI, MRI->getMF());
1377
1378 auto It = getOpVariableMBBIt(I);
1379 auto MIB = BuildMI(*It->getParent(), It, It->getDebugLoc(),
1380 TII.get(SPIRV::OpVariable))
1381 .addDef(PointerVReg)
1382 .addUse(GR.getSPIRVTypeID(PointerType))
1383 .addImm(static_cast<uint32_t>(SPIRV::StorageClass::Function))
1384 .constrainAllUses(TII, TRI, RBI);
1385
1386 MIB = MIB &
1387 BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
1388 .addDef(ResVReg)
1389 .addUse(GR.getSPIRVTypeID(ResType))
1390 .addImm(static_cast<uint32_t>(Ex.first))
1391 .addImm(Opcode)
1392 .add(I.getOperand(2))
1393 .addUse(PointerVReg)
1394 .constrainAllUses(TII, TRI, RBI);
1395
1396 MIB = MIB &
1397 BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpLoad))
1398 .addDef(I.getOperand(1).getReg())
1399 .addUse(GR.getSPIRVTypeID(PointeeTy))
1400 .addUse(PointerVReg)
1401 .constrainAllUses(TII, TRI, RBI);
1402 return MIB;
1403 }
1404 return false;
1405}
1406
1407bool SPIRVInstructionSelector::selectOpWithSrcs(Register ResVReg,
1408 const SPIRVType *ResType,
1409 MachineInstr &I,
1410 std::vector<Register> Srcs,
1411 unsigned Opcode) const {
1412 auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(Opcode))
1413 .addDef(ResVReg)
1414 .addUse(GR.getSPIRVTypeID(ResType));
1415 for (Register SReg : Srcs) {
1416 MIB.addUse(SReg);
1417 }
1418 return MIB.constrainAllUses(TII, TRI, RBI);
1419}
1420
1421bool SPIRVInstructionSelector::selectUnOp(Register ResVReg,
1422 const SPIRVType *ResType,
1423 MachineInstr &I,
1424 unsigned Opcode) const {
1425 if (STI.isPhysicalSPIRV() && I.getOperand(1).isReg()) {
1426 Register SrcReg = I.getOperand(1).getReg();
1427 bool IsGV = false;
1429 MRI->def_instr_begin(SrcReg);
1430 DefIt != MRI->def_instr_end(); DefIt = std::next(DefIt)) {
1431 unsigned DefOpCode = DefIt->getOpcode();
1432 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1433 // We need special handling to look through the type assignment or the
1434 // COPY pseudo-op and see if this is a constant or a global.
1435 if (auto *VRD = getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1436 DefOpCode = VRD->getOpcode();
1437 }
1438 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1439 DefOpCode == TargetOpcode::G_CONSTANT ||
1440 DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {
1441 IsGV = true;
1442 break;
1443 }
1444 }
1445 if (IsGV) {
1446 uint32_t SpecOpcode = 0;
1447 switch (Opcode) {
1448 case SPIRV::OpConvertPtrToU:
1449 SpecOpcode = static_cast<uint32_t>(SPIRV::Opcode::ConvertPtrToU);
1450 break;
1451 case SPIRV::OpConvertUToPtr:
1452 SpecOpcode = static_cast<uint32_t>(SPIRV::Opcode::ConvertUToPtr);
1453 break;
1454 }
1455 if (SpecOpcode)
1456 return BuildMI(*I.getParent(), I, I.getDebugLoc(),
1457 TII.get(SPIRV::OpSpecConstantOp))
1458 .addDef(ResVReg)
1459 .addUse(GR.getSPIRVTypeID(ResType))
1460 .addImm(SpecOpcode)
1461 .addUse(SrcReg)
1462 .constrainAllUses(TII, TRI, RBI);
1463 }
1464 }
1465 return selectOpWithSrcs(ResVReg, ResType, I, {I.getOperand(1).getReg()},
1466 Opcode);
1467}
1468
1469bool SPIRVInstructionSelector::selectBitcast(Register ResVReg,
1470 const SPIRVType *ResType,
1471 MachineInstr &I) const {
1472 Register OpReg = I.getOperand(1).getReg();
1473 SPIRVType *OpType = OpReg.isValid() ? GR.getSPIRVTypeForVReg(OpReg) : nullptr;
1474 if (!GR.isBitcastCompatible(ResType, OpType))
1475 report_fatal_error("incompatible result and operand types in a bitcast");
1476 return selectUnOp(ResVReg, ResType, I, SPIRV::OpBitcast);
1477}
1478
1481 MachineIRBuilder &MIRBuilder,
1482 SPIRVGlobalRegistry &GR) {
1483 uint32_t SpvMemOp = static_cast<uint32_t>(SPIRV::MemoryOperand::None);
1484 if (MemOp->isVolatile())
1485 SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1486 if (MemOp->isNonTemporal())
1487 SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1488 if (MemOp->getAlign().value())
1489 SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1490
1491 [[maybe_unused]] MachineInstr *AliasList = nullptr;
1492 [[maybe_unused]] MachineInstr *NoAliasList = nullptr;
1493 const SPIRVSubtarget *ST =
1494 static_cast<const SPIRVSubtarget *>(&MIRBuilder.getMF().getSubtarget());
1495 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1496 if (auto *MD = MemOp->getAAInfo().Scope) {
1497 AliasList = GR.getOrAddMemAliasingINTELInst(MIRBuilder, MD);
1498 if (AliasList)
1499 SpvMemOp |=
1500 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1501 }
1502 if (auto *MD = MemOp->getAAInfo().NoAlias) {
1503 NoAliasList = GR.getOrAddMemAliasingINTELInst(MIRBuilder, MD);
1504 if (NoAliasList)
1505 SpvMemOp |=
1506 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1507 }
1508 }
1509
1510 if (SpvMemOp != static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1511 MIB.addImm(SpvMemOp);
1512 if (SpvMemOp & static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1513 MIB.addImm(MemOp->getAlign().value());
1514 if (AliasList)
1515 MIB.addUse(AliasList->getOperand(0).getReg());
1516 if (NoAliasList)
1517 MIB.addUse(NoAliasList->getOperand(0).getReg());
1518 }
1519}
1520
1522 uint32_t SpvMemOp = static_cast<uint32_t>(SPIRV::MemoryOperand::None);
1524 SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1526 SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1527
1528 if (SpvMemOp != static_cast<uint32_t>(SPIRV::MemoryOperand::None))
1529 MIB.addImm(SpvMemOp);
1530}
1531
1532bool SPIRVInstructionSelector::selectLoad(Register ResVReg,
1533 const SPIRVType *ResType,
1534 MachineInstr &I) const {
1535 unsigned OpOffset = isa<GIntrinsic>(I) ? 1 : 0;
1536 Register Ptr = I.getOperand(1 + OpOffset).getReg();
1537
1538 auto *PtrDef = getVRegDef(*MRI, Ptr);
1539 auto *IntPtrDef = dyn_cast<GIntrinsic>(PtrDef);
1540 if (IntPtrDef &&
1541 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1542 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1543 SPIRVType *HandleType = GR.getSPIRVTypeForVReg(HandleReg);
1544 if (HandleType->getOpcode() == SPIRV::OpTypeImage) {
1545 Register NewHandleReg =
1546 MRI->createVirtualRegister(MRI->getRegClass(HandleReg));
1547 auto *HandleDef = cast<GIntrinsic>(getVRegDef(*MRI, HandleReg));
1548 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef, I)) {
1549 return false;
1550 }
1551
1552 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1553 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
1554 I.getDebugLoc(), I);
1555 }
1556 }
1557
1558 auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpLoad))
1559 .addDef(ResVReg)
1560 .addUse(GR.getSPIRVTypeID(ResType))
1561 .addUse(Ptr);
1562 if (!I.getNumMemOperands()) {
1563 assert(I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1564 I.getOpcode() ==
1565 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1566 addMemoryOperands(I.getOperand(2 + OpOffset).getImm(), MIB);
1567 } else {
1568 MachineIRBuilder MIRBuilder(I);
1569 addMemoryOperands(*I.memoperands_begin(), MIB, MIRBuilder, GR);
1570 }
1571 return MIB.constrainAllUses(TII, TRI, RBI);
1572}
1573
1574bool SPIRVInstructionSelector::selectStore(MachineInstr &I) const {
1575 unsigned OpOffset = isa<GIntrinsic>(I) ? 1 : 0;
1576 Register StoreVal = I.getOperand(0 + OpOffset).getReg();
1577 Register Ptr = I.getOperand(1 + OpOffset).getReg();
1578
1579 auto *PtrDef = getVRegDef(*MRI, Ptr);
1580 auto *IntPtrDef = dyn_cast<GIntrinsic>(PtrDef);
1581 if (IntPtrDef &&
1582 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1583 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1584 Register NewHandleReg =
1585 MRI->createVirtualRegister(MRI->getRegClass(HandleReg));
1586 auto *HandleDef = cast<GIntrinsic>(getVRegDef(*MRI, HandleReg));
1587 SPIRVType *HandleType = GR.getSPIRVTypeForVReg(HandleReg);
1588 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef, I)) {
1589 return false;
1590 }
1591
1592 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1593 if (HandleType->getOpcode() == SPIRV::OpTypeImage) {
1594 auto BMI = BuildMI(*I.getParent(), I, I.getDebugLoc(),
1595 TII.get(SPIRV::OpImageWrite))
1596 .addUse(NewHandleReg)
1597 .addUse(IdxReg)
1598 .addUse(StoreVal);
1599
1600 const llvm::Type *LLVMHandleType = GR.getTypeForSPIRVType(HandleType);
1601 if (sampledTypeIsSignedInteger(LLVMHandleType))
1602 BMI.addImm(0x1000); // SignExtend
1603
1604 return BMI.constrainAllUses(TII, TRI, RBI);
1605 }
1606 }
1607
1608 MachineBasicBlock &BB = *I.getParent();
1609 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpStore))
1610 .addUse(Ptr)
1611 .addUse(StoreVal);
1612 if (!I.getNumMemOperands()) {
1613 assert(I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1614 I.getOpcode() ==
1615 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1616 addMemoryOperands(I.getOperand(2 + OpOffset).getImm(), MIB);
1617 } else {
1618 MachineIRBuilder MIRBuilder(I);
1619 addMemoryOperands(*I.memoperands_begin(), MIB, MIRBuilder, GR);
1620 }
1621 return MIB.constrainAllUses(TII, TRI, RBI);
1622}
1623
1624bool SPIRVInstructionSelector::selectStackSave(Register ResVReg,
1625 const SPIRVType *ResType,
1626 MachineInstr &I) const {
1627 if (!STI.canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
1629 "llvm.stacksave intrinsic: this instruction requires the following "
1630 "SPIR-V extension: SPV_INTEL_variable_length_array",
1631 false);
1632 MachineBasicBlock &BB = *I.getParent();
1633 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSaveMemoryINTEL))
1634 .addDef(ResVReg)
1635 .addUse(GR.getSPIRVTypeID(ResType))
1636 .constrainAllUses(TII, TRI, RBI);
1637}
1638
1639bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &I) const {
1640 if (!STI.canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
1642 "llvm.stackrestore intrinsic: this instruction requires the following "
1643 "SPIR-V extension: SPV_INTEL_variable_length_array",
1644 false);
1645 if (!I.getOperand(0).isReg())
1646 return false;
1647 MachineBasicBlock &BB = *I.getParent();
1648 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpRestoreMemoryINTEL))
1649 .addUse(I.getOperand(0).getReg())
1650 .constrainAllUses(TII, TRI, RBI);
1651}
1652
1654SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &I) const {
1655 MachineIRBuilder MIRBuilder(I);
1656 assert(I.getOperand(1).isReg() && I.getOperand(2).isReg());
1657
1658 // TODO: check if we have such GV, add init, use buildGlobalVariable.
1659 unsigned Num = getIConstVal(I.getOperand(2).getReg(), MRI);
1660 Function &CurFunction = GR.CurMF->getFunction();
1661 Type *LLVMArrTy =
1662 ArrayType::get(IntegerType::get(CurFunction.getContext(), 8), Num);
1663 GlobalVariable *GV = new GlobalVariable(*CurFunction.getParent(), LLVMArrTy,
1665 Constant::getNullValue(LLVMArrTy));
1666
1667 Type *ValTy = Type::getInt8Ty(I.getMF()->getFunction().getContext());
1668 Type *ArrTy = ArrayType::get(ValTy, Num);
1670 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
1671
1672 SPIRVType *SpvArrTy = GR.getOrCreateSPIRVType(
1673 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None, false);
1674
1675 unsigned Val = getIConstVal(I.getOperand(1).getReg(), MRI);
1676 Register Const = GR.getOrCreateConstIntArray(Val, Num, I, SpvArrTy, TII);
1677
1678 Register VarReg = MRI->createGenericVirtualRegister(LLT::scalar(64));
1679 auto MIBVar =
1680 BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpVariable))
1681 .addDef(VarReg)
1682 .addUse(GR.getSPIRVTypeID(VarTy))
1683 .addImm(SPIRV::StorageClass::UniformConstant)
1684 .addUse(Const);
1685 if (!MIBVar.constrainAllUses(TII, TRI, RBI))
1686 return Register();
1687
1688 GR.add(GV, MIBVar);
1689 GR.addGlobalObject(GV, GR.CurMF, VarReg);
1690
1691 buildOpDecorate(VarReg, I, TII, SPIRV::Decoration::Constant, {});
1692 return VarReg;
1693}
1694
1695bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &I,
1696 Register SrcReg) const {
1697 MachineBasicBlock &BB = *I.getParent();
1698 Register DstReg = I.getOperand(0).getReg();
1699 SPIRVType *DstTy = GR.getSPIRVTypeForVReg(DstReg);
1700 SPIRVType *SrcTy = GR.getSPIRVTypeForVReg(SrcReg);
1701 if (GR.getPointeeType(DstTy) != GR.getPointeeType(SrcTy))
1702 report_fatal_error("OpCopyMemory requires operands to have the same type");
1703 uint64_t CopySize = getIConstVal(I.getOperand(2).getReg(), MRI);
1704 SPIRVType *PointeeTy = GR.getPointeeType(DstTy);
1705 const Type *LLVMPointeeTy = GR.getTypeForSPIRVType(PointeeTy);
1706 if (!LLVMPointeeTy)
1708 "Unable to determine pointee type size for OpCopyMemory");
1709 const DataLayout &DL = I.getMF()->getFunction().getDataLayout();
1710 if (CopySize != DL.getTypeStoreSize(const_cast<Type *>(LLVMPointeeTy)))
1712 "OpCopyMemory requires the size to match the pointee type size");
1713 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCopyMemory))
1714 .addUse(DstReg)
1715 .addUse(SrcReg);
1716 if (I.getNumMemOperands()) {
1717 MachineIRBuilder MIRBuilder(I);
1718 addMemoryOperands(*I.memoperands_begin(), MIB, MIRBuilder, GR);
1719 }
1720 return MIB.constrainAllUses(TII, TRI, RBI);
1721}
1722
1723bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &I,
1724 Register SrcReg) const {
1725 MachineBasicBlock &BB = *I.getParent();
1726 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCopyMemorySized))
1727 .addUse(I.getOperand(0).getReg())
1728 .addUse(SrcReg)
1729 .addUse(I.getOperand(2).getReg());
1730 if (I.getNumMemOperands()) {
1731 MachineIRBuilder MIRBuilder(I);
1732 addMemoryOperands(*I.memoperands_begin(), MIB, MIRBuilder, GR);
1733 }
1734 return MIB.constrainAllUses(TII, TRI, RBI);
1735}
1736
1737bool SPIRVInstructionSelector::selectMemOperation(Register ResVReg,
1738 MachineInstr &I) const {
1739 Register SrcReg = I.getOperand(1).getReg();
1740 bool Result = true;
1741 if (I.getOpcode() == TargetOpcode::G_MEMSET) {
1742 Register VarReg = getOrCreateMemSetGlobal(I);
1743 if (!VarReg.isValid())
1744 return false;
1745 Type *ValTy = Type::getInt8Ty(I.getMF()->getFunction().getContext());
1747 ValTy, I, SPIRV::StorageClass::UniformConstant);
1748 SrcReg = MRI->createGenericVirtualRegister(LLT::scalar(64));
1749 Result &= selectOpWithSrcs(SrcReg, SourceTy, I, {VarReg}, SPIRV::OpBitcast);
1750 }
1751 if (STI.isLogicalSPIRV()) {
1752 Result &= selectCopyMemory(I, SrcReg);
1753 } else {
1754 Result &= selectCopyMemorySized(I, SrcReg);
1755 }
1756 if (ResVReg.isValid() && ResVReg != I.getOperand(0).getReg())
1757 Result &= BuildCOPY(ResVReg, I.getOperand(0).getReg(), I);
1758 return Result;
1759}
1760
1761bool SPIRVInstructionSelector::selectAtomicRMW(Register ResVReg,
1762 const SPIRVType *ResType,
1763 MachineInstr &I,
1764 unsigned NewOpcode,
1765 unsigned NegateOpcode) const {
1766 bool Result = true;
1767 assert(I.hasOneMemOperand());
1768 const MachineMemOperand *MemOp = *I.memoperands_begin();
1769 uint32_t Scope = static_cast<uint32_t>(getMemScope(
1770 GR.CurMF->getFunction().getContext(), MemOp->getSyncScopeID()));
1771 auto ScopeConstant = buildI32Constant(Scope, I);
1772 Register ScopeReg = ScopeConstant.first;
1773 Result &= ScopeConstant.second;
1774
1775 Register Ptr = I.getOperand(1).getReg();
1776 // TODO: Changed as it's implemented in the translator. See test/atomicrmw.ll
1777 // auto ScSem =
1778 // getMemSemanticsForStorageClass(GR.getPointerStorageClass(Ptr));
1779 AtomicOrdering AO = MemOp->getSuccessOrdering();
1780 uint32_t MemSem = static_cast<uint32_t>(getMemSemantics(AO));
1781 auto MemSemConstant = buildI32Constant(MemSem /*| ScSem*/, I);
1782 Register MemSemReg = MemSemConstant.first;
1783 Result &= MemSemConstant.second;
1784
1785 Register ValueReg = I.getOperand(2).getReg();
1786 if (NegateOpcode != 0) {
1787 // Translation with negative value operand is requested
1788 Register TmpReg = createVirtualRegister(ResType, &GR, MRI, MRI->getMF());
1789 Result &= selectOpWithSrcs(TmpReg, ResType, I, {ValueReg}, NegateOpcode);
1790 ValueReg = TmpReg;
1791 }
1792
1793 return Result &&
1794 BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(NewOpcode))
1795 .addDef(ResVReg)
1796 .addUse(GR.getSPIRVTypeID(ResType))
1797 .addUse(Ptr)
1798 .addUse(ScopeReg)
1799 .addUse(MemSemReg)
1800 .addUse(ValueReg)
1801 .constrainAllUses(TII, TRI, RBI);
1802}
1803
1804bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &I) const {
1805 unsigned ArgI = I.getNumOperands() - 1;
1806 Register SrcReg =
1807 I.getOperand(ArgI).isReg() ? I.getOperand(ArgI).getReg() : Register(0);
1808 SPIRVType *SrcType =
1809 SrcReg.isValid() ? GR.getSPIRVTypeForVReg(SrcReg) : nullptr;
1810 if (!SrcType || SrcType->getOpcode() != SPIRV::OpTypeVector)
1812 "cannot select G_UNMERGE_VALUES with a non-vector argument");
1813
1814 SPIRVType *ScalarType =
1815 GR.getSPIRVTypeForVReg(SrcType->getOperand(1).getReg());
1816 MachineBasicBlock &BB = *I.getParent();
1817 bool Res = false;
1818 unsigned CurrentIndex = 0;
1819 for (unsigned i = 0; i < I.getNumDefs(); ++i) {
1820 Register ResVReg = I.getOperand(i).getReg();
1821 SPIRVType *ResType = GR.getSPIRVTypeForVReg(ResVReg);
1822 if (!ResType) {
1823 LLT ResLLT = MRI->getType(ResVReg);
1824 assert(ResLLT.isValid());
1825 if (ResLLT.isVector()) {
1826 ResType = GR.getOrCreateSPIRVVectorType(
1827 ScalarType, ResLLT.getNumElements(), I, TII);
1828 } else {
1829 ResType = ScalarType;
1830 }
1831 MRI->setRegClass(ResVReg, GR.getRegClass(ResType));
1832 GR.assignSPIRVTypeToVReg(ResType, ResVReg, *GR.CurMF);
1833 }
1834
1835 if (ResType->getOpcode() == SPIRV::OpTypeVector) {
1836 Register UndefReg = GR.getOrCreateUndef(I, SrcType, TII);
1837 auto MIB =
1838 BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpVectorShuffle))
1839 .addDef(ResVReg)
1840 .addUse(GR.getSPIRVTypeID(ResType))
1841 .addUse(SrcReg)
1842 .addUse(UndefReg);
1843 unsigned NumElements = GR.getScalarOrVectorComponentCount(ResType);
1844 for (unsigned j = 0; j < NumElements; ++j) {
1845 MIB.addImm(CurrentIndex + j);
1846 }
1847 CurrentIndex += NumElements;
1848 Res |= MIB.constrainAllUses(TII, TRI, RBI);
1849 } else {
1850 auto MIB =
1851 BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
1852 .addDef(ResVReg)
1853 .addUse(GR.getSPIRVTypeID(ResType))
1854 .addUse(SrcReg)
1855 .addImm(CurrentIndex);
1856 CurrentIndex++;
1857 Res |= MIB.constrainAllUses(TII, TRI, RBI);
1858 }
1859 }
1860 return Res;
1861}
1862
1863bool SPIRVInstructionSelector::selectFence(MachineInstr &I) const {
1864 AtomicOrdering AO = AtomicOrdering(I.getOperand(0).getImm());
1865 uint32_t MemSem = static_cast<uint32_t>(getMemSemantics(AO));
1866 auto MemSemConstant = buildI32Constant(MemSem, I);
1867 Register MemSemReg = MemSemConstant.first;
1868 bool Result = MemSemConstant.second;
1869 SyncScope::ID Ord = SyncScope::ID(I.getOperand(1).getImm());
1870 uint32_t Scope = static_cast<uint32_t>(
1871 getMemScope(GR.CurMF->getFunction().getContext(), Ord));
1872 auto ScopeConstant = buildI32Constant(Scope, I);
1873 Register ScopeReg = ScopeConstant.first;
1874 Result &= ScopeConstant.second;
1875 MachineBasicBlock &BB = *I.getParent();
1876 return Result &&
1877 BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpMemoryBarrier))
1878 .addUse(ScopeReg)
1879 .addUse(MemSemReg)
1880 .constrainAllUses(TII, TRI, RBI);
1881}
1882
1883bool SPIRVInstructionSelector::selectOverflowArith(Register ResVReg,
1884 const SPIRVType *ResType,
1885 MachineInstr &I,
1886 unsigned Opcode) const {
1887 Type *ResTy = nullptr;
1888 StringRef ResName;
1889 if (!GR.findValueAttrs(&I, ResTy, ResName))
1891 "Not enough info to select the arithmetic with overflow instruction");
1892 if (!ResTy || !ResTy->isStructTy())
1893 report_fatal_error("Expect struct type result for the arithmetic "
1894 "with overflow instruction");
1895 // "Result Type must be from OpTypeStruct. The struct must have two members,
1896 // and the two members must be the same type."
1897 Type *ResElemTy = cast<StructType>(ResTy)->getElementType(0);
1898 ResTy = StructType::get(ResElemTy, ResElemTy);
1899 // Build SPIR-V types and constant(s) if needed.
1900 MachineIRBuilder MIRBuilder(I);
1901 SPIRVType *StructType = GR.getOrCreateSPIRVType(
1902 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite, false);
1903 assert(I.getNumDefs() > 1 && "Not enought operands");
1904 SPIRVType *BoolType = GR.getOrCreateSPIRVBoolType(I, TII);
1905 unsigned N = GR.getScalarOrVectorComponentCount(ResType);
1906 if (N > 1)
1907 BoolType = GR.getOrCreateSPIRVVectorType(BoolType, N, I, TII);
1908 Register BoolTypeReg = GR.getSPIRVTypeID(BoolType);
1909 Register ZeroReg = buildZerosVal(ResType, I);
1910 // A new virtual register to store the result struct.
1911 Register StructVReg = MRI->createGenericVirtualRegister(LLT::scalar(64));
1912 MRI->setRegClass(StructVReg, &SPIRV::IDRegClass);
1913 // Build the result name if needed.
1914 if (ResName.size() > 0)
1915 buildOpName(StructVReg, ResName, MIRBuilder);
1916 // Build the arithmetic with overflow instruction.
1917 MachineBasicBlock &BB = *I.getParent();
1918 auto MIB =
1919 BuildMI(BB, MIRBuilder.getInsertPt(), I.getDebugLoc(), TII.get(Opcode))
1920 .addDef(StructVReg)
1921 .addUse(GR.getSPIRVTypeID(StructType));
1922 for (unsigned i = I.getNumDefs(); i < I.getNumOperands(); ++i)
1923 MIB.addUse(I.getOperand(i).getReg());
1924 bool Result = MIB.constrainAllUses(TII, TRI, RBI);
1925 // Build instructions to extract fields of the instruction's result.
1926 // A new virtual register to store the higher part of the result struct.
1927 Register HigherVReg = MRI->createGenericVirtualRegister(LLT::scalar(64));
1928 MRI->setRegClass(HigherVReg, &SPIRV::iIDRegClass);
1929 for (unsigned i = 0; i < I.getNumDefs(); ++i) {
1930 auto MIB =
1931 BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
1932 .addDef(i == 1 ? HigherVReg : I.getOperand(i).getReg())
1933 .addUse(GR.getSPIRVTypeID(ResType))
1934 .addUse(StructVReg)
1935 .addImm(i);
1936 Result &= MIB.constrainAllUses(TII, TRI, RBI);
1937 }
1938 // Build boolean value from the higher part.
1939 return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpINotEqual))
1940 .addDef(I.getOperand(1).getReg())
1941 .addUse(BoolTypeReg)
1942 .addUse(HigherVReg)
1943 .addUse(ZeroReg)
1944 .constrainAllUses(TII, TRI, RBI);
1945}
1946
1947bool SPIRVInstructionSelector::selectAtomicCmpXchg(Register ResVReg,
1948 const SPIRVType *ResType,
1949 MachineInstr &I) const {
1950 bool Result = true;
1951 Register ScopeReg;
1952 Register MemSemEqReg;
1953 Register MemSemNeqReg;
1954 Register Ptr = I.getOperand(2).getReg();
1955 if (!isa<GIntrinsic>(I)) {
1956 assert(I.hasOneMemOperand());
1957 const MachineMemOperand *MemOp = *I.memoperands_begin();
1958 unsigned Scope = static_cast<uint32_t>(getMemScope(
1959 GR.CurMF->getFunction().getContext(), MemOp->getSyncScopeID()));
1960 auto ScopeConstant = buildI32Constant(Scope, I);
1961 ScopeReg = ScopeConstant.first;
1962 Result &= ScopeConstant.second;
1963
1964 unsigned ScSem = static_cast<uint32_t>(
1966 AtomicOrdering AO = MemOp->getSuccessOrdering();
1967 unsigned MemSemEq = static_cast<uint32_t>(getMemSemantics(AO)) | ScSem;
1968 auto MemSemEqConstant = buildI32Constant(MemSemEq, I);
1969 MemSemEqReg = MemSemEqConstant.first;
1970 Result &= MemSemEqConstant.second;
1971 AtomicOrdering FO = MemOp->getFailureOrdering();
1972 unsigned MemSemNeq = static_cast<uint32_t>(getMemSemantics(FO)) | ScSem;
1973 if (MemSemEq == MemSemNeq)
1974 MemSemNeqReg = MemSemEqReg;
1975 else {
1976 auto MemSemNeqConstant = buildI32Constant(MemSemEq, I);
1977 MemSemNeqReg = MemSemNeqConstant.first;
1978 Result &= MemSemNeqConstant.second;
1979 }
1980 } else {
1981 ScopeReg = I.getOperand(5).getReg();
1982 MemSemEqReg = I.getOperand(6).getReg();
1983 MemSemNeqReg = I.getOperand(7).getReg();
1984 }
1985
1986 Register Cmp = I.getOperand(3).getReg();
1987 Register Val = I.getOperand(4).getReg();
1988 SPIRVType *SpvValTy = GR.getSPIRVTypeForVReg(Val);
1989 Register ACmpRes = createVirtualRegister(SpvValTy, &GR, MRI, *I.getMF());
1990 const DebugLoc &DL = I.getDebugLoc();
1991 Result &=
1992 BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpAtomicCompareExchange))
1993 .addDef(ACmpRes)
1994 .addUse(GR.getSPIRVTypeID(SpvValTy))
1995 .addUse(Ptr)
1996 .addUse(ScopeReg)
1997 .addUse(MemSemEqReg)
1998 .addUse(MemSemNeqReg)
1999 .addUse(Val)
2000 .addUse(Cmp)
2001 .constrainAllUses(TII, TRI, RBI);
2002 SPIRVType *BoolTy = GR.getOrCreateSPIRVBoolType(I, TII);
2003 Register CmpSuccReg = createVirtualRegister(BoolTy, &GR, MRI, *I.getMF());
2004 Result &= BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpIEqual))
2005 .addDef(CmpSuccReg)
2006 .addUse(GR.getSPIRVTypeID(BoolTy))
2007 .addUse(ACmpRes)
2008 .addUse(Cmp)
2009 .constrainAllUses(TII, TRI, RBI);
2010 Register TmpReg = createVirtualRegister(ResType, &GR, MRI, *I.getMF());
2011 Result &= BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpCompositeInsert))
2012 .addDef(TmpReg)
2013 .addUse(GR.getSPIRVTypeID(ResType))
2014 .addUse(ACmpRes)
2015 .addUse(GR.getOrCreateUndef(I, ResType, TII))
2016 .addImm(0)
2017 .constrainAllUses(TII, TRI, RBI);
2018 return Result &&
2019 BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpCompositeInsert))
2020 .addDef(ResVReg)
2021 .addUse(GR.getSPIRVTypeID(ResType))
2022 .addUse(CmpSuccReg)
2023 .addUse(TmpReg)
2024 .addImm(1)
2025 .constrainAllUses(TII, TRI, RBI);
2026}
2027
2028static bool isUSMStorageClass(SPIRV::StorageClass::StorageClass SC) {
2029 switch (SC) {
2030 case SPIRV::StorageClass::DeviceOnlyINTEL:
2031 case SPIRV::StorageClass::HostOnlyINTEL:
2032 return true;
2033 default:
2034 return false;
2035 }
2036}
2037
2038// Returns true ResVReg is referred only from global vars and OpName's.
2040 bool IsGRef = false;
2041 bool IsAllowedRefs =
2042 llvm::all_of(MRI->use_instructions(ResVReg), [&IsGRef](auto const &It) {
2043 unsigned Opcode = It.getOpcode();
2044 if (Opcode == SPIRV::OpConstantComposite ||
2045 Opcode == SPIRV::OpVariable ||
2046 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2047 return IsGRef = true;
2048 return Opcode == SPIRV::OpName;
2049 });
2050 return IsAllowedRefs && IsGRef;
2051}
2052
2053Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2054 MachineInstr &I, SPIRV::StorageClass::StorageClass SC) const {
2056 Type::getInt8Ty(I.getMF()->getFunction().getContext()), I, SC));
2057}
2058
2059MachineInstrBuilder
2060SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &I, Register Dest,
2061 Register Src, Register DestType,
2062 uint32_t Opcode) const {
2063 return BuildMI(*I.getParent(), I, I.getDebugLoc(),
2064 TII.get(SPIRV::OpSpecConstantOp))
2065 .addDef(Dest)
2066 .addUse(DestType)
2067 .addImm(Opcode)
2068 .addUse(Src);
2069}
2070
2071MachineInstrBuilder
2072SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &I, Register SrcPtr,
2073 SPIRVType *SrcPtrTy) const {
2074 SPIRVType *GenericPtrTy =
2075 GR.changePointerStorageClass(SrcPtrTy, SPIRV::StorageClass::Generic, I);
2076 Register Tmp = MRI->createVirtualRegister(&SPIRV::pIDRegClass);
2078 SPIRV::StorageClass::Generic),
2079 GR.getPointerSize()));
2080 MachineFunction *MF = I.getParent()->getParent();
2081 GR.assignSPIRVTypeToVReg(GenericPtrTy, Tmp, *MF);
2082 MachineInstrBuilder MIB = buildSpecConstantOp(
2083 I, Tmp, SrcPtr, GR.getSPIRVTypeID(GenericPtrTy),
2084 static_cast<uint32_t>(SPIRV::Opcode::PtrCastToGeneric));
2085 GR.add(MIB.getInstr(), MIB);
2086 return MIB;
2087}
2088
2089// In SPIR-V address space casting can only happen to and from the Generic
2090// storage class. We can also only cast Workgroup, CrossWorkgroup, or Function
2091// pointers to and from Generic pointers. As such, we can convert e.g. from
2092// Workgroup to Function by going via a Generic pointer as an intermediary. All
2093// other combinations can only be done by a bitcast, and are probably not safe.
2094bool SPIRVInstructionSelector::selectAddrSpaceCast(Register ResVReg,
2095 const SPIRVType *ResType,
2096 MachineInstr &I) const {
2097 MachineBasicBlock &BB = *I.getParent();
2098 const DebugLoc &DL = I.getDebugLoc();
2099
2100 Register SrcPtr = I.getOperand(1).getReg();
2101 SPIRVType *SrcPtrTy = GR.getSPIRVTypeForVReg(SrcPtr);
2102
2103 // don't generate a cast for a null that may be represented by OpTypeInt
2104 if (SrcPtrTy->getOpcode() != SPIRV::OpTypePointer ||
2105 ResType->getOpcode() != SPIRV::OpTypePointer)
2106 return BuildCOPY(ResVReg, SrcPtr, I);
2107
2108 SPIRV::StorageClass::StorageClass SrcSC = GR.getPointerStorageClass(SrcPtrTy);
2109 SPIRV::StorageClass::StorageClass DstSC = GR.getPointerStorageClass(ResType);
2110
2111 if (isASCastInGVar(MRI, ResVReg)) {
2112 // AddrSpaceCast uses within OpVariable and OpConstantComposite instructions
2113 // are expressed by OpSpecConstantOp with an Opcode.
2114 // TODO: maybe insert a check whether the Kernel capability was declared and
2115 // so PtrCastToGeneric/GenericCastToPtr are available.
2116 unsigned SpecOpcode =
2117 DstSC == SPIRV::StorageClass::Generic && isGenericCastablePtr(SrcSC)
2118 ? static_cast<uint32_t>(SPIRV::Opcode::PtrCastToGeneric)
2119 : (SrcSC == SPIRV::StorageClass::Generic &&
2121 ? static_cast<uint32_t>(SPIRV::Opcode::GenericCastToPtr)
2122 : 0);
2123 // TODO: OpConstantComposite expects i8*, so we are forced to forget a
2124 // correct value of ResType and use general i8* instead. Maybe this should
2125 // be addressed in the emit-intrinsic step to infer a correct
2126 // OpConstantComposite type.
2127 if (SpecOpcode) {
2128 return buildSpecConstantOp(I, ResVReg, SrcPtr,
2129 getUcharPtrTypeReg(I, DstSC), SpecOpcode)
2130 .constrainAllUses(TII, TRI, RBI);
2131 } else if (isGenericCastablePtr(SrcSC) && isGenericCastablePtr(DstSC)) {
2132 MachineInstrBuilder MIB = buildConstGenericPtr(I, SrcPtr, SrcPtrTy);
2133 return MIB.constrainAllUses(TII, TRI, RBI) &&
2134 buildSpecConstantOp(
2135 I, ResVReg, MIB->getOperand(0).getReg(),
2136 getUcharPtrTypeReg(I, DstSC),
2137 static_cast<uint32_t>(SPIRV::Opcode::GenericCastToPtr))
2138 .constrainAllUses(TII, TRI, RBI);
2139 }
2140 }
2141
2142 // don't generate a cast between identical storage classes
2143 if (SrcSC == DstSC)
2144 return BuildCOPY(ResVReg, SrcPtr, I);
2145
2146 if ((SrcSC == SPIRV::StorageClass::Function &&
2147 DstSC == SPIRV::StorageClass::Private) ||
2148 (DstSC == SPIRV::StorageClass::Function &&
2149 SrcSC == SPIRV::StorageClass::Private))
2150 return BuildCOPY(ResVReg, SrcPtr, I);
2151
2152 // Casting from an eligible pointer to Generic.
2153 if (DstSC == SPIRV::StorageClass::Generic && isGenericCastablePtr(SrcSC))
2154 return selectUnOp(ResVReg, ResType, I, SPIRV::OpPtrCastToGeneric);
2155 // Casting from Generic to an eligible pointer.
2156 if (SrcSC == SPIRV::StorageClass::Generic && isGenericCastablePtr(DstSC))
2157 return selectUnOp(ResVReg, ResType, I, SPIRV::OpGenericCastToPtr);
2158 // Casting between 2 eligible pointers using Generic as an intermediary.
2159 if (isGenericCastablePtr(SrcSC) && isGenericCastablePtr(DstSC)) {
2160 SPIRVType *GenericPtrTy =
2161 GR.changePointerStorageClass(SrcPtrTy, SPIRV::StorageClass::Generic, I);
2162 Register Tmp = createVirtualRegister(GenericPtrTy, &GR, MRI, MRI->getMF());
2163 bool Result = BuildMI(BB, I, DL, TII.get(SPIRV::OpPtrCastToGeneric))
2164 .addDef(Tmp)
2165 .addUse(GR.getSPIRVTypeID(GenericPtrTy))
2166 .addUse(SrcPtr)
2167 .constrainAllUses(TII, TRI, RBI);
2168 return Result && BuildMI(BB, I, DL, TII.get(SPIRV::OpGenericCastToPtr))
2169 .addDef(ResVReg)
2170 .addUse(GR.getSPIRVTypeID(ResType))
2171 .addUse(Tmp)
2172 .constrainAllUses(TII, TRI, RBI);
2173 }
2174
2175 // Check if instructions from the SPV_INTEL_usm_storage_classes extension may
2176 // be applied
2177 if (isUSMStorageClass(SrcSC) && DstSC == SPIRV::StorageClass::CrossWorkgroup)
2178 return selectUnOp(ResVReg, ResType, I,
2179 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2180 if (SrcSC == SPIRV::StorageClass::CrossWorkgroup && isUSMStorageClass(DstSC))
2181 return selectUnOp(ResVReg, ResType, I,
2182 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2183 if (isUSMStorageClass(SrcSC) && DstSC == SPIRV::StorageClass::Generic)
2184 return selectUnOp(ResVReg, ResType, I, SPIRV::OpPtrCastToGeneric);
2185 if (SrcSC == SPIRV::StorageClass::Generic && isUSMStorageClass(DstSC))
2186 return selectUnOp(ResVReg, ResType, I, SPIRV::OpGenericCastToPtr);
2187
2188 // Bitcast for pointers requires that the address spaces must match
2189 return false;
2190}
2191
2192static unsigned getFCmpOpcode(unsigned PredNum) {
2193 auto Pred = static_cast<CmpInst::Predicate>(PredNum);
2194 switch (Pred) {
2195 case CmpInst::FCMP_OEQ:
2196 return SPIRV::OpFOrdEqual;
2197 case CmpInst::FCMP_OGE:
2198 return SPIRV::OpFOrdGreaterThanEqual;
2199 case CmpInst::FCMP_OGT:
2200 return SPIRV::OpFOrdGreaterThan;
2201 case CmpInst::FCMP_OLE:
2202 return SPIRV::OpFOrdLessThanEqual;
2203 case CmpInst::FCMP_OLT:
2204 return SPIRV::OpFOrdLessThan;
2205 case CmpInst::FCMP_ONE:
2206 return SPIRV::OpFOrdNotEqual;
2207 case CmpInst::FCMP_ORD:
2208 return SPIRV::OpOrdered;
2209 case CmpInst::FCMP_UEQ:
2210 return SPIRV::OpFUnordEqual;
2211 case CmpInst::FCMP_UGE:
2212 return SPIRV::OpFUnordGreaterThanEqual;
2213 case CmpInst::FCMP_UGT:
2214 return SPIRV::OpFUnordGreaterThan;
2215 case CmpInst::FCMP_ULE:
2216 return SPIRV::OpFUnordLessThanEqual;
2217 case CmpInst::FCMP_ULT:
2218 return SPIRV::OpFUnordLessThan;
2219 case CmpInst::FCMP_UNE:
2220 return SPIRV::OpFUnordNotEqual;
2221 case CmpInst::FCMP_UNO:
2222 return SPIRV::OpUnordered;
2223 default:
2224 llvm_unreachable("Unknown predicate type for FCmp");
2225 }
2226}
2227
2228static unsigned getICmpOpcode(unsigned PredNum) {
2229 auto Pred = static_cast<CmpInst::Predicate>(PredNum);
2230 switch (Pred) {
2231 case CmpInst::ICMP_EQ:
2232 return SPIRV::OpIEqual;
2233 case CmpInst::ICMP_NE:
2234 return SPIRV::OpINotEqual;
2235 case CmpInst::ICMP_SGE:
2236 return SPIRV::OpSGreaterThanEqual;
2237 case CmpInst::ICMP_SGT:
2238 return SPIRV::OpSGreaterThan;
2239 case CmpInst::ICMP_SLE:
2240 return SPIRV::OpSLessThanEqual;
2241 case CmpInst::ICMP_SLT:
2242 return SPIRV::OpSLessThan;
2243 case CmpInst::ICMP_UGE:
2244 return SPIRV::OpUGreaterThanEqual;
2245 case CmpInst::ICMP_UGT:
2246 return SPIRV::OpUGreaterThan;
2247 case CmpInst::ICMP_ULE:
2248 return SPIRV::OpULessThanEqual;
2249 case CmpInst::ICMP_ULT:
2250 return SPIRV::OpULessThan;
2251 default:
2252 llvm_unreachable("Unknown predicate type for ICmp");
2253 }
2254}
2255
2256static unsigned getPtrCmpOpcode(unsigned Pred) {
2257 switch (static_cast<CmpInst::Predicate>(Pred)) {
2258 case CmpInst::ICMP_EQ:
2259 return SPIRV::OpPtrEqual;
2260 case CmpInst::ICMP_NE:
2261 return SPIRV::OpPtrNotEqual;
2262 default:
2263 llvm_unreachable("Unknown predicate type for pointer comparison");
2264 }
2265}
2266
2267// Return the logical operation, or abort if none exists.
2268static unsigned getBoolCmpOpcode(unsigned PredNum) {
2269 auto Pred = static_cast<CmpInst::Predicate>(PredNum);
2270 switch (Pred) {
2271 case CmpInst::ICMP_EQ:
2272 return SPIRV::OpLogicalEqual;
2273 case CmpInst::ICMP_NE:
2274 return SPIRV::OpLogicalNotEqual;
2275 default:
2276 llvm_unreachable("Unknown predicate type for Bool comparison");
2277 }
2278}
2279
2280static APFloat getZeroFP(const Type *LLVMFloatTy) {
2281 if (!LLVMFloatTy)
2283 switch (LLVMFloatTy->getScalarType()->getTypeID()) {
2284 case Type::HalfTyID:
2286 default:
2287 case Type::FloatTyID:
2289 case Type::DoubleTyID:
2291 }
2292}
2293
2294static APFloat getOneFP(const Type *LLVMFloatTy) {
2295 if (!LLVMFloatTy)
2297 switch (LLVMFloatTy->getScalarType()->getTypeID()) {
2298 case Type::HalfTyID:
2300 default:
2301 case Type::FloatTyID:
2303 case Type::DoubleTyID:
2305 }
2306}
2307
2308bool SPIRVInstructionSelector::selectAnyOrAll(Register ResVReg,
2309 const SPIRVType *ResType,
2310 MachineInstr &I,
2311 unsigned OpAnyOrAll) const {
2312 assert(I.getNumOperands() == 3);
2313 assert(I.getOperand(2).isReg());
2314 MachineBasicBlock &BB = *I.getParent();
2315 Register InputRegister = I.getOperand(2).getReg();
2316 SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);
2317
2318 if (!InputType)
2319 report_fatal_error("Input Type could not be determined.");
2320
2321 bool IsBoolTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeBool);
2322 bool IsVectorTy = InputType->getOpcode() == SPIRV::OpTypeVector;
2323 if (IsBoolTy && !IsVectorTy) {
2324 assert(ResVReg == I.getOperand(0).getReg());
2325 return BuildCOPY(ResVReg, InputRegister, I);
2326 }
2327
2328 bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);
2329 unsigned SpirvNotEqualId =
2330 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
2331 SPIRVType *SpvBoolScalarTy = GR.getOrCreateSPIRVBoolType(I, TII);
2332 SPIRVType *SpvBoolTy = SpvBoolScalarTy;
2333 Register NotEqualReg = ResVReg;
2334
2335 if (IsVectorTy) {
2336 NotEqualReg =
2337 IsBoolTy ? InputRegister
2338 : createVirtualRegister(SpvBoolTy, &GR, MRI, MRI->getMF());
2339 const unsigned NumElts = InputType->getOperand(2).getImm();
2340 SpvBoolTy = GR.getOrCreateSPIRVVectorType(SpvBoolTy, NumElts, I, TII);
2341 }
2342
2343 bool Result = true;
2344 if (!IsBoolTy) {
2345 Register ConstZeroReg =
2346 IsFloatTy ? buildZerosValF(InputType, I) : buildZerosVal(InputType, I);
2347
2348 Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SpirvNotEqualId))
2349 .addDef(NotEqualReg)
2350 .addUse(GR.getSPIRVTypeID(SpvBoolTy))
2351 .addUse(InputRegister)
2352 .addUse(ConstZeroReg)
2353 .constrainAllUses(TII, TRI, RBI);
2354 }
2355
2356 if (!IsVectorTy)
2357 return Result;
2358
2359 return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(OpAnyOrAll))
2360 .addDef(ResVReg)
2361 .addUse(GR.getSPIRVTypeID(SpvBoolScalarTy))
2362 .addUse(NotEqualReg)
2363 .constrainAllUses(TII, TRI, RBI);
2364}
2365
2366bool SPIRVInstructionSelector::selectAll(Register ResVReg,
2367 const SPIRVType *ResType,
2368 MachineInstr &I) const {
2369 return selectAnyOrAll(ResVReg, ResType, I, SPIRV::OpAll);
2370}
2371
2372bool SPIRVInstructionSelector::selectAny(Register ResVReg,
2373 const SPIRVType *ResType,
2374 MachineInstr &I) const {
2375 return selectAnyOrAll(ResVReg, ResType, I, SPIRV::OpAny);
2376}
2377
2378// Select the OpDot instruction for the given float dot
2379bool SPIRVInstructionSelector::selectFloatDot(Register ResVReg,
2380 const SPIRVType *ResType,
2381 MachineInstr &I) const {
2382 assert(I.getNumOperands() == 4);
2383 assert(I.getOperand(2).isReg());
2384 assert(I.getOperand(3).isReg());
2385
2386 [[maybe_unused]] SPIRVType *VecType =
2387 GR.getSPIRVTypeForVReg(I.getOperand(2).getReg());
2388
2389 assert(VecType->getOpcode() == SPIRV::OpTypeVector &&
2390 GR.getScalarOrVectorComponentCount(VecType) > 1 &&
2391 "dot product requires a vector of at least 2 components");
2392
2393 [[maybe_unused]] SPIRVType *EltType =
2394 GR.getSPIRVTypeForVReg(VecType->getOperand(1).getReg());
2395
2396 assert(EltType->getOpcode() == SPIRV::OpTypeFloat);
2397
2398 MachineBasicBlock &BB = *I.getParent();
2399 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpDot))
2400 .addDef(ResVReg)
2401 .addUse(GR.getSPIRVTypeID(ResType))
2402 .addUse(I.getOperand(2).getReg())
2403 .addUse(I.getOperand(3).getReg())
2404 .constrainAllUses(TII, TRI, RBI);
2405}
2406
2407bool SPIRVInstructionSelector::selectIntegerDot(Register ResVReg,
2408 const SPIRVType *ResType,
2409 MachineInstr &I,
2410 bool Signed) const {
2411 assert(I.getNumOperands() == 4);
2412 assert(I.getOperand(2).isReg());
2413 assert(I.getOperand(3).isReg());
2414 MachineBasicBlock &BB = *I.getParent();
2415
2416 auto DotOp = Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2417 return BuildMI(BB, I, I.getDebugLoc(), TII.get(DotOp))
2418 .addDef(ResVReg)
2419 .addUse(GR.getSPIRVTypeID(ResType))
2420 .addUse(I.getOperand(2).getReg())
2421 .addUse(I.getOperand(3).getReg())
2422 .constrainAllUses(TII, TRI, RBI);
2423}
2424
2425// Since pre-1.6 SPIRV has no integer dot implementation,
2426// expand by piecewise multiplying and adding the results
2427bool SPIRVInstructionSelector::selectIntegerDotExpansion(
2428 Register ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
2429 assert(I.getNumOperands() == 4);
2430 assert(I.getOperand(2).isReg());
2431 assert(I.getOperand(3).isReg());
2432 MachineBasicBlock &BB = *I.getParent();
2433
2434 // Multiply the vectors, then sum the results
2435 Register Vec0 = I.getOperand(2).getReg();
2436 Register Vec1 = I.getOperand(3).getReg();
2437 Register TmpVec = MRI->createVirtualRegister(GR.getRegClass(ResType));
2438 SPIRVType *VecType = GR.getSPIRVTypeForVReg(Vec0);
2439
2440 bool Result = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIMulV))
2441 .addDef(TmpVec)
2442 .addUse(GR.getSPIRVTypeID(VecType))
2443 .addUse(Vec0)
2444 .addUse(Vec1)
2445 .constrainAllUses(TII, TRI, RBI);
2446
2447 assert(VecType->getOpcode() == SPIRV::OpTypeVector &&
2448 GR.getScalarOrVectorComponentCount(VecType) > 1 &&
2449 "dot product requires a vector of at least 2 components");
2450
2451 Register Res = MRI->createVirtualRegister(GR.getRegClass(ResType));
2452 Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
2453 .addDef(Res)
2454 .addUse(GR.getSPIRVTypeID(ResType))
2455 .addUse(TmpVec)
2456 .addImm(0)
2457 .constrainAllUses(TII, TRI, RBI);
2458
2459 for (unsigned i = 1; i < GR.getScalarOrVectorComponentCount(VecType); i++) {
2460 Register Elt = MRI->createVirtualRegister(GR.getRegClass(ResType));
2461
2462 Result &=
2463 BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
2464 .addDef(Elt)
2465 .addUse(GR.getSPIRVTypeID(ResType))
2466 .addUse(TmpVec)
2467 .addImm(i)
2468 .constrainAllUses(TII, TRI, RBI);
2469
2470 Register Sum = i < GR.getScalarOrVectorComponentCount(VecType) - 1
2471 ? MRI->createVirtualRegister(GR.getRegClass(ResType))
2472 : ResVReg;
2473
2474 Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIAddS))
2475 .addDef(Sum)
2476 .addUse(GR.getSPIRVTypeID(ResType))
2477 .addUse(Res)
2478 .addUse(Elt)
2479 .constrainAllUses(TII, TRI, RBI);
2480 Res = Sum;
2481 }
2482
2483 return Result;
2484}
2485
2486bool SPIRVInstructionSelector::selectOpIsInf(Register ResVReg,
2487 const SPIRVType *ResType,
2488 MachineInstr &I) const {
2489 MachineBasicBlock &BB = *I.getParent();
2490 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIsInf))
2491 .addDef(ResVReg)
2492 .addUse(GR.getSPIRVTypeID(ResType))
2493 .addUse(I.getOperand(2).getReg())
2494 .constrainAllUses(TII, TRI, RBI);
2495}
2496
2497bool SPIRVInstructionSelector::selectOpIsNan(Register ResVReg,
2498 const SPIRVType *ResType,
2499 MachineInstr &I) const {
2500 MachineBasicBlock &BB = *I.getParent();
2501 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIsNan))
2502 .addDef(ResVReg)
2503 .addUse(GR.getSPIRVTypeID(ResType))
2504 .addUse(I.getOperand(2).getReg())
2505 .constrainAllUses(TII, TRI, RBI);
2506}
2507
2508template <bool Signed>
2509bool SPIRVInstructionSelector::selectDot4AddPacked(Register ResVReg,
2510 const SPIRVType *ResType,
2511 MachineInstr &I) const {
2512 assert(I.getNumOperands() == 5);
2513 assert(I.getOperand(2).isReg());
2514 assert(I.getOperand(3).isReg());
2515 assert(I.getOperand(4).isReg());
2516 MachineBasicBlock &BB = *I.getParent();
2517
2518 Register Acc = I.getOperand(2).getReg();
2519 Register X = I.getOperand(3).getReg();
2520 Register Y = I.getOperand(4).getReg();
2521
2522 auto DotOp = Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2523 Register Dot = MRI->createVirtualRegister(GR.getRegClass(ResType));
2524 bool Result = BuildMI(BB, I, I.getDebugLoc(), TII.get(DotOp))
2525 .addDef(Dot)
2526 .addUse(GR.getSPIRVTypeID(ResType))
2527 .addUse(X)
2528 .addUse(Y)
2529 .constrainAllUses(TII, TRI, RBI);
2530
2531 return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIAddS))
2532 .addDef(ResVReg)
2533 .addUse(GR.getSPIRVTypeID(ResType))
2534 .addUse(Dot)
2535 .addUse(Acc)
2536 .constrainAllUses(TII, TRI, RBI);
2537}
2538
2539// Since pre-1.6 SPIRV has no DotProductInput4x8BitPacked implementation,
2540// extract the elements of the packed inputs, multiply them and add the result
2541// to the accumulator.
2542template <bool Signed>
2543bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
2544 Register ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
2545 assert(I.getNumOperands() == 5);
2546 assert(I.getOperand(2).isReg());
2547 assert(I.getOperand(3).isReg());
2548 assert(I.getOperand(4).isReg());
2549 MachineBasicBlock &BB = *I.getParent();
2550
2551 bool Result = true;
2552
2553 Register Acc = I.getOperand(2).getReg();
2554 Register X = I.getOperand(3).getReg();
2555 Register Y = I.getOperand(4).getReg();
2556
2557 SPIRVType *EltType = GR.getOrCreateSPIRVIntegerType(8, I, TII);
2558 auto ExtractOp =
2559 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
2560
2561 bool ZeroAsNull = !STI.isShader();
2562 // Extract the i8 element, multiply and add it to the accumulator
2563 for (unsigned i = 0; i < 4; i++) {
2564 // A[i]
2565 Register AElt = MRI->createVirtualRegister(&SPIRV::IDRegClass);
2566 Result &=
2567 BuildMI(BB, I, I.getDebugLoc(), TII.get(ExtractOp))
2568 .addDef(AElt)
2569 .addUse(GR.getSPIRVTypeID(ResType))
2570 .addUse(X)
2571 .addUse(GR.getOrCreateConstInt(i * 8, I, EltType, TII, ZeroAsNull))
2572 .addUse(GR.getOrCreateConstInt(8, I, EltType, TII, ZeroAsNull))
2573 .constrainAllUses(TII, TRI, RBI);
2574
2575 // B[i]
2576 Register BElt = MRI->createVirtualRegister(&SPIRV::IDRegClass);
2577 Result &=
2578 BuildMI(BB, I, I.getDebugLoc(), TII.get(ExtractOp))
2579 .addDef(BElt)
2580 .addUse(GR.getSPIRVTypeID(ResType))
2581 .addUse(Y)
2582 .addUse(GR.getOrCreateConstInt(i * 8, I, EltType, TII, ZeroAsNull))
2583 .addUse(GR.getOrCreateConstInt(8, I, EltType, TII, ZeroAsNull))
2584 .constrainAllUses(TII, TRI, RBI);
2585
2586 // A[i] * B[i]
2587 Register Mul = MRI->createVirtualRegister(&SPIRV::IDRegClass);
2588 Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIMulS))
2589 .addDef(Mul)
2590 .addUse(GR.getSPIRVTypeID(ResType))
2591 .addUse(AElt)
2592 .addUse(BElt)
2593 .constrainAllUses(TII, TRI, RBI);
2594
2595 // Discard 24 highest-bits so that stored i32 register is i8 equivalent
2596 Register MaskMul = MRI->createVirtualRegister(&SPIRV::IDRegClass);
2597 Result &=
2598 BuildMI(BB, I, I.getDebugLoc(), TII.get(ExtractOp))
2599 .addDef(MaskMul)
2600 .addUse(GR.getSPIRVTypeID(ResType))
2601 .addUse(Mul)
2602 .addUse(GR.getOrCreateConstInt(0, I, EltType, TII, ZeroAsNull))
2603 .addUse(GR.getOrCreateConstInt(8, I, EltType, TII, ZeroAsNull))
2604 .constrainAllUses(TII, TRI, RBI);
2605
2606 // Acc = Acc + A[i] * B[i]
2607 Register Sum =
2608 i < 3 ? MRI->createVirtualRegister(&SPIRV::IDRegClass) : ResVReg;
2609 Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIAddS))
2610 .addDef(Sum)
2611 .addUse(GR.getSPIRVTypeID(ResType))
2612 .addUse(Acc)
2613 .addUse(MaskMul)
2614 .constrainAllUses(TII, TRI, RBI);
2615
2616 Acc = Sum;
2617 }
2618
2619 return Result;
2620}
2621
2622/// Transform saturate(x) to clamp(x, 0.0f, 1.0f) as SPIRV
2623/// does not have a saturate builtin.
2624bool SPIRVInstructionSelector::selectSaturate(Register ResVReg,
2625 const SPIRVType *ResType,
2626 MachineInstr &I) const {
2627 assert(I.getNumOperands() == 3);
2628 assert(I.getOperand(2).isReg());
2629 MachineBasicBlock &BB = *I.getParent();
2630 Register VZero = buildZerosValF(ResType, I);
2631 Register VOne = buildOnesValF(ResType, I);
2632
2633 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
2634 .addDef(ResVReg)
2635 .addUse(GR.getSPIRVTypeID(ResType))
2636 .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::GLSL_std_450))
2637 .addImm(GL::FClamp)
2638 .addUse(I.getOperand(2).getReg())
2639 .addUse(VZero)
2640 .addUse(VOne)
2641 .constrainAllUses(TII, TRI, RBI);
2642}
2643
2644bool SPIRVInstructionSelector::selectSign(Register ResVReg,
2645 const SPIRVType *ResType,
2646 MachineInstr &I) const {
2647 assert(I.getNumOperands() == 3);
2648 assert(I.getOperand(2).isReg());
2649 MachineBasicBlock &BB = *I.getParent();
2650 Register InputRegister = I.getOperand(2).getReg();
2651 SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);
2652 auto &DL = I.getDebugLoc();
2653
2654 if (!InputType)
2655 report_fatal_error("Input Type could not be determined.");
2656
2657 bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);
2658
2659 unsigned SignBitWidth = GR.getScalarOrVectorBitWidth(InputType);
2660 unsigned ResBitWidth = GR.getScalarOrVectorBitWidth(ResType);
2661
2662 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
2663
2664 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
2665 Register SignReg = NeedsConversion
2666 ? MRI->createVirtualRegister(&SPIRV::IDRegClass)
2667 : ResVReg;
2668
2669 bool Result =
2670 BuildMI(BB, I, DL, TII.get(SPIRV::OpExtInst))
2671 .addDef(SignReg)
2672 .addUse(GR.getSPIRVTypeID(InputType))
2673 .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::GLSL_std_450))
2674 .addImm(SignOpcode)
2675 .addUse(InputRegister)
2676 .constrainAllUses(TII, TRI, RBI);
2677
2678 if (NeedsConversion) {
2679 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
2680 Result &= BuildMI(*I.getParent(), I, DL, TII.get(ConvertOpcode))
2681 .addDef(ResVReg)
2682 .addUse(GR.getSPIRVTypeID(ResType))
2683 .addUse(SignReg)
2684 .constrainAllUses(TII, TRI, RBI);
2685 }
2686
2687 return Result;
2688}
2689
2690bool SPIRVInstructionSelector::selectWaveOpInst(Register ResVReg,
2691 const SPIRVType *ResType,
2692 MachineInstr &I,
2693 unsigned Opcode) const {
2694 MachineBasicBlock &BB = *I.getParent();
2695 SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);
2696
2697 auto BMI = BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
2698 .addDef(ResVReg)
2699 .addUse(GR.getSPIRVTypeID(ResType))
2700 .addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I,
2701 IntTy, TII, !STI.isShader()));
2702
2703 for (unsigned J = 2; J < I.getNumOperands(); J++) {
2704 BMI.addUse(I.getOperand(J).getReg());
2705 }
2706
2707 return BMI.constrainAllUses(TII, TRI, RBI);
2708}
2709
2710bool SPIRVInstructionSelector::selectWaveActiveCountBits(
2711 Register ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
2712
2713 SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);
2714 SPIRVType *BallotType = GR.getOrCreateSPIRVVectorType(IntTy, 4, I, TII);
2715 Register BallotReg = MRI->createVirtualRegister(GR.getRegClass(BallotType));
2716 bool Result = selectWaveOpInst(BallotReg, BallotType, I,
2717 SPIRV::OpGroupNonUniformBallot);
2718
2719 MachineBasicBlock &BB = *I.getParent();
2720 Result &= BuildMI(BB, I, I.getDebugLoc(),
2721 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
2722 .addDef(ResVReg)
2723 .addUse(GR.getSPIRVTypeID(ResType))
2724 .addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I, IntTy,
2725 TII, !STI.isShader()))
2726 .addImm(SPIRV::GroupOperation::Reduce)
2727 .addUse(BallotReg)
2728 .constrainAllUses(TII, TRI, RBI);
2729
2730 return Result;
2731}
2732
2733bool SPIRVInstructionSelector::selectWavePrefixBitCount(
2734 Register ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
2735
2736 assert(I.getNumOperands() == 3);
2737
2738 auto Op = I.getOperand(2);
2739 assert(Op.isReg());
2740
2741 MachineBasicBlock &BB = *I.getParent();
2742 DebugLoc DL = I.getDebugLoc();
2743
2744 Register InputRegister = Op.getReg();
2745 SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);
2746
2747 if (!InputType)
2748 report_fatal_error("Input Type could not be determined.");
2749
2750 if (InputType->getOpcode() != SPIRV::OpTypeBool)
2751 report_fatal_error("WavePrefixBitCount requires boolean input");
2752
2753 // Types
2755
2756 // Ballot result type: vector<uint32>
2757 // Match DXC: %v4uint for Subgroup size
2758 SPIRVType *BallotTy = GR.getOrCreateSPIRVVectorType(Int32Ty, 4, I, TII);
2759
2760 // Create a vreg for the ballot result
2761 Register BallotVReg = MRI->createVirtualRegister(&SPIRV::IDRegClass);
2762
2763 // 1. OpGroupNonUniformBallot
2764 BuildMI(BB, I, DL, TII.get(SPIRV::OpGroupNonUniformBallot))
2765 .addDef(BallotVReg)
2766 .addUse(GR.getSPIRVTypeID(BallotTy))
2767 .addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I, Int32Ty, TII))
2768 .addUse(InputRegister)
2769 .constrainAllUses(TII, TRI, RBI);
2770
2771 // 2. OpGroupNonUniformBallotBitCount
2772 BuildMI(BB, I, DL, TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
2773 .addDef(ResVReg)
2774 .addUse(GR.getSPIRVTypeID(ResType))
2775 .addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I, Int32Ty, TII))
2776 .addImm(SPIRV::GroupOperation::ExclusiveScan)
2777 .addUse(BallotVReg)
2778 .constrainAllUses(TII, TRI, RBI);
2779
2780 return true;
2781}
2782
2783bool SPIRVInstructionSelector::selectWaveReduceMax(Register ResVReg,
2784 const SPIRVType *ResType,
2785 MachineInstr &I,
2786 bool IsUnsigned) const {
2787 assert(I.getNumOperands() == 3);
2788 assert(I.getOperand(2).isReg());
2789 MachineBasicBlock &BB = *I.getParent();
2790 Register InputRegister = I.getOperand(2).getReg();
2791 SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);
2792
2793 if (!InputType)
2794 report_fatal_error("Input Type could not be determined.");
2795
2796 SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);
2797 // Retreive the operation to use based on input type
2798 bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);
2799 auto IntegerOpcodeType =
2800 IsUnsigned ? SPIRV::OpGroupNonUniformUMax : SPIRV::OpGroupNonUniformSMax;
2801 auto Opcode = IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntegerOpcodeType;
2802 return BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
2803 .addDef(ResVReg)
2804 .addUse(GR.getSPIRVTypeID(ResType))
2805 .addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I, IntTy, TII,
2806 !STI.isShader()))
2807 .addImm(SPIRV::GroupOperation::Reduce)
2808 .addUse(I.getOperand(2).getReg())
2809 .constrainAllUses(TII, TRI, RBI);
2810}
2811
2812bool SPIRVInstructionSelector::selectWaveReduceMin(Register ResVReg,
2813 const SPIRVType *ResType,
2814 MachineInstr &I,
2815 bool IsUnsigned) const {
2816 assert(I.getNumOperands() == 3);
2817 assert(I.getOperand(2).isReg());
2818 MachineBasicBlock &BB = *I.getParent();
2819 Register InputRegister = I.getOperand(2).getReg();
2820 SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);
2821
2822 if (!InputType)
2823 report_fatal_error("Input Type could not be determined.");
2824
2825 SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);
2826 // Retreive the operation to use based on input type
2827 bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);
2828 auto IntegerOpcodeType =
2829 IsUnsigned ? SPIRV::OpGroupNonUniformUMin : SPIRV::OpGroupNonUniformSMin;
2830 auto Opcode = IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntegerOpcodeType;
2831 return BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
2832 .addDef(ResVReg)
2833 .addUse(GR.getSPIRVTypeID(ResType))
2834 .addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I, IntTy, TII,
2835 !STI.isShader()))
2836 .addImm(SPIRV::GroupOperation::Reduce)
2837 .addUse(I.getOperand(2).getReg())
2838 .constrainAllUses(TII, TRI, RBI);
2839}
2840
2841bool SPIRVInstructionSelector::selectWaveReduceSum(Register ResVReg,
2842 const SPIRVType *ResType,
2843 MachineInstr &I) const {
2844 assert(I.getNumOperands() == 3);
2845 assert(I.getOperand(2).isReg());
2846 MachineBasicBlock &BB = *I.getParent();
2847 Register InputRegister = I.getOperand(2).getReg();
2848 SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);
2849
2850 if (!InputType)
2851 report_fatal_error("Input Type could not be determined.");
2852
2853 SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);
2854 // Retreive the operation to use based on input type
2855 bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);
2856 auto Opcode =
2857 IsFloatTy ? SPIRV::OpGroupNonUniformFAdd : SPIRV::OpGroupNonUniformIAdd;
2858 return BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
2859 .addDef(ResVReg)
2860 .addUse(GR.getSPIRVTypeID(ResType))
2861 .addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I, IntTy, TII,
2862 !STI.isShader()))
2863 .addImm(SPIRV::GroupOperation::Reduce)
2864 .addUse(I.getOperand(2).getReg());
2865}
2866
2867bool SPIRVInstructionSelector::selectBitreverse(Register ResVReg,
2868 const SPIRVType *ResType,
2869 MachineInstr &I) const {
2870 MachineBasicBlock &BB = *I.getParent();
2871 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpBitReverse))
2872 .addDef(ResVReg)
2873 .addUse(GR.getSPIRVTypeID(ResType))
2874 .addUse(I.getOperand(1).getReg())
2875 .constrainAllUses(TII, TRI, RBI);
2876}
2877
2878bool SPIRVInstructionSelector::selectFreeze(Register ResVReg,
2879 const SPIRVType *ResType,
2880 MachineInstr &I) const {
2881 // There is no way to implement `freeze` correctly without support on SPIR-V
2882 // standard side, but we may at least address a simple (static) case when
2883 // undef/poison value presence is obvious. The main benefit of even
2884 // incomplete `freeze` support is preventing of translation from crashing due
2885 // to lack of support on legalization and instruction selection steps.
2886 if (!I.getOperand(0).isReg() || !I.getOperand(1).isReg())
2887 return false;
2888 Register OpReg = I.getOperand(1).getReg();
2889 if (MachineInstr *Def = MRI->getVRegDef(OpReg)) {
2890 if (Def->getOpcode() == TargetOpcode::COPY)
2891 Def = MRI->getVRegDef(Def->getOperand(1).getReg());
2892 Register Reg;
2893 switch (Def->getOpcode()) {
2894 case SPIRV::ASSIGN_TYPE:
2895 if (MachineInstr *AssignToDef =
2896 MRI->getVRegDef(Def->getOperand(1).getReg())) {
2897 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
2898 Reg = Def->getOperand(2).getReg();
2899 }
2900 break;
2901 case SPIRV::OpUndef:
2902 Reg = Def->getOperand(1).getReg();
2903 break;
2904 }
2905 unsigned DestOpCode;
2906 if (Reg.isValid()) {
2907 DestOpCode = SPIRV::OpConstantNull;
2908 } else {
2909 DestOpCode = TargetOpcode::COPY;
2910 Reg = OpReg;
2911 }
2912 return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(DestOpCode))
2913 .addDef(I.getOperand(0).getReg())
2914 .addUse(Reg)
2915 .constrainAllUses(TII, TRI, RBI);
2916 }
2917 return false;
2918}
2919
2920bool SPIRVInstructionSelector::selectBuildVector(Register ResVReg,
2921 const SPIRVType *ResType,
2922 MachineInstr &I) const {
2923 unsigned N = 0;
2924 if (ResType->getOpcode() == SPIRV::OpTypeVector)
2925 N = GR.getScalarOrVectorComponentCount(ResType);
2926 else if (ResType->getOpcode() == SPIRV::OpTypeArray)
2927 N = getArrayComponentCount(MRI, ResType);
2928 else
2929 report_fatal_error("Cannot select G_BUILD_VECTOR with a non-vector result");
2930 if (I.getNumExplicitOperands() - I.getNumExplicitDefs() != N)
2931 report_fatal_error("G_BUILD_VECTOR and the result type are inconsistent");
2932
2933 // check if we may construct a constant vector
2934 bool IsConst = true;
2935 for (unsigned i = I.getNumExplicitDefs();
2936 i < I.getNumExplicitOperands() && IsConst; ++i)
2937 if (!isConstReg(MRI, I.getOperand(i).getReg()))
2938 IsConst = false;
2939
2940 if (!IsConst && N < 2)
2942 "There must be at least two constituent operands in a vector");
2943
2944 MRI->setRegClass(ResVReg, GR.getRegClass(ResType));
2945 auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
2946 TII.get(IsConst ? SPIRV::OpConstantComposite
2947 : SPIRV::OpCompositeConstruct))
2948 .addDef(ResVReg)
2949 .addUse(GR.getSPIRVTypeID(ResType));
2950 for (unsigned i = I.getNumExplicitDefs(); i < I.getNumExplicitOperands(); ++i)
2951 MIB.addUse(I.getOperand(i).getReg());
2952 return MIB.constrainAllUses(TII, TRI, RBI);
2953}
2954
2955bool SPIRVInstructionSelector::selectSplatVector(Register ResVReg,
2956 const SPIRVType *ResType,
2957 MachineInstr &I) const {
2958 unsigned N = 0;
2959 if (ResType->getOpcode() == SPIRV::OpTypeVector)
2960 N = GR.getScalarOrVectorComponentCount(ResType);
2961 else if (ResType->getOpcode() == SPIRV::OpTypeArray)
2962 N = getArrayComponentCount(MRI, ResType);
2963 else
2964 report_fatal_error("Cannot select G_SPLAT_VECTOR with a non-vector result");
2965
2966 unsigned OpIdx = I.getNumExplicitDefs();
2967 if (!I.getOperand(OpIdx).isReg())
2968 report_fatal_error("Unexpected argument in G_SPLAT_VECTOR");
2969
2970 // check if we may construct a constant vector
2971 Register OpReg = I.getOperand(OpIdx).getReg();
2972 bool IsConst = isConstReg(MRI, OpReg);
2973
2974 if (!IsConst && N < 2)
2976 "There must be at least two constituent operands in a vector");
2977
2978 MRI->setRegClass(ResVReg, GR.getRegClass(ResType));
2979 auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
2980 TII.get(IsConst ? SPIRV::OpConstantComposite
2981 : SPIRV::OpCompositeConstruct))
2982 .addDef(ResVReg)
2983 .addUse(GR.getSPIRVTypeID(ResType));
2984 for (unsigned i = 0; i < N; ++i)
2985 MIB.addUse(OpReg);
2986 return MIB.constrainAllUses(TII, TRI, RBI);
2987}
2988
2989bool SPIRVInstructionSelector::selectDiscard(Register ResVReg,
2990 const SPIRVType *ResType,
2991 MachineInstr &I) const {
2992
2993 unsigned Opcode;
2994
2995 if (STI.canUseExtension(
2996 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
2997 STI.isAtLeastSPIRVVer(llvm::VersionTuple(1, 6))) {
2998 Opcode = SPIRV::OpDemoteToHelperInvocation;
2999 } else {
3000 Opcode = SPIRV::OpKill;
3001 // OpKill must be the last operation of any basic block.
3002 if (MachineInstr *NextI = I.getNextNode()) {
3003 GR.invalidateMachineInstr(NextI);
3004 NextI->removeFromParent();
3005 }
3006 }
3007
3008 MachineBasicBlock &BB = *I.getParent();
3009 return BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
3010 .constrainAllUses(TII, TRI, RBI);
3011}
3012
3013bool SPIRVInstructionSelector::selectCmp(Register ResVReg,
3014 const SPIRVType *ResType,
3015 unsigned CmpOpc,
3016 MachineInstr &I) const {
3017 Register Cmp0 = I.getOperand(2).getReg();
3018 Register Cmp1 = I.getOperand(3).getReg();
3019 assert(GR.getSPIRVTypeForVReg(Cmp0)->getOpcode() ==
3020 GR.getSPIRVTypeForVReg(Cmp1)->getOpcode() &&
3021 "CMP operands should have the same type");
3022 return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(CmpOpc))
3023 .addDef(ResVReg)
3024 .addUse(GR.getSPIRVTypeID(ResType))
3025 .addUse(Cmp0)
3026 .addUse(Cmp1)
3027 .setMIFlags(I.getFlags())
3028 .constrainAllUses(TII, TRI, RBI);
3029}
3030
3031bool SPIRVInstructionSelector::selectICmp(Register ResVReg,
3032 const SPIRVType *ResType,
3033 MachineInstr &I) const {
3034 auto Pred = I.getOperand(1).getPredicate();
3035 unsigned CmpOpc;
3036
3037 Register CmpOperand = I.getOperand(2).getReg();
3038 if (GR.isScalarOfType(CmpOperand, SPIRV::OpTypePointer))
3039 CmpOpc = getPtrCmpOpcode(Pred);
3040 else if (GR.isScalarOrVectorOfType(CmpOperand, SPIRV::OpTypeBool))
3041 CmpOpc = getBoolCmpOpcode(Pred);
3042 else
3043 CmpOpc = getICmpOpcode(Pred);
3044 return selectCmp(ResVReg, ResType, CmpOpc, I);
3045}
3046
3047std::pair<Register, bool>
3048SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &I,
3049 const SPIRVType *ResType) const {
3050 Type *LLVMTy = IntegerType::get(GR.CurMF->getFunction().getContext(), 32);
3051 const SPIRVType *SpvI32Ty =
3052 ResType ? ResType : GR.getOrCreateSPIRVIntegerType(32, I, TII);
3053 // Find a constant in DT or build a new one.
3054 auto ConstInt = ConstantInt::get(LLVMTy, Val);
3055 Register NewReg = GR.find(ConstInt, GR.CurMF);
3056 bool Result = true;
3057 if (!NewReg.isValid()) {
3058 NewReg = MRI->createGenericVirtualRegister(LLT::scalar(64));
3059 MachineBasicBlock &BB = *I.getParent();
3060 MachineInstr *MI =
3061 Val == 0
3062 ? BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantNull))
3063 .addDef(NewReg)
3064 .addUse(GR.getSPIRVTypeID(SpvI32Ty))
3065 : BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantI))
3066 .addDef(NewReg)
3067 .addUse(GR.getSPIRVTypeID(SpvI32Ty))
3068 .addImm(APInt(32, Val).getZExtValue());
3070 GR.add(ConstInt, MI);
3071 }
3072 return {NewReg, Result};
3073}
3074
3075bool SPIRVInstructionSelector::selectFCmp(Register ResVReg,
3076 const SPIRVType *ResType,
3077 MachineInstr &I) const {
3078 unsigned CmpOp = getFCmpOpcode(I.getOperand(1).getPredicate());
3079 return selectCmp(ResVReg, ResType, CmpOp, I);
3080}
3081
3082Register SPIRVInstructionSelector::buildZerosVal(const SPIRVType *ResType,
3083 MachineInstr &I) const {
3084 // OpenCL uses nulls for Zero. In HLSL we don't use null constants.
3085 bool ZeroAsNull = !STI.isShader();
3086 if (ResType->getOpcode() == SPIRV::OpTypeVector)
3087 return GR.getOrCreateConstVector(0UL, I, ResType, TII, ZeroAsNull);
3088 return GR.getOrCreateConstInt(0, I, ResType, TII, ZeroAsNull);
3089}
3090
3091bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
3092 Register Reg) const {
3094 if (!Type)
3095 return false;
3097 if (!CompType || CompType->getOpcode() != SPIRV::OpTypeInt)
3098 return false;
3099
3100 auto IsZero = [this](Register Reg) {
3101 MachineInstr *Def = getDefInstrMaybeConstant(Reg, MRI);
3102 if (!Def)
3103 return false;
3104
3105 if (Def->getOpcode() == SPIRV::OpConstantNull)
3106 return true;
3107
3108 if (Def->getOpcode() == TargetOpcode::G_CONSTANT ||
3109 Def->getOpcode() == SPIRV::OpConstantI)
3110 return getIConstVal(Reg, MRI) == 0;
3111
3112 return false;
3113 };
3114
3115 if (IsZero(Reg))
3116 return true;
3117
3118 MachineInstr *Def = MRI->getVRegDef(Reg);
3119 if (!Def)
3120 return false;
3121
3122 if (Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
3123 (Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
3124 cast<GIntrinsic>(Def)->getIntrinsicID() ==
3125 Intrinsic::spv_const_composite)) {
3126 unsigned StartOp = Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
3127 for (unsigned i = StartOp; i < Def->getNumOperands(); ++i) {
3128 if (!IsZero(Def->getOperand(i).getReg()))
3129 return false;
3130 }
3131 return true;
3132 }
3133
3134 return false;
3135}
3136
3137Register SPIRVInstructionSelector::buildZerosValF(const SPIRVType *ResType,
3138 MachineInstr &I) const {
3139 // OpenCL uses nulls for Zero. In HLSL we don't use null constants.
3140 bool ZeroAsNull = !STI.isShader();
3141 APFloat VZero = getZeroFP(GR.getTypeForSPIRVType(ResType));
3142 if (ResType->getOpcode() == SPIRV::OpTypeVector)
3143 return GR.getOrCreateConstVector(VZero, I, ResType, TII, ZeroAsNull);
3144 return GR.getOrCreateConstFP(VZero, I, ResType, TII, ZeroAsNull);
3145}
3146
3147Register SPIRVInstructionSelector::buildOnesValF(const SPIRVType *ResType,
3148 MachineInstr &I) const {
3149 // OpenCL uses nulls for Zero. In HLSL we don't use null constants.
3150 bool ZeroAsNull = !STI.isShader();
3151 APFloat VOne = getOneFP(GR.getTypeForSPIRVType(ResType));
3152 if (ResType->getOpcode() == SPIRV::OpTypeVector)
3153 return GR.getOrCreateConstVector(VOne, I, ResType, TII, ZeroAsNull);
3154 return GR.getOrCreateConstFP(VOne, I, ResType, TII, ZeroAsNull);
3155}
3156
3157Register SPIRVInstructionSelector::buildOnesVal(bool AllOnes,
3158 const SPIRVType *ResType,
3159 MachineInstr &I) const {
3160 unsigned BitWidth = GR.getScalarOrVectorBitWidth(ResType);
3161 APInt One =
3162 AllOnes ? APInt::getAllOnes(BitWidth) : APInt::getOneBitSet(BitWidth, 0);
3163 if (ResType->getOpcode() == SPIRV::OpTypeVector)
3164 return GR.getOrCreateConstVector(One.getZExtValue(), I, ResType, TII);
3165 return GR.getOrCreateConstInt(One.getZExtValue(), I, ResType, TII);
3166}
3167
3168bool SPIRVInstructionSelector::selectSelect(Register ResVReg,
3169 const SPIRVType *ResType,
3170 MachineInstr &I) const {
3171 Register SelectFirstArg = I.getOperand(2).getReg();
3172 Register SelectSecondArg = I.getOperand(3).getReg();
3173 assert(ResType == GR.getSPIRVTypeForVReg(SelectFirstArg) &&
3174 ResType == GR.getSPIRVTypeForVReg(SelectSecondArg));
3175
3176 bool IsFloatTy =
3177 GR.isScalarOrVectorOfType(SelectFirstArg, SPIRV::OpTypeFloat);
3178 bool IsPtrTy =
3179 GR.isScalarOrVectorOfType(SelectFirstArg, SPIRV::OpTypePointer);
3180 bool IsVectorTy = GR.getSPIRVTypeForVReg(SelectFirstArg)->getOpcode() ==
3181 SPIRV::OpTypeVector;
3182
3183 bool IsScalarBool =
3184 GR.isScalarOfType(I.getOperand(1).getReg(), SPIRV::OpTypeBool);
3185 unsigned Opcode;
3186 if (IsVectorTy) {
3187 if (IsFloatTy) {
3188 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
3189 } else if (IsPtrTy) {
3190 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
3191 } else {
3192 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
3193 }
3194 } else {
3195 if (IsFloatTy) {
3196 Opcode = IsScalarBool ? SPIRV::OpSelectSFSCond : SPIRV::OpSelectVFVCond;
3197 } else if (IsPtrTy) {
3198 Opcode = IsScalarBool ? SPIRV::OpSelectSPSCond : SPIRV::OpSelectVPVCond;
3199 } else {
3200 Opcode = IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
3201 }
3202 }
3203 return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(Opcode))
3204 .addDef(ResVReg)
3205 .addUse(GR.getSPIRVTypeID(ResType))
3206 .addUse(I.getOperand(1).getReg())
3207 .addUse(SelectFirstArg)
3208 .addUse(SelectSecondArg)
3209 .constrainAllUses(TII, TRI, RBI);
3210}
3211
3212bool SPIRVInstructionSelector::selectSelectDefaultArgs(Register ResVReg,
3213 const SPIRVType *ResType,
3214 MachineInstr &I,
3215 bool IsSigned) const {
3216 // To extend a bool, we need to use OpSelect between constants.
3217 Register ZeroReg = buildZerosVal(ResType, I);
3218 Register OneReg = buildOnesVal(IsSigned, ResType, I);
3219 bool IsScalarBool =
3220 GR.isScalarOfType(I.getOperand(1).getReg(), SPIRV::OpTypeBool);
3221 unsigned Opcode =
3222 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
3223 return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(Opcode))
3224 .addDef(ResVReg)
3225 .addUse(GR.getSPIRVTypeID(ResType))
3226 .addUse(I.getOperand(1).getReg())
3227 .addUse(OneReg)
3228 .addUse(ZeroReg)
3229 .constrainAllUses(TII, TRI, RBI);
3230}
3231
3232bool SPIRVInstructionSelector::selectIToF(Register ResVReg,
3233 const SPIRVType *ResType,
3234 MachineInstr &I, bool IsSigned,
3235 unsigned Opcode) const {
3236 Register SrcReg = I.getOperand(1).getReg();
3237 // We can convert bool value directly to float type without OpConvert*ToF,
3238 // however the translator generates OpSelect+OpConvert*ToF, so we do the same.
3239 if (GR.isScalarOrVectorOfType(I.getOperand(1).getReg(), SPIRV::OpTypeBool)) {
3240 unsigned BitWidth = GR.getScalarOrVectorBitWidth(ResType);
3242 if (ResType->getOpcode() == SPIRV::OpTypeVector) {
3243 const unsigned NumElts = ResType->getOperand(2).getImm();
3244 TmpType = GR.getOrCreateSPIRVVectorType(TmpType, NumElts, I, TII);
3245 }
3246 SrcReg = createVirtualRegister(TmpType, &GR, MRI, MRI->getMF());
3247 selectSelectDefaultArgs(SrcReg, TmpType, I, false);
3248 }
3249 return selectOpWithSrcs(ResVReg, ResType, I, {SrcReg}, Opcode);
3250}
3251
3252bool SPIRVInstructionSelector::selectExt(Register ResVReg,
3253 const SPIRVType *ResType,
3254 MachineInstr &I, bool IsSigned) const {
3255 Register SrcReg = I.getOperand(1).getReg();
3256 if (GR.isScalarOrVectorOfType(SrcReg, SPIRV::OpTypeBool))
3257 return selectSelectDefaultArgs(ResVReg, ResType, I, IsSigned);
3258
3259 SPIRVType *SrcType = GR.getSPIRVTypeForVReg(SrcReg);
3260 if (SrcType == ResType)
3261 return BuildCOPY(ResVReg, SrcReg, I);
3262
3263 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
3264 return selectUnOp(ResVReg, ResType, I, Opcode);
3265}
3266
3267bool SPIRVInstructionSelector::selectSUCmp(Register ResVReg,
3268 const SPIRVType *ResType,
3269 MachineInstr &I,
3270 bool IsSigned) const {
3271 MachineIRBuilder MIRBuilder(I);
3272 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
3273 MachineBasicBlock &BB = *I.getParent();
3274 // Ensure we have bool.
3275 SPIRVType *BoolType = GR.getOrCreateSPIRVBoolType(I, TII);
3276 unsigned N = GR.getScalarOrVectorComponentCount(ResType);
3277 if (N > 1)
3278 BoolType = GR.getOrCreateSPIRVVectorType(BoolType, N, I, TII);
3279 Register BoolTypeReg = GR.getSPIRVTypeID(BoolType);
3280 // Build less-than-equal and less-than.
3281 // TODO: replace with one-liner createVirtualRegister() from
3282 // llvm/lib/Target/SPIRV/SPIRVUtils.cpp when PR #116609 is merged.
3283 Register IsLessEqReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
3284 MRI->setType(IsLessEqReg, LLT::scalar(64));
3285 GR.assignSPIRVTypeToVReg(ResType, IsLessEqReg, MIRBuilder.getMF());
3286 bool Result = BuildMI(BB, I, I.getDebugLoc(),
3287 TII.get(IsSigned ? SPIRV::OpSLessThanEqual
3288 : SPIRV::OpULessThanEqual))
3289 .addDef(IsLessEqReg)
3290 .addUse(BoolTypeReg)
3291 .addUse(I.getOperand(1).getReg())
3292 .addUse(I.getOperand(2).getReg())
3293 .constrainAllUses(TII, TRI, RBI);
3294 Register IsLessReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
3295 MRI->setType(IsLessReg, LLT::scalar(64));
3296 GR.assignSPIRVTypeToVReg(ResType, IsLessReg, MIRBuilder.getMF());
3297 Result &= BuildMI(BB, I, I.getDebugLoc(),
3298 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
3299 .addDef(IsLessReg)
3300 .addUse(BoolTypeReg)
3301 .addUse(I.getOperand(1).getReg())
3302 .addUse(I.getOperand(2).getReg())
3303 .constrainAllUses(TII, TRI, RBI);
3304 // Build selects.
3305 Register ResTypeReg = GR.getSPIRVTypeID(ResType);
3306 Register NegOneOrZeroReg =
3307 MRI->createVirtualRegister(GR.getRegClass(ResType));
3308 MRI->setType(NegOneOrZeroReg, LLT::scalar(64));
3309 GR.assignSPIRVTypeToVReg(ResType, NegOneOrZeroReg, MIRBuilder.getMF());
3310 unsigned SelectOpcode =
3311 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
3312 Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SelectOpcode))
3313 .addDef(NegOneOrZeroReg)
3314 .addUse(ResTypeReg)
3315 .addUse(IsLessReg)
3316 .addUse(buildOnesVal(true, ResType, I)) // -1
3317 .addUse(buildZerosVal(ResType, I))
3318 .constrainAllUses(TII, TRI, RBI);
3319 return Result & BuildMI(BB, I, I.getDebugLoc(), TII.get(SelectOpcode))
3320 .addDef(ResVReg)
3321 .addUse(ResTypeReg)
3322 .addUse(IsLessEqReg)
3323 .addUse(NegOneOrZeroReg) // -1 or 0
3324 .addUse(buildOnesVal(false, ResType, I))
3325 .constrainAllUses(TII, TRI, RBI);
3326}
3327
3328bool SPIRVInstructionSelector::selectIntToBool(Register IntReg,
3329 Register ResVReg,
3330 MachineInstr &I,
3331 const SPIRVType *IntTy,
3332 const SPIRVType *BoolTy) const {
3333 // To truncate to a bool, we use OpBitwiseAnd 1 and OpINotEqual to zero.
3334 Register BitIntReg = createVirtualRegister(IntTy, &GR, MRI, MRI->getMF());
3335 bool IsVectorTy = IntTy->getOpcode() == SPIRV::OpTypeVector;
3336 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
3337 Register Zero = buildZerosVal(IntTy, I);
3338 Register One = buildOnesVal(false, IntTy, I);
3339 MachineBasicBlock &BB = *I.getParent();
3340 bool Result = BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
3341 .addDef(BitIntReg)
3342 .addUse(GR.getSPIRVTypeID(IntTy))
3343 .addUse(IntReg)
3344 .addUse(One)
3345 .constrainAllUses(TII, TRI, RBI);
3346 return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpINotEqual))
3347 .addDef(ResVReg)
3348 .addUse(GR.getSPIRVTypeID(BoolTy))
3349 .addUse(BitIntReg)
3350 .addUse(Zero)
3351 .constrainAllUses(TII, TRI, RBI);
3352}
3353
3354bool SPIRVInstructionSelector::selectTrunc(Register ResVReg,
3355 const SPIRVType *ResType,
3356 MachineInstr &I) const {
3357 Register IntReg = I.getOperand(1).getReg();
3358 const SPIRVType *ArgType = GR.getSPIRVTypeForVReg(IntReg);
3359 if (GR.isScalarOrVectorOfType(ResVReg, SPIRV::OpTypeBool))
3360 return selectIntToBool(IntReg, ResVReg, I, ArgType, ResType);
3361 if (ArgType == ResType)
3362 return BuildCOPY(ResVReg, IntReg, I);
3363 bool IsSigned = GR.isScalarOrVectorSigned(ResType);
3364 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
3365 return selectUnOp(ResVReg, ResType, I, Opcode);
3366}
3367
3368bool SPIRVInstructionSelector::selectConst(Register ResVReg,
3369 const SPIRVType *ResType,
3370 MachineInstr &I) const {
3371 unsigned Opcode = I.getOpcode();
3372 unsigned TpOpcode = ResType->getOpcode();
3373 Register Reg;
3374 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
3375 assert(Opcode == TargetOpcode::G_CONSTANT &&
3376 I.getOperand(1).getCImm()->isZero());
3377 MachineBasicBlock &DepMBB = I.getMF()->front();
3378 MachineIRBuilder MIRBuilder(DepMBB, DepMBB.getFirstNonPHI());
3379 Reg = GR.getOrCreateConstNullPtr(MIRBuilder, ResType);
3380 } else if (Opcode == TargetOpcode::G_FCONSTANT) {
3381 Reg = GR.getOrCreateConstFP(I.getOperand(1).getFPImm()->getValue(), I,
3382 ResType, TII, !STI.isShader());
3383 } else {
3384 Reg = GR.getOrCreateConstInt(I.getOperand(1).getCImm()->getZExtValue(), I,
3385 ResType, TII, !STI.isShader());
3386 }
3387 return Reg == ResVReg ? true : BuildCOPY(ResVReg, Reg, I);
3388}
3389
3390bool SPIRVInstructionSelector::selectOpUndef(Register ResVReg,
3391 const SPIRVType *ResType,
3392 MachineInstr &I) const {
3393 return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpUndef))
3394 .addDef(ResVReg)
3395 .addUse(GR.getSPIRVTypeID(ResType))
3396 .constrainAllUses(TII, TRI, RBI);
3397}
3398
3399bool SPIRVInstructionSelector::selectInsertVal(Register ResVReg,
3400 const SPIRVType *ResType,
3401 MachineInstr &I) const {
3402 MachineBasicBlock &BB = *I.getParent();
3403 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeInsert))
3404 .addDef(ResVReg)
3405 .addUse(GR.getSPIRVTypeID(ResType))
3406 // object to insert
3407 .addUse(I.getOperand(3).getReg())
3408 // composite to insert into
3409 .addUse(I.getOperand(2).getReg());
3410 for (unsigned i = 4; i < I.getNumOperands(); i++)
3411 MIB.addImm(foldImm(I.getOperand(i), MRI));
3412 return MIB.constrainAllUses(TII, TRI, RBI);
3413}
3414
3415bool SPIRVInstructionSelector::selectExtractVal(Register ResVReg,
3416 const SPIRVType *ResType,
3417 MachineInstr &I) const {
3418 Type *MaybeResTy = nullptr;
3419 StringRef ResName;
3420 if (GR.findValueAttrs(&I, MaybeResTy, ResName) &&
3421 MaybeResTy != GR.getTypeForSPIRVType(ResType)) {
3422 assert(!MaybeResTy ||
3423 MaybeResTy->isAggregateType() &&
3424 "Expected aggregate type for extractv instruction");
3425 ResType = GR.getOrCreateSPIRVType(MaybeResTy, I,
3426 SPIRV::AccessQualifier::ReadWrite, false);
3427 GR.assignSPIRVTypeToVReg(ResType, ResVReg, *I.getMF());
3428 }
3429 MachineBasicBlock &BB = *I.getParent();
3430 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
3431 .addDef(ResVReg)
3432 .addUse(GR.getSPIRVTypeID(ResType))
3433 .addUse(I.getOperand(2).getReg());
3434 for (unsigned i = 3; i < I.getNumOperands(); i++)
3435 MIB.addImm(foldImm(I.getOperand(i), MRI));
3436 return MIB.constrainAllUses(TII, TRI, RBI);
3437}
3438
3439bool SPIRVInstructionSelector::selectInsertElt(Register ResVReg,
3440 const SPIRVType *ResType,
3441 MachineInstr &I) const {
3442 if (getImm(I.getOperand(4), MRI))
3443 return selectInsertVal(ResVReg, ResType, I);
3444 MachineBasicBlock &BB = *I.getParent();
3445 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpVectorInsertDynamic))
3446 .addDef(ResVReg)
3447 .addUse(GR.getSPIRVTypeID(ResType))
3448 .addUse(I.getOperand(2).getReg())
3449 .addUse(I.getOperand(3).getReg())
3450 .addUse(I.getOperand(4).getReg())
3451 .constrainAllUses(TII, TRI, RBI);
3452}
3453
3454bool SPIRVInstructionSelector::selectExtractElt(Register ResVReg,
3455 const SPIRVType *ResType,
3456 MachineInstr &I) const {
3457 if (getImm(I.getOperand(3), MRI))
3458 return selectExtractVal(ResVReg, ResType, I);
3459 MachineBasicBlock &BB = *I.getParent();
3460 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpVectorExtractDynamic))
3461 .addDef(ResVReg)
3462 .addUse(GR.getSPIRVTypeID(ResType))
3463 .addUse(I.getOperand(2).getReg())
3464 .addUse(I.getOperand(3).getReg())
3465 .constrainAllUses(TII, TRI, RBI);
3466}
3467
3468bool SPIRVInstructionSelector::selectGEP(Register ResVReg,
3469 const SPIRVType *ResType,
3470 MachineInstr &I) const {
3471 const bool IsGEPInBounds = I.getOperand(2).getImm();
3472
3473 // OpAccessChain could be used for OpenCL, but the SPIRV-LLVM Translator only
3474 // relies on PtrAccessChain, so we'll try not to deviate. For Vulkan however,
3475 // we have to use Op[InBounds]AccessChain.
3476 const unsigned Opcode = STI.isLogicalSPIRV()
3477 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
3478 : SPIRV::OpAccessChain)
3479 : (IsGEPInBounds ? SPIRV::OpInBoundsPtrAccessChain
3480 : SPIRV::OpPtrAccessChain);
3481
3482 auto Res = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(Opcode))
3483 .addDef(ResVReg)
3484 .addUse(GR.getSPIRVTypeID(ResType))
3485 // Object to get a pointer to.
3486 .addUse(I.getOperand(3).getReg());
3487 assert(
3488 (Opcode == SPIRV::OpPtrAccessChain ||
3489 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
3490 (getImm(I.getOperand(4), MRI) && foldImm(I.getOperand(4), MRI) == 0)) &&
3491 "Cannot translate GEP to OpAccessChain. First index must be 0.");
3492
3493 // Adding indices.
3494 const unsigned StartingIndex =
3495 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
3496 ? 5
3497 : 4;
3498 for (unsigned i = StartingIndex; i < I.getNumExplicitOperands(); ++i)
3499 Res.addUse(I.getOperand(i).getReg());
3500 return Res.constrainAllUses(TII, TRI, RBI);
3501}
3502
3503// Maybe wrap a value into OpSpecConstantOp
3504bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
3505 MachineInstr &I, SmallVector<Register> &CompositeArgs) const {
3506 bool Result = true;
3507 unsigned Lim = I.getNumExplicitOperands();
3508 for (unsigned i = I.getNumExplicitDefs() + 1; i < Lim; ++i) {
3509 Register OpReg = I.getOperand(i).getReg();
3510 MachineInstr *OpDefine = MRI->getVRegDef(OpReg);
3511 SPIRVType *OpType = GR.getSPIRVTypeForVReg(OpReg);
3512 SmallPtrSet<SPIRVType *, 4> Visited;
3513 if (!OpDefine || !OpType || isConstReg(MRI, OpDefine, Visited) ||
3514 OpDefine->getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
3515 OpDefine->getOpcode() == TargetOpcode::G_INTTOPTR ||
3516 GR.isAggregateType(OpType)) {
3517 // The case of G_ADDRSPACE_CAST inside spv_const_composite() is processed
3518 // by selectAddrSpaceCast(), and G_INTTOPTR is processed by selectUnOp()
3519 CompositeArgs.push_back(OpReg);
3520 continue;
3521 }
3522 MachineFunction *MF = I.getMF();
3523 Register WrapReg = GR.find(OpDefine, MF);
3524 if (WrapReg.isValid()) {
3525 CompositeArgs.push_back(WrapReg);
3526 continue;
3527 }
3528 // Create a new register for the wrapper
3529 WrapReg = MRI->createVirtualRegister(GR.getRegClass(OpType));
3530 CompositeArgs.push_back(WrapReg);
3531 // Decorate the wrapper register and generate a new instruction
3532 MRI->setType(WrapReg, LLT::pointer(0, 64));
3533 GR.assignSPIRVTypeToVReg(OpType, WrapReg, *MF);
3534 auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
3535 TII.get(SPIRV::OpSpecConstantOp))
3536 .addDef(WrapReg)
3537 .addUse(GR.getSPIRVTypeID(OpType))
3538 .addImm(static_cast<uint32_t>(SPIRV::Opcode::Bitcast))
3539 .addUse(OpReg);
3540 GR.add(OpDefine, MIB);
3541 Result = MIB.constrainAllUses(TII, TRI, RBI);
3542 if (!Result)
3543 break;
3544 }
3545 return Result;
3546}
3547
3548bool SPIRVInstructionSelector::selectDerivativeInst(
3549 Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
3550 const unsigned DPdOpCode) const {
3551 // TODO: This should check specifically for Fragment Execution Model, but STI
3552 // doesn't provide that information yet. See #167562
3553 errorIfInstrOutsideShader(I);
3554
3555 // If the arg/result types are half then we need to wrap the instr in
3556 // conversions to float
3557 // This case occurs because a half arg/result is legal in HLSL but not spirv.
3558 Register SrcReg = I.getOperand(2).getReg();
3559 SPIRVType *SrcType = GR.getSPIRVTypeForVReg(SrcReg);
3560 unsigned BitWidth = std::min(GR.getScalarOrVectorBitWidth(SrcType),
3561 GR.getScalarOrVectorBitWidth(ResType));
3562 if (BitWidth == 32)
3563 return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(DPdOpCode))
3564 .addDef(ResVReg)
3565 .addUse(GR.getSPIRVTypeID(ResType))
3566 .addUse(I.getOperand(2).getReg());
3567
3568 MachineIRBuilder MIRBuilder(I);
3569 unsigned componentCount = GR.getScalarOrVectorComponentCount(SrcType);
3570 SPIRVType *F32ConvertTy = GR.getOrCreateSPIRVFloatType(32, I, TII);
3571 if (componentCount != 1)
3572 F32ConvertTy = GR.getOrCreateSPIRVVectorType(F32ConvertTy, componentCount,
3573 MIRBuilder, false);
3574
3575 const TargetRegisterClass *RegClass = GR.getRegClass(SrcType);
3576 Register ConvertToVReg = MRI->createVirtualRegister(RegClass);
3577 Register DpdOpVReg = MRI->createVirtualRegister(RegClass);
3578
3579 bool Result =
3580 BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpFConvert))
3581 .addDef(ConvertToVReg)
3582 .addUse(GR.getSPIRVTypeID(F32ConvertTy))
3583 .addUse(SrcReg)
3584 .constrainAllUses(TII, TRI, RBI);
3585 Result &= BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(DPdOpCode))
3586 .addDef(DpdOpVReg)
3587 .addUse(GR.getSPIRVTypeID(F32ConvertTy))
3588 .addUse(ConvertToVReg)
3589 .constrainAllUses(TII, TRI, RBI);
3590 Result &=
3591 BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpFConvert))
3592 .addDef(ResVReg)
3593 .addUse(GR.getSPIRVTypeID(ResType))
3594 .addUse(DpdOpVReg)
3595 .constrainAllUses(TII, TRI, RBI);
3596 return Result;
3597}
3598
3599bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
3600 const SPIRVType *ResType,
3601 MachineInstr &I) const {
3602 MachineBasicBlock &BB = *I.getParent();
3603 Intrinsic::ID IID = cast<GIntrinsic>(I).getIntrinsicID();
3604 switch (IID) {
3605 case Intrinsic::spv_load:
3606 return selectLoad(ResVReg, ResType, I);
3607 case Intrinsic::spv_store:
3608 return selectStore(I);
3609 case Intrinsic::spv_extractv:
3610 return selectExtractVal(ResVReg, ResType, I);
3611 case Intrinsic::spv_insertv:
3612 return selectInsertVal(ResVReg, ResType, I);
3613 case Intrinsic::spv_extractelt:
3614 return selectExtractElt(ResVReg, ResType, I);
3615 case Intrinsic::spv_insertelt:
3616 return selectInsertElt(ResVReg, ResType, I);
3617 case Intrinsic::spv_gep:
3618 return selectGEP(ResVReg, ResType, I);
3619 case Intrinsic::spv_bitcast: {
3620 Register OpReg = I.getOperand(2).getReg();
3621 SPIRVType *OpType =
3622 OpReg.isValid() ? GR.getSPIRVTypeForVReg(OpReg) : nullptr;
3623 if (!GR.isBitcastCompatible(ResType, OpType))
3624 report_fatal_error("incompatible result and operand types in a bitcast");
3625 return selectOpWithSrcs(ResVReg, ResType, I, {OpReg}, SPIRV::OpBitcast);
3626 }
3627 case Intrinsic::spv_unref_global:
3628 case Intrinsic::spv_init_global: {
3629 MachineInstr *MI = MRI->getVRegDef(I.getOperand(1).getReg());
3630 MachineInstr *Init = I.getNumExplicitOperands() > 2
3631 ? MRI->getVRegDef(I.getOperand(2).getReg())
3632 : nullptr;
3633 assert(MI);
3634 Register GVarVReg = MI->getOperand(0).getReg();
3635 bool Res = selectGlobalValue(GVarVReg, *MI, Init);
3636 // We violate SSA form by inserting OpVariable and still having a gMIR
3637 // instruction %vreg = G_GLOBAL_VALUE @gvar. We need to fix this by erasing
3638 // the duplicated definition.
3639 if (MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
3641 MI->removeFromParent();
3642 }
3643 return Res;
3644 }
3645 case Intrinsic::spv_undef: {
3646 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpUndef))
3647 .addDef(ResVReg)
3648 .addUse(GR.getSPIRVTypeID(ResType));
3649 return MIB.constrainAllUses(TII, TRI, RBI);
3650 }
3651 case Intrinsic::spv_const_composite: {
3652 // If no values are attached, the composite is null constant.
3653 bool IsNull = I.getNumExplicitDefs() + 1 == I.getNumExplicitOperands();
3654 SmallVector<Register> CompositeArgs;
3655 MRI->setRegClass(ResVReg, GR.getRegClass(ResType));
3656
3657 // skip type MD node we already used when generated assign.type for this
3658 if (!IsNull) {
3659 if (!wrapIntoSpecConstantOp(I, CompositeArgs))
3660 return false;
3661 MachineIRBuilder MIR(I);
3662 SmallVector<MachineInstr *, 4> Instructions = createContinuedInstructions(
3663 MIR, SPIRV::OpConstantComposite, 3,
3664 SPIRV::OpConstantCompositeContinuedINTEL, CompositeArgs, ResVReg,
3665 GR.getSPIRVTypeID(ResType));
3666 for (auto *Instr : Instructions) {
3667 Instr->setDebugLoc(I.getDebugLoc());
3668 if (!constrainSelectedInstRegOperands(*Instr, TII, TRI, RBI))
3669 return false;
3670 }
3671 return true;
3672 } else {
3673 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantNull))
3674 .addDef(ResVReg)
3675 .addUse(GR.getSPIRVTypeID(ResType));
3676 return MIB.constrainAllUses(TII, TRI, RBI);
3677 }
3678 }
3679 case Intrinsic::spv_assign_name: {
3680 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpName));
3681 MIB.addUse(I.getOperand(I.getNumExplicitDefs() + 1).getReg());
3682 for (unsigned i = I.getNumExplicitDefs() + 2;
3683 i < I.getNumExplicitOperands(); ++i) {
3684 MIB.addImm(I.getOperand(i).getImm());
3685 }
3686 return MIB.constrainAllUses(TII, TRI, RBI);
3687 }
3688 case Intrinsic::spv_switch: {
3689 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSwitch));
3690 for (unsigned i = 1; i < I.getNumExplicitOperands(); ++i) {
3691 if (I.getOperand(i).isReg())
3692 MIB.addReg(I.getOperand(i).getReg());
3693 else if (I.getOperand(i).isCImm())
3694 addNumImm(I.getOperand(i).getCImm()->getValue(), MIB);
3695 else if (I.getOperand(i).isMBB())
3696 MIB.addMBB(I.getOperand(i).getMBB());
3697 else
3698 llvm_unreachable("Unexpected OpSwitch operand");
3699 }
3700 return MIB.constrainAllUses(TII, TRI, RBI);
3701 }
3702 case Intrinsic::spv_loop_merge: {
3703 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpLoopMerge));
3704 for (unsigned i = 1; i < I.getNumExplicitOperands(); ++i) {
3705 if (I.getOperand(i).isMBB())
3706 MIB.addMBB(I.getOperand(i).getMBB());
3707 else
3708 MIB.addImm(foldImm(I.getOperand(i), MRI));
3709 }
3710 return MIB.constrainAllUses(TII, TRI, RBI);
3711 }
3712 case Intrinsic::spv_selection_merge: {
3713 auto MIB =
3714 BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSelectionMerge));
3715 assert(I.getOperand(1).isMBB() &&
3716 "operand 1 to spv_selection_merge must be a basic block");
3717 MIB.addMBB(I.getOperand(1).getMBB());
3718 MIB.addImm(getSelectionOperandForImm(I.getOperand(2).getImm()));
3719 return MIB.constrainAllUses(TII, TRI, RBI);
3720 }
3721 case Intrinsic::spv_cmpxchg:
3722 return selectAtomicCmpXchg(ResVReg, ResType, I);
3723 case Intrinsic::spv_unreachable:
3724 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpUnreachable))
3725 .constrainAllUses(TII, TRI, RBI);
3726 case Intrinsic::spv_alloca:
3727 return selectFrameIndex(ResVReg, ResType, I);
3728 case Intrinsic::spv_alloca_array:
3729 return selectAllocaArray(ResVReg, ResType, I);
3730 case Intrinsic::spv_assume:
3731 if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_expect_assume))
3732 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpAssumeTrueKHR))
3733 .addUse(I.getOperand(1).getReg())
3734 .constrainAllUses(TII, TRI, RBI);
3735 break;
3736 case Intrinsic::spv_expect:
3737 if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_expect_assume))
3738 return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExpectKHR))
3739 .addDef(ResVReg)
3740 .addUse(GR.getSPIRVTypeID(ResType))
3741 .addUse(I.getOperand(2).getReg())
3742 .addUse(I.getOperand(3).getReg())
3743 .constrainAllUses(TII, TRI, RBI);
3744 break;
3745 case Intrinsic::arithmetic_fence:
3746 if (STI.canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence))
3747 return BuildMI(BB, I, I.getDebugLoc(),
3748 TII.get(SPIRV::OpArithmeticFenceEXT))
3749 .addDef(ResVReg)
3750 .addUse(GR.getSPIRVTypeID(ResType))
3751 .addUse(I.getOperand(2).getReg())
3752 .constrainAllUses(TII, TRI, RBI);
3753 else
3754 return BuildCOPY(ResVReg, I.getOperand(2).getReg(), I);
3755 break;
3756 case Intrinsic::spv_thread_id:
3757 // The HLSL SV_DispatchThreadID semantic is lowered to llvm.spv.thread.id
3758 // intrinsic in LLVM IR for SPIR-V backend.
3759 //
3760 // In SPIR-V backend, llvm.spv.thread.id is now correctly translated to a
3761 // `GlobalInvocationId` builtin variable
3762 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
3763 ResType, I);
3764 case Intrinsic::spv_thread_id_in_group:
3765 // The HLSL SV_GroupThreadId semantic is lowered to
3766 // llvm.spv.thread.id.in.group intrinsic in LLVM IR for SPIR-V backend.
3767 //
3768 // In SPIR-V backend, llvm.spv.thread.id.in.group is now correctly
3769 // translated to a `LocalInvocationId` builtin variable
3770 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
3771 ResType, I);
3772 case Intrinsic::spv_group_id:
3773 // The HLSL SV_GroupId semantic is lowered to
3774 // llvm.spv.group.id intrinsic in LLVM IR for SPIR-V backend.
3775 //
3776 // In SPIR-V backend, llvm.spv.group.id is now translated to a `WorkgroupId`
3777 // builtin variable
3778 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
3779 I);
3780 case Intrinsic::spv_flattened_thread_id_in_group:
3781 // The HLSL SV_GroupIndex semantic is lowered to
3782 // llvm.spv.flattened.thread.id.in.group() intrinsic in LLVM IR for SPIR-V
3783 // backend.
3784 //
3785 // In SPIR-V backend, llvm.spv.flattened.thread.id.in.group is translated to
3786 // a `LocalInvocationIndex` builtin variable
3787 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
3788 ResType, I);
3789 case Intrinsic::spv_workgroup_size:
3790 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
3791 ResType, I);
3792 case Intrinsic::spv_global_size:
3793 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
3794 I);
3795 case Intrinsic::spv_global_offset:
3796 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
3797 ResType, I);
3798 case Intrinsic::spv_num_workgroups:
3799 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
3800 ResType, I);
3801 case Intrinsic::spv_subgroup_size:
3802 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
3803 I);
3804 case Intrinsic::spv_num_subgroups:
3805 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
3806 I);
3807 case Intrinsic::spv_subgroup_id:
3808 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType, I);
3809 case Intrinsic::spv_subgroup_local_invocation_id:
3810 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
3811 ResVReg, ResType, I);
3812 case Intrinsic::spv_subgroup_max_size:
3813 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
3814 I);
3815 case Intrinsic::spv_fdot:
3816 return selectFloatDot(ResVReg, ResType, I);
3817 case Intrinsic::spv_udot:
3818 case Intrinsic::spv_sdot:
3819 if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
3820 STI.isAtLeastSPIRVVer(VersionTuple(1, 6)))
3821 return selectIntegerDot(ResVReg, ResType, I,
3822 /*Signed=*/IID == Intrinsic::spv_sdot);
3823 return selectIntegerDotExpansion(ResVReg, ResType, I);
3824 case Intrinsic::spv_dot4add_i8packed:
3825 if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
3826 STI.isAtLeastSPIRVVer(VersionTuple(1, 6)))
3827 return selectDot4AddPacked<true>(ResVReg, ResType, I);
3828 return selectDot4AddPackedExpansion<true>(ResVReg, ResType, I);
3829 case Intrinsic::spv_dot4add_u8packed:
3830 if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
3831 STI.isAtLeastSPIRVVer(VersionTuple(1, 6)))
3832 return selectDot4AddPacked<false>(ResVReg, ResType, I);
3833 return selectDot4AddPackedExpansion<false>(ResVReg, ResType, I);
3834 case Intrinsic::spv_all:
3835 return selectAll(ResVReg, ResType, I);
3836 case Intrinsic::spv_any:
3837 return selectAny(ResVReg, ResType, I);
3838 case Intrinsic::spv_cross:
3839 return selectExtInst(ResVReg, ResType, I, CL::cross, GL::Cross);
3840 case Intrinsic::spv_distance:
3841 return selectExtInst(ResVReg, ResType, I, CL::distance, GL::Distance);
3842 case Intrinsic::spv_lerp:
3843 return selectExtInst(ResVReg, ResType, I, CL::mix, GL::FMix);
3844 case Intrinsic::spv_length:
3845 return selectExtInst(ResVReg, ResType, I, CL::length, GL::Length);
3846 case Intrinsic::spv_degrees:
3847 return selectExtInst(ResVReg, ResType, I, CL::degrees, GL::Degrees);
3848 case Intrinsic::spv_faceforward:
3849 return selectExtInst(ResVReg, ResType, I, GL::FaceForward);
3850 case Intrinsic::spv_frac:
3851 return selectExtInst(ResVReg, ResType, I, CL::fract, GL::Fract);
3852 case Intrinsic::spv_isinf:
3853 return selectOpIsInf(ResVReg, ResType, I);
3854 case Intrinsic::spv_isnan:
3855 return selectOpIsNan(ResVReg, ResType, I);
3856 case Intrinsic::spv_normalize:
3857 return selectExtInst(ResVReg, ResType, I, CL::normalize, GL::Normalize);
3858 case Intrinsic::spv_refract:
3859 return selectExtInst(ResVReg, ResType, I, GL::Refract);
3860 case Intrinsic::spv_reflect:
3861 return selectExtInst(ResVReg, ResType, I, GL::Reflect);
3862 case Intrinsic::spv_rsqrt:
3863 return selectExtInst(ResVReg, ResType, I, CL::rsqrt, GL::InverseSqrt);
3864 case Intrinsic::spv_sign:
3865 return selectSign(ResVReg, ResType, I);
3866 case Intrinsic::spv_smoothstep:
3867 return selectExtInst(ResVReg, ResType, I, CL::smoothstep, GL::SmoothStep);
3868 case Intrinsic::spv_firstbituhigh: // There is no CL equivalent of FindUMsb
3869 return selectFirstBitHigh(ResVReg, ResType, I, /*IsSigned=*/false);
3870 case Intrinsic::spv_firstbitshigh: // There is no CL equivalent of FindSMsb
3871 return selectFirstBitHigh(ResVReg, ResType, I, /*IsSigned=*/true);
3872 case Intrinsic::spv_firstbitlow: // There is no CL equivlent of FindILsb
3873 return selectFirstBitLow(ResVReg, ResType, I);
3874 case Intrinsic::spv_group_memory_barrier_with_group_sync: {
3875 bool Result = true;
3876 auto MemSemConstant =
3877 buildI32Constant(SPIRV::MemorySemantics::SequentiallyConsistent, I);
3878 Register MemSemReg = MemSemConstant.first;
3879 Result &= MemSemConstant.second;
3880 auto ScopeConstant = buildI32Constant(SPIRV::Scope::Workgroup, I);
3881 Register ScopeReg = ScopeConstant.first;
3882 Result &= ScopeConstant.second;
3883 MachineBasicBlock &BB = *I.getParent();
3884 return Result &&
3885 BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpControlBarrier))
3886 .addUse(ScopeReg)
3887 .addUse(ScopeReg)
3888 .addUse(MemSemReg)
3889 .constrainAllUses(TII, TRI, RBI);
3890 }
3891 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
3892 Register PtrReg = I.getOperand(I.getNumExplicitDefs() + 1).getReg();
3893 SPIRV::StorageClass::StorageClass ResSC =
3894 GR.getPointerStorageClass(ResType);
3895 if (!isGenericCastablePtr(ResSC))
3896 report_fatal_error("The target storage class is not castable from the "
3897 "Generic storage class");
3898 return BuildMI(BB, I, I.getDebugLoc(),
3899 TII.get(SPIRV::OpGenericCastToPtrExplicit))
3900 .addDef(ResVReg)
3901 .addUse(GR.getSPIRVTypeID(ResType))
3902 .addUse(PtrReg)
3903 .addImm(ResSC)
3904 .constrainAllUses(TII, TRI, RBI);
3905 }
3906 case Intrinsic::spv_lifetime_start:
3907 case Intrinsic::spv_lifetime_end: {
3908 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
3909 : SPIRV::OpLifetimeStop;
3910 int64_t Size = I.getOperand(I.getNumExplicitDefs() + 1).getImm();
3911 Register PtrReg = I.getOperand(I.getNumExplicitDefs() + 2).getReg();
3912 if (Size == -1)
3913 Size = 0;
3914 return BuildMI(BB, I, I.getDebugLoc(), TII.get(Op))
3915 .addUse(PtrReg)
3916 .addImm(Size)
3917 .constrainAllUses(TII, TRI, RBI);
3918 }
3919 case Intrinsic::spv_saturate:
3920 return selectSaturate(ResVReg, ResType, I);
3921 case Intrinsic::spv_nclamp:
3922 return selectExtInst(ResVReg, ResType, I, CL::fclamp, GL::NClamp);
3923 case Intrinsic::spv_uclamp:
3924 return selectExtInst(ResVReg, ResType, I, CL::u_clamp, GL::UClamp);
3925 case Intrinsic::spv_sclamp:
3926 return selectExtInst(ResVReg, ResType, I, CL::s_clamp, GL::SClamp);
3927 case Intrinsic::spv_subgroup_prefix_bit_count:
3928 return selectWavePrefixBitCount(ResVReg, ResType, I);
3929 case Intrinsic::spv_wave_active_countbits:
3930 return selectWaveActiveCountBits(ResVReg, ResType, I);
3931 case Intrinsic::spv_wave_all:
3932 return selectWaveOpInst(ResVReg, ResType, I, SPIRV::OpGroupNonUniformAll);
3933 case Intrinsic::spv_wave_any:
3934 return selectWaveOpInst(ResVReg, ResType, I, SPIRV::OpGroupNonUniformAny);
3935 case Intrinsic::spv_subgroup_ballot:
3936 return selectWaveOpInst(ResVReg, ResType, I,
3937 SPIRV::OpGroupNonUniformBallot);
3938 case Intrinsic::spv_wave_is_first_lane:
3939 return selectWaveOpInst(ResVReg, ResType, I, SPIRV::OpGroupNonUniformElect);
3940 case Intrinsic::spv_wave_reduce_umax:
3941 return selectWaveReduceMax(ResVReg, ResType, I, /*IsUnsigned*/ true);
3942 case Intrinsic::spv_wave_reduce_max:
3943 return selectWaveReduceMax(ResVReg, ResType, I, /*IsUnsigned*/ false);
3944 case Intrinsic::spv_wave_reduce_umin:
3945 return selectWaveReduceMin(ResVReg, ResType, I, /*IsUnsigned*/ true);
3946 case Intrinsic::spv_wave_reduce_min:
3947 return selectWaveReduceMin(ResVReg, ResType, I, /*IsUnsigned*/ false);
3948 case Intrinsic::spv_wave_reduce_sum:
3949 return selectWaveReduceSum(ResVReg, ResType, I);
3950 case Intrinsic::spv_wave_readlane:
3951 return selectWaveOpInst(ResVReg, ResType, I,
3952 SPIRV::OpGroupNonUniformShuffle);
3953 case Intrinsic::spv_step:
3954 return selectExtInst(ResVReg, ResType, I, CL::step, GL::Step);
3955 case Intrinsic::spv_radians:
3956 return selectExtInst(ResVReg, ResType, I, CL::radians, GL::Radians);
3957 // Discard intrinsics which we do not expect to actually represent code after
3958 // lowering or intrinsics which are not implemented but should not crash when
3959 // found in a customer's LLVM IR input.
3960 case Intrinsic::instrprof_increment:
3961 case Intrinsic::instrprof_increment_step:
3962 case Intrinsic::instrprof_value_profile:
3963 break;
3964 // Discard internal intrinsics.
3965 case Intrinsic::spv_value_md:
3966 break;
3967 case Intrinsic::spv_resource_handlefrombinding: {
3968 return selectHandleFromBinding(ResVReg, ResType, I);
3969 }
3970 case Intrinsic::spv_resource_counterhandlefrombinding:
3971 return selectCounterHandleFromBinding(ResVReg, ResType, I);
3972 case Intrinsic::spv_resource_updatecounter:
3973 return selectUpdateCounter(ResVReg, ResType, I);
3974 case Intrinsic::spv_resource_store_typedbuffer: {
3975 return selectImageWriteIntrinsic(I);
3976 }
3977 case Intrinsic::spv_resource_load_typedbuffer: {
3978 return selectReadImageIntrinsic(ResVReg, ResType, I);
3979 }
3980 case Intrinsic::spv_resource_sample:
3981 case Intrinsic::spv_resource_sample_clamp: {
3982 return selectSampleIntrinsic(ResVReg, ResType, I);
3983 }
3984 case Intrinsic::spv_resource_getpointer: {
3985 return selectResourceGetPointer(ResVReg, ResType, I);
3986 }
3987 case Intrinsic::spv_pushconstant_getpointer: {
3988 return selectPushConstantGetPointer(ResVReg, ResType, I);
3989 }
3990 case Intrinsic::spv_discard: {
3991 return selectDiscard(ResVReg, ResType, I);
3992 }
3993 case Intrinsic::spv_resource_nonuniformindex: {
3994 return selectResourceNonUniformIndex(ResVReg, ResType, I);
3995 }
3996 case Intrinsic::spv_unpackhalf2x16: {
3997 return selectExtInst(ResVReg, ResType, I, GL::UnpackHalf2x16);
3998 }
3999 case Intrinsic::spv_packhalf2x16: {
4000 return selectExtInst(ResVReg, ResType, I, GL::PackHalf2x16);
4001 }
4002 case Intrinsic::spv_ddx:
4003 return selectDerivativeInst(ResVReg, ResType, I, SPIRV::OpDPdx);
4004 case Intrinsic::spv_ddy:
4005 return selectDerivativeInst(ResVReg, ResType, I, SPIRV::OpDPdy);
4006 case Intrinsic::spv_ddx_coarse:
4007 return selectDerivativeInst(ResVReg, ResType, I, SPIRV::OpDPdxCoarse);
4008 case Intrinsic::spv_ddy_coarse:
4009 return selectDerivativeInst(ResVReg, ResType, I, SPIRV::OpDPdyCoarse);
4010 case Intrinsic::spv_ddx_fine:
4011 return selectDerivativeInst(ResVReg, ResType, I, SPIRV::OpDPdxFine);
4012 case Intrinsic::spv_ddy_fine:
4013 return selectDerivativeInst(ResVReg, ResType, I, SPIRV::OpDPdyFine);
4014 case Intrinsic::spv_fwidth:
4015 return selectDerivativeInst(ResVReg, ResType, I, SPIRV::OpFwidth);
4016 default: {
4017 std::string DiagMsg;
4018 raw_string_ostream OS(DiagMsg);
4019 I.print(OS);
4020 DiagMsg = "Intrinsic selection not implemented: " + DiagMsg;
4021 report_fatal_error(DiagMsg.c_str(), false);
4022 }
4023 }
4024 return true;
4025}
4026
4027bool SPIRVInstructionSelector::selectHandleFromBinding(Register &ResVReg,
4028 const SPIRVType *ResType,
4029 MachineInstr &I) const {
4030 // The images need to be loaded in the same basic block as their use. We defer
4031 // loading the image to the intrinsic that uses it.
4032 if (ResType->getOpcode() == SPIRV::OpTypeImage)
4033 return true;
4034
4035 return loadHandleBeforePosition(ResVReg, GR.getSPIRVTypeForVReg(ResVReg),
4036 *cast<GIntrinsic>(&I), I);
4037}
4038
4039bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
4040 Register &ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
4041 auto &Intr = cast<GIntrinsic>(I);
4042 assert(Intr.getIntrinsicID() ==
4043 Intrinsic::spv_resource_counterhandlefrombinding);
4044
4045 // Extract information from the intrinsic call.
4046 Register MainHandleReg = Intr.getOperand(2).getReg();
4047 auto *MainHandleDef = cast<GIntrinsic>(getVRegDef(*MRI, MainHandleReg));
4048 assert(MainHandleDef->getIntrinsicID() ==
4049 Intrinsic::spv_resource_handlefrombinding);
4050
4051 uint32_t Set = getIConstVal(Intr.getOperand(4).getReg(), MRI);
4052 uint32_t Binding = getIConstVal(Intr.getOperand(3).getReg(), MRI);
4053 uint32_t ArraySize = getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
4054 Register IndexReg = MainHandleDef->getOperand(5).getReg();
4055 std::string CounterName =
4056 getStringValueFromReg(MainHandleDef->getOperand(6).getReg(), *MRI) +
4057 ".counter";
4058
4059 // Create the counter variable.
4060 MachineIRBuilder MIRBuilder(I);
4061 Register CounterVarReg = buildPointerToResource(
4062 GR.getPointeeType(ResType), GR.getPointerStorageClass(ResType), Set,
4063 Binding, ArraySize, IndexReg, CounterName, MIRBuilder);
4064
4065 return BuildCOPY(ResVReg, CounterVarReg, I);
4066}
4067
4068bool SPIRVInstructionSelector::selectUpdateCounter(Register &ResVReg,
4069 const SPIRVType *ResType,
4070 MachineInstr &I) const {
4071 auto &Intr = cast<GIntrinsic>(I);
4072 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
4073
4074 Register CounterHandleReg = Intr.getOperand(2).getReg();
4075 Register IncrReg = Intr.getOperand(3).getReg();
4076
4077 // The counter handle is a pointer to the counter variable (which is a struct
4078 // containing an i32). We need to get a pointer to that i32 member to do the
4079 // atomic operation.
4080#ifndef NDEBUG
4081 SPIRVType *CounterVarType = GR.getSPIRVTypeForVReg(CounterHandleReg);
4082 SPIRVType *CounterVarPointeeType = GR.getPointeeType(CounterVarType);
4083 assert(CounterVarPointeeType &&
4084 CounterVarPointeeType->getOpcode() == SPIRV::OpTypeStruct &&
4085 "Counter variable must be a struct");
4086 assert(GR.getPointerStorageClass(CounterVarType) ==
4087 SPIRV::StorageClass::StorageBuffer &&
4088 "Counter variable must be in the storage buffer storage class");
4089 assert(CounterVarPointeeType->getNumOperands() == 2 &&
4090 "Counter variable must have exactly 1 member in the struct");
4091 const SPIRVType *MemberType =
4092 GR.getSPIRVTypeForVReg(CounterVarPointeeType->getOperand(1).getReg());
4093 assert(MemberType->getOpcode() == SPIRV::OpTypeInt &&
4094 "Counter variable struct must have a single i32 member");
4095#endif
4096
4097 // The struct has a single i32 member.
4098 MachineIRBuilder MIRBuilder(I);
4099 const Type *LLVMIntType =
4100 Type::getInt32Ty(I.getMF()->getFunction().getContext());
4101
4102 SPIRVType *IntPtrType = GR.getOrCreateSPIRVPointerType(
4103 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
4104
4105 auto Zero = buildI32Constant(0, I);
4106 if (!Zero.second)
4107 return false;
4108
4109 Register PtrToCounter =
4110 MRI->createVirtualRegister(GR.getRegClass(IntPtrType));
4111 if (!BuildMI(*I.getParent(), I, I.getDebugLoc(),
4112 TII.get(SPIRV::OpAccessChain))
4113 .addDef(PtrToCounter)
4114 .addUse(GR.getSPIRVTypeID(IntPtrType))
4115 .addUse(CounterHandleReg)
4116 .addUse(Zero.first)
4117 .constrainAllUses(TII, TRI, RBI)) {
4118 return false;
4119 }
4120
4121 // For UAV/SSBO counters, the scope is Device. The counter variable is not
4122 // used as a flag. So the memory semantics can be None.
4123 auto Scope = buildI32Constant(SPIRV::Scope::Device, I);
4124 if (!Scope.second)
4125 return false;
4126 auto Semantics = buildI32Constant(SPIRV::MemorySemantics::None, I);
4127 if (!Semantics.second)
4128 return false;
4129
4130 int64_t IncrVal = getIConstValSext(IncrReg, MRI);
4131 auto Incr = buildI32Constant(static_cast<uint32_t>(IncrVal), I);
4132 if (!Incr.second)
4133 return false;
4134
4135 Register AtomicRes = MRI->createVirtualRegister(GR.getRegClass(ResType));
4136 if (!BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpAtomicIAdd))
4137 .addDef(AtomicRes)
4138 .addUse(GR.getSPIRVTypeID(ResType))
4139 .addUse(PtrToCounter)
4140 .addUse(Scope.first)
4141 .addUse(Semantics.first)
4142 .addUse(Incr.first)
4143 .constrainAllUses(TII, TRI, RBI)) {
4144 return false;
4145 }
4146 if (IncrVal >= 0) {
4147 return BuildCOPY(ResVReg, AtomicRes, I);
4148 }
4149
4150 // In HLSL, IncrementCounter returns the value *before* the increment, while
4151 // DecrementCounter returns the value *after* the decrement. Both are lowered
4152 // to the same atomic intrinsic which returns the value *before* the
4153 // operation. So for decrements (negative IncrVal), we must subtract the
4154 // increment value from the result to get the post-decrement value.
4155 return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpIAddS))
4156 .addDef(ResVReg)
4157 .addUse(GR.getSPIRVTypeID(ResType))
4158 .addUse(AtomicRes)
4159 .addUse(Incr.first)
4160 .constrainAllUses(TII, TRI, RBI);
4161}
4162bool SPIRVInstructionSelector::selectReadImageIntrinsic(
4163 Register &ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
4164
4165 // If the load of the image is in a different basic block, then
4166 // this will generate invalid code. A proper solution is to move
4167 // the OpLoad from selectHandleFromBinding here. However, to do
4168 // that we will need to change the return type of the intrinsic.
4169 // We will do that when we can, but for now trying to move forward with other
4170 // issues.
4171 Register ImageReg = I.getOperand(2).getReg();
4172 auto *ImageDef = cast<GIntrinsic>(getVRegDef(*MRI, ImageReg));
4173 Register NewImageReg = MRI->createVirtualRegister(MRI->getRegClass(ImageReg));
4174 if (!loadHandleBeforePosition(NewImageReg, GR.getSPIRVTypeForVReg(ImageReg),
4175 *ImageDef, I)) {
4176 return false;
4177 }
4178
4179 Register IdxReg = I.getOperand(3).getReg();
4180 DebugLoc Loc = I.getDebugLoc();
4181 MachineInstr &Pos = I;
4182
4183 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
4184 Pos);
4185}
4186
4187bool SPIRVInstructionSelector::selectSampleIntrinsic(Register &ResVReg,
4188 const SPIRVType *ResType,
4189 MachineInstr &I) const {
4190 Register ImageReg = I.getOperand(2).getReg();
4191 Register SamplerReg = I.getOperand(3).getReg();
4192 Register CoordinateReg = I.getOperand(4).getReg();
4193 std::optional<Register> OffsetReg;
4194 std::optional<Register> ClampReg;
4195
4196 if (I.getNumOperands() > 5)
4197 OffsetReg = I.getOperand(5).getReg();
4198 if (I.getNumOperands() > 6)
4199 ClampReg = I.getOperand(6).getReg();
4200
4201 DebugLoc Loc = I.getDebugLoc();
4202
4203 auto *ImageDef = cast<GIntrinsic>(getVRegDef(*MRI, ImageReg));
4204 Register NewImageReg = MRI->createVirtualRegister(MRI->getRegClass(ImageReg));
4205 if (!loadHandleBeforePosition(NewImageReg, GR.getSPIRVTypeForVReg(ImageReg),
4206 *ImageDef, I)) {
4207 return false;
4208 }
4209
4210 auto *SamplerDef = cast<GIntrinsic>(getVRegDef(*MRI, SamplerReg));
4211 Register NewSamplerReg =
4212 MRI->createVirtualRegister(MRI->getRegClass(SamplerReg));
4213 if (!loadHandleBeforePosition(
4214 NewSamplerReg, GR.getSPIRVTypeForVReg(SamplerReg), *SamplerDef, I)) {
4215 return false;
4216 }
4217
4218 MachineIRBuilder MIRBuilder(I);
4219 SPIRVType *SampledImageType = GR.getOrCreateOpTypeSampledImage(
4220 GR.getSPIRVTypeForVReg(ImageReg), MIRBuilder);
4221
4222 Register SampledImageReg =
4223 MRI->createVirtualRegister(GR.getRegClass(SampledImageType));
4224 bool Succeed = BuildMI(*I.getParent(), I, Loc, TII.get(SPIRV::OpSampledImage))
4225 .addDef(SampledImageReg)
4226 .addUse(GR.getSPIRVTypeID(SampledImageType))
4227 .addUse(NewImageReg)
4228 .addUse(NewSamplerReg)
4229 .constrainAllUses(TII, TRI, RBI);
4230 if (!Succeed)
4231 return false;
4232
4233 auto MIB =
4234 BuildMI(*I.getParent(), I, Loc, TII.get(SPIRV::OpImageSampleImplicitLod))
4235 .addDef(ResVReg)
4236 .addUse(GR.getSPIRVTypeID(ResType))
4237 .addUse(SampledImageReg)
4238 .addUse(CoordinateReg);
4239
4240 uint32_t ImageOperands = 0;
4241 if (OffsetReg && !isScalarOrVectorIntConstantZero(*OffsetReg)) {
4242 ImageOperands |= 0x8; // ConstOffset
4243 }
4244
4245 if (ClampReg) {
4246 ImageOperands |= 0x80; // MinLod
4247 }
4248
4249 if (ImageOperands != 0) {
4250 MIB.addImm(ImageOperands);
4251 if (ImageOperands & 0x8)
4252 MIB.addUse(*OffsetReg);
4253 if (ImageOperands & 0x80)
4254 MIB.addUse(*ClampReg);
4255 }
4256
4257 return MIB.constrainAllUses(TII, TRI, RBI);
4258}
4259
4260bool SPIRVInstructionSelector::generateImageReadOrFetch(
4261 Register &ResVReg, const SPIRVType *ResType, Register ImageReg,
4262 Register IdxReg, DebugLoc Loc, MachineInstr &Pos) const {
4263 SPIRVType *ImageType = GR.getSPIRVTypeForVReg(ImageReg);
4264 assert(ImageType && ImageType->getOpcode() == SPIRV::OpTypeImage &&
4265 "ImageReg is not an image type.");
4266
4267 bool IsSignedInteger =
4268 sampledTypeIsSignedInteger(GR.getTypeForSPIRVType(ImageType));
4269 // Check if the "sampled" operand of the image type is 1.
4270 // https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpImageFetch
4271 auto SampledOp = ImageType->getOperand(6);
4272 bool IsFetch = (SampledOp.getImm() == 1);
4273
4274 uint64_t ResultSize = GR.getScalarOrVectorComponentCount(ResType);
4275 if (ResultSize == 4) {
4276 auto BMI =
4277 BuildMI(*Pos.getParent(), Pos, Loc,
4278 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
4279 .addDef(ResVReg)
4280 .addUse(GR.getSPIRVTypeID(ResType))
4281 .addUse(ImageReg)
4282 .addUse(IdxReg);
4283
4284 if (IsSignedInteger)
4285 BMI.addImm(0x1000); // SignExtend
4286 return BMI.constrainAllUses(TII, TRI, RBI);
4287 }
4288
4289 SPIRVType *ReadType = widenTypeToVec4(ResType, Pos);
4290 Register ReadReg = MRI->createVirtualRegister(GR.getRegClass(ReadType));
4291 auto BMI =
4292 BuildMI(*Pos.getParent(), Pos, Loc,
4293 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
4294 .addDef(ReadReg)
4295 .addUse(GR.getSPIRVTypeID(ReadType))
4296 .addUse(ImageReg)
4297 .addUse(IdxReg);
4298 if (IsSignedInteger)
4299 BMI.addImm(0x1000); // SignExtend
4300 bool Succeed = BMI.constrainAllUses(TII, TRI, RBI);
4301 if (!Succeed)
4302 return false;
4303
4304 if (ResultSize == 1) {
4305 return BuildMI(*Pos.getParent(), Pos, Loc,
4306 TII.get(SPIRV::OpCompositeExtract))
4307 .addDef(ResVReg)
4308 .addUse(GR.getSPIRVTypeID(ResType))
4309 .addUse(ReadReg)
4310 .addImm(0)
4311 .constrainAllUses(TII, TRI, RBI);
4312 }
4313 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
4314}
4315
4316bool SPIRVInstructionSelector::selectResourceGetPointer(
4317 Register &ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
4318 Register ResourcePtr = I.getOperand(2).getReg();
4319 SPIRVType *RegType = GR.getSPIRVTypeForVReg(ResourcePtr, I.getMF());
4320 if (RegType->getOpcode() == SPIRV::OpTypeImage) {
4321 // For texel buffers, the index into the image is part of the OpImageRead or
4322 // OpImageWrite instructions. So we will do nothing in this case. This
4323 // intrinsic will be combined with the load or store when selecting the load
4324 // or store.
4325 return true;
4326 }
4327
4328 assert(ResType->getOpcode() == SPIRV::OpTypePointer);
4329 MachineIRBuilder MIRBuilder(I);
4330
4331 Register IndexReg = I.getOperand(3).getReg();
4332 Register ZeroReg =
4333 buildZerosVal(GR.getOrCreateSPIRVIntegerType(32, I, TII), I);
4334 return BuildMI(*I.getParent(), I, I.getDebugLoc(),
4335 TII.get(SPIRV::OpAccessChain))
4336 .addDef(ResVReg)
4337 .addUse(GR.getSPIRVTypeID(ResType))
4338 .addUse(ResourcePtr)
4339 .addUse(ZeroReg)
4340 .addUse(IndexReg)
4341 .constrainAllUses(TII, TRI, RBI);
4342}
4343
4344bool SPIRVInstructionSelector::selectPushConstantGetPointer(
4345 Register &ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
4346 MRI->replaceRegWith(ResVReg, I.getOperand(2).getReg());
4347 return true;
4348}
4349
4350bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
4351 Register &ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
4352 Register ObjReg = I.getOperand(2).getReg();
4353 if (!BuildCOPY(ResVReg, ObjReg, I))
4354 return false;
4355
4356 buildOpDecorate(ResVReg, I, TII, SPIRV::Decoration::NonUniformEXT, {});
4357 // Check for the registers that use the index marked as non-uniform
4358 // and recursively mark them as non-uniform.
4359 // Per the spec, it's necessary that the final argument used for
4360 // load/store/sample/atomic must be decorated, so we need to propagate the
4361 // decoration through access chains and copies.
4362 // https://docs.vulkan.org/samples/latest/samples/extensions/descriptor_indexing/README.html#_when_to_use_non_uniform_indexing_qualifier
4363 decorateUsesAsNonUniform(ResVReg);
4364 return true;
4365}
4366
4367void SPIRVInstructionSelector::decorateUsesAsNonUniform(
4368 Register &NonUniformReg) const {
4369 llvm::SmallVector<Register> WorkList = {NonUniformReg};
4370 while (WorkList.size() > 0) {
4371 Register CurrentReg = WorkList.back();
4372 WorkList.pop_back();
4373
4374 bool IsDecorated = false;
4375 for (MachineInstr &Use : MRI->use_instructions(CurrentReg)) {
4376 if (Use.getOpcode() == SPIRV::OpDecorate &&
4377 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
4378 IsDecorated = true;
4379 continue;
4380 }
4381 // Check if the instruction has the result register and add it to the
4382 // worklist.
4383 if (Use.getOperand(0).isReg() && Use.getOperand(0).isDef()) {
4384 Register ResultReg = Use.getOperand(0).getReg();
4385 if (ResultReg == CurrentReg)
4386 continue;
4387 WorkList.push_back(ResultReg);
4388 }
4389 }
4390
4391 if (!IsDecorated) {
4392 buildOpDecorate(CurrentReg, *MRI->getVRegDef(CurrentReg), TII,
4393 SPIRV::Decoration::NonUniformEXT, {});
4394 }
4395 }
4396}
4397
4398bool SPIRVInstructionSelector::extractSubvector(
4399 Register &ResVReg, const SPIRVType *ResType, Register &ReadReg,
4400 MachineInstr &InsertionPoint) const {
4401 SPIRVType *InputType = GR.getResultType(ReadReg);
4402 [[maybe_unused]] uint64_t InputSize =
4403 GR.getScalarOrVectorComponentCount(InputType);
4404 uint64_t ResultSize = GR.getScalarOrVectorComponentCount(ResType);
4405 assert(InputSize > 1 && "The input must be a vector.");
4406 assert(ResultSize > 1 && "The result must be a vector.");
4407 assert(ResultSize < InputSize &&
4408 "Cannot extract more element than there are in the input.");
4409 SmallVector<Register> ComponentRegisters;
4410 SPIRVType *ScalarType = GR.getScalarOrVectorComponentType(ResType);
4411 const TargetRegisterClass *ScalarRegClass = GR.getRegClass(ScalarType);
4412 for (uint64_t I = 0; I < ResultSize; I++) {
4413 Register ComponentReg = MRI->createVirtualRegister(ScalarRegClass);
4414 bool Succeed = BuildMI(*InsertionPoint.getParent(), InsertionPoint,
4415 InsertionPoint.getDebugLoc(),
4416 TII.get(SPIRV::OpCompositeExtract))
4417 .addDef(ComponentReg)
4418 .addUse(ScalarType->getOperand(0).getReg())
4419 .addUse(ReadReg)
4420 .addImm(I)
4421 .constrainAllUses(TII, TRI, RBI);
4422 if (!Succeed)
4423 return false;
4424 ComponentRegisters.emplace_back(ComponentReg);
4425 }
4426
4427 MachineInstrBuilder MIB = BuildMI(*InsertionPoint.getParent(), InsertionPoint,
4428 InsertionPoint.getDebugLoc(),
4429 TII.get(SPIRV::OpCompositeConstruct))
4430 .addDef(ResVReg)
4431 .addUse(GR.getSPIRVTypeID(ResType));
4432
4433 for (Register ComponentReg : ComponentRegisters)
4434 MIB.addUse(ComponentReg);
4435 return MIB.constrainAllUses(TII, TRI, RBI);
4436}
4437
4438bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
4439 MachineInstr &I) const {
4440 // If the load of the image is in a different basic block, then
4441 // this will generate invalid code. A proper solution is to move
4442 // the OpLoad from selectHandleFromBinding here. However, to do
4443 // that we will need to change the return type of the intrinsic.
4444 // We will do that when we can, but for now trying to move forward with other
4445 // issues.
4446 Register ImageReg = I.getOperand(1).getReg();
4447 auto *ImageDef = cast<GIntrinsic>(getVRegDef(*MRI, ImageReg));
4448 Register NewImageReg = MRI->createVirtualRegister(MRI->getRegClass(ImageReg));
4449 if (!loadHandleBeforePosition(NewImageReg, GR.getSPIRVTypeForVReg(ImageReg),
4450 *ImageDef, I)) {
4451 return false;
4452 }
4453
4454 Register CoordinateReg = I.getOperand(2).getReg();
4455 Register DataReg = I.getOperand(3).getReg();
4456 assert(GR.getResultType(DataReg)->getOpcode() == SPIRV::OpTypeVector);
4458 return BuildMI(*I.getParent(), I, I.getDebugLoc(),
4459 TII.get(SPIRV::OpImageWrite))
4460 .addUse(NewImageReg)
4461 .addUse(CoordinateReg)
4462 .addUse(DataReg)
4463 .constrainAllUses(TII, TRI, RBI);
4464}
4465
4466Register SPIRVInstructionSelector::buildPointerToResource(
4467 const SPIRVType *SpirvResType, SPIRV::StorageClass::StorageClass SC,
4468 uint32_t Set, uint32_t Binding, uint32_t ArraySize, Register IndexReg,
4469 StringRef Name, MachineIRBuilder MIRBuilder) const {
4470 const Type *ResType = GR.getTypeForSPIRVType(SpirvResType);
4471 if (ArraySize == 1) {
4472 SPIRVType *PtrType =
4473 GR.getOrCreateSPIRVPointerType(ResType, MIRBuilder, SC);
4474 assert(GR.getPointeeType(PtrType) == SpirvResType &&
4475 "SpirvResType did not have an explicit layout.");
4476 return GR.getOrCreateGlobalVariableWithBinding(PtrType, Set, Binding, Name,
4477 MIRBuilder);
4478 }
4479
4480 const Type *VarType = ArrayType::get(const_cast<Type *>(ResType), ArraySize);
4481 SPIRVType *VarPointerType =
4482 GR.getOrCreateSPIRVPointerType(VarType, MIRBuilder, SC);
4484 VarPointerType, Set, Binding, Name, MIRBuilder);
4485
4486 SPIRVType *ResPointerType =
4487 GR.getOrCreateSPIRVPointerType(ResType, MIRBuilder, SC);
4488 Register AcReg = MRI->createVirtualRegister(GR.getRegClass(ResPointerType));
4489
4490 MIRBuilder.buildInstr(SPIRV::OpAccessChain)
4491 .addDef(AcReg)
4492 .addUse(GR.getSPIRVTypeID(ResPointerType))
4493 .addUse(VarReg)
4494 .addUse(IndexReg);
4495
4496 return AcReg;
4497}
4498
4499bool SPIRVInstructionSelector::selectFirstBitSet16(
4500 Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
4501 unsigned ExtendOpcode, unsigned BitSetOpcode) const {
4502 Register ExtReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
4503 bool Result = selectOpWithSrcs(ExtReg, ResType, I, {I.getOperand(2).getReg()},
4504 ExtendOpcode);
4505
4506 return Result &&
4507 selectFirstBitSet32(ResVReg, ResType, I, ExtReg, BitSetOpcode);
4508}
4509
4510bool SPIRVInstructionSelector::selectFirstBitSet32(
4511 Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
4512 Register SrcReg, unsigned BitSetOpcode) const {
4513 return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
4514 .addDef(ResVReg)
4515 .addUse(GR.getSPIRVTypeID(ResType))
4516 .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::GLSL_std_450))
4517 .addImm(BitSetOpcode)
4518 .addUse(SrcReg)
4519 .constrainAllUses(TII, TRI, RBI);
4520}
4521
4522bool SPIRVInstructionSelector::selectFirstBitSet64Overflow(
4523 Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
4524 Register SrcReg, unsigned BitSetOpcode, bool SwapPrimarySide) const {
4525
4526 // SPIR-V allow vectors of size 2,3,4 only. Calling with a larger vectors
4527 // requires creating a param register and return register with an invalid
4528 // vector size. If that is resolved, then this function can be used for
4529 // vectors of any component size.
4530 unsigned ComponentCount = GR.getScalarOrVectorComponentCount(ResType);
4531 assert(ComponentCount < 5 && "Vec 5+ will generate invalid SPIR-V ops");
4532
4533 MachineIRBuilder MIRBuilder(I);
4535 SPIRVType *I64Type = GR.getOrCreateSPIRVIntegerType(64, MIRBuilder);
4536 SPIRVType *I64x2Type =
4537 GR.getOrCreateSPIRVVectorType(I64Type, 2, MIRBuilder, false);
4538 SPIRVType *Vec2ResType =
4539 GR.getOrCreateSPIRVVectorType(BaseType, 2, MIRBuilder, false);
4540
4541 std::vector<Register> PartialRegs;
4542
4543 // Loops 0, 2, 4, ... but stops one loop early when ComponentCount is odd
4544 unsigned CurrentComponent = 0;
4545 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
4546 // This register holds the firstbitX result for each of the i64x2 vectors
4547 // extracted from SrcReg
4548 Register BitSetResult =
4549 MRI->createVirtualRegister(GR.getRegClass(I64x2Type));
4550
4551 auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
4552 TII.get(SPIRV::OpVectorShuffle))
4553 .addDef(BitSetResult)
4554 .addUse(GR.getSPIRVTypeID(I64x2Type))
4555 .addUse(SrcReg)
4556 .addUse(SrcReg)
4557 .addImm(CurrentComponent)
4558 .addImm(CurrentComponent + 1);
4559
4560 if (!MIB.constrainAllUses(TII, TRI, RBI))
4561 return false;
4562
4563 Register SubVecBitSetReg =
4564 MRI->createVirtualRegister(GR.getRegClass(Vec2ResType));
4565
4566 if (!selectFirstBitSet64(SubVecBitSetReg, Vec2ResType, I, BitSetResult,
4567 BitSetOpcode, SwapPrimarySide))
4568 return false;
4569
4570 PartialRegs.push_back(SubVecBitSetReg);
4571 }
4572
4573 // On odd component counts we need to handle one more component
4574 if (CurrentComponent != ComponentCount) {
4575 bool ZeroAsNull = !STI.isShader();
4576 Register FinalElemReg = MRI->createVirtualRegister(GR.getRegClass(I64Type));
4577 Register ConstIntLastIdx = GR.getOrCreateConstInt(
4578 ComponentCount - 1, I, BaseType, TII, ZeroAsNull);
4579
4580 if (!selectOpWithSrcs(FinalElemReg, I64Type, I, {SrcReg, ConstIntLastIdx},
4581 SPIRV::OpVectorExtractDynamic))
4582 return false;
4583
4584 Register FinalElemBitSetReg =
4585 MRI->createVirtualRegister(GR.getRegClass(BaseType));
4586
4587 if (!selectFirstBitSet64(FinalElemBitSetReg, BaseType, I, FinalElemReg,
4588 BitSetOpcode, SwapPrimarySide))
4589 return false;
4590
4591 PartialRegs.push_back(FinalElemBitSetReg);
4592 }
4593
4594 // Join all the resulting registers back into the return type in order
4595 // (ie i32x2, i32x2, i32x1 -> i32x5)
4596 return selectOpWithSrcs(ResVReg, ResType, I, std::move(PartialRegs),
4597 SPIRV::OpCompositeConstruct);
4598}
4599
4600bool SPIRVInstructionSelector::selectFirstBitSet64(
4601 Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
4602 Register SrcReg, unsigned BitSetOpcode, bool SwapPrimarySide) const {
4603 unsigned ComponentCount = GR.getScalarOrVectorComponentCount(ResType);
4605 bool ZeroAsNull = !STI.isShader();
4606 Register ConstIntZero =
4607 GR.getOrCreateConstInt(0, I, BaseType, TII, ZeroAsNull);
4608 Register ConstIntOne =
4609 GR.getOrCreateConstInt(1, I, BaseType, TII, ZeroAsNull);
4610
4611 // SPIRV doesn't support vectors with more than 4 components. Since the
4612 // algoritm below converts i64 -> i32x2 and i64x4 -> i32x8 it can only
4613 // operate on vectors with 2 or less components. When largers vectors are
4614 // seen. Split them, recurse, then recombine them.
4615 if (ComponentCount > 2) {
4616 return selectFirstBitSet64Overflow(ResVReg, ResType, I, SrcReg,
4617 BitSetOpcode, SwapPrimarySide);
4618 }
4619
4620 // 1. Split int64 into 2 pieces using a bitcast
4621 MachineIRBuilder MIRBuilder(I);
4622 SPIRVType *PostCastType = GR.getOrCreateSPIRVVectorType(
4623 BaseType, 2 * ComponentCount, MIRBuilder, false);
4624 Register BitcastReg =
4625 MRI->createVirtualRegister(GR.getRegClass(PostCastType));
4626
4627 if (!selectOpWithSrcs(BitcastReg, PostCastType, I, {SrcReg},
4628 SPIRV::OpBitcast))
4629 return false;
4630
4631 // 2. Find the first set bit from the primary side for all the pieces in #1
4632 Register FBSReg = MRI->createVirtualRegister(GR.getRegClass(PostCastType));
4633 if (!selectFirstBitSet32(FBSReg, PostCastType, I, BitcastReg, BitSetOpcode))
4634 return false;
4635
4636 // 3. Split result vector into high bits and low bits
4637 Register HighReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
4638 Register LowReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
4639
4640 bool IsScalarRes = ResType->getOpcode() != SPIRV::OpTypeVector;
4641 if (IsScalarRes) {
4642 // if scalar do a vector extract
4643 if (!selectOpWithSrcs(HighReg, ResType, I, {FBSReg, ConstIntZero},
4644 SPIRV::OpVectorExtractDynamic))
4645 return false;
4646 if (!selectOpWithSrcs(LowReg, ResType, I, {FBSReg, ConstIntOne},
4647 SPIRV::OpVectorExtractDynamic))
4648 return false;
4649 } else {
4650 // if vector do a shufflevector
4651 auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
4652 TII.get(SPIRV::OpVectorShuffle))
4653 .addDef(HighReg)
4654 .addUse(GR.getSPIRVTypeID(ResType))
4655 .addUse(FBSReg)
4656 // Per the spec, repeat the vector if only one vec is needed
4657 .addUse(FBSReg);
4658
4659 // high bits are stored in even indexes. Extract them from FBSReg
4660 for (unsigned J = 0; J < ComponentCount * 2; J += 2) {
4661 MIB.addImm(J);
4662 }
4663
4664 if (!MIB.constrainAllUses(TII, TRI, RBI))
4665 return false;
4666
4667 MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
4668 TII.get(SPIRV::OpVectorShuffle))
4669 .addDef(LowReg)
4670 .addUse(GR.getSPIRVTypeID(ResType))
4671 .addUse(FBSReg)
4672 // Per the spec, repeat the vector if only one vec is needed
4673 .addUse(FBSReg);
4674
4675 // low bits are stored in odd indexes. Extract them from FBSReg
4676 for (unsigned J = 1; J < ComponentCount * 2; J += 2) {
4677 MIB.addImm(J);
4678 }
4679 if (!MIB.constrainAllUses(TII, TRI, RBI))
4680 return false;
4681 }
4682
4683 // 4. Check the result. When primary bits == -1 use secondary, otherwise use
4684 // primary
4685 SPIRVType *BoolType = GR.getOrCreateSPIRVBoolType(I, TII);
4686 Register NegOneReg;
4687 Register Reg0;
4688 Register Reg32;
4689 unsigned SelectOp;
4690 unsigned AddOp;
4691
4692 if (IsScalarRes) {
4693 NegOneReg =
4694 GR.getOrCreateConstInt((unsigned)-1, I, ResType, TII, ZeroAsNull);
4695 Reg0 = GR.getOrCreateConstInt(0, I, ResType, TII, ZeroAsNull);
4696 Reg32 = GR.getOrCreateConstInt(32, I, ResType, TII, ZeroAsNull);
4697 SelectOp = SPIRV::OpSelectSISCond;
4698 AddOp = SPIRV::OpIAddS;
4699 } else {
4700 BoolType = GR.getOrCreateSPIRVVectorType(BoolType, ComponentCount,
4701 MIRBuilder, false);
4702 NegOneReg =
4703 GR.getOrCreateConstVector((unsigned)-1, I, ResType, TII, ZeroAsNull);
4704 Reg0 = GR.getOrCreateConstVector(0, I, ResType, TII, ZeroAsNull);
4705 Reg32 = GR.getOrCreateConstVector(32, I, ResType, TII, ZeroAsNull);
4706 SelectOp = SPIRV::OpSelectVIVCond;
4707 AddOp = SPIRV::OpIAddV;
4708 }
4709
4710 Register PrimaryReg = HighReg;
4711 Register SecondaryReg = LowReg;
4712 Register PrimaryShiftReg = Reg32;
4713 Register SecondaryShiftReg = Reg0;
4714
4715 // By default the emitted opcodes check for the set bit from the MSB side.
4716 // Setting SwapPrimarySide checks the set bit from the LSB side
4717 if (SwapPrimarySide) {
4718 PrimaryReg = LowReg;
4719 SecondaryReg = HighReg;
4720 PrimaryShiftReg = Reg0;
4721 SecondaryShiftReg = Reg32;
4722 }
4723
4724 // Check if the primary bits are == -1
4725 Register BReg = MRI->createVirtualRegister(GR.getRegClass(BoolType));
4726 if (!selectOpWithSrcs(BReg, BoolType, I, {PrimaryReg, NegOneReg},
4727 SPIRV::OpIEqual))
4728 return false;
4729
4730 // Select secondary bits if true in BReg, otherwise primary bits
4731 Register TmpReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
4732 if (!selectOpWithSrcs(TmpReg, ResType, I, {BReg, SecondaryReg, PrimaryReg},
4733 SelectOp))
4734 return false;
4735
4736 // 5. Add 32 when high bits are used, otherwise 0 for low bits
4737 Register ValReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
4738 if (!selectOpWithSrcs(ValReg, ResType, I,
4739 {BReg, SecondaryShiftReg, PrimaryShiftReg}, SelectOp))
4740 return false;
4741
4742 return selectOpWithSrcs(ResVReg, ResType, I, {ValReg, TmpReg}, AddOp);
4743}
4744
4745bool SPIRVInstructionSelector::selectFirstBitHigh(Register ResVReg,
4746 const SPIRVType *ResType,
4747 MachineInstr &I,
4748 bool IsSigned) const {
4749 // FindUMsb and FindSMsb intrinsics only support 32 bit integers
4750 Register OpReg = I.getOperand(2).getReg();
4751 SPIRVType *OpType = GR.getSPIRVTypeForVReg(OpReg);
4752 // zero or sign extend
4753 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4754 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
4755
4756 switch (GR.getScalarOrVectorBitWidth(OpType)) {
4757 case 16:
4758 return selectFirstBitSet16(ResVReg, ResType, I, ExtendOpcode, BitSetOpcode);
4759 case 32:
4760 return selectFirstBitSet32(ResVReg, ResType, I, OpReg, BitSetOpcode);
4761 case 64:
4762 return selectFirstBitSet64(ResVReg, ResType, I, OpReg, BitSetOpcode,
4763 /*SwapPrimarySide=*/false);
4764 default:
4766 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
4767 }
4768}
4769
4770bool SPIRVInstructionSelector::selectFirstBitLow(Register ResVReg,
4771 const SPIRVType *ResType,
4772 MachineInstr &I) const {
4773 // FindILsb intrinsic only supports 32 bit integers
4774 Register OpReg = I.getOperand(2).getReg();
4775 SPIRVType *OpType = GR.getSPIRVTypeForVReg(OpReg);
4776 // OpUConvert treats the operand bits as an unsigned i16 and zero extends it
4777 // to an unsigned i32. As this leaves all the least significant bits unchanged
4778 // so the first set bit from the LSB side doesn't change.
4779 unsigned ExtendOpcode = SPIRV::OpUConvert;
4780 unsigned BitSetOpcode = GL::FindILsb;
4781
4782 switch (GR.getScalarOrVectorBitWidth(OpType)) {
4783 case 16:
4784 return selectFirstBitSet16(ResVReg, ResType, I, ExtendOpcode, BitSetOpcode);
4785 case 32:
4786 return selectFirstBitSet32(ResVReg, ResType, I, OpReg, BitSetOpcode);
4787 case 64:
4788 return selectFirstBitSet64(ResVReg, ResType, I, OpReg, BitSetOpcode,
4789 /*SwapPrimarySide=*/true);
4790 default:
4791 report_fatal_error("spv_firstbitlow only supports 16,32,64 bits.");
4792 }
4793}
4794
4795bool SPIRVInstructionSelector::selectAllocaArray(Register ResVReg,
4796 const SPIRVType *ResType,
4797 MachineInstr &I) const {
4798 // there was an allocation size parameter to the allocation instruction
4799 // that is not 1
4800 MachineBasicBlock &BB = *I.getParent();
4801 bool Res = BuildMI(BB, I, I.getDebugLoc(),
4802 TII.get(SPIRV::OpVariableLengthArrayINTEL))
4803 .addDef(ResVReg)
4804 .addUse(GR.getSPIRVTypeID(ResType))
4805 .addUse(I.getOperand(2).getReg())
4806 .constrainAllUses(TII, TRI, RBI);
4807 if (!STI.isShader()) {
4808 unsigned Alignment = I.getOperand(3).getImm();
4809 buildOpDecorate(ResVReg, I, TII, SPIRV::Decoration::Alignment, {Alignment});
4810 }
4811 return Res;
4812}
4813
4814bool SPIRVInstructionSelector::selectFrameIndex(Register ResVReg,
4815 const SPIRVType *ResType,
4816 MachineInstr &I) const {
4817 // Change order of instructions if needed: all OpVariable instructions in a
4818 // function must be the first instructions in the first block
4819 auto It = getOpVariableMBBIt(I);
4820 bool Res = BuildMI(*It->getParent(), It, It->getDebugLoc(),
4821 TII.get(SPIRV::OpVariable))
4822 .addDef(ResVReg)
4823 .addUse(GR.getSPIRVTypeID(ResType))
4824 .addImm(static_cast<uint32_t>(SPIRV::StorageClass::Function))
4825 .constrainAllUses(TII, TRI, RBI);
4826 if (!STI.isShader()) {
4827 unsigned Alignment = I.getOperand(2).getImm();
4828 buildOpDecorate(ResVReg, *It, TII, SPIRV::Decoration::Alignment,
4829 {Alignment});
4830 }
4831 return Res;
4832}
4833
4834bool SPIRVInstructionSelector::selectBranch(MachineInstr &I) const {
4835 // InstructionSelector walks backwards through the instructions. We can use
4836 // both a G_BR and a G_BRCOND to create an OpBranchConditional. We hit G_BR
4837 // first, so can generate an OpBranchConditional here. If there is no
4838 // G_BRCOND, we just use OpBranch for a regular unconditional branch.
4839 const MachineInstr *PrevI = I.getPrevNode();
4840 MachineBasicBlock &MBB = *I.getParent();
4841 if (PrevI != nullptr && PrevI->getOpcode() == TargetOpcode::G_BRCOND) {
4842 return BuildMI(MBB, I, I.getDebugLoc(), TII.get(SPIRV::OpBranchConditional))
4843 .addUse(PrevI->getOperand(0).getReg())
4844 .addMBB(PrevI->getOperand(1).getMBB())
4845 .addMBB(I.getOperand(0).getMBB())
4846 .constrainAllUses(TII, TRI, RBI);
4847 }
4848 return BuildMI(MBB, I, I.getDebugLoc(), TII.get(SPIRV::OpBranch))
4849 .addMBB(I.getOperand(0).getMBB())
4850 .constrainAllUses(TII, TRI, RBI);
4851}
4852
4853bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &I) const {
4854 // InstructionSelector walks backwards through the instructions. For an
4855 // explicit conditional branch with no fallthrough, we use both a G_BR and a
4856 // G_BRCOND to create an OpBranchConditional. We should hit G_BR first, and
4857 // generate the OpBranchConditional in selectBranch above.
4858 //
4859 // If an OpBranchConditional has been generated, we simply return, as the work
4860 // is alread done. If there is no OpBranchConditional, LLVM must be relying on
4861 // implicit fallthrough to the next basic block, so we need to create an
4862 // OpBranchConditional with an explicit "false" argument pointing to the next
4863 // basic block that LLVM would fall through to.
4864 const MachineInstr *NextI = I.getNextNode();
4865 // Check if this has already been successfully selected.
4866 if (NextI != nullptr && NextI->getOpcode() == SPIRV::OpBranchConditional)
4867 return true;
4868 // Must be relying on implicit block fallthrough, so generate an
4869 // OpBranchConditional with the "next" basic block as the "false" target.
4870 MachineBasicBlock &MBB = *I.getParent();
4871 unsigned NextMBBNum = MBB.getNextNode()->getNumber();
4872 MachineBasicBlock *NextMBB = I.getMF()->getBlockNumbered(NextMBBNum);
4873 return BuildMI(MBB, I, I.getDebugLoc(), TII.get(SPIRV::OpBranchConditional))
4874 .addUse(I.getOperand(0).getReg())
4875 .addMBB(I.getOperand(1).getMBB())
4876 .addMBB(NextMBB)
4877 .constrainAllUses(TII, TRI, RBI);
4878}
4879
4880bool SPIRVInstructionSelector::selectPhi(Register ResVReg,
4881 const SPIRVType *ResType,
4882 MachineInstr &I) const {
4883 auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpPhi))
4884 .addDef(ResVReg)
4885 .addUse(GR.getSPIRVTypeID(ResType));
4886 const unsigned NumOps = I.getNumOperands();
4887 for (unsigned i = 1; i < NumOps; i += 2) {
4888 MIB.addUse(I.getOperand(i + 0).getReg());
4889 MIB.addMBB(I.getOperand(i + 1).getMBB());
4890 }
4891 bool Res = MIB.constrainAllUses(TII, TRI, RBI);
4892 MIB->setDesc(TII.get(TargetOpcode::PHI));
4893 MIB->removeOperand(1);
4894 return Res;
4895}
4896
4897bool SPIRVInstructionSelector::selectGlobalValue(
4898 Register ResVReg, MachineInstr &I, const MachineInstr *Init) const {
4899 // FIXME: don't use MachineIRBuilder here, replace it with BuildMI.
4900 MachineIRBuilder MIRBuilder(I);
4901 const GlobalValue *GV = I.getOperand(1).getGlobal();
4903
4904 std::string GlobalIdent;
4905 if (!GV->hasName()) {
4906 unsigned &ID = UnnamedGlobalIDs[GV];
4907 if (ID == 0)
4908 ID = UnnamedGlobalIDs.size();
4909 GlobalIdent = "__unnamed_" + Twine(ID).str();
4910 } else {
4911 GlobalIdent = GV->getName();
4912 }
4913
4914 // Behaviour of functions as operands depends on availability of the
4915 // corresponding extension (SPV_INTEL_function_pointers):
4916 // - If there is an extension to operate with functions as operands:
4917 // We create a proper constant operand and evaluate a correct type for a
4918 // function pointer.
4919 // - Without the required extension:
4920 // We have functions as operands in tests with blocks of instruction e.g. in
4921 // transcoding/global_block.ll. These operands are not used and should be
4922 // substituted by zero constants. Their type is expected to be always
4923 // OpTypePointer Function %uchar.
4924 if (isa<Function>(GV)) {
4925 const Constant *ConstVal = GV;
4926 MachineBasicBlock &BB = *I.getParent();
4927 Register NewReg = GR.find(ConstVal, GR.CurMF);
4928 if (!NewReg.isValid()) {
4929 Register NewReg = ResVReg;
4930 const Function *GVFun =
4931 STI.canUseExtension(SPIRV::Extension::SPV_INTEL_function_pointers)
4932 ? dyn_cast<Function>(GV)
4933 : nullptr;
4935 GVType, I,
4936 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
4938 if (GVFun) {
4939 // References to a function via function pointers generate virtual
4940 // registers without a definition. We will resolve it later, during
4941 // module analysis stage.
4942 Register ResTypeReg = GR.getSPIRVTypeID(ResType);
4943 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
4944 Register FuncVReg =
4945 MRI->createGenericVirtualRegister(GR.getRegType(ResType));
4946 MRI->setRegClass(FuncVReg, &SPIRV::pIDRegClass);
4947 MachineInstrBuilder MIB1 =
4948 BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpUndef))
4949 .addDef(FuncVReg)
4950 .addUse(ResTypeReg);
4951 MachineInstrBuilder MIB2 =
4952 BuildMI(BB, I, I.getDebugLoc(),
4953 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
4954 .addDef(NewReg)
4955 .addUse(ResTypeReg)
4956 .addUse(FuncVReg);
4957 GR.add(ConstVal, MIB2);
4958 // mapping the function pointer to the used Function
4959 GR.recordFunctionPointer(&MIB2.getInstr()->getOperand(2), GVFun);
4960 return MIB1.constrainAllUses(TII, TRI, RBI) &&
4961 MIB2.constrainAllUses(TII, TRI, RBI);
4962 }
4963 MachineInstrBuilder MIB3 =
4964 BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantNull))
4965 .addDef(NewReg)
4966 .addUse(GR.getSPIRVTypeID(ResType));
4967 GR.add(ConstVal, MIB3);
4968 return MIB3.constrainAllUses(TII, TRI, RBI);
4969 }
4970 assert(NewReg != ResVReg);
4971 return BuildCOPY(ResVReg, NewReg, I);
4972 }
4974 assert(GlobalVar->getName() != "llvm.global.annotations");
4975
4976 // Skip empty declaration for GVs with initializers till we get the decl with
4977 // passed initializer.
4978 if (hasInitializer(GlobalVar) && !Init)
4979 return true;
4980
4981 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
4982 getSpirvLinkageTypeFor(STI, *GV);
4983
4984 const unsigned AddrSpace = GV->getAddressSpace();
4985 SPIRV::StorageClass::StorageClass StorageClass =
4986 addressSpaceToStorageClass(AddrSpace, STI);
4987 SPIRVType *ResType = GR.getOrCreateSPIRVPointerType(GVType, I, StorageClass);
4989 ResVReg, ResType, GlobalIdent, GV, StorageClass, Init,
4990 GlobalVar->isConstant(), LnkType, MIRBuilder, true);
4991 // TODO: For AMDGCN, we pipe externally_initialized through via
4992 // HostAccessINTEL, with ReadWrite (3) access, which is we then handle during
4993 // reverse translation. We should remove this once SPIR-V gains the ability to
4994 // express the concept.
4995 if (GlobalVar->isExternallyInitialized() &&
4996 STI.getTargetTriple().getVendor() == Triple::AMD) {
4997 constexpr unsigned ReadWriteINTEL = 3u;
4998 buildOpDecorate(Reg, MIRBuilder, SPIRV::Decoration::HostAccessINTEL,
4999 {ReadWriteINTEL});
5000 MachineInstrBuilder MIB(*MF, --MIRBuilder.getInsertPt());
5001 addStringImm(GV->getName(), MIB);
5002 }
5003 return Reg.isValid();
5004}
5005
5006bool SPIRVInstructionSelector::selectLog10(Register ResVReg,
5007 const SPIRVType *ResType,
5008 MachineInstr &I) const {
5009 if (STI.canUseExtInstSet(SPIRV::InstructionSet::OpenCL_std)) {
5010 return selectExtInst(ResVReg, ResType, I, CL::log10);
5011 }
5012
5013 // There is no log10 instruction in the GLSL Extended Instruction set, so it
5014 // is implemented as:
5015 // log10(x) = log2(x) * (1 / log2(10))
5016 // = log2(x) * 0.30103
5017
5018 MachineIRBuilder MIRBuilder(I);
5019 MachineBasicBlock &BB = *I.getParent();
5020
5021 // Build log2(x).
5022 Register VarReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
5023 bool Result =
5024 BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
5025 .addDef(VarReg)
5026 .addUse(GR.getSPIRVTypeID(ResType))
5027 .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::GLSL_std_450))
5028 .addImm(GL::Log2)
5029 .add(I.getOperand(1))
5030 .constrainAllUses(TII, TRI, RBI);
5031
5032 // Build 0.30103.
5033 assert(ResType->getOpcode() == SPIRV::OpTypeVector ||
5034 ResType->getOpcode() == SPIRV::OpTypeFloat);
5035 // TODO: Add matrix implementation once supported by the HLSL frontend.
5036 const SPIRVType *SpirvScalarType =
5037 ResType->getOpcode() == SPIRV::OpTypeVector
5038 ? GR.getSPIRVTypeForVReg(ResType->getOperand(1).getReg())
5039 : ResType;
5040 Register ScaleReg =
5041 GR.buildConstantFP(APFloat(0.30103f), MIRBuilder, SpirvScalarType);
5042
5043 // Multiply log2(x) by 0.30103 to get log10(x) result.
5044 auto Opcode = ResType->getOpcode() == SPIRV::OpTypeVector
5045 ? SPIRV::OpVectorTimesScalar
5046 : SPIRV::OpFMulS;
5047 return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
5048 .addDef(ResVReg)
5049 .addUse(GR.getSPIRVTypeID(ResType))
5050 .addUse(VarReg)
5051 .addUse(ScaleReg)
5052 .constrainAllUses(TII, TRI, RBI);
5053}
5054
5055bool SPIRVInstructionSelector::selectModf(Register ResVReg,
5056 const SPIRVType *ResType,
5057 MachineInstr &I) const {
5058 // llvm.modf has a single arg --the number to be decomposed-- and returns a
5059 // struct { restype, restype }, while OpenCLLIB::modf has two args --the
5060 // number to be decomposed and a pointer--, returns the fractional part and
5061 // the integral part is stored in the pointer argument. Therefore, we can't
5062 // use directly the OpenCLLIB::modf intrinsic. However, we can do some
5063 // scaffolding to make it work. The idea is to create an alloca instruction
5064 // to get a ptr, pass this ptr to OpenCL::modf, and then load the value
5065 // from this ptr to place it in the struct. llvm.modf returns the fractional
5066 // part as the first element of the result, and the integral part as the
5067 // second element of the result.
5068
5069 // At this point, the return type is not a struct anymore, but rather two
5070 // independent elements of SPIRVResType. We can get each independent element
5071 // from I.getDefs() or I.getOperands().
5072 if (STI.canUseExtInstSet(SPIRV::InstructionSet::OpenCL_std)) {
5073 MachineIRBuilder MIRBuilder(I);
5074 // Get pointer type for alloca variable.
5075 const SPIRVType *PtrType = GR.getOrCreateSPIRVPointerType(
5076 ResType, MIRBuilder, SPIRV::StorageClass::Function);
5077 // Create new register for the pointer type of alloca variable.
5078 Register PtrTyReg =
5079 MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::iIDRegClass);
5080 MIRBuilder.getMRI()->setType(
5081 PtrTyReg,
5082 LLT::pointer(storageClassToAddressSpace(SPIRV::StorageClass::Function),
5083 GR.getPointerSize()));
5084
5085 // Assign SPIR-V type of the pointer type of the alloca variable to the
5086 // new register.
5087 GR.assignSPIRVTypeToVReg(PtrType, PtrTyReg, MIRBuilder.getMF());
5088 MachineBasicBlock &EntryBB = I.getMF()->front();
5091 auto AllocaMIB =
5092 BuildMI(EntryBB, VarPos, I.getDebugLoc(), TII.get(SPIRV::OpVariable))
5093 .addDef(PtrTyReg)
5094 .addUse(GR.getSPIRVTypeID(PtrType))
5095 .addImm(static_cast<uint32_t>(SPIRV::StorageClass::Function));
5096 Register Variable = AllocaMIB->getOperand(0).getReg();
5097
5098 MachineBasicBlock &BB = *I.getParent();
5099 // Create the OpenCLLIB::modf instruction.
5100 auto MIB =
5101 BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
5102 .addDef(ResVReg)
5103 .addUse(GR.getSPIRVTypeID(ResType))
5104 .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::OpenCL_std))
5105 .addImm(CL::modf)
5106 .setMIFlags(I.getFlags())
5107 .add(I.getOperand(I.getNumExplicitDefs())) // Floating point value.
5108 .addUse(Variable); // Pointer to integral part.
5109 // Assign the integral part stored in the ptr to the second element of the
5110 // result.
5111 Register IntegralPartReg = I.getOperand(1).getReg();
5112 if (IntegralPartReg.isValid()) {
5113 // Load the value from the pointer to integral part.
5114 auto LoadMIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpLoad))
5115 .addDef(IntegralPartReg)
5116 .addUse(GR.getSPIRVTypeID(ResType))
5117 .addUse(Variable);
5118 return LoadMIB.constrainAllUses(TII, TRI, RBI);
5119 }
5120
5121 return MIB.constrainAllUses(TII, TRI, RBI);
5122 } else if (STI.canUseExtInstSet(SPIRV::InstructionSet::GLSL_std_450)) {
5123 assert(false && "GLSL::Modf is deprecated.");
5124 // FIXME: GL::Modf is deprecated, use Modfstruct instead.
5125 return false;
5126 }
5127 return false;
5128}
5129
5130// Generate the instructions to load 3-element vector builtin input
5131// IDs/Indices.
5132// Like: GlobalInvocationId, LocalInvocationId, etc....
5133
5134bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
5135 SPIRV::BuiltIn::BuiltIn BuiltInValue, Register ResVReg,
5136 const SPIRVType *ResType, MachineInstr &I) const {
5137 MachineIRBuilder MIRBuilder(I);
5138 const SPIRVType *Vec3Ty =
5139 GR.getOrCreateSPIRVVectorType(ResType, 3, MIRBuilder, false);
5140 const SPIRVType *PtrType = GR.getOrCreateSPIRVPointerType(
5141 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
5142
5143 // Create new register for the input ID builtin variable.
5144 Register NewRegister =
5145 MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::iIDRegClass);
5146 MIRBuilder.getMRI()->setType(NewRegister, LLT::pointer(0, 64));
5147 GR.assignSPIRVTypeToVReg(PtrType, NewRegister, MIRBuilder.getMF());
5148
5149 // Build global variable with the necessary decorations for the input ID
5150 // builtin variable.
5152 NewRegister, PtrType, getLinkStringForBuiltIn(BuiltInValue), nullptr,
5153 SPIRV::StorageClass::Input, nullptr, true, std::nullopt, MIRBuilder,
5154 false);
5155
5156 // Create new register for loading value.
5157 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
5158 Register LoadedRegister = MRI->createVirtualRegister(&SPIRV::iIDRegClass);
5159 MIRBuilder.getMRI()->setType(LoadedRegister, LLT::pointer(0, 64));
5160 GR.assignSPIRVTypeToVReg(Vec3Ty, LoadedRegister, MIRBuilder.getMF());
5161
5162 // Load v3uint value from the global variable.
5163 bool Result =
5164 BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpLoad))
5165 .addDef(LoadedRegister)
5166 .addUse(GR.getSPIRVTypeID(Vec3Ty))
5167 .addUse(Variable);
5168
5169 // Get the input ID index. Expecting operand is a constant immediate value,
5170 // wrapped in a type assignment.
5171 assert(I.getOperand(2).isReg());
5172 const uint32_t ThreadId = foldImm(I.getOperand(2), MRI);
5173
5174 // Extract the input ID from the loaded vector value.
5175 MachineBasicBlock &BB = *I.getParent();
5176 auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
5177 .addDef(ResVReg)
5178 .addUse(GR.getSPIRVTypeID(ResType))
5179 .addUse(LoadedRegister)
5180 .addImm(ThreadId);
5181 return Result && MIB.constrainAllUses(TII, TRI, RBI);
5182}
5183
5184// Generate the instructions to load 32-bit integer builtin input IDs/Indices.
5185// Like LocalInvocationIndex
5186bool SPIRVInstructionSelector::loadBuiltinInputID(
5187 SPIRV::BuiltIn::BuiltIn BuiltInValue, Register ResVReg,
5188 const SPIRVType *ResType, MachineInstr &I) const {
5189 MachineIRBuilder MIRBuilder(I);
5190 const SPIRVType *PtrType = GR.getOrCreateSPIRVPointerType(
5191 ResType, MIRBuilder, SPIRV::StorageClass::Input);
5192
5193 // Create new register for the input ID builtin variable.
5194 Register NewRegister =
5195 MIRBuilder.getMRI()->createVirtualRegister(GR.getRegClass(PtrType));
5196 MIRBuilder.getMRI()->setType(
5197 NewRegister,
5198 LLT::pointer(storageClassToAddressSpace(SPIRV::StorageClass::Input),
5199 GR.getPointerSize()));
5200 GR.assignSPIRVTypeToVReg(PtrType, NewRegister, MIRBuilder.getMF());
5201
5202 // Build global variable with the necessary decorations for the input ID
5203 // builtin variable.
5205 NewRegister, PtrType, getLinkStringForBuiltIn(BuiltInValue), nullptr,
5206 SPIRV::StorageClass::Input, nullptr, true, std::nullopt, MIRBuilder,
5207 false);
5208
5209 // Load uint value from the global variable.
5210 auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpLoad))
5211 .addDef(ResVReg)
5212 .addUse(GR.getSPIRVTypeID(ResType))
5213 .addUse(Variable);
5214
5215 return MIB.constrainAllUses(TII, TRI, RBI);
5216}
5217
5218SPIRVType *SPIRVInstructionSelector::widenTypeToVec4(const SPIRVType *Type,
5219 MachineInstr &I) const {
5220 MachineIRBuilder MIRBuilder(I);
5221 if (Type->getOpcode() != SPIRV::OpTypeVector)
5222 return GR.getOrCreateSPIRVVectorType(Type, 4, MIRBuilder, false);
5223
5224 uint64_t VectorSize = Type->getOperand(2).getImm();
5225 if (VectorSize == 4)
5226 return Type;
5227
5228 Register ScalarTypeReg = Type->getOperand(1).getReg();
5229 const SPIRVType *ScalarType = GR.getSPIRVTypeForVReg(ScalarTypeReg);
5230 return GR.getOrCreateSPIRVVectorType(ScalarType, 4, MIRBuilder, false);
5231}
5232
5233bool SPIRVInstructionSelector::loadHandleBeforePosition(
5234 Register &HandleReg, const SPIRVType *ResType, GIntrinsic &HandleDef,
5235 MachineInstr &Pos) const {
5236
5237 assert(HandleDef.getIntrinsicID() ==
5238 Intrinsic::spv_resource_handlefrombinding);
5239 uint32_t Set = foldImm(HandleDef.getOperand(2), MRI);
5240 uint32_t Binding = foldImm(HandleDef.getOperand(3), MRI);
5241 uint32_t ArraySize = foldImm(HandleDef.getOperand(4), MRI);
5242 Register IndexReg = HandleDef.getOperand(5).getReg();
5243 std::string Name =
5244 getStringValueFromReg(HandleDef.getOperand(6).getReg(), *MRI);
5245
5246 bool IsStructuredBuffer = ResType->getOpcode() == SPIRV::OpTypePointer;
5247 MachineIRBuilder MIRBuilder(HandleDef);
5248 SPIRVType *VarType = ResType;
5249 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
5250
5251 if (IsStructuredBuffer) {
5252 VarType = GR.getPointeeType(ResType);
5253 SC = GR.getPointerStorageClass(ResType);
5254 }
5255
5256 Register VarReg = buildPointerToResource(VarType, SC, Set, Binding, ArraySize,
5257 IndexReg, Name, MIRBuilder);
5258
5259 // The handle for the buffer is the pointer to the resource. For an image, the
5260 // handle is the image object. So images get an extra load.
5261 uint32_t LoadOpcode =
5262 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
5263 GR.assignSPIRVTypeToVReg(ResType, HandleReg, *Pos.getMF());
5264 return BuildMI(*Pos.getParent(), Pos, HandleDef.getDebugLoc(),
5265 TII.get(LoadOpcode))
5266 .addDef(HandleReg)
5267 .addUse(GR.getSPIRVTypeID(ResType))
5268 .addUse(VarReg)
5269 .constrainAllUses(TII, TRI, RBI);
5270}
5271
5272void SPIRVInstructionSelector::errorIfInstrOutsideShader(
5273 MachineInstr &I) const {
5274 if (!STI.isShader()) {
5275 std::string DiagMsg;
5276 raw_string_ostream OS(DiagMsg);
5277 I.print(OS, true, false, false, false);
5278 DiagMsg += " is only supported in shaders.\n";
5279 report_fatal_error(DiagMsg.c_str(), false);
5280 }
5281}
5282
5283namespace llvm {
5284InstructionSelector *
5286 const SPIRVSubtarget &Subtarget,
5287 const RegisterBankInfo &RBI) {
5288 return new SPIRVInstructionSelector(TM, Subtarget, RBI);
5289}
5290} // namespace llvm
unsigned const MachineRegisterInfo * MRI
MachineInstrBuilder & UseMI
#define GET_GLOBALISEL_PREDICATES_INIT
#define GET_GLOBALISEL_TEMPORARIES_INIT
@ Generic
assert(UImm &&(UImm !=~static_cast< T >(0)) &&"Invalid immediate!")
This file declares a class to represent arbitrary precision floating point values and provide a varie...
static bool selectUnmergeValues(MachineInstrBuilder &MIB, const ARMBaseInstrInfo &TII, MachineRegisterInfo &MRI, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI)
MachineBasicBlock & MBB
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
basic Basic Alias true
static GCRegistry::Add< CoreCLRGC > E("coreclr", "CoreCLR-compatible GC")
DXIL Resource Implicit Binding
#define DEBUG_TYPE
Declares convenience wrapper classes for interpreting MachineInstr instances as specific generic oper...
const HexagonInstrInfo * TII
IRTranslator LLVM IR MI
LLVMTypeRef LLVMIntType(unsigned NumBits)
Definition Core.cpp:701
const size_t AbstractManglingParser< Derived, Alloc >::NumOps
#define I(x, y, z)
Definition MD5.cpp:57
Register Reg
Register const TargetRegisterInfo * TRI
Promote Memory to Register
Definition Mem2Reg.cpp:110
MachineInstr unsigned OpIdx
uint64_t IntrinsicInst * II
static StringRef getName(Value *V)
static unsigned getFCmpOpcode(CmpInst::Predicate Pred, unsigned Size)
static APFloat getOneFP(const Type *LLVMFloatTy)
static bool isUSMStorageClass(SPIRV::StorageClass::StorageClass SC)
static bool isASCastInGVar(MachineRegisterInfo *MRI, Register ResVReg)
static bool mayApplyGenericSelection(unsigned Opcode)
static APFloat getZeroFP(const Type *LLVMFloatTy)
std::vector< std::pair< SPIRV::InstructionSet::InstructionSet, uint32_t > > ExtInstList
static bool intrinsicHasSideEffects(Intrinsic::ID ID)
static unsigned getBoolCmpOpcode(unsigned PredNum)
static unsigned getICmpOpcode(unsigned PredNum)
static bool isOpcodeWithNoSideEffects(unsigned Opcode)
static void addMemoryOperands(MachineMemOperand *MemOp, MachineInstrBuilder &MIB, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry &GR)
static bool isConstReg(MachineRegisterInfo *MRI, MachineInstr *OpDef, SmallPtrSet< SPIRVType *, 4 > &Visited)
static unsigned getPtrCmpOpcode(unsigned Pred)
bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
spirv structurize SPIRV
BaseType
A given derived pointer can have multiple base pointers through phi/selects.
This file contains some functions that are useful when dealing with strings.
#define LLVM_DEBUG(...)
Definition Debug.h:114
static TableGen::Emitter::Opt Y("gen-skeleton-entry", EmitSkeleton, "Generate example skeleton entry")
static TableGen::Emitter::OptClass< SkeletonEmitter > X("gen-skeleton-class", "Generate example skeleton class")
BinaryOperator * Mul
static const fltSemantics & IEEEsingle()
Definition APFloat.h:296
static const fltSemantics & IEEEdouble()
Definition APFloat.h:297
static const fltSemantics & IEEEhalf()
Definition APFloat.h:294
static APFloat getOne(const fltSemantics &Sem, bool Negative=false)
Factory for Positive and Negative One.
Definition APFloat.h:1151
static APFloat getZero(const fltSemantics &Sem, bool Negative=false)
Factory for Positive and Negative Zero.
Definition APFloat.h:1142
static APInt getAllOnes(unsigned numBits)
Return an APInt of a specified width with all bits set.
Definition APInt.h:235
uint64_t getZExtValue() const
Get zero extended value.
Definition APInt.h:1549
BlockFrequencyInfo pass uses BlockFrequencyInfoImpl implementation to estimate IR basic block frequen...
Predicate
This enumeration lists the possible predicates for CmpInst subclasses.
Definition InstrTypes.h:676
@ FCMP_OEQ
0 0 0 1 True if ordered and equal
Definition InstrTypes.h:679
@ ICMP_SLT
signed less than
Definition InstrTypes.h:705
@ ICMP_SLE
signed less or equal
Definition InstrTypes.h:706
@ FCMP_OLT
0 1 0 0 True if ordered and less than
Definition InstrTypes.h:682
@ FCMP_ULE
1 1 0 1 True if unordered, less than, or equal
Definition InstrTypes.h:691
@ FCMP_OGT
0 0 1 0 True if ordered and greater than
Definition InstrTypes.h:680
@ FCMP_OGE
0 0 1 1 True if ordered and greater than or equal
Definition InstrTypes.h:681
@ ICMP_UGE
unsigned greater or equal
Definition InstrTypes.h:700
@ ICMP_UGT
unsigned greater than
Definition InstrTypes.h:699
@ ICMP_SGT
signed greater than
Definition InstrTypes.h:703
@ FCMP_ULT
1 1 0 0 True if unordered or less than
Definition InstrTypes.h:690
@ FCMP_ONE
0 1 1 0 True if ordered and operands are unequal
Definition InstrTypes.h:684
@ FCMP_UEQ
1 0 0 1 True if unordered or equal
Definition InstrTypes.h:687
@ ICMP_ULT
unsigned less than
Definition InstrTypes.h:701
@ FCMP_UGT
1 0 1 0 True if unordered or greater than
Definition InstrTypes.h:688
@ FCMP_OLE
0 1 0 1 True if ordered and less than or equal
Definition InstrTypes.h:683
@ FCMP_ORD
0 1 1 1 True if ordered (no nans)
Definition InstrTypes.h:685
@ ICMP_NE
not equal
Definition InstrTypes.h:698
@ ICMP_SGE
signed greater or equal
Definition InstrTypes.h:704
@ FCMP_UNE
1 1 1 0 True if unordered or not equal
Definition InstrTypes.h:692
@ ICMP_ULE
unsigned less or equal
Definition InstrTypes.h:702
@ FCMP_UGE
1 0 1 1 True if unordered, greater than, or equal
Definition InstrTypes.h:689
@ FCMP_UNO
1 0 0 0 True if unordered: isnan(X) | isnan(Y)
Definition InstrTypes.h:686
static LLVM_ABI Constant * getNullValue(Type *Ty)
Constructor to create a '0' constant of arbitrary type.
A debug info location.
Definition DebugLoc.h:123
LLVMContext & getContext() const
getContext - Return a reference to the LLVMContext associated with this function.
Definition Function.cpp:358
Represents a call to an intrinsic.
Intrinsic::ID getIntrinsicID() const
unsigned getAddressSpace() const
Module * getParent()
Get the module that this global value is contained inside of...
@ InternalLinkage
Rename collisions when linking (static functions).
Definition GlobalValue.h:60
static LLVM_ABI IntegerType * get(LLVMContext &C, unsigned NumBits)
This static method is the primary way of constructing an IntegerType.
Definition Type.cpp:318
constexpr bool isScalar() const
static constexpr LLT scalar(unsigned SizeInBits)
Get a low-level scalar or aggregate "bag of bits".
constexpr bool isValid() const
constexpr uint16_t getNumElements() const
Returns the number of elements in a vector LLT.
constexpr bool isVector() const
static constexpr LLT pointer(unsigned AddressSpace, unsigned SizeInBits)
Get a low-level pointer in the given address space.
constexpr bool isPointer() const
static constexpr LLT fixed_vector(unsigned NumElements, unsigned ScalarSizeInBits)
Get a low-level fixed-width vector of some number of elements and element width.
int getNumber() const
MachineBasicBlocks are uniquely numbered at the function level, unless they're not in a MachineFuncti...
LLVM_ABI iterator getFirstNonPHI()
Returns a pointer to the first instruction in this block that is not a PHINode instruction.
const MachineFunction * getParent() const
Return the MachineFunction containing this basic block.
MachineInstrBundleIterator< MachineInstr > iterator
MachineRegisterInfo & getRegInfo()
getRegInfo - Return information about the registers currently in use.
Function & getFunction()
Return the LLVM function that this machine code represents.
Helper class to build MachineInstr.
MachineBasicBlock::iterator getInsertPt()
Current insertion point for new instructions.
MachineInstrBuilder buildInstr(unsigned Opcode)
Build and insert <empty> = Opcode <empty>.
MachineFunction & getMF()
Getter for the function we currently build.
MachineRegisterInfo * getMRI()
Getter for MRI.
const MachineInstrBuilder & addUse(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a virtual register use operand.
const MachineInstrBuilder & addReg(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a new virtual register operand.
const MachineInstrBuilder & addImm(int64_t Val) const
Add a new immediate operand.
const MachineInstrBuilder & add(const MachineOperand &MO) const
bool constrainAllUses(const TargetInstrInfo &TII, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI) const
const MachineInstrBuilder & addMBB(MachineBasicBlock *MBB, unsigned TargetFlags=0) const
const MachineInstrBuilder & addDef(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a virtual register definition operand.
const MachineInstrBuilder & setMIFlags(unsigned Flags) const
MachineInstr * getInstr() const
If conversion operators fail, use this method to get the MachineInstr explicitly.
Representation of each machine instruction.
unsigned getOpcode() const
Returns the opcode of this MachineInstr.
const MachineBasicBlock * getParent() const
unsigned getNumOperands() const
Retuns the total number of operands.
LLVM_ABI void setDesc(const MCInstrDesc &TID)
Replace the instruction descriptor (thus opcode) of the current instruction with a new one.
LLVM_ABI unsigned getNumExplicitDefs() const
Returns the number of non-implicit definitions.
LLVM_ABI const MachineFunction * getMF() const
Return the function that contains the basic block that this instruction belongs to.
const DebugLoc & getDebugLoc() const
Returns the debug location id of this MachineInstr.
LLVM_ABI void removeOperand(unsigned OpNo)
Erase an operand from an instruction, leaving it with one fewer operand than it started with.
const MachineOperand & getOperand(unsigned i) const
A description of a memory reference used in the backend.
@ MOVolatile
The memory access is volatile.
@ MONonTemporal
The memory access is non-temporal.
int64_t getImm() const
bool isReg() const
isReg - Tests if this is a MO_Register operand.
MachineBasicBlock * getMBB() const
Register getReg() const
getReg - Returns the register number.
MachineRegisterInfo - Keep track of information for virtual and physical registers,...
defusechain_instr_iterator< true, false, false, true > use_instr_iterator
use_instr_iterator/use_instr_begin/use_instr_end - Walk all uses of the specified register,...
defusechain_instr_iterator< false, true, false, true > def_instr_iterator
def_instr_iterator/def_instr_begin/def_instr_end - Walk all defs of the specified register,...
LLVM_ABI Register createVirtualRegister(const TargetRegisterClass *RegClass, StringRef Name="")
createVirtualRegister - Create and return a new virtual register in the function with the specified r...
LLVM_ABI void setType(Register VReg, LLT Ty)
Set the low-level type of VReg to Ty.
Analysis providing profile information.
Holds all the information related to register banks.
Wrapper class representing virtual and physical registers.
Definition Register.h:20
constexpr bool isValid() const
Definition Register.h:112
constexpr bool isPhysical() const
Return true if the specified register number is in the physical register namespace.
Definition Register.h:83
SPIRVType * getSPIRVTypeForVReg(Register VReg, const MachineFunction *MF=nullptr) const
Register getOrCreateConstInt(uint64_t Val, MachineInstr &I, SPIRVType *SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
SPIRVType * getResultType(Register VReg, MachineFunction *MF=nullptr)
SPIRVType * getOrCreateSPIRVBoolType(MachineIRBuilder &MIRBuilder, bool EmitIR)
MachineInstr * getOrAddMemAliasingINTELInst(MachineIRBuilder &MIRBuilder, const MDNode *AliasingListMD)
void assignSPIRVTypeToVReg(SPIRVType *Type, Register VReg, const MachineFunction &MF)
Register getOrCreateUndef(MachineInstr &I, SPIRVType *SpvType, const SPIRVInstrInfo &TII)
Register buildGlobalVariable(Register Reg, SPIRVType *BaseType, StringRef Name, const GlobalValue *GV, SPIRV::StorageClass::StorageClass Storage, const MachineInstr *Init, bool IsConst, const std::optional< SPIRV::LinkageType::LinkageType > &LinkageType, MachineIRBuilder &MIRBuilder, bool IsInstSelector)
SPIRVType * changePointerStorageClass(SPIRVType *PtrType, SPIRV::StorageClass::StorageClass SC, MachineInstr &I)
const Type * getTypeForSPIRVType(const SPIRVType *Ty) const
bool isBitcastCompatible(const SPIRVType *Type1, const SPIRVType *Type2) const
unsigned getScalarOrVectorComponentCount(Register VReg) const
SPIRVType * getOrCreateSPIRVFloatType(unsigned BitWidth, MachineInstr &I, const SPIRVInstrInfo &TII)
bool isScalarOrVectorSigned(const SPIRVType *Type) const
Register getOrCreateGlobalVariableWithBinding(const SPIRVType *VarType, uint32_t Set, uint32_t Binding, StringRef Name, MachineIRBuilder &MIRBuilder)
SPIRVType * getOrCreateSPIRVType(const Type *Type, MachineInstr &I, SPIRV::AccessQualifier::AccessQualifier AQ, bool EmitIR)
SPIRVType * getOrCreateSPIRVPointerType(const Type *BaseType, MachineIRBuilder &MIRBuilder, SPIRV::StorageClass::StorageClass SC)
Register buildConstantFP(APFloat Val, MachineIRBuilder &MIRBuilder, SPIRVType *SpvType=nullptr)
SPIRVType * getPointeeType(SPIRVType *PtrType)
void invalidateMachineInstr(MachineInstr *MI)
Register getSPIRVTypeID(const SPIRVType *SpirvType) const
bool isScalarOfType(Register VReg, unsigned TypeOpcode) const
bool findValueAttrs(const MachineInstr *Key, Type *&Ty, StringRef &Name)
SPIRVType * getOrCreateOpTypeSampledImage(SPIRVType *ImageType, MachineIRBuilder &MIRBuilder)
void addGlobalObject(const Value *V, const MachineFunction *MF, Register R)
SPIRVType * getScalarOrVectorComponentType(Register VReg) const
void recordFunctionPointer(const MachineOperand *MO, const Function *F)
bool isAggregateType(SPIRVType *Type) const
const TargetRegisterClass * getRegClass(SPIRVType *SpvType) const
SPIRVType * getOrCreateSPIRVVectorType(SPIRVType *BaseType, unsigned NumElements, MachineIRBuilder &MIRBuilder, bool EmitIR)
bool isScalarOrVectorOfType(Register VReg, unsigned TypeOpcode) const
Register getOrCreateConstIntArray(uint64_t Val, size_t Num, MachineInstr &I, SPIRVType *SpvType, const SPIRVInstrInfo &TII)
MachineFunction * setCurrentFunc(MachineFunction &MF)
Register getOrCreateConstVector(uint64_t Val, MachineInstr &I, SPIRVType *SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
SPIRVType * getOrCreateSPIRVIntegerType(unsigned BitWidth, MachineIRBuilder &MIRBuilder)
Type * getDeducedGlobalValueType(const GlobalValue *Global)
LLT getRegType(SPIRVType *SpvType) const
SPIRV::StorageClass::StorageClass getPointerStorageClass(Register VReg) const
Register getOrCreateConstFP(APFloat Val, MachineInstr &I, SPIRVType *SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
Register getOrCreateConstNullPtr(MachineIRBuilder &MIRBuilder, SPIRVType *SpvType)
unsigned getScalarOrVectorBitWidth(const SPIRVType *Type) const
const SPIRVType * retrieveScalarOrVectorIntType(const SPIRVType *Type) const
bool erase(const MachineInstr *MI)
bool add(SPIRV::IRHandle Handle, const MachineInstr *MI)
Register find(SPIRV::IRHandle Handle, const MachineFunction *MF)
bool isPhysicalSPIRV() const
bool isAtLeastSPIRVVer(VersionTuple VerToCompareTo) const
bool canUseExtInstSet(SPIRV::InstructionSet::InstructionSet E) const
bool isLogicalSPIRV() const
bool canUseExtension(SPIRV::Extension::Extension E) const
std::pair< iterator, bool > insert(PtrType Ptr)
Inserts Ptr if and only if there is no element in the container equal to Ptr.
bool contains(ConstPtrType Ptr) const
SmallPtrSet - This class implements a set which is optimized for holding SmallSize or less elements.
reference emplace_back(ArgTypes &&... Args)
void push_back(const T &Elt)
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
StringRef - Represent a constant reference to a string, i.e.
Definition StringRef.h:55
constexpr size_t size() const
size - Get the string size.
Definition StringRef.h:146
static LLVM_ABI StructType * get(LLVMContext &Context, ArrayRef< Type * > Elements, bool isPacked=false)
This static method is the primary way to create a literal StructType.
Definition Type.cpp:413
The instances of the Type class are immutable: once they are created, they are never changed.
Definition Type.h:45
@ HalfTyID
16-bit floating point type
Definition Type.h:56
@ FloatTyID
32-bit floating point type
Definition Type.h:58
@ DoubleTyID
64-bit floating point type
Definition Type.h:59
Type * getScalarType() const
If this is a vector type, return the element type, otherwise return 'this'.
Definition Type.h:352
bool isStructTy() const
True if this is an instance of StructType.
Definition Type.h:261
bool isAggregateType() const
Return true if the type is an aggregate type.
Definition Type.h:304
TypeID getTypeID() const
Return the type id for the type.
Definition Type.h:136
Value * getOperand(unsigned i) const
Definition User.h:207
bool hasName() const
Definition Value.h:262
LLVM_ABI StringRef getName() const
Return a constant reference to the value's name.
Definition Value.cpp:322
NodeTy * getNextNode()
Get the next node, or nullptr for the list tail.
Definition ilist_node.h:348
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
constexpr char IsConst[]
Key for Kernel::Arg::Metadata::mIsConst.
unsigned ID
LLVM IR allows to use arbitrary numbers as calling convention identifiers.
Definition CallingConv.h:24
NodeAddr< DefNode * > Def
Definition RDFGraph.h:384
NodeAddr< InstrNode * > Instr
Definition RDFGraph.h:389
NodeAddr< UseNode * > Use
Definition RDFGraph.h:385
This is an optimization pass for GlobalISel generic memory operations.
Definition Types.h:26
void buildOpName(Register Target, const StringRef &Name, MachineIRBuilder &MIRBuilder)
bool all_of(R &&range, UnaryPredicate P)
Provide wrappers to std::all_of which take ranges instead of having to pass begin/end explicitly.
Definition STLExtras.h:1737
int64_t getIConstValSext(Register ConstReg, const MachineRegisterInfo *MRI)
MachineInstrBuilder BuildMI(MachineFunction &MF, const MIMetadata &MIMD, const MCInstrDesc &MCID)
Builder interface. Specify how to create the initial instruction itself.
bool isTypeFoldingSupported(unsigned Opcode)
decltype(auto) dyn_cast(const From &Val)
dyn_cast<X> - Return the argument parameter cast to the specified type.
Definition Casting.h:643
FunctionAddr VTableAddr uintptr_t uintptr_t Int32Ty
Definition InstrProf.h:296
void addNumImm(const APInt &Imm, MachineInstrBuilder &MIB)
LLVM_ABI void salvageDebugInfo(const MachineRegisterInfo &MRI, MachineInstr &MI)
Assuming the instruction MI is going to be deleted, attempt to salvage debug users of MI by writing t...
Definition Utils.cpp:1731
LLVM_ABI bool constrainSelectedInstRegOperands(MachineInstr &I, const TargetInstrInfo &TII, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI)
Mutate the newly-selected instruction I to constrain its (possibly generic) virtual register operands...
Definition Utils.cpp:155
bool isPreISelGenericOpcode(unsigned Opcode)
Check whether the given Opcode is a generic opcode that is not supposed to appear after ISel.
unsigned getArrayComponentCount(const MachineRegisterInfo *MRI, const MachineInstr *ResType)
uint64_t getIConstVal(Register ConstReg, const MachineRegisterInfo *MRI)
SmallVector< MachineInstr *, 4 > createContinuedInstructions(MachineIRBuilder &MIRBuilder, unsigned Opcode, unsigned MinWC, unsigned ContinuedOpcode, ArrayRef< Register > Args, Register ReturnRegister, Register TypeID)
SPIRV::MemorySemantics::MemorySemantics getMemSemanticsForStorageClass(SPIRV::StorageClass::StorageClass SC)
constexpr unsigned storageClassToAddressSpace(SPIRV::StorageClass::StorageClass SC)
Definition SPIRVUtils.h:244
MachineBasicBlock::iterator getFirstValidInstructionInsertPoint(MachineBasicBlock &BB)
void buildOpDecorate(Register Reg, MachineIRBuilder &MIRBuilder, SPIRV::Decoration::Decoration Dec, const std::vector< uint32_t > &DecArgs, StringRef StrImm)
MachineBasicBlock::iterator getOpVariableMBBIt(MachineInstr &I)
Register createVirtualRegister(SPIRVType *SpvType, SPIRVGlobalRegistry *GR, MachineRegisterInfo *MRI, const MachineFunction &MF)
MachineInstr * getImm(const MachineOperand &MO, const MachineRegisterInfo *MRI)
Type * toTypedPointer(Type *Ty)
Definition SPIRVUtils.h:458
LLVM_ABI raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
Definition Debug.cpp:207
LLVM_ABI void report_fatal_error(Error Err, bool gen_crash_diag=true)
Definition Error.cpp:163
const MachineInstr SPIRVType
constexpr bool isGenericCastablePtr(SPIRV::StorageClass::StorageClass SC)
Definition SPIRVUtils.h:229
class LLVM_GSL_OWNER SmallVector
Forward declaration of SmallVector so that calculateSmallVectorDefaultInlinedElements can reference s...
MachineInstr * passCopy(MachineInstr *Def, const MachineRegisterInfo *MRI)
bool isa(const From &Val)
isa<X> - Return true if the parameter to the template is an instance of one of the template type argu...
Definition Casting.h:547
std::optional< SPIRV::LinkageType::LinkageType > getSpirvLinkageTypeFor(const SPIRVSubtarget &ST, const GlobalValue &GV)
LLVM_ABI raw_fd_ostream & errs()
This returns a reference to a raw_ostream for standard error.
SPIRV::StorageClass::StorageClass addressSpaceToStorageClass(unsigned AddrSpace, const SPIRVSubtarget &STI)
AtomicOrdering
Atomic ordering for LLVM's memory model.
SPIRV::Scope::Scope getMemScope(LLVMContext &Ctx, SyncScope::ID Id)
InstructionSelector * createSPIRVInstructionSelector(const SPIRVTargetMachine &TM, const SPIRVSubtarget &Subtarget, const RegisterBankInfo &RBI)
std::string getStringValueFromReg(Register Reg, MachineRegisterInfo &MRI)
int64_t foldImm(const MachineOperand &MO, const MachineRegisterInfo *MRI)
DWARFExpression::Operation Op
MachineInstr * getDefInstrMaybeConstant(Register &ConstReg, const MachineRegisterInfo *MRI)
constexpr unsigned BitWidth
decltype(auto) cast(const From &Val)
cast<X> - Return the argument parameter cast to the specified type.
Definition Casting.h:559
bool hasInitializer(const GlobalVariable *GV)
Definition SPIRVUtils.h:346
void addStringImm(const StringRef &Str, MCInst &Inst)
MachineInstr * getVRegDef(MachineRegisterInfo &MRI, Register Reg)
SPIRV::MemorySemantics::MemorySemantics getMemSemantics(AtomicOrdering Ord)
std::string getLinkStringForBuiltIn(SPIRV::BuiltIn::BuiltIn BuiltInValue)
LLVM_ABI bool isTriviallyDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
Check whether an instruction MI is dead: it only defines dead virtual registers, and doesn't have oth...
Definition Utils.cpp:222
#define N