1bdd1243dSDimitry Andric //===- SPIRVBuiltins.cpp - SPIR-V Built-in Functions ------------*- C++ -*-===//
2bdd1243dSDimitry Andric //
3bdd1243dSDimitry Andric // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4bdd1243dSDimitry Andric // See https://llvm.org/LICENSE.txt for license information.
5bdd1243dSDimitry Andric // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6bdd1243dSDimitry Andric //
7bdd1243dSDimitry Andric //===----------------------------------------------------------------------===//
8bdd1243dSDimitry Andric //
9bdd1243dSDimitry Andric // This file implements lowering builtin function calls and types using their
10bdd1243dSDimitry Andric // demangled names and TableGen records.
11bdd1243dSDimitry Andric //
12bdd1243dSDimitry Andric //===----------------------------------------------------------------------===//
13bdd1243dSDimitry Andric
14bdd1243dSDimitry Andric #include "SPIRVBuiltins.h"
15bdd1243dSDimitry Andric #include "SPIRV.h"
16bdd1243dSDimitry Andric #include "SPIRVUtils.h"
1706c3fb27SDimitry Andric #include "llvm/ADT/StringExtras.h"
18bdd1243dSDimitry Andric #include "llvm/Analysis/ValueTracking.h"
19bdd1243dSDimitry Andric #include "llvm/IR/IntrinsicsSPIRV.h"
20bdd1243dSDimitry Andric #include <string>
21bdd1243dSDimitry Andric #include <tuple>
22bdd1243dSDimitry Andric
23bdd1243dSDimitry Andric #define DEBUG_TYPE "spirv-builtins"
24bdd1243dSDimitry Andric
25bdd1243dSDimitry Andric namespace llvm {
26bdd1243dSDimitry Andric namespace SPIRV {
27bdd1243dSDimitry Andric #define GET_BuiltinGroup_DECL
28bdd1243dSDimitry Andric #include "SPIRVGenTables.inc"
29bdd1243dSDimitry Andric
30bdd1243dSDimitry Andric struct DemangledBuiltin {
31bdd1243dSDimitry Andric StringRef Name;
32bdd1243dSDimitry Andric InstructionSet::InstructionSet Set;
33bdd1243dSDimitry Andric BuiltinGroup Group;
34bdd1243dSDimitry Andric uint8_t MinNumArgs;
35bdd1243dSDimitry Andric uint8_t MaxNumArgs;
36bdd1243dSDimitry Andric };
37bdd1243dSDimitry Andric
38bdd1243dSDimitry Andric #define GET_DemangledBuiltins_DECL
39bdd1243dSDimitry Andric #define GET_DemangledBuiltins_IMPL
40bdd1243dSDimitry Andric
41bdd1243dSDimitry Andric struct IncomingCall {
42bdd1243dSDimitry Andric const std::string BuiltinName;
43bdd1243dSDimitry Andric const DemangledBuiltin *Builtin;
44bdd1243dSDimitry Andric
45bdd1243dSDimitry Andric const Register ReturnRegister;
46bdd1243dSDimitry Andric const SPIRVType *ReturnType;
47bdd1243dSDimitry Andric const SmallVectorImpl<Register> &Arguments;
48bdd1243dSDimitry Andric
IncomingCallllvm::SPIRV::IncomingCall49bdd1243dSDimitry Andric IncomingCall(const std::string BuiltinName, const DemangledBuiltin *Builtin,
50bdd1243dSDimitry Andric const Register ReturnRegister, const SPIRVType *ReturnType,
51bdd1243dSDimitry Andric const SmallVectorImpl<Register> &Arguments)
52bdd1243dSDimitry Andric : BuiltinName(BuiltinName), Builtin(Builtin),
53bdd1243dSDimitry Andric ReturnRegister(ReturnRegister), ReturnType(ReturnType),
54bdd1243dSDimitry Andric Arguments(Arguments) {}
55bdd1243dSDimitry Andric };
56bdd1243dSDimitry Andric
57bdd1243dSDimitry Andric struct NativeBuiltin {
58bdd1243dSDimitry Andric StringRef Name;
59bdd1243dSDimitry Andric InstructionSet::InstructionSet Set;
60bdd1243dSDimitry Andric uint32_t Opcode;
61bdd1243dSDimitry Andric };
62bdd1243dSDimitry Andric
63bdd1243dSDimitry Andric #define GET_NativeBuiltins_DECL
64bdd1243dSDimitry Andric #define GET_NativeBuiltins_IMPL
65bdd1243dSDimitry Andric
66bdd1243dSDimitry Andric struct GroupBuiltin {
67bdd1243dSDimitry Andric StringRef Name;
68bdd1243dSDimitry Andric uint32_t Opcode;
69bdd1243dSDimitry Andric uint32_t GroupOperation;
70bdd1243dSDimitry Andric bool IsElect;
71bdd1243dSDimitry Andric bool IsAllOrAny;
72bdd1243dSDimitry Andric bool IsAllEqual;
73bdd1243dSDimitry Andric bool IsBallot;
74bdd1243dSDimitry Andric bool IsInverseBallot;
75bdd1243dSDimitry Andric bool IsBallotBitExtract;
76bdd1243dSDimitry Andric bool IsBallotFindBit;
77bdd1243dSDimitry Andric bool IsLogical;
78bdd1243dSDimitry Andric bool NoGroupOperation;
79bdd1243dSDimitry Andric bool HasBoolArg;
80bdd1243dSDimitry Andric };
81bdd1243dSDimitry Andric
82bdd1243dSDimitry Andric #define GET_GroupBuiltins_DECL
83bdd1243dSDimitry Andric #define GET_GroupBuiltins_IMPL
84bdd1243dSDimitry Andric
85bdd1243dSDimitry Andric struct GetBuiltin {
86bdd1243dSDimitry Andric StringRef Name;
87bdd1243dSDimitry Andric InstructionSet::InstructionSet Set;
88bdd1243dSDimitry Andric BuiltIn::BuiltIn Value;
89bdd1243dSDimitry Andric };
90bdd1243dSDimitry Andric
91bdd1243dSDimitry Andric using namespace BuiltIn;
92bdd1243dSDimitry Andric #define GET_GetBuiltins_DECL
93bdd1243dSDimitry Andric #define GET_GetBuiltins_IMPL
94bdd1243dSDimitry Andric
95bdd1243dSDimitry Andric struct ImageQueryBuiltin {
96bdd1243dSDimitry Andric StringRef Name;
97bdd1243dSDimitry Andric InstructionSet::InstructionSet Set;
98bdd1243dSDimitry Andric uint32_t Component;
99bdd1243dSDimitry Andric };
100bdd1243dSDimitry Andric
101bdd1243dSDimitry Andric #define GET_ImageQueryBuiltins_DECL
102bdd1243dSDimitry Andric #define GET_ImageQueryBuiltins_IMPL
103bdd1243dSDimitry Andric
104bdd1243dSDimitry Andric struct ConvertBuiltin {
105bdd1243dSDimitry Andric StringRef Name;
106bdd1243dSDimitry Andric InstructionSet::InstructionSet Set;
107bdd1243dSDimitry Andric bool IsDestinationSigned;
108bdd1243dSDimitry Andric bool IsSaturated;
109bdd1243dSDimitry Andric bool IsRounded;
110bdd1243dSDimitry Andric FPRoundingMode::FPRoundingMode RoundingMode;
111bdd1243dSDimitry Andric };
112bdd1243dSDimitry Andric
113bdd1243dSDimitry Andric struct VectorLoadStoreBuiltin {
114bdd1243dSDimitry Andric StringRef Name;
115bdd1243dSDimitry Andric InstructionSet::InstructionSet Set;
116bdd1243dSDimitry Andric uint32_t Number;
117bdd1243dSDimitry Andric bool IsRounded;
118bdd1243dSDimitry Andric FPRoundingMode::FPRoundingMode RoundingMode;
119bdd1243dSDimitry Andric };
120bdd1243dSDimitry Andric
121bdd1243dSDimitry Andric using namespace FPRoundingMode;
122bdd1243dSDimitry Andric #define GET_ConvertBuiltins_DECL
123bdd1243dSDimitry Andric #define GET_ConvertBuiltins_IMPL
124bdd1243dSDimitry Andric
125bdd1243dSDimitry Andric using namespace InstructionSet;
126bdd1243dSDimitry Andric #define GET_VectorLoadStoreBuiltins_DECL
127bdd1243dSDimitry Andric #define GET_VectorLoadStoreBuiltins_IMPL
128bdd1243dSDimitry Andric
129bdd1243dSDimitry Andric #define GET_CLMemoryScope_DECL
130bdd1243dSDimitry Andric #define GET_CLSamplerAddressingMode_DECL
131bdd1243dSDimitry Andric #define GET_CLMemoryFenceFlags_DECL
132bdd1243dSDimitry Andric #define GET_ExtendedBuiltins_DECL
133bdd1243dSDimitry Andric #include "SPIRVGenTables.inc"
134bdd1243dSDimitry Andric } // namespace SPIRV
135bdd1243dSDimitry Andric
136bdd1243dSDimitry Andric //===----------------------------------------------------------------------===//
137bdd1243dSDimitry Andric // Misc functions for looking up builtins and veryfying requirements using
138bdd1243dSDimitry Andric // TableGen records
139bdd1243dSDimitry Andric //===----------------------------------------------------------------------===//
140bdd1243dSDimitry Andric
141bdd1243dSDimitry Andric /// Looks up the demangled builtin call in the SPIRVBuiltins.td records using
142bdd1243dSDimitry Andric /// the provided \p DemangledCall and specified \p Set.
143bdd1243dSDimitry Andric ///
144bdd1243dSDimitry Andric /// The lookup follows the following algorithm, returning the first successful
145bdd1243dSDimitry Andric /// match:
146bdd1243dSDimitry Andric /// 1. Search with the plain demangled name (expecting a 1:1 match).
147bdd1243dSDimitry Andric /// 2. Search with the prefix before or suffix after the demangled name
148bdd1243dSDimitry Andric /// signyfying the type of the first argument.
149bdd1243dSDimitry Andric ///
150bdd1243dSDimitry Andric /// \returns Wrapper around the demangled call and found builtin definition.
151bdd1243dSDimitry Andric static std::unique_ptr<const SPIRV::IncomingCall>
lookupBuiltin(StringRef DemangledCall,SPIRV::InstructionSet::InstructionSet Set,Register ReturnRegister,const SPIRVType * ReturnType,const SmallVectorImpl<Register> & Arguments)152bdd1243dSDimitry Andric lookupBuiltin(StringRef DemangledCall,
153bdd1243dSDimitry Andric SPIRV::InstructionSet::InstructionSet Set,
154bdd1243dSDimitry Andric Register ReturnRegister, const SPIRVType *ReturnType,
155bdd1243dSDimitry Andric const SmallVectorImpl<Register> &Arguments) {
156bdd1243dSDimitry Andric // Extract the builtin function name and types of arguments from the call
157bdd1243dSDimitry Andric // skeleton.
158bdd1243dSDimitry Andric std::string BuiltinName =
159bdd1243dSDimitry Andric DemangledCall.substr(0, DemangledCall.find('(')).str();
160bdd1243dSDimitry Andric
161bdd1243dSDimitry Andric // Check if the extracted name contains type information between angle
162bdd1243dSDimitry Andric // brackets. If so, the builtin is an instantiated template - needs to have
163bdd1243dSDimitry Andric // the information after angle brackets and return type removed.
164bdd1243dSDimitry Andric if (BuiltinName.find('<') && BuiltinName.back() == '>') {
165bdd1243dSDimitry Andric BuiltinName = BuiltinName.substr(0, BuiltinName.find('<'));
1665f757f3fSDimitry Andric BuiltinName = BuiltinName.substr(BuiltinName.find_last_of(' ') + 1);
167bdd1243dSDimitry Andric }
168bdd1243dSDimitry Andric
169bdd1243dSDimitry Andric // Check if the extracted name begins with "__spirv_ImageSampleExplicitLod"
170bdd1243dSDimitry Andric // contains return type information at the end "_R<type>", if so extract the
171bdd1243dSDimitry Andric // plain builtin name without the type information.
172bdd1243dSDimitry Andric if (StringRef(BuiltinName).contains("__spirv_ImageSampleExplicitLod") &&
173bdd1243dSDimitry Andric StringRef(BuiltinName).contains("_R")) {
174bdd1243dSDimitry Andric BuiltinName = BuiltinName.substr(0, BuiltinName.find("_R"));
175bdd1243dSDimitry Andric }
176bdd1243dSDimitry Andric
177bdd1243dSDimitry Andric SmallVector<StringRef, 10> BuiltinArgumentTypes;
178bdd1243dSDimitry Andric StringRef BuiltinArgs =
179bdd1243dSDimitry Andric DemangledCall.slice(DemangledCall.find('(') + 1, DemangledCall.find(')'));
180bdd1243dSDimitry Andric BuiltinArgs.split(BuiltinArgumentTypes, ',', -1, false);
181bdd1243dSDimitry Andric
182bdd1243dSDimitry Andric // Look up the builtin in the defined set. Start with the plain demangled
183bdd1243dSDimitry Andric // name, expecting a 1:1 match in the defined builtin set.
184bdd1243dSDimitry Andric const SPIRV::DemangledBuiltin *Builtin;
185bdd1243dSDimitry Andric if ((Builtin = SPIRV::lookupBuiltin(BuiltinName, Set)))
186bdd1243dSDimitry Andric return std::make_unique<SPIRV::IncomingCall>(
187bdd1243dSDimitry Andric BuiltinName, Builtin, ReturnRegister, ReturnType, Arguments);
188bdd1243dSDimitry Andric
189bdd1243dSDimitry Andric // If the initial look up was unsuccessful and the demangled call takes at
190bdd1243dSDimitry Andric // least 1 argument, add a prefix or suffix signifying the type of the first
191bdd1243dSDimitry Andric // argument and repeat the search.
192bdd1243dSDimitry Andric if (BuiltinArgumentTypes.size() >= 1) {
193bdd1243dSDimitry Andric char FirstArgumentType = BuiltinArgumentTypes[0][0];
194bdd1243dSDimitry Andric // Prefix to be added to the builtin's name for lookup.
195bdd1243dSDimitry Andric // For example, OpenCL "abs" taking an unsigned value has a prefix "u_".
196bdd1243dSDimitry Andric std::string Prefix;
197bdd1243dSDimitry Andric
198bdd1243dSDimitry Andric switch (FirstArgumentType) {
199bdd1243dSDimitry Andric // Unsigned:
200bdd1243dSDimitry Andric case 'u':
201bdd1243dSDimitry Andric if (Set == SPIRV::InstructionSet::OpenCL_std)
202bdd1243dSDimitry Andric Prefix = "u_";
203bdd1243dSDimitry Andric else if (Set == SPIRV::InstructionSet::GLSL_std_450)
204bdd1243dSDimitry Andric Prefix = "u";
205bdd1243dSDimitry Andric break;
206bdd1243dSDimitry Andric // Signed:
207bdd1243dSDimitry Andric case 'c':
208bdd1243dSDimitry Andric case 's':
209bdd1243dSDimitry Andric case 'i':
210bdd1243dSDimitry Andric case 'l':
211bdd1243dSDimitry Andric if (Set == SPIRV::InstructionSet::OpenCL_std)
212bdd1243dSDimitry Andric Prefix = "s_";
213bdd1243dSDimitry Andric else if (Set == SPIRV::InstructionSet::GLSL_std_450)
214bdd1243dSDimitry Andric Prefix = "s";
215bdd1243dSDimitry Andric break;
216bdd1243dSDimitry Andric // Floating-point:
217bdd1243dSDimitry Andric case 'f':
218bdd1243dSDimitry Andric case 'd':
219bdd1243dSDimitry Andric case 'h':
220bdd1243dSDimitry Andric if (Set == SPIRV::InstructionSet::OpenCL_std ||
221bdd1243dSDimitry Andric Set == SPIRV::InstructionSet::GLSL_std_450)
222bdd1243dSDimitry Andric Prefix = "f";
223bdd1243dSDimitry Andric break;
224bdd1243dSDimitry Andric }
225bdd1243dSDimitry Andric
226bdd1243dSDimitry Andric // If argument-type name prefix was added, look up the builtin again.
227bdd1243dSDimitry Andric if (!Prefix.empty() &&
228bdd1243dSDimitry Andric (Builtin = SPIRV::lookupBuiltin(Prefix + BuiltinName, Set)))
229bdd1243dSDimitry Andric return std::make_unique<SPIRV::IncomingCall>(
230bdd1243dSDimitry Andric BuiltinName, Builtin, ReturnRegister, ReturnType, Arguments);
231bdd1243dSDimitry Andric
232bdd1243dSDimitry Andric // If lookup with a prefix failed, find a suffix to be added to the
233bdd1243dSDimitry Andric // builtin's name for lookup. For example, OpenCL "group_reduce_max" taking
234bdd1243dSDimitry Andric // an unsigned value has a suffix "u".
235bdd1243dSDimitry Andric std::string Suffix;
236bdd1243dSDimitry Andric
237bdd1243dSDimitry Andric switch (FirstArgumentType) {
238bdd1243dSDimitry Andric // Unsigned:
239bdd1243dSDimitry Andric case 'u':
240bdd1243dSDimitry Andric Suffix = "u";
241bdd1243dSDimitry Andric break;
242bdd1243dSDimitry Andric // Signed:
243bdd1243dSDimitry Andric case 'c':
244bdd1243dSDimitry Andric case 's':
245bdd1243dSDimitry Andric case 'i':
246bdd1243dSDimitry Andric case 'l':
247bdd1243dSDimitry Andric Suffix = "s";
248bdd1243dSDimitry Andric break;
249bdd1243dSDimitry Andric // Floating-point:
250bdd1243dSDimitry Andric case 'f':
251bdd1243dSDimitry Andric case 'd':
252bdd1243dSDimitry Andric case 'h':
253bdd1243dSDimitry Andric Suffix = "f";
254bdd1243dSDimitry Andric break;
255bdd1243dSDimitry Andric }
256bdd1243dSDimitry Andric
257bdd1243dSDimitry Andric // If argument-type name suffix was added, look up the builtin again.
258bdd1243dSDimitry Andric if (!Suffix.empty() &&
259bdd1243dSDimitry Andric (Builtin = SPIRV::lookupBuiltin(BuiltinName + Suffix, Set)))
260bdd1243dSDimitry Andric return std::make_unique<SPIRV::IncomingCall>(
261bdd1243dSDimitry Andric BuiltinName, Builtin, ReturnRegister, ReturnType, Arguments);
262bdd1243dSDimitry Andric }
263bdd1243dSDimitry Andric
264bdd1243dSDimitry Andric // No builtin with such name was found in the set.
265bdd1243dSDimitry Andric return nullptr;
266bdd1243dSDimitry Andric }
267bdd1243dSDimitry Andric
268bdd1243dSDimitry Andric //===----------------------------------------------------------------------===//
269bdd1243dSDimitry Andric // Helper functions for building misc instructions
270bdd1243dSDimitry Andric //===----------------------------------------------------------------------===//
271bdd1243dSDimitry Andric
272bdd1243dSDimitry Andric /// Helper function building either a resulting scalar or vector bool register
273bdd1243dSDimitry Andric /// depending on the expected \p ResultType.
274bdd1243dSDimitry Andric ///
275bdd1243dSDimitry Andric /// \returns Tuple of the resulting register and its type.
276bdd1243dSDimitry Andric static std::tuple<Register, SPIRVType *>
buildBoolRegister(MachineIRBuilder & MIRBuilder,const SPIRVType * ResultType,SPIRVGlobalRegistry * GR)277bdd1243dSDimitry Andric buildBoolRegister(MachineIRBuilder &MIRBuilder, const SPIRVType *ResultType,
278bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
279bdd1243dSDimitry Andric LLT Type;
280bdd1243dSDimitry Andric SPIRVType *BoolType = GR->getOrCreateSPIRVBoolType(MIRBuilder);
281bdd1243dSDimitry Andric
282bdd1243dSDimitry Andric if (ResultType->getOpcode() == SPIRV::OpTypeVector) {
283bdd1243dSDimitry Andric unsigned VectorElements = ResultType->getOperand(2).getImm();
284bdd1243dSDimitry Andric BoolType =
285bdd1243dSDimitry Andric GR->getOrCreateSPIRVVectorType(BoolType, VectorElements, MIRBuilder);
286bdd1243dSDimitry Andric const FixedVectorType *LLVMVectorType =
287bdd1243dSDimitry Andric cast<FixedVectorType>(GR->getTypeForSPIRVType(BoolType));
288bdd1243dSDimitry Andric Type = LLT::vector(LLVMVectorType->getElementCount(), 1);
289bdd1243dSDimitry Andric } else {
290bdd1243dSDimitry Andric Type = LLT::scalar(1);
291bdd1243dSDimitry Andric }
292bdd1243dSDimitry Andric
293bdd1243dSDimitry Andric Register ResultRegister =
294bdd1243dSDimitry Andric MIRBuilder.getMRI()->createGenericVirtualRegister(Type);
29506c3fb27SDimitry Andric MIRBuilder.getMRI()->setRegClass(ResultRegister, &SPIRV::IDRegClass);
296bdd1243dSDimitry Andric GR->assignSPIRVTypeToVReg(BoolType, ResultRegister, MIRBuilder.getMF());
297bdd1243dSDimitry Andric return std::make_tuple(ResultRegister, BoolType);
298bdd1243dSDimitry Andric }
299bdd1243dSDimitry Andric
300bdd1243dSDimitry Andric /// Helper function for building either a vector or scalar select instruction
301bdd1243dSDimitry Andric /// depending on the expected \p ResultType.
buildSelectInst(MachineIRBuilder & MIRBuilder,Register ReturnRegister,Register SourceRegister,const SPIRVType * ReturnType,SPIRVGlobalRegistry * GR)302bdd1243dSDimitry Andric static bool buildSelectInst(MachineIRBuilder &MIRBuilder,
303bdd1243dSDimitry Andric Register ReturnRegister, Register SourceRegister,
304bdd1243dSDimitry Andric const SPIRVType *ReturnType,
305bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
306bdd1243dSDimitry Andric Register TrueConst, FalseConst;
307bdd1243dSDimitry Andric
308bdd1243dSDimitry Andric if (ReturnType->getOpcode() == SPIRV::OpTypeVector) {
309bdd1243dSDimitry Andric unsigned Bits = GR->getScalarOrVectorBitWidth(ReturnType);
31006c3fb27SDimitry Andric uint64_t AllOnes = APInt::getAllOnes(Bits).getZExtValue();
311bdd1243dSDimitry Andric TrueConst = GR->getOrCreateConsIntVector(AllOnes, MIRBuilder, ReturnType);
312bdd1243dSDimitry Andric FalseConst = GR->getOrCreateConsIntVector(0, MIRBuilder, ReturnType);
313bdd1243dSDimitry Andric } else {
314bdd1243dSDimitry Andric TrueConst = GR->buildConstantInt(1, MIRBuilder, ReturnType);
315bdd1243dSDimitry Andric FalseConst = GR->buildConstantInt(0, MIRBuilder, ReturnType);
316bdd1243dSDimitry Andric }
317bdd1243dSDimitry Andric return MIRBuilder.buildSelect(ReturnRegister, SourceRegister, TrueConst,
318bdd1243dSDimitry Andric FalseConst);
319bdd1243dSDimitry Andric }
320bdd1243dSDimitry Andric
321bdd1243dSDimitry Andric /// Helper function for building a load instruction loading into the
322bdd1243dSDimitry Andric /// \p DestinationReg.
buildLoadInst(SPIRVType * BaseType,Register PtrRegister,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR,LLT LowLevelType,Register DestinationReg=Register (0))323bdd1243dSDimitry Andric static Register buildLoadInst(SPIRVType *BaseType, Register PtrRegister,
324bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
325bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR, LLT LowLevelType,
326bdd1243dSDimitry Andric Register DestinationReg = Register(0)) {
327bdd1243dSDimitry Andric MachineRegisterInfo *MRI = MIRBuilder.getMRI();
328bdd1243dSDimitry Andric if (!DestinationReg.isValid()) {
329bdd1243dSDimitry Andric DestinationReg = MRI->createVirtualRegister(&SPIRV::IDRegClass);
330bdd1243dSDimitry Andric MRI->setType(DestinationReg, LLT::scalar(32));
331bdd1243dSDimitry Andric GR->assignSPIRVTypeToVReg(BaseType, DestinationReg, MIRBuilder.getMF());
332bdd1243dSDimitry Andric }
333bdd1243dSDimitry Andric // TODO: consider using correct address space and alignment (p0 is canonical
334bdd1243dSDimitry Andric // type for selection though).
335bdd1243dSDimitry Andric MachinePointerInfo PtrInfo = MachinePointerInfo();
336bdd1243dSDimitry Andric MIRBuilder.buildLoad(DestinationReg, PtrRegister, PtrInfo, Align());
337bdd1243dSDimitry Andric return DestinationReg;
338bdd1243dSDimitry Andric }
339bdd1243dSDimitry Andric
340bdd1243dSDimitry Andric /// Helper function for building a load instruction for loading a builtin global
341bdd1243dSDimitry Andric /// variable of \p BuiltinValue value.
buildBuiltinVariableLoad(MachineIRBuilder & MIRBuilder,SPIRVType * VariableType,SPIRVGlobalRegistry * GR,SPIRV::BuiltIn::BuiltIn BuiltinValue,LLT LLType,Register Reg=Register (0))342bdd1243dSDimitry Andric static Register buildBuiltinVariableLoad(MachineIRBuilder &MIRBuilder,
343bdd1243dSDimitry Andric SPIRVType *VariableType,
344bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR,
345bdd1243dSDimitry Andric SPIRV::BuiltIn::BuiltIn BuiltinValue,
346bdd1243dSDimitry Andric LLT LLType,
347bdd1243dSDimitry Andric Register Reg = Register(0)) {
348bdd1243dSDimitry Andric Register NewRegister =
349bdd1243dSDimitry Andric MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::IDRegClass);
350bdd1243dSDimitry Andric MIRBuilder.getMRI()->setType(NewRegister,
351bdd1243dSDimitry Andric LLT::pointer(0, GR->getPointerSize()));
352bdd1243dSDimitry Andric SPIRVType *PtrType = GR->getOrCreateSPIRVPointerType(
353bdd1243dSDimitry Andric VariableType, MIRBuilder, SPIRV::StorageClass::Input);
354bdd1243dSDimitry Andric GR->assignSPIRVTypeToVReg(PtrType, NewRegister, MIRBuilder.getMF());
355bdd1243dSDimitry Andric
356bdd1243dSDimitry Andric // Set up the global OpVariable with the necessary builtin decorations.
357bdd1243dSDimitry Andric Register Variable = GR->buildGlobalVariable(
358bdd1243dSDimitry Andric NewRegister, PtrType, getLinkStringForBuiltIn(BuiltinValue), nullptr,
359bdd1243dSDimitry Andric SPIRV::StorageClass::Input, nullptr, true, true,
360bdd1243dSDimitry Andric SPIRV::LinkageType::Import, MIRBuilder, false);
361bdd1243dSDimitry Andric
362bdd1243dSDimitry Andric // Load the value from the global variable.
363bdd1243dSDimitry Andric Register LoadedRegister =
364bdd1243dSDimitry Andric buildLoadInst(VariableType, Variable, MIRBuilder, GR, LLType, Reg);
365bdd1243dSDimitry Andric MIRBuilder.getMRI()->setType(LoadedRegister, LLType);
366bdd1243dSDimitry Andric return LoadedRegister;
367bdd1243dSDimitry Andric }
368bdd1243dSDimitry Andric
369bdd1243dSDimitry Andric /// Helper external function for inserting ASSIGN_TYPE instuction between \p Reg
370bdd1243dSDimitry Andric /// and its definition, set the new register as a destination of the definition,
371bdd1243dSDimitry Andric /// assign SPIRVType to both registers. If SpirvTy is provided, use it as
372bdd1243dSDimitry Andric /// SPIRVType in ASSIGN_TYPE, otherwise create it from \p Ty. Defined in
373bdd1243dSDimitry Andric /// SPIRVPreLegalizer.cpp.
374bdd1243dSDimitry Andric extern Register insertAssignInstr(Register Reg, Type *Ty, SPIRVType *SpirvTy,
375bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR,
376bdd1243dSDimitry Andric MachineIRBuilder &MIB,
377bdd1243dSDimitry Andric MachineRegisterInfo &MRI);
378bdd1243dSDimitry Andric
379bdd1243dSDimitry Andric // TODO: Move to TableGen.
380bdd1243dSDimitry Andric static SPIRV::MemorySemantics::MemorySemantics
getSPIRVMemSemantics(std::memory_order MemOrder)381bdd1243dSDimitry Andric getSPIRVMemSemantics(std::memory_order MemOrder) {
382bdd1243dSDimitry Andric switch (MemOrder) {
383bdd1243dSDimitry Andric case std::memory_order::memory_order_relaxed:
384bdd1243dSDimitry Andric return SPIRV::MemorySemantics::None;
385bdd1243dSDimitry Andric case std::memory_order::memory_order_acquire:
386bdd1243dSDimitry Andric return SPIRV::MemorySemantics::Acquire;
387bdd1243dSDimitry Andric case std::memory_order::memory_order_release:
388bdd1243dSDimitry Andric return SPIRV::MemorySemantics::Release;
389bdd1243dSDimitry Andric case std::memory_order::memory_order_acq_rel:
390bdd1243dSDimitry Andric return SPIRV::MemorySemantics::AcquireRelease;
391bdd1243dSDimitry Andric case std::memory_order::memory_order_seq_cst:
392bdd1243dSDimitry Andric return SPIRV::MemorySemantics::SequentiallyConsistent;
393bdd1243dSDimitry Andric default:
394bdd1243dSDimitry Andric llvm_unreachable("Unknown CL memory scope");
395bdd1243dSDimitry Andric }
396bdd1243dSDimitry Andric }
397bdd1243dSDimitry Andric
getSPIRVScope(SPIRV::CLMemoryScope ClScope)398bdd1243dSDimitry Andric static SPIRV::Scope::Scope getSPIRVScope(SPIRV::CLMemoryScope ClScope) {
399bdd1243dSDimitry Andric switch (ClScope) {
400bdd1243dSDimitry Andric case SPIRV::CLMemoryScope::memory_scope_work_item:
401bdd1243dSDimitry Andric return SPIRV::Scope::Invocation;
402bdd1243dSDimitry Andric case SPIRV::CLMemoryScope::memory_scope_work_group:
403bdd1243dSDimitry Andric return SPIRV::Scope::Workgroup;
404bdd1243dSDimitry Andric case SPIRV::CLMemoryScope::memory_scope_device:
405bdd1243dSDimitry Andric return SPIRV::Scope::Device;
406bdd1243dSDimitry Andric case SPIRV::CLMemoryScope::memory_scope_all_svm_devices:
407bdd1243dSDimitry Andric return SPIRV::Scope::CrossDevice;
408bdd1243dSDimitry Andric case SPIRV::CLMemoryScope::memory_scope_sub_group:
409bdd1243dSDimitry Andric return SPIRV::Scope::Subgroup;
410bdd1243dSDimitry Andric }
411bdd1243dSDimitry Andric llvm_unreachable("Unknown CL memory scope");
412bdd1243dSDimitry Andric }
413bdd1243dSDimitry Andric
buildConstantIntReg(uint64_t Val,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR,unsigned BitWidth=32)414bdd1243dSDimitry Andric static Register buildConstantIntReg(uint64_t Val, MachineIRBuilder &MIRBuilder,
415bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR,
416bdd1243dSDimitry Andric unsigned BitWidth = 32) {
417bdd1243dSDimitry Andric SPIRVType *IntType = GR->getOrCreateSPIRVIntegerType(BitWidth, MIRBuilder);
418bdd1243dSDimitry Andric return GR->buildConstantInt(Val, MIRBuilder, IntType);
419bdd1243dSDimitry Andric }
420bdd1243dSDimitry Andric
buildScopeReg(Register CLScopeRegister,SPIRV::Scope::Scope Scope,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR,MachineRegisterInfo * MRI)421bdd1243dSDimitry Andric static Register buildScopeReg(Register CLScopeRegister,
42206c3fb27SDimitry Andric SPIRV::Scope::Scope Scope,
423bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
424bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR,
42506c3fb27SDimitry Andric MachineRegisterInfo *MRI) {
42606c3fb27SDimitry Andric if (CLScopeRegister.isValid()) {
427bdd1243dSDimitry Andric auto CLScope =
428bdd1243dSDimitry Andric static_cast<SPIRV::CLMemoryScope>(getIConstVal(CLScopeRegister, MRI));
42906c3fb27SDimitry Andric Scope = getSPIRVScope(CLScope);
430bdd1243dSDimitry Andric
43106c3fb27SDimitry Andric if (CLScope == static_cast<unsigned>(Scope)) {
43206c3fb27SDimitry Andric MRI->setRegClass(CLScopeRegister, &SPIRV::IDRegClass);
433bdd1243dSDimitry Andric return CLScopeRegister;
43406c3fb27SDimitry Andric }
43506c3fb27SDimitry Andric }
436bdd1243dSDimitry Andric return buildConstantIntReg(Scope, MIRBuilder, GR);
437bdd1243dSDimitry Andric }
438bdd1243dSDimitry Andric
buildMemSemanticsReg(Register SemanticsRegister,Register PtrRegister,unsigned & Semantics,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)439bdd1243dSDimitry Andric static Register buildMemSemanticsReg(Register SemanticsRegister,
44006c3fb27SDimitry Andric Register PtrRegister, unsigned &Semantics,
44106c3fb27SDimitry Andric MachineIRBuilder &MIRBuilder,
442bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
44306c3fb27SDimitry Andric if (SemanticsRegister.isValid()) {
44406c3fb27SDimitry Andric MachineRegisterInfo *MRI = MIRBuilder.getMRI();
445bdd1243dSDimitry Andric std::memory_order Order =
446bdd1243dSDimitry Andric static_cast<std::memory_order>(getIConstVal(SemanticsRegister, MRI));
44706c3fb27SDimitry Andric Semantics =
448bdd1243dSDimitry Andric getSPIRVMemSemantics(Order) |
449bdd1243dSDimitry Andric getMemSemanticsForStorageClass(GR->getPointerStorageClass(PtrRegister));
450bdd1243dSDimitry Andric
45106c3fb27SDimitry Andric if (Order == Semantics) {
45206c3fb27SDimitry Andric MRI->setRegClass(SemanticsRegister, &SPIRV::IDRegClass);
453bdd1243dSDimitry Andric return SemanticsRegister;
45406c3fb27SDimitry Andric }
45506c3fb27SDimitry Andric }
45606c3fb27SDimitry Andric return buildConstantIntReg(Semantics, MIRBuilder, GR);
457bdd1243dSDimitry Andric }
458bdd1243dSDimitry Andric
459bdd1243dSDimitry Andric /// Helper function for translating atomic init to OpStore.
buildAtomicInitInst(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder)460bdd1243dSDimitry Andric static bool buildAtomicInitInst(const SPIRV::IncomingCall *Call,
461bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder) {
462bdd1243dSDimitry Andric assert(Call->Arguments.size() == 2 &&
463bdd1243dSDimitry Andric "Need 2 arguments for atomic init translation");
46406c3fb27SDimitry Andric MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
46506c3fb27SDimitry Andric MIRBuilder.getMRI()->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
466bdd1243dSDimitry Andric MIRBuilder.buildInstr(SPIRV::OpStore)
467bdd1243dSDimitry Andric .addUse(Call->Arguments[0])
468bdd1243dSDimitry Andric .addUse(Call->Arguments[1]);
469bdd1243dSDimitry Andric return true;
470bdd1243dSDimitry Andric }
471bdd1243dSDimitry Andric
472bdd1243dSDimitry Andric /// Helper function for building an atomic load instruction.
buildAtomicLoadInst(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)473bdd1243dSDimitry Andric static bool buildAtomicLoadInst(const SPIRV::IncomingCall *Call,
474bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
475bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
476bdd1243dSDimitry Andric Register PtrRegister = Call->Arguments[0];
47706c3fb27SDimitry Andric MIRBuilder.getMRI()->setRegClass(PtrRegister, &SPIRV::IDRegClass);
478bdd1243dSDimitry Andric // TODO: if true insert call to __translate_ocl_memory_sccope before
479bdd1243dSDimitry Andric // OpAtomicLoad and the function implementation. We can use Translator's
480bdd1243dSDimitry Andric // output for transcoding/atomic_explicit_arguments.cl as an example.
481bdd1243dSDimitry Andric Register ScopeRegister;
48206c3fb27SDimitry Andric if (Call->Arguments.size() > 1) {
483bdd1243dSDimitry Andric ScopeRegister = Call->Arguments[1];
48406c3fb27SDimitry Andric MIRBuilder.getMRI()->setRegClass(ScopeRegister, &SPIRV::IDRegClass);
48506c3fb27SDimitry Andric } else
486bdd1243dSDimitry Andric ScopeRegister = buildConstantIntReg(SPIRV::Scope::Device, MIRBuilder, GR);
487bdd1243dSDimitry Andric
488bdd1243dSDimitry Andric Register MemSemanticsReg;
489bdd1243dSDimitry Andric if (Call->Arguments.size() > 2) {
490bdd1243dSDimitry Andric // TODO: Insert call to __translate_ocl_memory_order before OpAtomicLoad.
491bdd1243dSDimitry Andric MemSemanticsReg = Call->Arguments[2];
49206c3fb27SDimitry Andric MIRBuilder.getMRI()->setRegClass(MemSemanticsReg, &SPIRV::IDRegClass);
493bdd1243dSDimitry Andric } else {
494bdd1243dSDimitry Andric int Semantics =
495bdd1243dSDimitry Andric SPIRV::MemorySemantics::SequentiallyConsistent |
496bdd1243dSDimitry Andric getMemSemanticsForStorageClass(GR->getPointerStorageClass(PtrRegister));
497bdd1243dSDimitry Andric MemSemanticsReg = buildConstantIntReg(Semantics, MIRBuilder, GR);
498bdd1243dSDimitry Andric }
499bdd1243dSDimitry Andric
500bdd1243dSDimitry Andric MIRBuilder.buildInstr(SPIRV::OpAtomicLoad)
501bdd1243dSDimitry Andric .addDef(Call->ReturnRegister)
502bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(Call->ReturnType))
503bdd1243dSDimitry Andric .addUse(PtrRegister)
504bdd1243dSDimitry Andric .addUse(ScopeRegister)
505bdd1243dSDimitry Andric .addUse(MemSemanticsReg);
506bdd1243dSDimitry Andric return true;
507bdd1243dSDimitry Andric }
508bdd1243dSDimitry Andric
509bdd1243dSDimitry Andric /// Helper function for building an atomic store instruction.
buildAtomicStoreInst(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)510bdd1243dSDimitry Andric static bool buildAtomicStoreInst(const SPIRV::IncomingCall *Call,
511bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
512bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
513bdd1243dSDimitry Andric Register ScopeRegister =
514bdd1243dSDimitry Andric buildConstantIntReg(SPIRV::Scope::Device, MIRBuilder, GR);
515bdd1243dSDimitry Andric Register PtrRegister = Call->Arguments[0];
51606c3fb27SDimitry Andric MIRBuilder.getMRI()->setRegClass(PtrRegister, &SPIRV::IDRegClass);
517bdd1243dSDimitry Andric int Semantics =
518bdd1243dSDimitry Andric SPIRV::MemorySemantics::SequentiallyConsistent |
519bdd1243dSDimitry Andric getMemSemanticsForStorageClass(GR->getPointerStorageClass(PtrRegister));
520bdd1243dSDimitry Andric Register MemSemanticsReg = buildConstantIntReg(Semantics, MIRBuilder, GR);
52106c3fb27SDimitry Andric MIRBuilder.getMRI()->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
522bdd1243dSDimitry Andric MIRBuilder.buildInstr(SPIRV::OpAtomicStore)
523bdd1243dSDimitry Andric .addUse(PtrRegister)
524bdd1243dSDimitry Andric .addUse(ScopeRegister)
525bdd1243dSDimitry Andric .addUse(MemSemanticsReg)
526bdd1243dSDimitry Andric .addUse(Call->Arguments[1]);
527bdd1243dSDimitry Andric return true;
528bdd1243dSDimitry Andric }
529bdd1243dSDimitry Andric
530bdd1243dSDimitry Andric /// Helper function for building an atomic compare-exchange instruction.
buildAtomicCompareExchangeInst(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)531bdd1243dSDimitry Andric static bool buildAtomicCompareExchangeInst(const SPIRV::IncomingCall *Call,
532bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
533bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
534bdd1243dSDimitry Andric const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
535bdd1243dSDimitry Andric unsigned Opcode =
536bdd1243dSDimitry Andric SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
537bdd1243dSDimitry Andric bool IsCmpxchg = Call->Builtin->Name.contains("cmpxchg");
538bdd1243dSDimitry Andric MachineRegisterInfo *MRI = MIRBuilder.getMRI();
539bdd1243dSDimitry Andric
540bdd1243dSDimitry Andric Register ObjectPtr = Call->Arguments[0]; // Pointer (volatile A *object.)
541bdd1243dSDimitry Andric Register ExpectedArg = Call->Arguments[1]; // Comparator (C* expected).
542bdd1243dSDimitry Andric Register Desired = Call->Arguments[2]; // Value (C Desired).
54306c3fb27SDimitry Andric MRI->setRegClass(ObjectPtr, &SPIRV::IDRegClass);
54406c3fb27SDimitry Andric MRI->setRegClass(ExpectedArg, &SPIRV::IDRegClass);
54506c3fb27SDimitry Andric MRI->setRegClass(Desired, &SPIRV::IDRegClass);
546bdd1243dSDimitry Andric SPIRVType *SpvDesiredTy = GR->getSPIRVTypeForVReg(Desired);
547bdd1243dSDimitry Andric LLT DesiredLLT = MRI->getType(Desired);
548bdd1243dSDimitry Andric
549bdd1243dSDimitry Andric assert(GR->getSPIRVTypeForVReg(ObjectPtr)->getOpcode() ==
550bdd1243dSDimitry Andric SPIRV::OpTypePointer);
551bdd1243dSDimitry Andric unsigned ExpectedType = GR->getSPIRVTypeForVReg(ExpectedArg)->getOpcode();
552bdd1243dSDimitry Andric assert(IsCmpxchg ? ExpectedType == SPIRV::OpTypeInt
553bdd1243dSDimitry Andric : ExpectedType == SPIRV::OpTypePointer);
554bdd1243dSDimitry Andric assert(GR->isScalarOfType(Desired, SPIRV::OpTypeInt));
555bdd1243dSDimitry Andric
556bdd1243dSDimitry Andric SPIRVType *SpvObjectPtrTy = GR->getSPIRVTypeForVReg(ObjectPtr);
557bdd1243dSDimitry Andric assert(SpvObjectPtrTy->getOperand(2).isReg() && "SPIRV type is expected");
558bdd1243dSDimitry Andric auto StorageClass = static_cast<SPIRV::StorageClass::StorageClass>(
559bdd1243dSDimitry Andric SpvObjectPtrTy->getOperand(1).getImm());
560bdd1243dSDimitry Andric auto MemSemStorage = getMemSemanticsForStorageClass(StorageClass);
561bdd1243dSDimitry Andric
562bdd1243dSDimitry Andric Register MemSemEqualReg;
563bdd1243dSDimitry Andric Register MemSemUnequalReg;
564bdd1243dSDimitry Andric uint64_t MemSemEqual =
565bdd1243dSDimitry Andric IsCmpxchg
566bdd1243dSDimitry Andric ? SPIRV::MemorySemantics::None
567bdd1243dSDimitry Andric : SPIRV::MemorySemantics::SequentiallyConsistent | MemSemStorage;
568bdd1243dSDimitry Andric uint64_t MemSemUnequal =
569bdd1243dSDimitry Andric IsCmpxchg
570bdd1243dSDimitry Andric ? SPIRV::MemorySemantics::None
571bdd1243dSDimitry Andric : SPIRV::MemorySemantics::SequentiallyConsistent | MemSemStorage;
572bdd1243dSDimitry Andric if (Call->Arguments.size() >= 4) {
573bdd1243dSDimitry Andric assert(Call->Arguments.size() >= 5 &&
574bdd1243dSDimitry Andric "Need 5+ args for explicit atomic cmpxchg");
575bdd1243dSDimitry Andric auto MemOrdEq =
576bdd1243dSDimitry Andric static_cast<std::memory_order>(getIConstVal(Call->Arguments[3], MRI));
577bdd1243dSDimitry Andric auto MemOrdNeq =
578bdd1243dSDimitry Andric static_cast<std::memory_order>(getIConstVal(Call->Arguments[4], MRI));
579bdd1243dSDimitry Andric MemSemEqual = getSPIRVMemSemantics(MemOrdEq) | MemSemStorage;
580bdd1243dSDimitry Andric MemSemUnequal = getSPIRVMemSemantics(MemOrdNeq) | MemSemStorage;
581bdd1243dSDimitry Andric if (MemOrdEq == MemSemEqual)
582bdd1243dSDimitry Andric MemSemEqualReg = Call->Arguments[3];
583bdd1243dSDimitry Andric if (MemOrdNeq == MemSemEqual)
584bdd1243dSDimitry Andric MemSemUnequalReg = Call->Arguments[4];
58506c3fb27SDimitry Andric MRI->setRegClass(Call->Arguments[3], &SPIRV::IDRegClass);
58606c3fb27SDimitry Andric MRI->setRegClass(Call->Arguments[4], &SPIRV::IDRegClass);
587bdd1243dSDimitry Andric }
588bdd1243dSDimitry Andric if (!MemSemEqualReg.isValid())
589bdd1243dSDimitry Andric MemSemEqualReg = buildConstantIntReg(MemSemEqual, MIRBuilder, GR);
590bdd1243dSDimitry Andric if (!MemSemUnequalReg.isValid())
591bdd1243dSDimitry Andric MemSemUnequalReg = buildConstantIntReg(MemSemUnequal, MIRBuilder, GR);
592bdd1243dSDimitry Andric
593bdd1243dSDimitry Andric Register ScopeReg;
594bdd1243dSDimitry Andric auto Scope = IsCmpxchg ? SPIRV::Scope::Workgroup : SPIRV::Scope::Device;
595bdd1243dSDimitry Andric if (Call->Arguments.size() >= 6) {
596bdd1243dSDimitry Andric assert(Call->Arguments.size() == 6 &&
597bdd1243dSDimitry Andric "Extra args for explicit atomic cmpxchg");
598bdd1243dSDimitry Andric auto ClScope = static_cast<SPIRV::CLMemoryScope>(
599bdd1243dSDimitry Andric getIConstVal(Call->Arguments[5], MRI));
600bdd1243dSDimitry Andric Scope = getSPIRVScope(ClScope);
601bdd1243dSDimitry Andric if (ClScope == static_cast<unsigned>(Scope))
602bdd1243dSDimitry Andric ScopeReg = Call->Arguments[5];
60306c3fb27SDimitry Andric MRI->setRegClass(Call->Arguments[5], &SPIRV::IDRegClass);
604bdd1243dSDimitry Andric }
605bdd1243dSDimitry Andric if (!ScopeReg.isValid())
606bdd1243dSDimitry Andric ScopeReg = buildConstantIntReg(Scope, MIRBuilder, GR);
607bdd1243dSDimitry Andric
608bdd1243dSDimitry Andric Register Expected = IsCmpxchg
609bdd1243dSDimitry Andric ? ExpectedArg
610bdd1243dSDimitry Andric : buildLoadInst(SpvDesiredTy, ExpectedArg, MIRBuilder,
611bdd1243dSDimitry Andric GR, LLT::scalar(32));
612bdd1243dSDimitry Andric MRI->setType(Expected, DesiredLLT);
613bdd1243dSDimitry Andric Register Tmp = !IsCmpxchg ? MRI->createGenericVirtualRegister(DesiredLLT)
614bdd1243dSDimitry Andric : Call->ReturnRegister;
61506c3fb27SDimitry Andric if (!MRI->getRegClassOrNull(Tmp))
61606c3fb27SDimitry Andric MRI->setRegClass(Tmp, &SPIRV::IDRegClass);
617bdd1243dSDimitry Andric GR->assignSPIRVTypeToVReg(SpvDesiredTy, Tmp, MIRBuilder.getMF());
618bdd1243dSDimitry Andric
619bdd1243dSDimitry Andric SPIRVType *IntTy = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
620bdd1243dSDimitry Andric MIRBuilder.buildInstr(Opcode)
621bdd1243dSDimitry Andric .addDef(Tmp)
622bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(IntTy))
623bdd1243dSDimitry Andric .addUse(ObjectPtr)
624bdd1243dSDimitry Andric .addUse(ScopeReg)
625bdd1243dSDimitry Andric .addUse(MemSemEqualReg)
626bdd1243dSDimitry Andric .addUse(MemSemUnequalReg)
627bdd1243dSDimitry Andric .addUse(Desired)
628bdd1243dSDimitry Andric .addUse(Expected);
629bdd1243dSDimitry Andric if (!IsCmpxchg) {
630bdd1243dSDimitry Andric MIRBuilder.buildInstr(SPIRV::OpStore).addUse(ExpectedArg).addUse(Tmp);
631bdd1243dSDimitry Andric MIRBuilder.buildICmp(CmpInst::ICMP_EQ, Call->ReturnRegister, Tmp, Expected);
632bdd1243dSDimitry Andric }
633bdd1243dSDimitry Andric return true;
634bdd1243dSDimitry Andric }
635bdd1243dSDimitry Andric
636bdd1243dSDimitry Andric /// Helper function for building an atomic load instruction.
buildAtomicRMWInst(const SPIRV::IncomingCall * Call,unsigned Opcode,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)637bdd1243dSDimitry Andric static bool buildAtomicRMWInst(const SPIRV::IncomingCall *Call, unsigned Opcode,
638bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
639bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
64006c3fb27SDimitry Andric MachineRegisterInfo *MRI = MIRBuilder.getMRI();
64106c3fb27SDimitry Andric Register ScopeRegister =
64206c3fb27SDimitry Andric Call->Arguments.size() >= 4 ? Call->Arguments[3] : Register();
643bdd1243dSDimitry Andric
64406c3fb27SDimitry Andric assert(Call->Arguments.size() <= 4 &&
645bdd1243dSDimitry Andric "Too many args for explicit atomic RMW");
64606c3fb27SDimitry Andric ScopeRegister = buildScopeReg(ScopeRegister, SPIRV::Scope::Workgroup,
64706c3fb27SDimitry Andric MIRBuilder, GR, MRI);
648bdd1243dSDimitry Andric
649bdd1243dSDimitry Andric Register PtrRegister = Call->Arguments[0];
650bdd1243dSDimitry Andric unsigned Semantics = SPIRV::MemorySemantics::None;
65106c3fb27SDimitry Andric MRI->setRegClass(PtrRegister, &SPIRV::IDRegClass);
65206c3fb27SDimitry Andric Register MemSemanticsReg =
65306c3fb27SDimitry Andric Call->Arguments.size() >= 3 ? Call->Arguments[2] : Register();
65406c3fb27SDimitry Andric MemSemanticsReg = buildMemSemanticsReg(MemSemanticsReg, PtrRegister,
65506c3fb27SDimitry Andric Semantics, MIRBuilder, GR);
65606c3fb27SDimitry Andric MRI->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
657bdd1243dSDimitry Andric MIRBuilder.buildInstr(Opcode)
658bdd1243dSDimitry Andric .addDef(Call->ReturnRegister)
659bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(Call->ReturnType))
660bdd1243dSDimitry Andric .addUse(PtrRegister)
661bdd1243dSDimitry Andric .addUse(ScopeRegister)
662bdd1243dSDimitry Andric .addUse(MemSemanticsReg)
663bdd1243dSDimitry Andric .addUse(Call->Arguments[1]);
664bdd1243dSDimitry Andric return true;
665bdd1243dSDimitry Andric }
666bdd1243dSDimitry Andric
667bdd1243dSDimitry Andric /// Helper function for building atomic flag instructions (e.g.
668bdd1243dSDimitry Andric /// OpAtomicFlagTestAndSet).
buildAtomicFlagInst(const SPIRV::IncomingCall * Call,unsigned Opcode,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)669bdd1243dSDimitry Andric static bool buildAtomicFlagInst(const SPIRV::IncomingCall *Call,
670bdd1243dSDimitry Andric unsigned Opcode, MachineIRBuilder &MIRBuilder,
671bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
67206c3fb27SDimitry Andric MachineRegisterInfo *MRI = MIRBuilder.getMRI();
673bdd1243dSDimitry Andric Register PtrRegister = Call->Arguments[0];
674bdd1243dSDimitry Andric unsigned Semantics = SPIRV::MemorySemantics::SequentiallyConsistent;
67506c3fb27SDimitry Andric Register MemSemanticsReg =
67606c3fb27SDimitry Andric Call->Arguments.size() >= 2 ? Call->Arguments[1] : Register();
67706c3fb27SDimitry Andric MemSemanticsReg = buildMemSemanticsReg(MemSemanticsReg, PtrRegister,
67806c3fb27SDimitry Andric Semantics, MIRBuilder, GR);
679bdd1243dSDimitry Andric
680bdd1243dSDimitry Andric assert((Opcode != SPIRV::OpAtomicFlagClear ||
681bdd1243dSDimitry Andric (Semantics != SPIRV::MemorySemantics::Acquire &&
682bdd1243dSDimitry Andric Semantics != SPIRV::MemorySemantics::AcquireRelease)) &&
683bdd1243dSDimitry Andric "Invalid memory order argument!");
684bdd1243dSDimitry Andric
68506c3fb27SDimitry Andric Register ScopeRegister =
68606c3fb27SDimitry Andric Call->Arguments.size() >= 3 ? Call->Arguments[2] : Register();
68706c3fb27SDimitry Andric ScopeRegister =
68806c3fb27SDimitry Andric buildScopeReg(ScopeRegister, SPIRV::Scope::Device, MIRBuilder, GR, MRI);
689bdd1243dSDimitry Andric
690bdd1243dSDimitry Andric auto MIB = MIRBuilder.buildInstr(Opcode);
691bdd1243dSDimitry Andric if (Opcode == SPIRV::OpAtomicFlagTestAndSet)
692bdd1243dSDimitry Andric MIB.addDef(Call->ReturnRegister)
693bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(Call->ReturnType));
694bdd1243dSDimitry Andric
695bdd1243dSDimitry Andric MIB.addUse(PtrRegister).addUse(ScopeRegister).addUse(MemSemanticsReg);
696bdd1243dSDimitry Andric return true;
697bdd1243dSDimitry Andric }
698bdd1243dSDimitry Andric
699bdd1243dSDimitry Andric /// Helper function for building barriers, i.e., memory/control ordering
700bdd1243dSDimitry Andric /// operations.
buildBarrierInst(const SPIRV::IncomingCall * Call,unsigned Opcode,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)701bdd1243dSDimitry Andric static bool buildBarrierInst(const SPIRV::IncomingCall *Call, unsigned Opcode,
702bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
703bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
70406c3fb27SDimitry Andric MachineRegisterInfo *MRI = MIRBuilder.getMRI();
705bdd1243dSDimitry Andric unsigned MemFlags = getIConstVal(Call->Arguments[0], MRI);
706bdd1243dSDimitry Andric unsigned MemSemantics = SPIRV::MemorySemantics::None;
707bdd1243dSDimitry Andric
708bdd1243dSDimitry Andric if (MemFlags & SPIRV::CLK_LOCAL_MEM_FENCE)
709bdd1243dSDimitry Andric MemSemantics |= SPIRV::MemorySemantics::WorkgroupMemory;
710bdd1243dSDimitry Andric
711bdd1243dSDimitry Andric if (MemFlags & SPIRV::CLK_GLOBAL_MEM_FENCE)
712bdd1243dSDimitry Andric MemSemantics |= SPIRV::MemorySemantics::CrossWorkgroupMemory;
713bdd1243dSDimitry Andric
714bdd1243dSDimitry Andric if (MemFlags & SPIRV::CLK_IMAGE_MEM_FENCE)
715bdd1243dSDimitry Andric MemSemantics |= SPIRV::MemorySemantics::ImageMemory;
716bdd1243dSDimitry Andric
717bdd1243dSDimitry Andric if (Opcode == SPIRV::OpMemoryBarrier) {
718bdd1243dSDimitry Andric std::memory_order MemOrder =
719bdd1243dSDimitry Andric static_cast<std::memory_order>(getIConstVal(Call->Arguments[1], MRI));
720bdd1243dSDimitry Andric MemSemantics = getSPIRVMemSemantics(MemOrder) | MemSemantics;
721bdd1243dSDimitry Andric } else {
722bdd1243dSDimitry Andric MemSemantics |= SPIRV::MemorySemantics::SequentiallyConsistent;
723bdd1243dSDimitry Andric }
724bdd1243dSDimitry Andric
725bdd1243dSDimitry Andric Register MemSemanticsReg;
72606c3fb27SDimitry Andric if (MemFlags == MemSemantics) {
727bdd1243dSDimitry Andric MemSemanticsReg = Call->Arguments[0];
72806c3fb27SDimitry Andric MRI->setRegClass(MemSemanticsReg, &SPIRV::IDRegClass);
72906c3fb27SDimitry Andric } else
730bdd1243dSDimitry Andric MemSemanticsReg = buildConstantIntReg(MemSemantics, MIRBuilder, GR);
731bdd1243dSDimitry Andric
732bdd1243dSDimitry Andric Register ScopeReg;
733bdd1243dSDimitry Andric SPIRV::Scope::Scope Scope = SPIRV::Scope::Workgroup;
734bdd1243dSDimitry Andric SPIRV::Scope::Scope MemScope = Scope;
735bdd1243dSDimitry Andric if (Call->Arguments.size() >= 2) {
736bdd1243dSDimitry Andric assert(
737bdd1243dSDimitry Andric ((Opcode != SPIRV::OpMemoryBarrier && Call->Arguments.size() == 2) ||
738bdd1243dSDimitry Andric (Opcode == SPIRV::OpMemoryBarrier && Call->Arguments.size() == 3)) &&
739bdd1243dSDimitry Andric "Extra args for explicitly scoped barrier");
740bdd1243dSDimitry Andric Register ScopeArg = (Opcode == SPIRV::OpMemoryBarrier) ? Call->Arguments[2]
741bdd1243dSDimitry Andric : Call->Arguments[1];
742bdd1243dSDimitry Andric SPIRV::CLMemoryScope CLScope =
743bdd1243dSDimitry Andric static_cast<SPIRV::CLMemoryScope>(getIConstVal(ScopeArg, MRI));
744bdd1243dSDimitry Andric MemScope = getSPIRVScope(CLScope);
745bdd1243dSDimitry Andric if (!(MemFlags & SPIRV::CLK_LOCAL_MEM_FENCE) ||
746bdd1243dSDimitry Andric (Opcode == SPIRV::OpMemoryBarrier))
747bdd1243dSDimitry Andric Scope = MemScope;
748bdd1243dSDimitry Andric
74906c3fb27SDimitry Andric if (CLScope == static_cast<unsigned>(Scope)) {
750bdd1243dSDimitry Andric ScopeReg = Call->Arguments[1];
75106c3fb27SDimitry Andric MRI->setRegClass(ScopeReg, &SPIRV::IDRegClass);
75206c3fb27SDimitry Andric }
753bdd1243dSDimitry Andric }
754bdd1243dSDimitry Andric
755bdd1243dSDimitry Andric if (!ScopeReg.isValid())
756bdd1243dSDimitry Andric ScopeReg = buildConstantIntReg(Scope, MIRBuilder, GR);
757bdd1243dSDimitry Andric
758bdd1243dSDimitry Andric auto MIB = MIRBuilder.buildInstr(Opcode).addUse(ScopeReg);
759bdd1243dSDimitry Andric if (Opcode != SPIRV::OpMemoryBarrier)
760bdd1243dSDimitry Andric MIB.addUse(buildConstantIntReg(MemScope, MIRBuilder, GR));
761bdd1243dSDimitry Andric MIB.addUse(MemSemanticsReg);
762bdd1243dSDimitry Andric return true;
763bdd1243dSDimitry Andric }
764bdd1243dSDimitry Andric
getNumComponentsForDim(SPIRV::Dim::Dim dim)765bdd1243dSDimitry Andric static unsigned getNumComponentsForDim(SPIRV::Dim::Dim dim) {
766bdd1243dSDimitry Andric switch (dim) {
767bdd1243dSDimitry Andric case SPIRV::Dim::DIM_1D:
768bdd1243dSDimitry Andric case SPIRV::Dim::DIM_Buffer:
769bdd1243dSDimitry Andric return 1;
770bdd1243dSDimitry Andric case SPIRV::Dim::DIM_2D:
771bdd1243dSDimitry Andric case SPIRV::Dim::DIM_Cube:
772bdd1243dSDimitry Andric case SPIRV::Dim::DIM_Rect:
773bdd1243dSDimitry Andric return 2;
774bdd1243dSDimitry Andric case SPIRV::Dim::DIM_3D:
775bdd1243dSDimitry Andric return 3;
776bdd1243dSDimitry Andric default:
777bdd1243dSDimitry Andric llvm_unreachable("Cannot get num components for given Dim");
778bdd1243dSDimitry Andric }
779bdd1243dSDimitry Andric }
780bdd1243dSDimitry Andric
781bdd1243dSDimitry Andric /// Helper function for obtaining the number of size components.
getNumSizeComponents(SPIRVType * imgType)782bdd1243dSDimitry Andric static unsigned getNumSizeComponents(SPIRVType *imgType) {
783bdd1243dSDimitry Andric assert(imgType->getOpcode() == SPIRV::OpTypeImage);
784bdd1243dSDimitry Andric auto dim = static_cast<SPIRV::Dim::Dim>(imgType->getOperand(2).getImm());
785bdd1243dSDimitry Andric unsigned numComps = getNumComponentsForDim(dim);
786bdd1243dSDimitry Andric bool arrayed = imgType->getOperand(4).getImm() == 1;
787bdd1243dSDimitry Andric return arrayed ? numComps + 1 : numComps;
788bdd1243dSDimitry Andric }
789bdd1243dSDimitry Andric
790bdd1243dSDimitry Andric //===----------------------------------------------------------------------===//
791bdd1243dSDimitry Andric // Implementation functions for each builtin group
792bdd1243dSDimitry Andric //===----------------------------------------------------------------------===//
793bdd1243dSDimitry Andric
generateExtInst(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)794bdd1243dSDimitry Andric static bool generateExtInst(const SPIRV::IncomingCall *Call,
795bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
796bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
797bdd1243dSDimitry Andric // Lookup the extended instruction number in the TableGen records.
798bdd1243dSDimitry Andric const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
799bdd1243dSDimitry Andric uint32_t Number =
800bdd1243dSDimitry Andric SPIRV::lookupExtendedBuiltin(Builtin->Name, Builtin->Set)->Number;
801bdd1243dSDimitry Andric
802bdd1243dSDimitry Andric // Build extended instruction.
803bdd1243dSDimitry Andric auto MIB =
804bdd1243dSDimitry Andric MIRBuilder.buildInstr(SPIRV::OpExtInst)
805bdd1243dSDimitry Andric .addDef(Call->ReturnRegister)
806bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(Call->ReturnType))
807bdd1243dSDimitry Andric .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::OpenCL_std))
808bdd1243dSDimitry Andric .addImm(Number);
809bdd1243dSDimitry Andric
810bdd1243dSDimitry Andric for (auto Argument : Call->Arguments)
811bdd1243dSDimitry Andric MIB.addUse(Argument);
812bdd1243dSDimitry Andric return true;
813bdd1243dSDimitry Andric }
814bdd1243dSDimitry Andric
generateRelationalInst(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)815bdd1243dSDimitry Andric static bool generateRelationalInst(const SPIRV::IncomingCall *Call,
816bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
817bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
818bdd1243dSDimitry Andric // Lookup the instruction opcode in the TableGen records.
819bdd1243dSDimitry Andric const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
820bdd1243dSDimitry Andric unsigned Opcode =
821bdd1243dSDimitry Andric SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
822bdd1243dSDimitry Andric
823bdd1243dSDimitry Andric Register CompareRegister;
824bdd1243dSDimitry Andric SPIRVType *RelationType;
825bdd1243dSDimitry Andric std::tie(CompareRegister, RelationType) =
826bdd1243dSDimitry Andric buildBoolRegister(MIRBuilder, Call->ReturnType, GR);
827bdd1243dSDimitry Andric
828bdd1243dSDimitry Andric // Build relational instruction.
829bdd1243dSDimitry Andric auto MIB = MIRBuilder.buildInstr(Opcode)
830bdd1243dSDimitry Andric .addDef(CompareRegister)
831bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(RelationType));
832bdd1243dSDimitry Andric
833bdd1243dSDimitry Andric for (auto Argument : Call->Arguments)
834bdd1243dSDimitry Andric MIB.addUse(Argument);
835bdd1243dSDimitry Andric
836bdd1243dSDimitry Andric // Build select instruction.
837bdd1243dSDimitry Andric return buildSelectInst(MIRBuilder, Call->ReturnRegister, CompareRegister,
838bdd1243dSDimitry Andric Call->ReturnType, GR);
839bdd1243dSDimitry Andric }
840bdd1243dSDimitry Andric
generateGroupInst(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)841bdd1243dSDimitry Andric static bool generateGroupInst(const SPIRV::IncomingCall *Call,
842bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
843bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
844bdd1243dSDimitry Andric const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
845bdd1243dSDimitry Andric const SPIRV::GroupBuiltin *GroupBuiltin =
846bdd1243dSDimitry Andric SPIRV::lookupGroupBuiltin(Builtin->Name);
84706c3fb27SDimitry Andric MachineRegisterInfo *MRI = MIRBuilder.getMRI();
848bdd1243dSDimitry Andric Register Arg0;
849bdd1243dSDimitry Andric if (GroupBuiltin->HasBoolArg) {
850bdd1243dSDimitry Andric Register ConstRegister = Call->Arguments[0];
851bdd1243dSDimitry Andric auto ArgInstruction = getDefInstrMaybeConstant(ConstRegister, MRI);
852bdd1243dSDimitry Andric // TODO: support non-constant bool values.
853bdd1243dSDimitry Andric assert(ArgInstruction->getOpcode() == TargetOpcode::G_CONSTANT &&
854bdd1243dSDimitry Andric "Only constant bool value args are supported");
855bdd1243dSDimitry Andric if (GR->getSPIRVTypeForVReg(Call->Arguments[0])->getOpcode() !=
856bdd1243dSDimitry Andric SPIRV::OpTypeBool)
857bdd1243dSDimitry Andric Arg0 = GR->buildConstantInt(getIConstVal(ConstRegister, MRI), MIRBuilder,
858bdd1243dSDimitry Andric GR->getOrCreateSPIRVBoolType(MIRBuilder));
859bdd1243dSDimitry Andric }
860bdd1243dSDimitry Andric
861bdd1243dSDimitry Andric Register GroupResultRegister = Call->ReturnRegister;
862bdd1243dSDimitry Andric SPIRVType *GroupResultType = Call->ReturnType;
863bdd1243dSDimitry Andric
864bdd1243dSDimitry Andric // TODO: maybe we need to check whether the result type is already boolean
865bdd1243dSDimitry Andric // and in this case do not insert select instruction.
866bdd1243dSDimitry Andric const bool HasBoolReturnTy =
867bdd1243dSDimitry Andric GroupBuiltin->IsElect || GroupBuiltin->IsAllOrAny ||
868bdd1243dSDimitry Andric GroupBuiltin->IsAllEqual || GroupBuiltin->IsLogical ||
869bdd1243dSDimitry Andric GroupBuiltin->IsInverseBallot || GroupBuiltin->IsBallotBitExtract;
870bdd1243dSDimitry Andric
871bdd1243dSDimitry Andric if (HasBoolReturnTy)
872bdd1243dSDimitry Andric std::tie(GroupResultRegister, GroupResultType) =
873bdd1243dSDimitry Andric buildBoolRegister(MIRBuilder, Call->ReturnType, GR);
874bdd1243dSDimitry Andric
8755f757f3fSDimitry Andric auto Scope = Builtin->Name.starts_with("sub_group") ? SPIRV::Scope::Subgroup
876bdd1243dSDimitry Andric : SPIRV::Scope::Workgroup;
877bdd1243dSDimitry Andric Register ScopeRegister = buildConstantIntReg(Scope, MIRBuilder, GR);
878bdd1243dSDimitry Andric
879bdd1243dSDimitry Andric // Build work/sub group instruction.
880bdd1243dSDimitry Andric auto MIB = MIRBuilder.buildInstr(GroupBuiltin->Opcode)
881bdd1243dSDimitry Andric .addDef(GroupResultRegister)
882bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(GroupResultType))
883bdd1243dSDimitry Andric .addUse(ScopeRegister);
884bdd1243dSDimitry Andric
885bdd1243dSDimitry Andric if (!GroupBuiltin->NoGroupOperation)
886bdd1243dSDimitry Andric MIB.addImm(GroupBuiltin->GroupOperation);
887bdd1243dSDimitry Andric if (Call->Arguments.size() > 0) {
888bdd1243dSDimitry Andric MIB.addUse(Arg0.isValid() ? Arg0 : Call->Arguments[0]);
88906c3fb27SDimitry Andric MRI->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
89006c3fb27SDimitry Andric for (unsigned i = 1; i < Call->Arguments.size(); i++) {
891bdd1243dSDimitry Andric MIB.addUse(Call->Arguments[i]);
89206c3fb27SDimitry Andric MRI->setRegClass(Call->Arguments[i], &SPIRV::IDRegClass);
89306c3fb27SDimitry Andric }
894bdd1243dSDimitry Andric }
895bdd1243dSDimitry Andric
896bdd1243dSDimitry Andric // Build select instruction.
897bdd1243dSDimitry Andric if (HasBoolReturnTy)
898bdd1243dSDimitry Andric buildSelectInst(MIRBuilder, Call->ReturnRegister, GroupResultRegister,
899bdd1243dSDimitry Andric Call->ReturnType, GR);
900bdd1243dSDimitry Andric return true;
901bdd1243dSDimitry Andric }
902bdd1243dSDimitry Andric
903bdd1243dSDimitry Andric // These queries ask for a single size_t result for a given dimension index, e.g
904bdd1243dSDimitry Andric // size_t get_global_id(uint dimindex). In SPIR-V, the builtins corresonding to
905bdd1243dSDimitry Andric // these values are all vec3 types, so we need to extract the correct index or
906bdd1243dSDimitry Andric // return defaultVal (0 or 1 depending on the query). We also handle extending
907bdd1243dSDimitry Andric // or tuncating in case size_t does not match the expected result type's
908bdd1243dSDimitry Andric // bitwidth.
909bdd1243dSDimitry Andric //
910bdd1243dSDimitry Andric // For a constant index >= 3 we generate:
911bdd1243dSDimitry Andric // %res = OpConstant %SizeT 0
912bdd1243dSDimitry Andric //
913bdd1243dSDimitry Andric // For other indices we generate:
914bdd1243dSDimitry Andric // %g = OpVariable %ptr_V3_SizeT Input
915bdd1243dSDimitry Andric // OpDecorate %g BuiltIn XXX
916bdd1243dSDimitry Andric // OpDecorate %g LinkageAttributes "__spirv_BuiltInXXX"
917bdd1243dSDimitry Andric // OpDecorate %g Constant
918bdd1243dSDimitry Andric // %loadedVec = OpLoad %V3_SizeT %g
919bdd1243dSDimitry Andric //
920bdd1243dSDimitry Andric // Then, if the index is constant < 3, we generate:
921bdd1243dSDimitry Andric // %res = OpCompositeExtract %SizeT %loadedVec idx
922bdd1243dSDimitry Andric // If the index is dynamic, we generate:
923bdd1243dSDimitry Andric // %tmp = OpVectorExtractDynamic %SizeT %loadedVec %idx
924bdd1243dSDimitry Andric // %cmp = OpULessThan %bool %idx %const_3
925bdd1243dSDimitry Andric // %res = OpSelect %SizeT %cmp %tmp %const_0
926bdd1243dSDimitry Andric //
927bdd1243dSDimitry Andric // If the bitwidth of %res does not match the expected return type, we add an
928bdd1243dSDimitry Andric // extend or truncate.
genWorkgroupQuery(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR,SPIRV::BuiltIn::BuiltIn BuiltinValue,uint64_t DefaultValue)929bdd1243dSDimitry Andric static bool genWorkgroupQuery(const SPIRV::IncomingCall *Call,
930bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
931bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR,
932bdd1243dSDimitry Andric SPIRV::BuiltIn::BuiltIn BuiltinValue,
933bdd1243dSDimitry Andric uint64_t DefaultValue) {
934bdd1243dSDimitry Andric Register IndexRegister = Call->Arguments[0];
935bdd1243dSDimitry Andric const unsigned ResultWidth = Call->ReturnType->getOperand(1).getImm();
936bdd1243dSDimitry Andric const unsigned PointerSize = GR->getPointerSize();
937bdd1243dSDimitry Andric const SPIRVType *PointerSizeType =
938bdd1243dSDimitry Andric GR->getOrCreateSPIRVIntegerType(PointerSize, MIRBuilder);
939bdd1243dSDimitry Andric MachineRegisterInfo *MRI = MIRBuilder.getMRI();
940bdd1243dSDimitry Andric auto IndexInstruction = getDefInstrMaybeConstant(IndexRegister, MRI);
941bdd1243dSDimitry Andric
942bdd1243dSDimitry Andric // Set up the final register to do truncation or extension on at the end.
943bdd1243dSDimitry Andric Register ToTruncate = Call->ReturnRegister;
944bdd1243dSDimitry Andric
945bdd1243dSDimitry Andric // If the index is constant, we can statically determine if it is in range.
946bdd1243dSDimitry Andric bool IsConstantIndex =
947bdd1243dSDimitry Andric IndexInstruction->getOpcode() == TargetOpcode::G_CONSTANT;
948bdd1243dSDimitry Andric
949bdd1243dSDimitry Andric // If it's out of range (max dimension is 3), we can just return the constant
950bdd1243dSDimitry Andric // default value (0 or 1 depending on which query function).
951bdd1243dSDimitry Andric if (IsConstantIndex && getIConstVal(IndexRegister, MRI) >= 3) {
95206c3fb27SDimitry Andric Register DefaultReg = Call->ReturnRegister;
953bdd1243dSDimitry Andric if (PointerSize != ResultWidth) {
95406c3fb27SDimitry Andric DefaultReg = MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));
95506c3fb27SDimitry Andric MRI->setRegClass(DefaultReg, &SPIRV::IDRegClass);
95606c3fb27SDimitry Andric GR->assignSPIRVTypeToVReg(PointerSizeType, DefaultReg,
957bdd1243dSDimitry Andric MIRBuilder.getMF());
95806c3fb27SDimitry Andric ToTruncate = DefaultReg;
959bdd1243dSDimitry Andric }
960bdd1243dSDimitry Andric auto NewRegister =
961bdd1243dSDimitry Andric GR->buildConstantInt(DefaultValue, MIRBuilder, PointerSizeType);
96206c3fb27SDimitry Andric MIRBuilder.buildCopy(DefaultReg, NewRegister);
963bdd1243dSDimitry Andric } else { // If it could be in range, we need to load from the given builtin.
964bdd1243dSDimitry Andric auto Vec3Ty =
965bdd1243dSDimitry Andric GR->getOrCreateSPIRVVectorType(PointerSizeType, 3, MIRBuilder);
966bdd1243dSDimitry Andric Register LoadedVector =
967bdd1243dSDimitry Andric buildBuiltinVariableLoad(MIRBuilder, Vec3Ty, GR, BuiltinValue,
968bdd1243dSDimitry Andric LLT::fixed_vector(3, PointerSize));
969bdd1243dSDimitry Andric // Set up the vreg to extract the result to (possibly a new temporary one).
970bdd1243dSDimitry Andric Register Extracted = Call->ReturnRegister;
971bdd1243dSDimitry Andric if (!IsConstantIndex || PointerSize != ResultWidth) {
972bdd1243dSDimitry Andric Extracted = MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));
97306c3fb27SDimitry Andric MRI->setRegClass(Extracted, &SPIRV::IDRegClass);
974bdd1243dSDimitry Andric GR->assignSPIRVTypeToVReg(PointerSizeType, Extracted, MIRBuilder.getMF());
975bdd1243dSDimitry Andric }
976bdd1243dSDimitry Andric // Use Intrinsic::spv_extractelt so dynamic vs static extraction is
977bdd1243dSDimitry Andric // handled later: extr = spv_extractelt LoadedVector, IndexRegister.
978bdd1243dSDimitry Andric MachineInstrBuilder ExtractInst = MIRBuilder.buildIntrinsic(
9795f757f3fSDimitry Andric Intrinsic::spv_extractelt, ArrayRef<Register>{Extracted}, true, false);
980bdd1243dSDimitry Andric ExtractInst.addUse(LoadedVector).addUse(IndexRegister);
981bdd1243dSDimitry Andric
982bdd1243dSDimitry Andric // If the index is dynamic, need check if it's < 3, and then use a select.
983bdd1243dSDimitry Andric if (!IsConstantIndex) {
984bdd1243dSDimitry Andric insertAssignInstr(Extracted, nullptr, PointerSizeType, GR, MIRBuilder,
985bdd1243dSDimitry Andric *MRI);
986bdd1243dSDimitry Andric
987bdd1243dSDimitry Andric auto IndexType = GR->getSPIRVTypeForVReg(IndexRegister);
988bdd1243dSDimitry Andric auto BoolType = GR->getOrCreateSPIRVBoolType(MIRBuilder);
989bdd1243dSDimitry Andric
990bdd1243dSDimitry Andric Register CompareRegister =
991bdd1243dSDimitry Andric MRI->createGenericVirtualRegister(LLT::scalar(1));
99206c3fb27SDimitry Andric MRI->setRegClass(CompareRegister, &SPIRV::IDRegClass);
993bdd1243dSDimitry Andric GR->assignSPIRVTypeToVReg(BoolType, CompareRegister, MIRBuilder.getMF());
994bdd1243dSDimitry Andric
995bdd1243dSDimitry Andric // Use G_ICMP to check if idxVReg < 3.
996bdd1243dSDimitry Andric MIRBuilder.buildICmp(CmpInst::ICMP_ULT, CompareRegister, IndexRegister,
997bdd1243dSDimitry Andric GR->buildConstantInt(3, MIRBuilder, IndexType));
998bdd1243dSDimitry Andric
999bdd1243dSDimitry Andric // Get constant for the default value (0 or 1 depending on which
1000bdd1243dSDimitry Andric // function).
1001bdd1243dSDimitry Andric Register DefaultRegister =
1002bdd1243dSDimitry Andric GR->buildConstantInt(DefaultValue, MIRBuilder, PointerSizeType);
1003bdd1243dSDimitry Andric
1004bdd1243dSDimitry Andric // Get a register for the selection result (possibly a new temporary one).
1005bdd1243dSDimitry Andric Register SelectionResult = Call->ReturnRegister;
1006bdd1243dSDimitry Andric if (PointerSize != ResultWidth) {
1007bdd1243dSDimitry Andric SelectionResult =
1008bdd1243dSDimitry Andric MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));
100906c3fb27SDimitry Andric MRI->setRegClass(SelectionResult, &SPIRV::IDRegClass);
1010bdd1243dSDimitry Andric GR->assignSPIRVTypeToVReg(PointerSizeType, SelectionResult,
1011bdd1243dSDimitry Andric MIRBuilder.getMF());
1012bdd1243dSDimitry Andric }
1013bdd1243dSDimitry Andric // Create the final G_SELECT to return the extracted value or the default.
1014bdd1243dSDimitry Andric MIRBuilder.buildSelect(SelectionResult, CompareRegister, Extracted,
1015bdd1243dSDimitry Andric DefaultRegister);
1016bdd1243dSDimitry Andric ToTruncate = SelectionResult;
1017bdd1243dSDimitry Andric } else {
1018bdd1243dSDimitry Andric ToTruncate = Extracted;
1019bdd1243dSDimitry Andric }
1020bdd1243dSDimitry Andric }
1021bdd1243dSDimitry Andric // Alter the result's bitwidth if it does not match the SizeT value extracted.
1022bdd1243dSDimitry Andric if (PointerSize != ResultWidth)
1023bdd1243dSDimitry Andric MIRBuilder.buildZExtOrTrunc(Call->ReturnRegister, ToTruncate);
1024bdd1243dSDimitry Andric return true;
1025bdd1243dSDimitry Andric }
1026bdd1243dSDimitry Andric
generateBuiltinVar(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)1027bdd1243dSDimitry Andric static bool generateBuiltinVar(const SPIRV::IncomingCall *Call,
1028bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
1029bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
1030bdd1243dSDimitry Andric // Lookup the builtin variable record.
1031bdd1243dSDimitry Andric const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1032bdd1243dSDimitry Andric SPIRV::BuiltIn::BuiltIn Value =
1033bdd1243dSDimitry Andric SPIRV::lookupGetBuiltin(Builtin->Name, Builtin->Set)->Value;
1034bdd1243dSDimitry Andric
1035bdd1243dSDimitry Andric if (Value == SPIRV::BuiltIn::GlobalInvocationId)
1036bdd1243dSDimitry Andric return genWorkgroupQuery(Call, MIRBuilder, GR, Value, 0);
1037bdd1243dSDimitry Andric
1038bdd1243dSDimitry Andric // Build a load instruction for the builtin variable.
1039bdd1243dSDimitry Andric unsigned BitWidth = GR->getScalarOrVectorBitWidth(Call->ReturnType);
1040bdd1243dSDimitry Andric LLT LLType;
1041bdd1243dSDimitry Andric if (Call->ReturnType->getOpcode() == SPIRV::OpTypeVector)
1042bdd1243dSDimitry Andric LLType =
1043bdd1243dSDimitry Andric LLT::fixed_vector(Call->ReturnType->getOperand(2).getImm(), BitWidth);
1044bdd1243dSDimitry Andric else
1045bdd1243dSDimitry Andric LLType = LLT::scalar(BitWidth);
1046bdd1243dSDimitry Andric
1047bdd1243dSDimitry Andric return buildBuiltinVariableLoad(MIRBuilder, Call->ReturnType, GR, Value,
1048bdd1243dSDimitry Andric LLType, Call->ReturnRegister);
1049bdd1243dSDimitry Andric }
1050bdd1243dSDimitry Andric
generateAtomicInst(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)1051bdd1243dSDimitry Andric static bool generateAtomicInst(const SPIRV::IncomingCall *Call,
1052bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
1053bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
1054bdd1243dSDimitry Andric // Lookup the instruction opcode in the TableGen records.
1055bdd1243dSDimitry Andric const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1056bdd1243dSDimitry Andric unsigned Opcode =
1057bdd1243dSDimitry Andric SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1058bdd1243dSDimitry Andric
1059bdd1243dSDimitry Andric switch (Opcode) {
1060bdd1243dSDimitry Andric case SPIRV::OpStore:
1061bdd1243dSDimitry Andric return buildAtomicInitInst(Call, MIRBuilder);
1062bdd1243dSDimitry Andric case SPIRV::OpAtomicLoad:
1063bdd1243dSDimitry Andric return buildAtomicLoadInst(Call, MIRBuilder, GR);
1064bdd1243dSDimitry Andric case SPIRV::OpAtomicStore:
1065bdd1243dSDimitry Andric return buildAtomicStoreInst(Call, MIRBuilder, GR);
1066bdd1243dSDimitry Andric case SPIRV::OpAtomicCompareExchange:
1067bdd1243dSDimitry Andric case SPIRV::OpAtomicCompareExchangeWeak:
1068bdd1243dSDimitry Andric return buildAtomicCompareExchangeInst(Call, MIRBuilder, GR);
1069bdd1243dSDimitry Andric case SPIRV::OpAtomicIAdd:
1070bdd1243dSDimitry Andric case SPIRV::OpAtomicISub:
1071bdd1243dSDimitry Andric case SPIRV::OpAtomicOr:
1072bdd1243dSDimitry Andric case SPIRV::OpAtomicXor:
1073bdd1243dSDimitry Andric case SPIRV::OpAtomicAnd:
1074bdd1243dSDimitry Andric case SPIRV::OpAtomicExchange:
1075bdd1243dSDimitry Andric return buildAtomicRMWInst(Call, Opcode, MIRBuilder, GR);
1076bdd1243dSDimitry Andric case SPIRV::OpMemoryBarrier:
1077bdd1243dSDimitry Andric return buildBarrierInst(Call, SPIRV::OpMemoryBarrier, MIRBuilder, GR);
1078bdd1243dSDimitry Andric case SPIRV::OpAtomicFlagTestAndSet:
1079bdd1243dSDimitry Andric case SPIRV::OpAtomicFlagClear:
1080bdd1243dSDimitry Andric return buildAtomicFlagInst(Call, Opcode, MIRBuilder, GR);
1081bdd1243dSDimitry Andric default:
1082bdd1243dSDimitry Andric return false;
1083bdd1243dSDimitry Andric }
1084bdd1243dSDimitry Andric }
1085bdd1243dSDimitry Andric
generateBarrierInst(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)1086bdd1243dSDimitry Andric static bool generateBarrierInst(const SPIRV::IncomingCall *Call,
1087bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
1088bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
1089bdd1243dSDimitry Andric // Lookup the instruction opcode in the TableGen records.
1090bdd1243dSDimitry Andric const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1091bdd1243dSDimitry Andric unsigned Opcode =
1092bdd1243dSDimitry Andric SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1093bdd1243dSDimitry Andric
1094bdd1243dSDimitry Andric return buildBarrierInst(Call, Opcode, MIRBuilder, GR);
1095bdd1243dSDimitry Andric }
1096bdd1243dSDimitry Andric
generateDotOrFMulInst(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)1097bdd1243dSDimitry Andric static bool generateDotOrFMulInst(const SPIRV::IncomingCall *Call,
1098bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
1099bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
1100bdd1243dSDimitry Andric unsigned Opcode = GR->getSPIRVTypeForVReg(Call->Arguments[0])->getOpcode();
1101bdd1243dSDimitry Andric bool IsVec = Opcode == SPIRV::OpTypeVector;
1102bdd1243dSDimitry Andric // Use OpDot only in case of vector args and OpFMul in case of scalar args.
1103bdd1243dSDimitry Andric MIRBuilder.buildInstr(IsVec ? SPIRV::OpDot : SPIRV::OpFMulS)
1104bdd1243dSDimitry Andric .addDef(Call->ReturnRegister)
1105bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1106bdd1243dSDimitry Andric .addUse(Call->Arguments[0])
1107bdd1243dSDimitry Andric .addUse(Call->Arguments[1]);
1108bdd1243dSDimitry Andric return true;
1109bdd1243dSDimitry Andric }
1110bdd1243dSDimitry Andric
generateGetQueryInst(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)1111bdd1243dSDimitry Andric static bool generateGetQueryInst(const SPIRV::IncomingCall *Call,
1112bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
1113bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
1114bdd1243dSDimitry Andric // Lookup the builtin record.
1115bdd1243dSDimitry Andric SPIRV::BuiltIn::BuiltIn Value =
1116bdd1243dSDimitry Andric SPIRV::lookupGetBuiltin(Call->Builtin->Name, Call->Builtin->Set)->Value;
1117bdd1243dSDimitry Andric uint64_t IsDefault = (Value == SPIRV::BuiltIn::GlobalSize ||
1118bdd1243dSDimitry Andric Value == SPIRV::BuiltIn::WorkgroupSize ||
1119bdd1243dSDimitry Andric Value == SPIRV::BuiltIn::EnqueuedWorkgroupSize);
1120bdd1243dSDimitry Andric return genWorkgroupQuery(Call, MIRBuilder, GR, Value, IsDefault ? 1 : 0);
1121bdd1243dSDimitry Andric }
1122bdd1243dSDimitry Andric
generateImageSizeQueryInst(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)1123bdd1243dSDimitry Andric static bool generateImageSizeQueryInst(const SPIRV::IncomingCall *Call,
1124bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
1125bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
1126bdd1243dSDimitry Andric // Lookup the image size query component number in the TableGen records.
1127bdd1243dSDimitry Andric const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1128bdd1243dSDimitry Andric uint32_t Component =
1129bdd1243dSDimitry Andric SPIRV::lookupImageQueryBuiltin(Builtin->Name, Builtin->Set)->Component;
1130bdd1243dSDimitry Andric // Query result may either be a vector or a scalar. If return type is not a
1131bdd1243dSDimitry Andric // vector, expect only a single size component. Otherwise get the number of
1132bdd1243dSDimitry Andric // expected components.
1133bdd1243dSDimitry Andric SPIRVType *RetTy = Call->ReturnType;
1134bdd1243dSDimitry Andric unsigned NumExpectedRetComponents = RetTy->getOpcode() == SPIRV::OpTypeVector
1135bdd1243dSDimitry Andric ? RetTy->getOperand(2).getImm()
1136bdd1243dSDimitry Andric : 1;
1137bdd1243dSDimitry Andric // Get the actual number of query result/size components.
1138bdd1243dSDimitry Andric SPIRVType *ImgType = GR->getSPIRVTypeForVReg(Call->Arguments[0]);
1139bdd1243dSDimitry Andric unsigned NumActualRetComponents = getNumSizeComponents(ImgType);
1140bdd1243dSDimitry Andric Register QueryResult = Call->ReturnRegister;
1141bdd1243dSDimitry Andric SPIRVType *QueryResultType = Call->ReturnType;
1142bdd1243dSDimitry Andric if (NumExpectedRetComponents != NumActualRetComponents) {
1143bdd1243dSDimitry Andric QueryResult = MIRBuilder.getMRI()->createGenericVirtualRegister(
1144bdd1243dSDimitry Andric LLT::fixed_vector(NumActualRetComponents, 32));
114506c3fb27SDimitry Andric MIRBuilder.getMRI()->setRegClass(QueryResult, &SPIRV::IDRegClass);
1146bdd1243dSDimitry Andric SPIRVType *IntTy = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
1147bdd1243dSDimitry Andric QueryResultType = GR->getOrCreateSPIRVVectorType(
1148bdd1243dSDimitry Andric IntTy, NumActualRetComponents, MIRBuilder);
1149bdd1243dSDimitry Andric GR->assignSPIRVTypeToVReg(QueryResultType, QueryResult, MIRBuilder.getMF());
1150bdd1243dSDimitry Andric }
1151bdd1243dSDimitry Andric bool IsDimBuf = ImgType->getOperand(2).getImm() == SPIRV::Dim::DIM_Buffer;
1152bdd1243dSDimitry Andric unsigned Opcode =
1153bdd1243dSDimitry Andric IsDimBuf ? SPIRV::OpImageQuerySize : SPIRV::OpImageQuerySizeLod;
115406c3fb27SDimitry Andric MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
1155bdd1243dSDimitry Andric auto MIB = MIRBuilder.buildInstr(Opcode)
1156bdd1243dSDimitry Andric .addDef(QueryResult)
1157bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(QueryResultType))
1158bdd1243dSDimitry Andric .addUse(Call->Arguments[0]);
1159bdd1243dSDimitry Andric if (!IsDimBuf)
1160bdd1243dSDimitry Andric MIB.addUse(buildConstantIntReg(0, MIRBuilder, GR)); // Lod id.
1161bdd1243dSDimitry Andric if (NumExpectedRetComponents == NumActualRetComponents)
1162bdd1243dSDimitry Andric return true;
1163bdd1243dSDimitry Andric if (NumExpectedRetComponents == 1) {
1164bdd1243dSDimitry Andric // Only 1 component is expected, build OpCompositeExtract instruction.
1165bdd1243dSDimitry Andric unsigned ExtractedComposite =
1166bdd1243dSDimitry Andric Component == 3 ? NumActualRetComponents - 1 : Component;
1167bdd1243dSDimitry Andric assert(ExtractedComposite < NumActualRetComponents &&
1168bdd1243dSDimitry Andric "Invalid composite index!");
1169bdd1243dSDimitry Andric MIRBuilder.buildInstr(SPIRV::OpCompositeExtract)
1170bdd1243dSDimitry Andric .addDef(Call->ReturnRegister)
1171bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1172bdd1243dSDimitry Andric .addUse(QueryResult)
1173bdd1243dSDimitry Andric .addImm(ExtractedComposite);
1174bdd1243dSDimitry Andric } else {
1175bdd1243dSDimitry Andric // More than 1 component is expected, fill a new vector.
1176bdd1243dSDimitry Andric auto MIB = MIRBuilder.buildInstr(SPIRV::OpVectorShuffle)
1177bdd1243dSDimitry Andric .addDef(Call->ReturnRegister)
1178bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1179bdd1243dSDimitry Andric .addUse(QueryResult)
1180bdd1243dSDimitry Andric .addUse(QueryResult);
1181bdd1243dSDimitry Andric for (unsigned i = 0; i < NumExpectedRetComponents; ++i)
1182bdd1243dSDimitry Andric MIB.addImm(i < NumActualRetComponents ? i : 0xffffffff);
1183bdd1243dSDimitry Andric }
1184bdd1243dSDimitry Andric return true;
1185bdd1243dSDimitry Andric }
1186bdd1243dSDimitry Andric
generateImageMiscQueryInst(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)1187bdd1243dSDimitry Andric static bool generateImageMiscQueryInst(const SPIRV::IncomingCall *Call,
1188bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
1189bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
1190bdd1243dSDimitry Andric assert(Call->ReturnType->getOpcode() == SPIRV::OpTypeInt &&
1191bdd1243dSDimitry Andric "Image samples query result must be of int type!");
1192bdd1243dSDimitry Andric
1193bdd1243dSDimitry Andric // Lookup the instruction opcode in the TableGen records.
1194bdd1243dSDimitry Andric const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1195bdd1243dSDimitry Andric unsigned Opcode =
1196bdd1243dSDimitry Andric SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1197bdd1243dSDimitry Andric
1198bdd1243dSDimitry Andric Register Image = Call->Arguments[0];
119906c3fb27SDimitry Andric MIRBuilder.getMRI()->setRegClass(Image, &SPIRV::IDRegClass);
1200bdd1243dSDimitry Andric SPIRV::Dim::Dim ImageDimensionality = static_cast<SPIRV::Dim::Dim>(
1201bdd1243dSDimitry Andric GR->getSPIRVTypeForVReg(Image)->getOperand(2).getImm());
1202bdd1243dSDimitry Andric
1203bdd1243dSDimitry Andric switch (Opcode) {
1204bdd1243dSDimitry Andric case SPIRV::OpImageQuerySamples:
1205bdd1243dSDimitry Andric assert(ImageDimensionality == SPIRV::Dim::DIM_2D &&
1206bdd1243dSDimitry Andric "Image must be of 2D dimensionality");
1207bdd1243dSDimitry Andric break;
1208bdd1243dSDimitry Andric case SPIRV::OpImageQueryLevels:
1209bdd1243dSDimitry Andric assert((ImageDimensionality == SPIRV::Dim::DIM_1D ||
1210bdd1243dSDimitry Andric ImageDimensionality == SPIRV::Dim::DIM_2D ||
1211bdd1243dSDimitry Andric ImageDimensionality == SPIRV::Dim::DIM_3D ||
1212bdd1243dSDimitry Andric ImageDimensionality == SPIRV::Dim::DIM_Cube) &&
1213bdd1243dSDimitry Andric "Image must be of 1D/2D/3D/Cube dimensionality");
1214bdd1243dSDimitry Andric break;
1215bdd1243dSDimitry Andric }
1216bdd1243dSDimitry Andric
1217bdd1243dSDimitry Andric MIRBuilder.buildInstr(Opcode)
1218bdd1243dSDimitry Andric .addDef(Call->ReturnRegister)
1219bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1220bdd1243dSDimitry Andric .addUse(Image);
1221bdd1243dSDimitry Andric return true;
1222bdd1243dSDimitry Andric }
1223bdd1243dSDimitry Andric
1224bdd1243dSDimitry Andric // TODO: Move to TableGen.
1225bdd1243dSDimitry Andric static SPIRV::SamplerAddressingMode::SamplerAddressingMode
getSamplerAddressingModeFromBitmask(unsigned Bitmask)1226bdd1243dSDimitry Andric getSamplerAddressingModeFromBitmask(unsigned Bitmask) {
1227bdd1243dSDimitry Andric switch (Bitmask & SPIRV::CLK_ADDRESS_MODE_MASK) {
1228bdd1243dSDimitry Andric case SPIRV::CLK_ADDRESS_CLAMP:
1229bdd1243dSDimitry Andric return SPIRV::SamplerAddressingMode::Clamp;
1230bdd1243dSDimitry Andric case SPIRV::CLK_ADDRESS_CLAMP_TO_EDGE:
1231bdd1243dSDimitry Andric return SPIRV::SamplerAddressingMode::ClampToEdge;
1232bdd1243dSDimitry Andric case SPIRV::CLK_ADDRESS_REPEAT:
1233bdd1243dSDimitry Andric return SPIRV::SamplerAddressingMode::Repeat;
1234bdd1243dSDimitry Andric case SPIRV::CLK_ADDRESS_MIRRORED_REPEAT:
1235bdd1243dSDimitry Andric return SPIRV::SamplerAddressingMode::RepeatMirrored;
1236bdd1243dSDimitry Andric case SPIRV::CLK_ADDRESS_NONE:
1237bdd1243dSDimitry Andric return SPIRV::SamplerAddressingMode::None;
1238bdd1243dSDimitry Andric default:
1239bdd1243dSDimitry Andric llvm_unreachable("Unknown CL address mode");
1240bdd1243dSDimitry Andric }
1241bdd1243dSDimitry Andric }
1242bdd1243dSDimitry Andric
getSamplerParamFromBitmask(unsigned Bitmask)1243bdd1243dSDimitry Andric static unsigned getSamplerParamFromBitmask(unsigned Bitmask) {
1244bdd1243dSDimitry Andric return (Bitmask & SPIRV::CLK_NORMALIZED_COORDS_TRUE) ? 1 : 0;
1245bdd1243dSDimitry Andric }
1246bdd1243dSDimitry Andric
1247bdd1243dSDimitry Andric static SPIRV::SamplerFilterMode::SamplerFilterMode
getSamplerFilterModeFromBitmask(unsigned Bitmask)1248bdd1243dSDimitry Andric getSamplerFilterModeFromBitmask(unsigned Bitmask) {
1249bdd1243dSDimitry Andric if (Bitmask & SPIRV::CLK_FILTER_LINEAR)
1250bdd1243dSDimitry Andric return SPIRV::SamplerFilterMode::Linear;
1251bdd1243dSDimitry Andric if (Bitmask & SPIRV::CLK_FILTER_NEAREST)
1252bdd1243dSDimitry Andric return SPIRV::SamplerFilterMode::Nearest;
1253bdd1243dSDimitry Andric return SPIRV::SamplerFilterMode::Nearest;
1254bdd1243dSDimitry Andric }
1255bdd1243dSDimitry Andric
generateReadImageInst(const StringRef DemangledCall,const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)1256bdd1243dSDimitry Andric static bool generateReadImageInst(const StringRef DemangledCall,
1257bdd1243dSDimitry Andric const SPIRV::IncomingCall *Call,
1258bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
1259bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
1260bdd1243dSDimitry Andric Register Image = Call->Arguments[0];
1261bdd1243dSDimitry Andric MachineRegisterInfo *MRI = MIRBuilder.getMRI();
126206c3fb27SDimitry Andric MRI->setRegClass(Image, &SPIRV::IDRegClass);
126306c3fb27SDimitry Andric MRI->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
126406c3fb27SDimitry Andric bool HasOclSampler = DemangledCall.contains_insensitive("ocl_sampler");
126506c3fb27SDimitry Andric bool HasMsaa = DemangledCall.contains_insensitive("msaa");
126606c3fb27SDimitry Andric if (HasOclSampler || HasMsaa)
126706c3fb27SDimitry Andric MRI->setRegClass(Call->Arguments[2], &SPIRV::IDRegClass);
126806c3fb27SDimitry Andric if (HasOclSampler) {
1269bdd1243dSDimitry Andric Register Sampler = Call->Arguments[1];
1270bdd1243dSDimitry Andric
1271bdd1243dSDimitry Andric if (!GR->isScalarOfType(Sampler, SPIRV::OpTypeSampler) &&
1272bdd1243dSDimitry Andric getDefInstrMaybeConstant(Sampler, MRI)->getOperand(1).isCImm()) {
1273bdd1243dSDimitry Andric uint64_t SamplerMask = getIConstVal(Sampler, MRI);
1274bdd1243dSDimitry Andric Sampler = GR->buildConstantSampler(
1275bdd1243dSDimitry Andric Register(), getSamplerAddressingModeFromBitmask(SamplerMask),
1276bdd1243dSDimitry Andric getSamplerParamFromBitmask(SamplerMask),
1277bdd1243dSDimitry Andric getSamplerFilterModeFromBitmask(SamplerMask), MIRBuilder,
1278bdd1243dSDimitry Andric GR->getSPIRVTypeForVReg(Sampler));
1279bdd1243dSDimitry Andric }
1280bdd1243dSDimitry Andric SPIRVType *ImageType = GR->getSPIRVTypeForVReg(Image);
1281bdd1243dSDimitry Andric SPIRVType *SampledImageType =
1282bdd1243dSDimitry Andric GR->getOrCreateOpTypeSampledImage(ImageType, MIRBuilder);
1283bdd1243dSDimitry Andric Register SampledImage = MRI->createVirtualRegister(&SPIRV::IDRegClass);
1284bdd1243dSDimitry Andric
1285bdd1243dSDimitry Andric MIRBuilder.buildInstr(SPIRV::OpSampledImage)
1286bdd1243dSDimitry Andric .addDef(SampledImage)
1287bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(SampledImageType))
1288bdd1243dSDimitry Andric .addUse(Image)
1289bdd1243dSDimitry Andric .addUse(Sampler);
1290bdd1243dSDimitry Andric
1291bdd1243dSDimitry Andric Register Lod = GR->buildConstantFP(APFloat::getZero(APFloat::IEEEsingle()),
1292bdd1243dSDimitry Andric MIRBuilder);
1293bdd1243dSDimitry Andric SPIRVType *TempType = Call->ReturnType;
1294bdd1243dSDimitry Andric bool NeedsExtraction = false;
1295bdd1243dSDimitry Andric if (TempType->getOpcode() != SPIRV::OpTypeVector) {
1296bdd1243dSDimitry Andric TempType =
1297bdd1243dSDimitry Andric GR->getOrCreateSPIRVVectorType(Call->ReturnType, 4, MIRBuilder);
1298bdd1243dSDimitry Andric NeedsExtraction = true;
1299bdd1243dSDimitry Andric }
1300bdd1243dSDimitry Andric LLT LLType = LLT::scalar(GR->getScalarOrVectorBitWidth(TempType));
1301bdd1243dSDimitry Andric Register TempRegister = MRI->createGenericVirtualRegister(LLType);
130206c3fb27SDimitry Andric MRI->setRegClass(TempRegister, &SPIRV::IDRegClass);
1303bdd1243dSDimitry Andric GR->assignSPIRVTypeToVReg(TempType, TempRegister, MIRBuilder.getMF());
1304bdd1243dSDimitry Andric
1305bdd1243dSDimitry Andric MIRBuilder.buildInstr(SPIRV::OpImageSampleExplicitLod)
1306bdd1243dSDimitry Andric .addDef(NeedsExtraction ? TempRegister : Call->ReturnRegister)
1307bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(TempType))
1308bdd1243dSDimitry Andric .addUse(SampledImage)
1309bdd1243dSDimitry Andric .addUse(Call->Arguments[2]) // Coordinate.
1310bdd1243dSDimitry Andric .addImm(SPIRV::ImageOperand::Lod)
1311bdd1243dSDimitry Andric .addUse(Lod);
1312bdd1243dSDimitry Andric
1313bdd1243dSDimitry Andric if (NeedsExtraction)
1314bdd1243dSDimitry Andric MIRBuilder.buildInstr(SPIRV::OpCompositeExtract)
1315bdd1243dSDimitry Andric .addDef(Call->ReturnRegister)
1316bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1317bdd1243dSDimitry Andric .addUse(TempRegister)
1318bdd1243dSDimitry Andric .addImm(0);
131906c3fb27SDimitry Andric } else if (HasMsaa) {
1320bdd1243dSDimitry Andric MIRBuilder.buildInstr(SPIRV::OpImageRead)
1321bdd1243dSDimitry Andric .addDef(Call->ReturnRegister)
1322bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1323bdd1243dSDimitry Andric .addUse(Image)
1324bdd1243dSDimitry Andric .addUse(Call->Arguments[1]) // Coordinate.
1325bdd1243dSDimitry Andric .addImm(SPIRV::ImageOperand::Sample)
1326bdd1243dSDimitry Andric .addUse(Call->Arguments[2]);
1327bdd1243dSDimitry Andric } else {
1328bdd1243dSDimitry Andric MIRBuilder.buildInstr(SPIRV::OpImageRead)
1329bdd1243dSDimitry Andric .addDef(Call->ReturnRegister)
1330bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1331bdd1243dSDimitry Andric .addUse(Image)
1332bdd1243dSDimitry Andric .addUse(Call->Arguments[1]); // Coordinate.
1333bdd1243dSDimitry Andric }
1334bdd1243dSDimitry Andric return true;
1335bdd1243dSDimitry Andric }
1336bdd1243dSDimitry Andric
generateWriteImageInst(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)1337bdd1243dSDimitry Andric static bool generateWriteImageInst(const SPIRV::IncomingCall *Call,
1338bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
1339bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
134006c3fb27SDimitry Andric MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
134106c3fb27SDimitry Andric MIRBuilder.getMRI()->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
134206c3fb27SDimitry Andric MIRBuilder.getMRI()->setRegClass(Call->Arguments[2], &SPIRV::IDRegClass);
1343bdd1243dSDimitry Andric MIRBuilder.buildInstr(SPIRV::OpImageWrite)
1344bdd1243dSDimitry Andric .addUse(Call->Arguments[0]) // Image.
1345bdd1243dSDimitry Andric .addUse(Call->Arguments[1]) // Coordinate.
1346bdd1243dSDimitry Andric .addUse(Call->Arguments[2]); // Texel.
1347bdd1243dSDimitry Andric return true;
1348bdd1243dSDimitry Andric }
1349bdd1243dSDimitry Andric
generateSampleImageInst(const StringRef DemangledCall,const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)1350bdd1243dSDimitry Andric static bool generateSampleImageInst(const StringRef DemangledCall,
1351bdd1243dSDimitry Andric const SPIRV::IncomingCall *Call,
1352bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
1353bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
135406c3fb27SDimitry Andric MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1355bdd1243dSDimitry Andric if (Call->Builtin->Name.contains_insensitive(
1356bdd1243dSDimitry Andric "__translate_sampler_initializer")) {
1357bdd1243dSDimitry Andric // Build sampler literal.
135806c3fb27SDimitry Andric uint64_t Bitmask = getIConstVal(Call->Arguments[0], MRI);
1359bdd1243dSDimitry Andric Register Sampler = GR->buildConstantSampler(
1360bdd1243dSDimitry Andric Call->ReturnRegister, getSamplerAddressingModeFromBitmask(Bitmask),
1361bdd1243dSDimitry Andric getSamplerParamFromBitmask(Bitmask),
1362bdd1243dSDimitry Andric getSamplerFilterModeFromBitmask(Bitmask), MIRBuilder, Call->ReturnType);
1363bdd1243dSDimitry Andric return Sampler.isValid();
1364bdd1243dSDimitry Andric } else if (Call->Builtin->Name.contains_insensitive("__spirv_SampledImage")) {
1365bdd1243dSDimitry Andric // Create OpSampledImage.
1366bdd1243dSDimitry Andric Register Image = Call->Arguments[0];
1367bdd1243dSDimitry Andric SPIRVType *ImageType = GR->getSPIRVTypeForVReg(Image);
1368bdd1243dSDimitry Andric SPIRVType *SampledImageType =
1369bdd1243dSDimitry Andric GR->getOrCreateOpTypeSampledImage(ImageType, MIRBuilder);
1370bdd1243dSDimitry Andric Register SampledImage =
1371bdd1243dSDimitry Andric Call->ReturnRegister.isValid()
1372bdd1243dSDimitry Andric ? Call->ReturnRegister
137306c3fb27SDimitry Andric : MRI->createVirtualRegister(&SPIRV::IDRegClass);
1374bdd1243dSDimitry Andric MIRBuilder.buildInstr(SPIRV::OpSampledImage)
1375bdd1243dSDimitry Andric .addDef(SampledImage)
1376bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(SampledImageType))
1377bdd1243dSDimitry Andric .addUse(Image)
1378bdd1243dSDimitry Andric .addUse(Call->Arguments[1]); // Sampler.
1379bdd1243dSDimitry Andric return true;
1380bdd1243dSDimitry Andric } else if (Call->Builtin->Name.contains_insensitive(
1381bdd1243dSDimitry Andric "__spirv_ImageSampleExplicitLod")) {
1382bdd1243dSDimitry Andric // Sample an image using an explicit level of detail.
1383bdd1243dSDimitry Andric std::string ReturnType = DemangledCall.str();
1384bdd1243dSDimitry Andric if (DemangledCall.contains("_R")) {
1385bdd1243dSDimitry Andric ReturnType = ReturnType.substr(ReturnType.find("_R") + 2);
1386bdd1243dSDimitry Andric ReturnType = ReturnType.substr(0, ReturnType.find('('));
1387bdd1243dSDimitry Andric }
1388bdd1243dSDimitry Andric SPIRVType *Type = GR->getOrCreateSPIRVTypeByName(ReturnType, MIRBuilder);
138906c3fb27SDimitry Andric MRI->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
139006c3fb27SDimitry Andric MRI->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
139106c3fb27SDimitry Andric MRI->setRegClass(Call->Arguments[3], &SPIRV::IDRegClass);
139206c3fb27SDimitry Andric
1393bdd1243dSDimitry Andric MIRBuilder.buildInstr(SPIRV::OpImageSampleExplicitLod)
1394bdd1243dSDimitry Andric .addDef(Call->ReturnRegister)
1395bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(Type))
1396bdd1243dSDimitry Andric .addUse(Call->Arguments[0]) // Image.
1397bdd1243dSDimitry Andric .addUse(Call->Arguments[1]) // Coordinate.
1398bdd1243dSDimitry Andric .addImm(SPIRV::ImageOperand::Lod)
1399bdd1243dSDimitry Andric .addUse(Call->Arguments[3]);
1400bdd1243dSDimitry Andric return true;
1401bdd1243dSDimitry Andric }
1402bdd1243dSDimitry Andric return false;
1403bdd1243dSDimitry Andric }
1404bdd1243dSDimitry Andric
generateSelectInst(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder)1405bdd1243dSDimitry Andric static bool generateSelectInst(const SPIRV::IncomingCall *Call,
1406bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder) {
1407bdd1243dSDimitry Andric MIRBuilder.buildSelect(Call->ReturnRegister, Call->Arguments[0],
1408bdd1243dSDimitry Andric Call->Arguments[1], Call->Arguments[2]);
1409bdd1243dSDimitry Andric return true;
1410bdd1243dSDimitry Andric }
1411bdd1243dSDimitry Andric
generateSpecConstantInst(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)1412bdd1243dSDimitry Andric static bool generateSpecConstantInst(const SPIRV::IncomingCall *Call,
1413bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
1414bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
1415bdd1243dSDimitry Andric // Lookup the instruction opcode in the TableGen records.
1416bdd1243dSDimitry Andric const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1417bdd1243dSDimitry Andric unsigned Opcode =
1418bdd1243dSDimitry Andric SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1419bdd1243dSDimitry Andric const MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1420bdd1243dSDimitry Andric
1421bdd1243dSDimitry Andric switch (Opcode) {
1422bdd1243dSDimitry Andric case SPIRV::OpSpecConstant: {
1423bdd1243dSDimitry Andric // Build the SpecID decoration.
1424bdd1243dSDimitry Andric unsigned SpecId =
1425bdd1243dSDimitry Andric static_cast<unsigned>(getIConstVal(Call->Arguments[0], MRI));
1426bdd1243dSDimitry Andric buildOpDecorate(Call->ReturnRegister, MIRBuilder, SPIRV::Decoration::SpecId,
1427bdd1243dSDimitry Andric {SpecId});
1428bdd1243dSDimitry Andric // Determine the constant MI.
1429bdd1243dSDimitry Andric Register ConstRegister = Call->Arguments[1];
1430bdd1243dSDimitry Andric const MachineInstr *Const = getDefInstrMaybeConstant(ConstRegister, MRI);
1431bdd1243dSDimitry Andric assert(Const &&
1432bdd1243dSDimitry Andric (Const->getOpcode() == TargetOpcode::G_CONSTANT ||
1433bdd1243dSDimitry Andric Const->getOpcode() == TargetOpcode::G_FCONSTANT) &&
1434bdd1243dSDimitry Andric "Argument should be either an int or floating-point constant");
1435bdd1243dSDimitry Andric // Determine the opcode and built the OpSpec MI.
1436bdd1243dSDimitry Andric const MachineOperand &ConstOperand = Const->getOperand(1);
1437bdd1243dSDimitry Andric if (Call->ReturnType->getOpcode() == SPIRV::OpTypeBool) {
1438bdd1243dSDimitry Andric assert(ConstOperand.isCImm() && "Int constant operand is expected");
1439bdd1243dSDimitry Andric Opcode = ConstOperand.getCImm()->getValue().getZExtValue()
1440bdd1243dSDimitry Andric ? SPIRV::OpSpecConstantTrue
1441bdd1243dSDimitry Andric : SPIRV::OpSpecConstantFalse;
1442bdd1243dSDimitry Andric }
1443bdd1243dSDimitry Andric auto MIB = MIRBuilder.buildInstr(Opcode)
1444bdd1243dSDimitry Andric .addDef(Call->ReturnRegister)
1445bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(Call->ReturnType));
1446bdd1243dSDimitry Andric
1447bdd1243dSDimitry Andric if (Call->ReturnType->getOpcode() != SPIRV::OpTypeBool) {
1448bdd1243dSDimitry Andric if (Const->getOpcode() == TargetOpcode::G_CONSTANT)
1449bdd1243dSDimitry Andric addNumImm(ConstOperand.getCImm()->getValue(), MIB);
1450bdd1243dSDimitry Andric else
1451bdd1243dSDimitry Andric addNumImm(ConstOperand.getFPImm()->getValueAPF().bitcastToAPInt(), MIB);
1452bdd1243dSDimitry Andric }
1453bdd1243dSDimitry Andric return true;
1454bdd1243dSDimitry Andric }
1455bdd1243dSDimitry Andric case SPIRV::OpSpecConstantComposite: {
1456bdd1243dSDimitry Andric auto MIB = MIRBuilder.buildInstr(Opcode)
1457bdd1243dSDimitry Andric .addDef(Call->ReturnRegister)
1458bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(Call->ReturnType));
1459bdd1243dSDimitry Andric for (unsigned i = 0; i < Call->Arguments.size(); i++)
1460bdd1243dSDimitry Andric MIB.addUse(Call->Arguments[i]);
1461bdd1243dSDimitry Andric return true;
1462bdd1243dSDimitry Andric }
1463bdd1243dSDimitry Andric default:
1464bdd1243dSDimitry Andric return false;
1465bdd1243dSDimitry Andric }
1466bdd1243dSDimitry Andric }
1467bdd1243dSDimitry Andric
buildNDRange(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)146806c3fb27SDimitry Andric static bool buildNDRange(const SPIRV::IncomingCall *Call,
146906c3fb27SDimitry Andric MachineIRBuilder &MIRBuilder,
147006c3fb27SDimitry Andric SPIRVGlobalRegistry *GR) {
147106c3fb27SDimitry Andric MachineRegisterInfo *MRI = MIRBuilder.getMRI();
147206c3fb27SDimitry Andric MRI->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
147306c3fb27SDimitry Andric SPIRVType *PtrType = GR->getSPIRVTypeForVReg(Call->Arguments[0]);
147406c3fb27SDimitry Andric assert(PtrType->getOpcode() == SPIRV::OpTypePointer &&
147506c3fb27SDimitry Andric PtrType->getOperand(2).isReg());
147606c3fb27SDimitry Andric Register TypeReg = PtrType->getOperand(2).getReg();
147706c3fb27SDimitry Andric SPIRVType *StructType = GR->getSPIRVTypeForVReg(TypeReg);
147806c3fb27SDimitry Andric MachineFunction &MF = MIRBuilder.getMF();
147906c3fb27SDimitry Andric Register TmpReg = MRI->createVirtualRegister(&SPIRV::IDRegClass);
148006c3fb27SDimitry Andric GR->assignSPIRVTypeToVReg(StructType, TmpReg, MF);
148106c3fb27SDimitry Andric // Skip the first arg, it's the destination pointer. OpBuildNDRange takes
148206c3fb27SDimitry Andric // three other arguments, so pass zero constant on absence.
148306c3fb27SDimitry Andric unsigned NumArgs = Call->Arguments.size();
148406c3fb27SDimitry Andric assert(NumArgs >= 2);
148506c3fb27SDimitry Andric Register GlobalWorkSize = Call->Arguments[NumArgs < 4 ? 1 : 2];
148606c3fb27SDimitry Andric MRI->setRegClass(GlobalWorkSize, &SPIRV::IDRegClass);
148706c3fb27SDimitry Andric Register LocalWorkSize =
148806c3fb27SDimitry Andric NumArgs == 2 ? Register(0) : Call->Arguments[NumArgs < 4 ? 2 : 3];
148906c3fb27SDimitry Andric if (LocalWorkSize.isValid())
149006c3fb27SDimitry Andric MRI->setRegClass(LocalWorkSize, &SPIRV::IDRegClass);
149106c3fb27SDimitry Andric Register GlobalWorkOffset = NumArgs <= 3 ? Register(0) : Call->Arguments[1];
149206c3fb27SDimitry Andric if (GlobalWorkOffset.isValid())
149306c3fb27SDimitry Andric MRI->setRegClass(GlobalWorkOffset, &SPIRV::IDRegClass);
149406c3fb27SDimitry Andric if (NumArgs < 4) {
149506c3fb27SDimitry Andric Register Const;
149606c3fb27SDimitry Andric SPIRVType *SpvTy = GR->getSPIRVTypeForVReg(GlobalWorkSize);
149706c3fb27SDimitry Andric if (SpvTy->getOpcode() == SPIRV::OpTypePointer) {
149806c3fb27SDimitry Andric MachineInstr *DefInstr = MRI->getUniqueVRegDef(GlobalWorkSize);
149906c3fb27SDimitry Andric assert(DefInstr && isSpvIntrinsic(*DefInstr, Intrinsic::spv_gep) &&
150006c3fb27SDimitry Andric DefInstr->getOperand(3).isReg());
150106c3fb27SDimitry Andric Register GWSPtr = DefInstr->getOperand(3).getReg();
150206c3fb27SDimitry Andric if (!MRI->getRegClassOrNull(GWSPtr))
150306c3fb27SDimitry Andric MRI->setRegClass(GWSPtr, &SPIRV::IDRegClass);
150406c3fb27SDimitry Andric // TODO: Maybe simplify generation of the type of the fields.
150506c3fb27SDimitry Andric unsigned Size = Call->Builtin->Name.equals("ndrange_3D") ? 3 : 2;
150606c3fb27SDimitry Andric unsigned BitWidth = GR->getPointerSize() == 64 ? 64 : 32;
150706c3fb27SDimitry Andric Type *BaseTy = IntegerType::get(MF.getFunction().getContext(), BitWidth);
150806c3fb27SDimitry Andric Type *FieldTy = ArrayType::get(BaseTy, Size);
150906c3fb27SDimitry Andric SPIRVType *SpvFieldTy = GR->getOrCreateSPIRVType(FieldTy, MIRBuilder);
151006c3fb27SDimitry Andric GlobalWorkSize = MRI->createVirtualRegister(&SPIRV::IDRegClass);
151106c3fb27SDimitry Andric GR->assignSPIRVTypeToVReg(SpvFieldTy, GlobalWorkSize, MF);
151206c3fb27SDimitry Andric MIRBuilder.buildInstr(SPIRV::OpLoad)
151306c3fb27SDimitry Andric .addDef(GlobalWorkSize)
151406c3fb27SDimitry Andric .addUse(GR->getSPIRVTypeID(SpvFieldTy))
151506c3fb27SDimitry Andric .addUse(GWSPtr);
151606c3fb27SDimitry Andric Const = GR->getOrCreateConsIntArray(0, MIRBuilder, SpvFieldTy);
151706c3fb27SDimitry Andric } else {
151806c3fb27SDimitry Andric Const = GR->buildConstantInt(0, MIRBuilder, SpvTy);
151906c3fb27SDimitry Andric }
152006c3fb27SDimitry Andric if (!LocalWorkSize.isValid())
152106c3fb27SDimitry Andric LocalWorkSize = Const;
152206c3fb27SDimitry Andric if (!GlobalWorkOffset.isValid())
152306c3fb27SDimitry Andric GlobalWorkOffset = Const;
152406c3fb27SDimitry Andric }
152506c3fb27SDimitry Andric assert(LocalWorkSize.isValid() && GlobalWorkOffset.isValid());
152606c3fb27SDimitry Andric MIRBuilder.buildInstr(SPIRV::OpBuildNDRange)
152706c3fb27SDimitry Andric .addDef(TmpReg)
152806c3fb27SDimitry Andric .addUse(TypeReg)
152906c3fb27SDimitry Andric .addUse(GlobalWorkSize)
153006c3fb27SDimitry Andric .addUse(LocalWorkSize)
153106c3fb27SDimitry Andric .addUse(GlobalWorkOffset);
153206c3fb27SDimitry Andric return MIRBuilder.buildInstr(SPIRV::OpStore)
153306c3fb27SDimitry Andric .addUse(Call->Arguments[0])
153406c3fb27SDimitry Andric .addUse(TmpReg);
153506c3fb27SDimitry Andric }
153606c3fb27SDimitry Andric
getBlockStructInstr(Register ParamReg,MachineRegisterInfo * MRI)1537bdd1243dSDimitry Andric static MachineInstr *getBlockStructInstr(Register ParamReg,
1538bdd1243dSDimitry Andric MachineRegisterInfo *MRI) {
1539bdd1243dSDimitry Andric // We expect the following sequence of instructions:
1540bdd1243dSDimitry Andric // %0:_(pN) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.alloca)
1541bdd1243dSDimitry Andric // or = G_GLOBAL_VALUE @block_literal_global
1542bdd1243dSDimitry Andric // %1:_(pN) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.bitcast), %0
1543bdd1243dSDimitry Andric // %2:_(p4) = G_ADDRSPACE_CAST %1:_(pN)
1544bdd1243dSDimitry Andric MachineInstr *MI = MRI->getUniqueVRegDef(ParamReg);
1545bdd1243dSDimitry Andric assert(MI->getOpcode() == TargetOpcode::G_ADDRSPACE_CAST &&
1546bdd1243dSDimitry Andric MI->getOperand(1).isReg());
1547bdd1243dSDimitry Andric Register BitcastReg = MI->getOperand(1).getReg();
1548bdd1243dSDimitry Andric MachineInstr *BitcastMI = MRI->getUniqueVRegDef(BitcastReg);
1549bdd1243dSDimitry Andric assert(isSpvIntrinsic(*BitcastMI, Intrinsic::spv_bitcast) &&
1550bdd1243dSDimitry Andric BitcastMI->getOperand(2).isReg());
1551bdd1243dSDimitry Andric Register ValueReg = BitcastMI->getOperand(2).getReg();
1552bdd1243dSDimitry Andric MachineInstr *ValueMI = MRI->getUniqueVRegDef(ValueReg);
1553bdd1243dSDimitry Andric return ValueMI;
1554bdd1243dSDimitry Andric }
1555bdd1243dSDimitry Andric
1556bdd1243dSDimitry Andric // Return an integer constant corresponding to the given register and
1557bdd1243dSDimitry Andric // defined in spv_track_constant.
1558bdd1243dSDimitry Andric // TODO: maybe unify with prelegalizer pass.
getConstFromIntrinsic(Register Reg,MachineRegisterInfo * MRI)1559bdd1243dSDimitry Andric static unsigned getConstFromIntrinsic(Register Reg, MachineRegisterInfo *MRI) {
1560bdd1243dSDimitry Andric MachineInstr *DefMI = MRI->getUniqueVRegDef(Reg);
1561bdd1243dSDimitry Andric assert(isSpvIntrinsic(*DefMI, Intrinsic::spv_track_constant) &&
1562bdd1243dSDimitry Andric DefMI->getOperand(2).isReg());
1563bdd1243dSDimitry Andric MachineInstr *DefMI2 = MRI->getUniqueVRegDef(DefMI->getOperand(2).getReg());
1564bdd1243dSDimitry Andric assert(DefMI2->getOpcode() == TargetOpcode::G_CONSTANT &&
1565bdd1243dSDimitry Andric DefMI2->getOperand(1).isCImm());
1566bdd1243dSDimitry Andric return DefMI2->getOperand(1).getCImm()->getValue().getZExtValue();
1567bdd1243dSDimitry Andric }
1568bdd1243dSDimitry Andric
1569bdd1243dSDimitry Andric // Return type of the instruction result from spv_assign_type intrinsic.
1570bdd1243dSDimitry Andric // TODO: maybe unify with prelegalizer pass.
getMachineInstrType(MachineInstr * MI)1571bdd1243dSDimitry Andric static const Type *getMachineInstrType(MachineInstr *MI) {
1572bdd1243dSDimitry Andric MachineInstr *NextMI = MI->getNextNode();
1573bdd1243dSDimitry Andric if (isSpvIntrinsic(*NextMI, Intrinsic::spv_assign_name))
1574bdd1243dSDimitry Andric NextMI = NextMI->getNextNode();
1575bdd1243dSDimitry Andric Register ValueReg = MI->getOperand(0).getReg();
1576bdd1243dSDimitry Andric if (!isSpvIntrinsic(*NextMI, Intrinsic::spv_assign_type) ||
1577bdd1243dSDimitry Andric NextMI->getOperand(1).getReg() != ValueReg)
1578bdd1243dSDimitry Andric return nullptr;
1579bdd1243dSDimitry Andric Type *Ty = getMDOperandAsType(NextMI->getOperand(2).getMetadata(), 0);
1580bdd1243dSDimitry Andric assert(Ty && "Type is expected");
1581bdd1243dSDimitry Andric return getTypedPtrEltType(Ty);
1582bdd1243dSDimitry Andric }
1583bdd1243dSDimitry Andric
getBlockStructType(Register ParamReg,MachineRegisterInfo * MRI)1584bdd1243dSDimitry Andric static const Type *getBlockStructType(Register ParamReg,
1585bdd1243dSDimitry Andric MachineRegisterInfo *MRI) {
1586bdd1243dSDimitry Andric // In principle, this information should be passed to us from Clang via
1587bdd1243dSDimitry Andric // an elementtype attribute. However, said attribute requires that
1588bdd1243dSDimitry Andric // the function call be an intrinsic, which is not. Instead, we rely on being
1589bdd1243dSDimitry Andric // able to trace this to the declaration of a variable: OpenCL C specification
1590bdd1243dSDimitry Andric // section 6.12.5 should guarantee that we can do this.
1591bdd1243dSDimitry Andric MachineInstr *MI = getBlockStructInstr(ParamReg, MRI);
1592bdd1243dSDimitry Andric if (MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
1593bdd1243dSDimitry Andric return getTypedPtrEltType(MI->getOperand(1).getGlobal()->getType());
1594bdd1243dSDimitry Andric assert(isSpvIntrinsic(*MI, Intrinsic::spv_alloca) &&
1595bdd1243dSDimitry Andric "Blocks in OpenCL C must be traceable to allocation site");
1596bdd1243dSDimitry Andric return getMachineInstrType(MI);
1597bdd1243dSDimitry Andric }
1598bdd1243dSDimitry Andric
1599bdd1243dSDimitry Andric // TODO: maybe move to the global register.
1600bdd1243dSDimitry Andric static SPIRVType *
getOrCreateSPIRVDeviceEventPointer(MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)1601bdd1243dSDimitry Andric getOrCreateSPIRVDeviceEventPointer(MachineIRBuilder &MIRBuilder,
1602bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
1603bdd1243dSDimitry Andric LLVMContext &Context = MIRBuilder.getMF().getFunction().getContext();
1604bdd1243dSDimitry Andric Type *OpaqueType = StructType::getTypeByName(Context, "spirv.DeviceEvent");
1605bdd1243dSDimitry Andric if (!OpaqueType)
1606bdd1243dSDimitry Andric OpaqueType = StructType::getTypeByName(Context, "opencl.clk_event_t");
1607bdd1243dSDimitry Andric if (!OpaqueType)
1608bdd1243dSDimitry Andric OpaqueType = StructType::create(Context, "spirv.DeviceEvent");
1609bdd1243dSDimitry Andric unsigned SC0 = storageClassToAddressSpace(SPIRV::StorageClass::Function);
1610bdd1243dSDimitry Andric unsigned SC1 = storageClassToAddressSpace(SPIRV::StorageClass::Generic);
1611bdd1243dSDimitry Andric Type *PtrType = PointerType::get(PointerType::get(OpaqueType, SC0), SC1);
1612bdd1243dSDimitry Andric return GR->getOrCreateSPIRVType(PtrType, MIRBuilder);
1613bdd1243dSDimitry Andric }
1614bdd1243dSDimitry Andric
buildEnqueueKernel(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)1615bdd1243dSDimitry Andric static bool buildEnqueueKernel(const SPIRV::IncomingCall *Call,
1616bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
1617bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
1618bdd1243dSDimitry Andric MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1619bdd1243dSDimitry Andric const DataLayout &DL = MIRBuilder.getDataLayout();
1620*cb14a3feSDimitry Andric bool HasEvents = Call->Builtin->Name.contains("events");
1621bdd1243dSDimitry Andric const SPIRVType *Int32Ty = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
1622bdd1243dSDimitry Andric
1623bdd1243dSDimitry Andric // Make vararg instructions before OpEnqueueKernel.
1624bdd1243dSDimitry Andric // Local sizes arguments: Sizes of block invoke arguments. Clang generates
1625bdd1243dSDimitry Andric // local size operands as an array, so we need to unpack them.
1626bdd1243dSDimitry Andric SmallVector<Register, 16> LocalSizes;
1627bdd1243dSDimitry Andric if (Call->Builtin->Name.find("_varargs") != StringRef::npos) {
1628bdd1243dSDimitry Andric const unsigned LocalSizeArrayIdx = HasEvents ? 9 : 6;
1629bdd1243dSDimitry Andric Register GepReg = Call->Arguments[LocalSizeArrayIdx];
1630bdd1243dSDimitry Andric MachineInstr *GepMI = MRI->getUniqueVRegDef(GepReg);
1631bdd1243dSDimitry Andric assert(isSpvIntrinsic(*GepMI, Intrinsic::spv_gep) &&
1632bdd1243dSDimitry Andric GepMI->getOperand(3).isReg());
1633bdd1243dSDimitry Andric Register ArrayReg = GepMI->getOperand(3).getReg();
1634bdd1243dSDimitry Andric MachineInstr *ArrayMI = MRI->getUniqueVRegDef(ArrayReg);
1635bdd1243dSDimitry Andric const Type *LocalSizeTy = getMachineInstrType(ArrayMI);
1636bdd1243dSDimitry Andric assert(LocalSizeTy && "Local size type is expected");
1637bdd1243dSDimitry Andric const uint64_t LocalSizeNum =
1638bdd1243dSDimitry Andric cast<ArrayType>(LocalSizeTy)->getNumElements();
1639bdd1243dSDimitry Andric unsigned SC = storageClassToAddressSpace(SPIRV::StorageClass::Generic);
1640bdd1243dSDimitry Andric const LLT LLType = LLT::pointer(SC, GR->getPointerSize());
1641bdd1243dSDimitry Andric const SPIRVType *PointerSizeTy = GR->getOrCreateSPIRVPointerType(
1642bdd1243dSDimitry Andric Int32Ty, MIRBuilder, SPIRV::StorageClass::Function);
1643bdd1243dSDimitry Andric for (unsigned I = 0; I < LocalSizeNum; ++I) {
164406c3fb27SDimitry Andric Register Reg = MRI->createVirtualRegister(&SPIRV::IDRegClass);
164506c3fb27SDimitry Andric MRI->setType(Reg, LLType);
1646bdd1243dSDimitry Andric GR->assignSPIRVTypeToVReg(PointerSizeTy, Reg, MIRBuilder.getMF());
16475f757f3fSDimitry Andric auto GEPInst = MIRBuilder.buildIntrinsic(
16485f757f3fSDimitry Andric Intrinsic::spv_gep, ArrayRef<Register>{Reg}, true, false);
1649bdd1243dSDimitry Andric GEPInst
1650bdd1243dSDimitry Andric .addImm(GepMI->getOperand(2).getImm()) // In bound.
1651bdd1243dSDimitry Andric .addUse(ArrayMI->getOperand(0).getReg()) // Alloca.
1652bdd1243dSDimitry Andric .addUse(buildConstantIntReg(0, MIRBuilder, GR)) // Indices.
1653bdd1243dSDimitry Andric .addUse(buildConstantIntReg(I, MIRBuilder, GR));
1654bdd1243dSDimitry Andric LocalSizes.push_back(Reg);
1655bdd1243dSDimitry Andric }
1656bdd1243dSDimitry Andric }
1657bdd1243dSDimitry Andric
1658bdd1243dSDimitry Andric // SPIRV OpEnqueueKernel instruction has 10+ arguments.
1659bdd1243dSDimitry Andric auto MIB = MIRBuilder.buildInstr(SPIRV::OpEnqueueKernel)
1660bdd1243dSDimitry Andric .addDef(Call->ReturnRegister)
1661bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(Int32Ty));
1662bdd1243dSDimitry Andric
1663bdd1243dSDimitry Andric // Copy all arguments before block invoke function pointer.
1664bdd1243dSDimitry Andric const unsigned BlockFIdx = HasEvents ? 6 : 3;
1665bdd1243dSDimitry Andric for (unsigned i = 0; i < BlockFIdx; i++)
1666bdd1243dSDimitry Andric MIB.addUse(Call->Arguments[i]);
1667bdd1243dSDimitry Andric
1668bdd1243dSDimitry Andric // If there are no event arguments in the original call, add dummy ones.
1669bdd1243dSDimitry Andric if (!HasEvents) {
1670bdd1243dSDimitry Andric MIB.addUse(buildConstantIntReg(0, MIRBuilder, GR)); // Dummy num events.
1671bdd1243dSDimitry Andric Register NullPtr = GR->getOrCreateConstNullPtr(
1672bdd1243dSDimitry Andric MIRBuilder, getOrCreateSPIRVDeviceEventPointer(MIRBuilder, GR));
1673bdd1243dSDimitry Andric MIB.addUse(NullPtr); // Dummy wait events.
1674bdd1243dSDimitry Andric MIB.addUse(NullPtr); // Dummy ret event.
1675bdd1243dSDimitry Andric }
1676bdd1243dSDimitry Andric
1677bdd1243dSDimitry Andric MachineInstr *BlockMI = getBlockStructInstr(Call->Arguments[BlockFIdx], MRI);
1678bdd1243dSDimitry Andric assert(BlockMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE);
1679bdd1243dSDimitry Andric // Invoke: Pointer to invoke function.
1680bdd1243dSDimitry Andric MIB.addGlobalAddress(BlockMI->getOperand(1).getGlobal());
1681bdd1243dSDimitry Andric
1682bdd1243dSDimitry Andric Register BlockLiteralReg = Call->Arguments[BlockFIdx + 1];
1683bdd1243dSDimitry Andric // Param: Pointer to block literal.
1684bdd1243dSDimitry Andric MIB.addUse(BlockLiteralReg);
1685bdd1243dSDimitry Andric
1686bdd1243dSDimitry Andric Type *PType = const_cast<Type *>(getBlockStructType(BlockLiteralReg, MRI));
1687bdd1243dSDimitry Andric // TODO: these numbers should be obtained from block literal structure.
1688bdd1243dSDimitry Andric // Param Size: Size of block literal structure.
1689bdd1243dSDimitry Andric MIB.addUse(buildConstantIntReg(DL.getTypeStoreSize(PType), MIRBuilder, GR));
1690bdd1243dSDimitry Andric // Param Aligment: Aligment of block literal structure.
1691bdd1243dSDimitry Andric MIB.addUse(
169206c3fb27SDimitry Andric buildConstantIntReg(DL.getPrefTypeAlign(PType).value(), MIRBuilder, GR));
1693bdd1243dSDimitry Andric
1694bdd1243dSDimitry Andric for (unsigned i = 0; i < LocalSizes.size(); i++)
1695bdd1243dSDimitry Andric MIB.addUse(LocalSizes[i]);
1696bdd1243dSDimitry Andric return true;
1697bdd1243dSDimitry Andric }
1698bdd1243dSDimitry Andric
generateEnqueueInst(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)1699bdd1243dSDimitry Andric static bool generateEnqueueInst(const SPIRV::IncomingCall *Call,
1700bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
1701bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
1702bdd1243dSDimitry Andric // Lookup the instruction opcode in the TableGen records.
1703bdd1243dSDimitry Andric const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1704bdd1243dSDimitry Andric unsigned Opcode =
1705bdd1243dSDimitry Andric SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1706bdd1243dSDimitry Andric
1707bdd1243dSDimitry Andric switch (Opcode) {
1708bdd1243dSDimitry Andric case SPIRV::OpRetainEvent:
1709bdd1243dSDimitry Andric case SPIRV::OpReleaseEvent:
171006c3fb27SDimitry Andric MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
1711bdd1243dSDimitry Andric return MIRBuilder.buildInstr(Opcode).addUse(Call->Arguments[0]);
1712bdd1243dSDimitry Andric case SPIRV::OpCreateUserEvent:
1713bdd1243dSDimitry Andric case SPIRV::OpGetDefaultQueue:
1714bdd1243dSDimitry Andric return MIRBuilder.buildInstr(Opcode)
1715bdd1243dSDimitry Andric .addDef(Call->ReturnRegister)
1716bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(Call->ReturnType));
1717bdd1243dSDimitry Andric case SPIRV::OpIsValidEvent:
171806c3fb27SDimitry Andric MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
1719bdd1243dSDimitry Andric return MIRBuilder.buildInstr(Opcode)
1720bdd1243dSDimitry Andric .addDef(Call->ReturnRegister)
1721bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1722bdd1243dSDimitry Andric .addUse(Call->Arguments[0]);
1723bdd1243dSDimitry Andric case SPIRV::OpSetUserEventStatus:
172406c3fb27SDimitry Andric MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
172506c3fb27SDimitry Andric MIRBuilder.getMRI()->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
1726bdd1243dSDimitry Andric return MIRBuilder.buildInstr(Opcode)
1727bdd1243dSDimitry Andric .addUse(Call->Arguments[0])
1728bdd1243dSDimitry Andric .addUse(Call->Arguments[1]);
1729bdd1243dSDimitry Andric case SPIRV::OpCaptureEventProfilingInfo:
173006c3fb27SDimitry Andric MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
173106c3fb27SDimitry Andric MIRBuilder.getMRI()->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
173206c3fb27SDimitry Andric MIRBuilder.getMRI()->setRegClass(Call->Arguments[2], &SPIRV::IDRegClass);
1733bdd1243dSDimitry Andric return MIRBuilder.buildInstr(Opcode)
1734bdd1243dSDimitry Andric .addUse(Call->Arguments[0])
1735bdd1243dSDimitry Andric .addUse(Call->Arguments[1])
1736bdd1243dSDimitry Andric .addUse(Call->Arguments[2]);
173706c3fb27SDimitry Andric case SPIRV::OpBuildNDRange:
173806c3fb27SDimitry Andric return buildNDRange(Call, MIRBuilder, GR);
1739bdd1243dSDimitry Andric case SPIRV::OpEnqueueKernel:
1740bdd1243dSDimitry Andric return buildEnqueueKernel(Call, MIRBuilder, GR);
1741bdd1243dSDimitry Andric default:
1742bdd1243dSDimitry Andric return false;
1743bdd1243dSDimitry Andric }
1744bdd1243dSDimitry Andric }
1745bdd1243dSDimitry Andric
generateAsyncCopy(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)1746bdd1243dSDimitry Andric static bool generateAsyncCopy(const SPIRV::IncomingCall *Call,
1747bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
1748bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
1749bdd1243dSDimitry Andric // Lookup the instruction opcode in the TableGen records.
1750bdd1243dSDimitry Andric const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1751bdd1243dSDimitry Andric unsigned Opcode =
1752bdd1243dSDimitry Andric SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1753bdd1243dSDimitry Andric auto Scope = buildConstantIntReg(SPIRV::Scope::Workgroup, MIRBuilder, GR);
1754bdd1243dSDimitry Andric
1755bdd1243dSDimitry Andric switch (Opcode) {
1756bdd1243dSDimitry Andric case SPIRV::OpGroupAsyncCopy:
1757bdd1243dSDimitry Andric return MIRBuilder.buildInstr(Opcode)
1758bdd1243dSDimitry Andric .addDef(Call->ReturnRegister)
1759bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1760bdd1243dSDimitry Andric .addUse(Scope)
1761bdd1243dSDimitry Andric .addUse(Call->Arguments[0])
1762bdd1243dSDimitry Andric .addUse(Call->Arguments[1])
1763bdd1243dSDimitry Andric .addUse(Call->Arguments[2])
1764bdd1243dSDimitry Andric .addUse(buildConstantIntReg(1, MIRBuilder, GR))
1765bdd1243dSDimitry Andric .addUse(Call->Arguments[3]);
1766bdd1243dSDimitry Andric case SPIRV::OpGroupWaitEvents:
1767bdd1243dSDimitry Andric return MIRBuilder.buildInstr(Opcode)
1768bdd1243dSDimitry Andric .addUse(Scope)
1769bdd1243dSDimitry Andric .addUse(Call->Arguments[0])
1770bdd1243dSDimitry Andric .addUse(Call->Arguments[1]);
1771bdd1243dSDimitry Andric default:
1772bdd1243dSDimitry Andric return false;
1773bdd1243dSDimitry Andric }
1774bdd1243dSDimitry Andric }
1775bdd1243dSDimitry Andric
generateConvertInst(const StringRef DemangledCall,const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)1776bdd1243dSDimitry Andric static bool generateConvertInst(const StringRef DemangledCall,
1777bdd1243dSDimitry Andric const SPIRV::IncomingCall *Call,
1778bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
1779bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
1780bdd1243dSDimitry Andric // Lookup the conversion builtin in the TableGen records.
1781bdd1243dSDimitry Andric const SPIRV::ConvertBuiltin *Builtin =
1782bdd1243dSDimitry Andric SPIRV::lookupConvertBuiltin(Call->Builtin->Name, Call->Builtin->Set);
1783bdd1243dSDimitry Andric
1784bdd1243dSDimitry Andric if (Builtin->IsSaturated)
1785bdd1243dSDimitry Andric buildOpDecorate(Call->ReturnRegister, MIRBuilder,
1786bdd1243dSDimitry Andric SPIRV::Decoration::SaturatedConversion, {});
1787bdd1243dSDimitry Andric if (Builtin->IsRounded)
1788bdd1243dSDimitry Andric buildOpDecorate(Call->ReturnRegister, MIRBuilder,
178906c3fb27SDimitry Andric SPIRV::Decoration::FPRoundingMode,
179006c3fb27SDimitry Andric {(unsigned)Builtin->RoundingMode});
1791bdd1243dSDimitry Andric
1792bdd1243dSDimitry Andric unsigned Opcode = SPIRV::OpNop;
1793bdd1243dSDimitry Andric if (GR->isScalarOrVectorOfType(Call->Arguments[0], SPIRV::OpTypeInt)) {
1794bdd1243dSDimitry Andric // Int -> ...
1795bdd1243dSDimitry Andric if (GR->isScalarOrVectorOfType(Call->ReturnRegister, SPIRV::OpTypeInt)) {
1796bdd1243dSDimitry Andric // Int -> Int
1797bdd1243dSDimitry Andric if (Builtin->IsSaturated)
1798bdd1243dSDimitry Andric Opcode = Builtin->IsDestinationSigned ? SPIRV::OpSatConvertUToS
1799bdd1243dSDimitry Andric : SPIRV::OpSatConvertSToU;
1800bdd1243dSDimitry Andric else
1801bdd1243dSDimitry Andric Opcode = Builtin->IsDestinationSigned ? SPIRV::OpUConvert
1802bdd1243dSDimitry Andric : SPIRV::OpSConvert;
1803bdd1243dSDimitry Andric } else if (GR->isScalarOrVectorOfType(Call->ReturnRegister,
1804bdd1243dSDimitry Andric SPIRV::OpTypeFloat)) {
1805bdd1243dSDimitry Andric // Int -> Float
1806bdd1243dSDimitry Andric bool IsSourceSigned =
1807bdd1243dSDimitry Andric DemangledCall[DemangledCall.find_first_of('(') + 1] != 'u';
1808bdd1243dSDimitry Andric Opcode = IsSourceSigned ? SPIRV::OpConvertSToF : SPIRV::OpConvertUToF;
1809bdd1243dSDimitry Andric }
1810bdd1243dSDimitry Andric } else if (GR->isScalarOrVectorOfType(Call->Arguments[0],
1811bdd1243dSDimitry Andric SPIRV::OpTypeFloat)) {
1812bdd1243dSDimitry Andric // Float -> ...
1813bdd1243dSDimitry Andric if (GR->isScalarOrVectorOfType(Call->ReturnRegister, SPIRV::OpTypeInt))
1814bdd1243dSDimitry Andric // Float -> Int
1815bdd1243dSDimitry Andric Opcode = Builtin->IsDestinationSigned ? SPIRV::OpConvertFToS
1816bdd1243dSDimitry Andric : SPIRV::OpConvertFToU;
1817bdd1243dSDimitry Andric else if (GR->isScalarOrVectorOfType(Call->ReturnRegister,
1818bdd1243dSDimitry Andric SPIRV::OpTypeFloat))
1819bdd1243dSDimitry Andric // Float -> Float
1820bdd1243dSDimitry Andric Opcode = SPIRV::OpFConvert;
1821bdd1243dSDimitry Andric }
1822bdd1243dSDimitry Andric
1823bdd1243dSDimitry Andric assert(Opcode != SPIRV::OpNop &&
1824bdd1243dSDimitry Andric "Conversion between the types not implemented!");
1825bdd1243dSDimitry Andric
1826bdd1243dSDimitry Andric MIRBuilder.buildInstr(Opcode)
1827bdd1243dSDimitry Andric .addDef(Call->ReturnRegister)
1828bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1829bdd1243dSDimitry Andric .addUse(Call->Arguments[0]);
1830bdd1243dSDimitry Andric return true;
1831bdd1243dSDimitry Andric }
1832bdd1243dSDimitry Andric
generateVectorLoadStoreInst(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)1833bdd1243dSDimitry Andric static bool generateVectorLoadStoreInst(const SPIRV::IncomingCall *Call,
1834bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
1835bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
1836bdd1243dSDimitry Andric // Lookup the vector load/store builtin in the TableGen records.
1837bdd1243dSDimitry Andric const SPIRV::VectorLoadStoreBuiltin *Builtin =
1838bdd1243dSDimitry Andric SPIRV::lookupVectorLoadStoreBuiltin(Call->Builtin->Name,
1839bdd1243dSDimitry Andric Call->Builtin->Set);
1840bdd1243dSDimitry Andric // Build extended instruction.
1841bdd1243dSDimitry Andric auto MIB =
1842bdd1243dSDimitry Andric MIRBuilder.buildInstr(SPIRV::OpExtInst)
1843bdd1243dSDimitry Andric .addDef(Call->ReturnRegister)
1844bdd1243dSDimitry Andric .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1845bdd1243dSDimitry Andric .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::OpenCL_std))
1846bdd1243dSDimitry Andric .addImm(Builtin->Number);
1847bdd1243dSDimitry Andric for (auto Argument : Call->Arguments)
1848bdd1243dSDimitry Andric MIB.addUse(Argument);
1849bdd1243dSDimitry Andric
1850bdd1243dSDimitry Andric // Rounding mode should be passed as a last argument in the MI for builtins
1851bdd1243dSDimitry Andric // like "vstorea_halfn_r".
1852bdd1243dSDimitry Andric if (Builtin->IsRounded)
1853bdd1243dSDimitry Andric MIB.addImm(static_cast<uint32_t>(Builtin->RoundingMode));
1854bdd1243dSDimitry Andric return true;
1855bdd1243dSDimitry Andric }
1856bdd1243dSDimitry Andric
generateLoadStoreInst(const SPIRV::IncomingCall * Call,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)1857bdd1243dSDimitry Andric static bool generateLoadStoreInst(const SPIRV::IncomingCall *Call,
1858bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
1859bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
1860bdd1243dSDimitry Andric // Lookup the instruction opcode in the TableGen records.
1861bdd1243dSDimitry Andric const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1862bdd1243dSDimitry Andric unsigned Opcode =
1863bdd1243dSDimitry Andric SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1864bdd1243dSDimitry Andric bool IsLoad = Opcode == SPIRV::OpLoad;
1865bdd1243dSDimitry Andric // Build the instruction.
1866bdd1243dSDimitry Andric auto MIB = MIRBuilder.buildInstr(Opcode);
1867bdd1243dSDimitry Andric if (IsLoad) {
1868bdd1243dSDimitry Andric MIB.addDef(Call->ReturnRegister);
1869bdd1243dSDimitry Andric MIB.addUse(GR->getSPIRVTypeID(Call->ReturnType));
1870bdd1243dSDimitry Andric }
1871bdd1243dSDimitry Andric // Add a pointer to the value to load/store.
1872bdd1243dSDimitry Andric MIB.addUse(Call->Arguments[0]);
1873bdd1243dSDimitry Andric MachineRegisterInfo *MRI = MIRBuilder.getMRI();
187406c3fb27SDimitry Andric MRI->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
187506c3fb27SDimitry Andric // Add a value to store.
187606c3fb27SDimitry Andric if (!IsLoad) {
187706c3fb27SDimitry Andric MIB.addUse(Call->Arguments[1]);
187806c3fb27SDimitry Andric MRI->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
187906c3fb27SDimitry Andric }
188006c3fb27SDimitry Andric // Add optional memory attributes and an alignment.
1881bdd1243dSDimitry Andric unsigned NumArgs = Call->Arguments.size();
188206c3fb27SDimitry Andric if ((IsLoad && NumArgs >= 2) || NumArgs >= 3) {
1883bdd1243dSDimitry Andric MIB.addImm(getConstFromIntrinsic(Call->Arguments[IsLoad ? 1 : 2], MRI));
188406c3fb27SDimitry Andric MRI->setRegClass(Call->Arguments[IsLoad ? 1 : 2], &SPIRV::IDRegClass);
188506c3fb27SDimitry Andric }
188606c3fb27SDimitry Andric if ((IsLoad && NumArgs >= 3) || NumArgs >= 4) {
1887bdd1243dSDimitry Andric MIB.addImm(getConstFromIntrinsic(Call->Arguments[IsLoad ? 2 : 3], MRI));
188806c3fb27SDimitry Andric MRI->setRegClass(Call->Arguments[IsLoad ? 2 : 3], &SPIRV::IDRegClass);
188906c3fb27SDimitry Andric }
1890bdd1243dSDimitry Andric return true;
1891bdd1243dSDimitry Andric }
1892bdd1243dSDimitry Andric
1893bdd1243dSDimitry Andric /// Lowers a builtin funtion call using the provided \p DemangledCall skeleton
1894bdd1243dSDimitry Andric /// and external instruction \p Set.
1895bdd1243dSDimitry Andric namespace SPIRV {
lowerBuiltin(const StringRef DemangledCall,SPIRV::InstructionSet::InstructionSet Set,MachineIRBuilder & MIRBuilder,const Register OrigRet,const Type * OrigRetTy,const SmallVectorImpl<Register> & Args,SPIRVGlobalRegistry * GR)1896bdd1243dSDimitry Andric std::optional<bool> lowerBuiltin(const StringRef DemangledCall,
1897bdd1243dSDimitry Andric SPIRV::InstructionSet::InstructionSet Set,
1898bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
1899bdd1243dSDimitry Andric const Register OrigRet, const Type *OrigRetTy,
1900bdd1243dSDimitry Andric const SmallVectorImpl<Register> &Args,
1901bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
1902bdd1243dSDimitry Andric LLVM_DEBUG(dbgs() << "Lowering builtin call: " << DemangledCall << "\n");
1903bdd1243dSDimitry Andric
1904bdd1243dSDimitry Andric // SPIR-V type and return register.
1905bdd1243dSDimitry Andric Register ReturnRegister = OrigRet;
1906bdd1243dSDimitry Andric SPIRVType *ReturnType = nullptr;
1907bdd1243dSDimitry Andric if (OrigRetTy && !OrigRetTy->isVoidTy()) {
1908bdd1243dSDimitry Andric ReturnType = GR->assignTypeToVReg(OrigRetTy, OrigRet, MIRBuilder);
190906c3fb27SDimitry Andric if (!MIRBuilder.getMRI()->getRegClassOrNull(ReturnRegister))
191006c3fb27SDimitry Andric MIRBuilder.getMRI()->setRegClass(ReturnRegister, &SPIRV::IDRegClass);
1911bdd1243dSDimitry Andric } else if (OrigRetTy && OrigRetTy->isVoidTy()) {
1912bdd1243dSDimitry Andric ReturnRegister = MIRBuilder.getMRI()->createVirtualRegister(&IDRegClass);
1913bdd1243dSDimitry Andric MIRBuilder.getMRI()->setType(ReturnRegister, LLT::scalar(32));
1914bdd1243dSDimitry Andric ReturnType = GR->assignTypeToVReg(OrigRetTy, ReturnRegister, MIRBuilder);
1915bdd1243dSDimitry Andric }
1916bdd1243dSDimitry Andric
1917bdd1243dSDimitry Andric // Lookup the builtin in the TableGen records.
1918bdd1243dSDimitry Andric std::unique_ptr<const IncomingCall> Call =
1919bdd1243dSDimitry Andric lookupBuiltin(DemangledCall, Set, ReturnRegister, ReturnType, Args);
1920bdd1243dSDimitry Andric
1921bdd1243dSDimitry Andric if (!Call) {
1922bdd1243dSDimitry Andric LLVM_DEBUG(dbgs() << "Builtin record was not found!\n");
1923bdd1243dSDimitry Andric return std::nullopt;
1924bdd1243dSDimitry Andric }
1925bdd1243dSDimitry Andric
1926bdd1243dSDimitry Andric // TODO: check if the provided args meet the builtin requirments.
1927bdd1243dSDimitry Andric assert(Args.size() >= Call->Builtin->MinNumArgs &&
1928bdd1243dSDimitry Andric "Too few arguments to generate the builtin");
1929bdd1243dSDimitry Andric if (Call->Builtin->MaxNumArgs && Args.size() > Call->Builtin->MaxNumArgs)
1930bdd1243dSDimitry Andric LLVM_DEBUG(dbgs() << "More arguments provided than required!\n");
1931bdd1243dSDimitry Andric
1932bdd1243dSDimitry Andric // Match the builtin with implementation based on the grouping.
1933bdd1243dSDimitry Andric switch (Call->Builtin->Group) {
1934bdd1243dSDimitry Andric case SPIRV::Extended:
1935bdd1243dSDimitry Andric return generateExtInst(Call.get(), MIRBuilder, GR);
1936bdd1243dSDimitry Andric case SPIRV::Relational:
1937bdd1243dSDimitry Andric return generateRelationalInst(Call.get(), MIRBuilder, GR);
1938bdd1243dSDimitry Andric case SPIRV::Group:
1939bdd1243dSDimitry Andric return generateGroupInst(Call.get(), MIRBuilder, GR);
1940bdd1243dSDimitry Andric case SPIRV::Variable:
1941bdd1243dSDimitry Andric return generateBuiltinVar(Call.get(), MIRBuilder, GR);
1942bdd1243dSDimitry Andric case SPIRV::Atomic:
1943bdd1243dSDimitry Andric return generateAtomicInst(Call.get(), MIRBuilder, GR);
1944bdd1243dSDimitry Andric case SPIRV::Barrier:
1945bdd1243dSDimitry Andric return generateBarrierInst(Call.get(), MIRBuilder, GR);
1946bdd1243dSDimitry Andric case SPIRV::Dot:
1947bdd1243dSDimitry Andric return generateDotOrFMulInst(Call.get(), MIRBuilder, GR);
1948bdd1243dSDimitry Andric case SPIRV::GetQuery:
1949bdd1243dSDimitry Andric return generateGetQueryInst(Call.get(), MIRBuilder, GR);
1950bdd1243dSDimitry Andric case SPIRV::ImageSizeQuery:
1951bdd1243dSDimitry Andric return generateImageSizeQueryInst(Call.get(), MIRBuilder, GR);
1952bdd1243dSDimitry Andric case SPIRV::ImageMiscQuery:
1953bdd1243dSDimitry Andric return generateImageMiscQueryInst(Call.get(), MIRBuilder, GR);
1954bdd1243dSDimitry Andric case SPIRV::ReadImage:
1955bdd1243dSDimitry Andric return generateReadImageInst(DemangledCall, Call.get(), MIRBuilder, GR);
1956bdd1243dSDimitry Andric case SPIRV::WriteImage:
1957bdd1243dSDimitry Andric return generateWriteImageInst(Call.get(), MIRBuilder, GR);
1958bdd1243dSDimitry Andric case SPIRV::SampleImage:
1959bdd1243dSDimitry Andric return generateSampleImageInst(DemangledCall, Call.get(), MIRBuilder, GR);
1960bdd1243dSDimitry Andric case SPIRV::Select:
1961bdd1243dSDimitry Andric return generateSelectInst(Call.get(), MIRBuilder);
1962bdd1243dSDimitry Andric case SPIRV::SpecConstant:
1963bdd1243dSDimitry Andric return generateSpecConstantInst(Call.get(), MIRBuilder, GR);
1964bdd1243dSDimitry Andric case SPIRV::Enqueue:
1965bdd1243dSDimitry Andric return generateEnqueueInst(Call.get(), MIRBuilder, GR);
1966bdd1243dSDimitry Andric case SPIRV::AsyncCopy:
1967bdd1243dSDimitry Andric return generateAsyncCopy(Call.get(), MIRBuilder, GR);
1968bdd1243dSDimitry Andric case SPIRV::Convert:
1969bdd1243dSDimitry Andric return generateConvertInst(DemangledCall, Call.get(), MIRBuilder, GR);
1970bdd1243dSDimitry Andric case SPIRV::VectorLoadStore:
1971bdd1243dSDimitry Andric return generateVectorLoadStoreInst(Call.get(), MIRBuilder, GR);
1972bdd1243dSDimitry Andric case SPIRV::LoadStore:
1973bdd1243dSDimitry Andric return generateLoadStoreInst(Call.get(), MIRBuilder, GR);
1974bdd1243dSDimitry Andric }
1975bdd1243dSDimitry Andric return false;
1976bdd1243dSDimitry Andric }
1977bdd1243dSDimitry Andric
197806c3fb27SDimitry Andric struct BuiltinType {
1979bdd1243dSDimitry Andric StringRef Name;
1980bdd1243dSDimitry Andric uint32_t Opcode;
1981bdd1243dSDimitry Andric };
1982bdd1243dSDimitry Andric
198306c3fb27SDimitry Andric #define GET_BuiltinTypes_DECL
198406c3fb27SDimitry Andric #define GET_BuiltinTypes_IMPL
1985bdd1243dSDimitry Andric
198606c3fb27SDimitry Andric struct OpenCLType {
1987bdd1243dSDimitry Andric StringRef Name;
198806c3fb27SDimitry Andric StringRef SpirvTypeLiteral;
1989bdd1243dSDimitry Andric };
1990bdd1243dSDimitry Andric
199106c3fb27SDimitry Andric #define GET_OpenCLTypes_DECL
199206c3fb27SDimitry Andric #define GET_OpenCLTypes_IMPL
1993bdd1243dSDimitry Andric
1994bdd1243dSDimitry Andric #include "SPIRVGenTables.inc"
1995bdd1243dSDimitry Andric } // namespace SPIRV
1996bdd1243dSDimitry Andric
1997bdd1243dSDimitry Andric //===----------------------------------------------------------------------===//
199806c3fb27SDimitry Andric // Misc functions for parsing builtin types.
1999bdd1243dSDimitry Andric //===----------------------------------------------------------------------===//
2000bdd1243dSDimitry Andric
parseTypeString(const StringRef Name,LLVMContext & Context)200106c3fb27SDimitry Andric static Type *parseTypeString(const StringRef Name, LLVMContext &Context) {
20025f757f3fSDimitry Andric if (Name.starts_with("void"))
200306c3fb27SDimitry Andric return Type::getVoidTy(Context);
20045f757f3fSDimitry Andric else if (Name.starts_with("int") || Name.starts_with("uint"))
200506c3fb27SDimitry Andric return Type::getInt32Ty(Context);
20065f757f3fSDimitry Andric else if (Name.starts_with("float"))
200706c3fb27SDimitry Andric return Type::getFloatTy(Context);
20085f757f3fSDimitry Andric else if (Name.starts_with("half"))
200906c3fb27SDimitry Andric return Type::getHalfTy(Context);
201006c3fb27SDimitry Andric llvm_unreachable("Unable to recognize type!");
2011bdd1243dSDimitry Andric }
2012bdd1243dSDimitry Andric
2013bdd1243dSDimitry Andric //===----------------------------------------------------------------------===//
2014bdd1243dSDimitry Andric // Implementation functions for builtin types.
2015bdd1243dSDimitry Andric //===----------------------------------------------------------------------===//
2016bdd1243dSDimitry Andric
getNonParameterizedType(const TargetExtType * ExtensionType,const SPIRV::BuiltinType * TypeRecord,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)201706c3fb27SDimitry Andric static SPIRVType *getNonParameterizedType(const TargetExtType *ExtensionType,
201806c3fb27SDimitry Andric const SPIRV::BuiltinType *TypeRecord,
2019bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
2020bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
2021bdd1243dSDimitry Andric unsigned Opcode = TypeRecord->Opcode;
2022bdd1243dSDimitry Andric // Create or get an existing type from GlobalRegistry.
202306c3fb27SDimitry Andric return GR->getOrCreateOpTypeByOpcode(ExtensionType, MIRBuilder, Opcode);
2024bdd1243dSDimitry Andric }
2025bdd1243dSDimitry Andric
getSamplerType(MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)2026bdd1243dSDimitry Andric static SPIRVType *getSamplerType(MachineIRBuilder &MIRBuilder,
2027bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
2028bdd1243dSDimitry Andric // Create or get an existing type from GlobalRegistry.
2029bdd1243dSDimitry Andric return GR->getOrCreateOpTypeSampler(MIRBuilder);
2030bdd1243dSDimitry Andric }
2031bdd1243dSDimitry Andric
getPipeType(const TargetExtType * ExtensionType,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)203206c3fb27SDimitry Andric static SPIRVType *getPipeType(const TargetExtType *ExtensionType,
2033bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
2034bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
203506c3fb27SDimitry Andric assert(ExtensionType->getNumIntParameters() == 1 &&
203606c3fb27SDimitry Andric "Invalid number of parameters for SPIR-V pipe builtin!");
2037bdd1243dSDimitry Andric // Create or get an existing type from GlobalRegistry.
203806c3fb27SDimitry Andric return GR->getOrCreateOpTypePipe(MIRBuilder,
203906c3fb27SDimitry Andric SPIRV::AccessQualifier::AccessQualifier(
204006c3fb27SDimitry Andric ExtensionType->getIntParameter(0)));
2041bdd1243dSDimitry Andric }
2042bdd1243dSDimitry Andric
2043bdd1243dSDimitry Andric static SPIRVType *
getImageType(const TargetExtType * ExtensionType,const SPIRV::AccessQualifier::AccessQualifier Qualifier,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)204406c3fb27SDimitry Andric getImageType(const TargetExtType *ExtensionType,
204506c3fb27SDimitry Andric const SPIRV::AccessQualifier::AccessQualifier Qualifier,
2046bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR) {
204706c3fb27SDimitry Andric assert(ExtensionType->getNumTypeParameters() == 1 &&
204806c3fb27SDimitry Andric "SPIR-V image builtin type must have sampled type parameter!");
204906c3fb27SDimitry Andric const SPIRVType *SampledType =
205006c3fb27SDimitry Andric GR->getOrCreateSPIRVType(ExtensionType->getTypeParameter(0), MIRBuilder);
205106c3fb27SDimitry Andric assert(ExtensionType->getNumIntParameters() == 7 &&
205206c3fb27SDimitry Andric "Invalid number of parameters for SPIR-V image builtin!");
205306c3fb27SDimitry Andric // Create or get an existing type from GlobalRegistry.
2054bdd1243dSDimitry Andric return GR->getOrCreateOpTypeImage(
205506c3fb27SDimitry Andric MIRBuilder, SampledType,
205606c3fb27SDimitry Andric SPIRV::Dim::Dim(ExtensionType->getIntParameter(0)),
205706c3fb27SDimitry Andric ExtensionType->getIntParameter(1), ExtensionType->getIntParameter(2),
205806c3fb27SDimitry Andric ExtensionType->getIntParameter(3), ExtensionType->getIntParameter(4),
205906c3fb27SDimitry Andric SPIRV::ImageFormat::ImageFormat(ExtensionType->getIntParameter(5)),
206006c3fb27SDimitry Andric Qualifier == SPIRV::AccessQualifier::WriteOnly
2061bdd1243dSDimitry Andric ? SPIRV::AccessQualifier::WriteOnly
206206c3fb27SDimitry Andric : SPIRV::AccessQualifier::AccessQualifier(
206306c3fb27SDimitry Andric ExtensionType->getIntParameter(6)));
2064bdd1243dSDimitry Andric }
2065bdd1243dSDimitry Andric
getSampledImageType(const TargetExtType * OpaqueType,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)206606c3fb27SDimitry Andric static SPIRVType *getSampledImageType(const TargetExtType *OpaqueType,
2067bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
2068bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
206906c3fb27SDimitry Andric SPIRVType *OpaqueImageType = getImageType(
207006c3fb27SDimitry Andric OpaqueType, SPIRV::AccessQualifier::ReadOnly, MIRBuilder, GR);
207106c3fb27SDimitry Andric // Create or get an existing type from GlobalRegistry.
207206c3fb27SDimitry Andric return GR->getOrCreateOpTypeSampledImage(OpaqueImageType, MIRBuilder);
2073bdd1243dSDimitry Andric }
2074bdd1243dSDimitry Andric
2075bdd1243dSDimitry Andric namespace SPIRV {
20765f757f3fSDimitry Andric const TargetExtType *
parseBuiltinTypeNameToTargetExtType(std::string TypeName,MachineIRBuilder & MIRBuilder)20775f757f3fSDimitry Andric parseBuiltinTypeNameToTargetExtType(std::string TypeName,
20785f757f3fSDimitry Andric MachineIRBuilder &MIRBuilder) {
20795f757f3fSDimitry Andric StringRef NameWithParameters = TypeName;
20805f757f3fSDimitry Andric
20815f757f3fSDimitry Andric // Pointers-to-opaque-structs representing OpenCL types are first translated
20825f757f3fSDimitry Andric // to equivalent SPIR-V types. OpenCL builtin type names should have the
20835f757f3fSDimitry Andric // following format: e.g. %opencl.event_t
20845f757f3fSDimitry Andric if (NameWithParameters.starts_with("opencl.")) {
20855f757f3fSDimitry Andric const SPIRV::OpenCLType *OCLTypeRecord =
20865f757f3fSDimitry Andric SPIRV::lookupOpenCLType(NameWithParameters);
20875f757f3fSDimitry Andric if (!OCLTypeRecord)
20885f757f3fSDimitry Andric report_fatal_error("Missing TableGen record for OpenCL type: " +
20895f757f3fSDimitry Andric NameWithParameters);
20905f757f3fSDimitry Andric NameWithParameters = OCLTypeRecord->SpirvTypeLiteral;
20915f757f3fSDimitry Andric // Continue with the SPIR-V builtin type...
20925f757f3fSDimitry Andric }
20935f757f3fSDimitry Andric
20945f757f3fSDimitry Andric // Names of the opaque structs representing a SPIR-V builtins without
20955f757f3fSDimitry Andric // parameters should have the following format: e.g. %spirv.Event
20965f757f3fSDimitry Andric assert(NameWithParameters.starts_with("spirv.") &&
20975f757f3fSDimitry Andric "Unknown builtin opaque type!");
20985f757f3fSDimitry Andric
20995f757f3fSDimitry Andric // Parameterized SPIR-V builtins names follow this format:
21005f757f3fSDimitry Andric // e.g. %spirv.Image._void_1_0_0_0_0_0_0, %spirv.Pipe._0
2101*cb14a3feSDimitry Andric if (!NameWithParameters.contains('_'))
21025f757f3fSDimitry Andric return TargetExtType::get(MIRBuilder.getContext(), NameWithParameters);
21035f757f3fSDimitry Andric
21045f757f3fSDimitry Andric SmallVector<StringRef> Parameters;
21055f757f3fSDimitry Andric unsigned BaseNameLength = NameWithParameters.find('_') - 1;
21065f757f3fSDimitry Andric SplitString(NameWithParameters.substr(BaseNameLength + 1), Parameters, "_");
21075f757f3fSDimitry Andric
21085f757f3fSDimitry Andric SmallVector<Type *, 1> TypeParameters;
21095f757f3fSDimitry Andric bool HasTypeParameter = !isDigit(Parameters[0][0]);
21105f757f3fSDimitry Andric if (HasTypeParameter)
21115f757f3fSDimitry Andric TypeParameters.push_back(parseTypeString(
21125f757f3fSDimitry Andric Parameters[0], MIRBuilder.getMF().getFunction().getContext()));
21135f757f3fSDimitry Andric SmallVector<unsigned> IntParameters;
21145f757f3fSDimitry Andric for (unsigned i = HasTypeParameter ? 1 : 0; i < Parameters.size(); i++) {
21155f757f3fSDimitry Andric unsigned IntParameter = 0;
21165f757f3fSDimitry Andric bool ValidLiteral = !Parameters[i].getAsInteger(10, IntParameter);
21175f757f3fSDimitry Andric assert(ValidLiteral &&
21185f757f3fSDimitry Andric "Invalid format of SPIR-V builtin parameter literal!");
21195f757f3fSDimitry Andric IntParameters.push_back(IntParameter);
21205f757f3fSDimitry Andric }
21215f757f3fSDimitry Andric return TargetExtType::get(MIRBuilder.getContext(),
21225f757f3fSDimitry Andric NameWithParameters.substr(0, BaseNameLength),
21235f757f3fSDimitry Andric TypeParameters, IntParameters);
21245f757f3fSDimitry Andric }
21255f757f3fSDimitry Andric
lowerBuiltinType(const Type * OpaqueType,SPIRV::AccessQualifier::AccessQualifier AccessQual,MachineIRBuilder & MIRBuilder,SPIRVGlobalRegistry * GR)212606c3fb27SDimitry Andric SPIRVType *lowerBuiltinType(const Type *OpaqueType,
2127bdd1243dSDimitry Andric SPIRV::AccessQualifier::AccessQualifier AccessQual,
2128bdd1243dSDimitry Andric MachineIRBuilder &MIRBuilder,
2129bdd1243dSDimitry Andric SPIRVGlobalRegistry *GR) {
213006c3fb27SDimitry Andric // In LLVM IR, SPIR-V and OpenCL builtin types are represented as either
213106c3fb27SDimitry Andric // target(...) target extension types or pointers-to-opaque-structs. The
213206c3fb27SDimitry Andric // approach relying on structs is deprecated and works only in the non-opaque
213306c3fb27SDimitry Andric // pointer mode (-opaque-pointers=0).
213406c3fb27SDimitry Andric // In order to maintain compatibility with LLVM IR generated by older versions
213506c3fb27SDimitry Andric // of Clang and LLVM/SPIR-V Translator, the pointers-to-opaque-structs are
213606c3fb27SDimitry Andric // "translated" to target extension types. This translation is temporary and
213706c3fb27SDimitry Andric // will be removed in the future release of LLVM.
213806c3fb27SDimitry Andric const TargetExtType *BuiltinType = dyn_cast<TargetExtType>(OpaqueType);
213906c3fb27SDimitry Andric if (!BuiltinType)
21405f757f3fSDimitry Andric BuiltinType = parseBuiltinTypeNameToTargetExtType(
21415f757f3fSDimitry Andric OpaqueType->getStructName().str(), MIRBuilder);
214206c3fb27SDimitry Andric
2143bdd1243dSDimitry Andric unsigned NumStartingVRegs = MIRBuilder.getMRI()->getNumVirtRegs();
2144bdd1243dSDimitry Andric
214506c3fb27SDimitry Andric const StringRef Name = BuiltinType->getName();
2146bdd1243dSDimitry Andric LLVM_DEBUG(dbgs() << "Lowering builtin type: " << Name << "\n");
2147bdd1243dSDimitry Andric
2148bdd1243dSDimitry Andric // Lookup the demangled builtin type in the TableGen records.
214906c3fb27SDimitry Andric const SPIRV::BuiltinType *TypeRecord = SPIRV::lookupBuiltinType(Name);
2150bdd1243dSDimitry Andric if (!TypeRecord)
2151bdd1243dSDimitry Andric report_fatal_error("Missing TableGen record for builtin type: " + Name);
2152bdd1243dSDimitry Andric
2153bdd1243dSDimitry Andric // "Lower" the BuiltinType into TargetType. The following get<...>Type methods
215406c3fb27SDimitry Andric // use the implementation details from TableGen records or TargetExtType
215506c3fb27SDimitry Andric // parameters to either create a new OpType<...> machine instruction or get an
215606c3fb27SDimitry Andric // existing equivalent SPIRVType from GlobalRegistry.
2157bdd1243dSDimitry Andric SPIRVType *TargetType;
2158bdd1243dSDimitry Andric switch (TypeRecord->Opcode) {
2159bdd1243dSDimitry Andric case SPIRV::OpTypeImage:
216006c3fb27SDimitry Andric TargetType = getImageType(BuiltinType, AccessQual, MIRBuilder, GR);
2161bdd1243dSDimitry Andric break;
2162bdd1243dSDimitry Andric case SPIRV::OpTypePipe:
216306c3fb27SDimitry Andric TargetType = getPipeType(BuiltinType, MIRBuilder, GR);
2164bdd1243dSDimitry Andric break;
2165bdd1243dSDimitry Andric case SPIRV::OpTypeDeviceEvent:
2166bdd1243dSDimitry Andric TargetType = GR->getOrCreateOpTypeDeviceEvent(MIRBuilder);
2167bdd1243dSDimitry Andric break;
2168bdd1243dSDimitry Andric case SPIRV::OpTypeSampler:
2169bdd1243dSDimitry Andric TargetType = getSamplerType(MIRBuilder, GR);
2170bdd1243dSDimitry Andric break;
2171bdd1243dSDimitry Andric case SPIRV::OpTypeSampledImage:
217206c3fb27SDimitry Andric TargetType = getSampledImageType(BuiltinType, MIRBuilder, GR);
2173bdd1243dSDimitry Andric break;
2174bdd1243dSDimitry Andric default:
217506c3fb27SDimitry Andric TargetType =
217606c3fb27SDimitry Andric getNonParameterizedType(BuiltinType, TypeRecord, MIRBuilder, GR);
2177bdd1243dSDimitry Andric break;
2178bdd1243dSDimitry Andric }
2179bdd1243dSDimitry Andric
2180bdd1243dSDimitry Andric // Emit OpName instruction if a new OpType<...> instruction was added
2181bdd1243dSDimitry Andric // (equivalent type was not found in GlobalRegistry).
2182bdd1243dSDimitry Andric if (NumStartingVRegs < MIRBuilder.getMRI()->getNumVirtRegs())
218306c3fb27SDimitry Andric buildOpName(GR->getSPIRVTypeID(TargetType), Name, MIRBuilder);
2184bdd1243dSDimitry Andric
2185bdd1243dSDimitry Andric return TargetType;
2186bdd1243dSDimitry Andric }
2187bdd1243dSDimitry Andric } // namespace SPIRV
2188bdd1243dSDimitry Andric } // namespace llvm
2189