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