Movatterモバイル変換

Go to the documentation of this file.

1//===- SPIRVBuiltins.cpp - SPIR-V Built-in Functions ------------*- C++ -*-===//

2//

3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.

4// See https://llvm.org/LICENSE.txt for license information.

5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception

6//

7//===----------------------------------------------------------------------===//

8//

9// This file implements lowering builtin function calls and types using their

10// demangled names and TableGen records.

11//

12//===----------------------------------------------------------------------===//

14#include "SPIRVBuiltins.h"

15#include "SPIRV.h"

16#include "SPIRVSubtarget.h"

17#include "SPIRVUtils.h"

18#include "llvm/ADT/StringExtras.h"

19#include "llvm/Analysis/ValueTracking.h"

20#include "llvm/IR/IntrinsicsSPIRV.h"

21#include <regex>

22#include <string>

23#include <tuple>

25#define DEBUG_TYPE "spirv-builtins"

27namespacellvm {

28namespaceSPIRV {

29#define GET_BuiltinGroup_DECL

30#include "SPIRVGenTables.inc"

32structDemangledBuiltin {

33StringRef Name;

34 InstructionSet::InstructionSetSet;

35 BuiltinGroupGroup;

36uint8_t MinNumArgs;

37uint8_t MaxNumArgs;

38};

40#define GET_DemangledBuiltins_DECL

41#define GET_DemangledBuiltins_IMPL

43structIncomingCall {

44const std::stringBuiltinName;

45constDemangledBuiltin *Builtin;

47constRegister ReturnRegister;

48constSPIRVType *ReturnType;

49constSmallVectorImpl<Register> &Arguments;

51IncomingCall(const std::stringBuiltinName,constDemangledBuiltin *Builtin,

52constRegister ReturnRegister,constSPIRVType *ReturnType,

53constSmallVectorImpl<Register> &Arguments)

54 :BuiltinName(BuiltinName),Builtin(Builtin),

55ReturnRegister(ReturnRegister),ReturnType(ReturnType),

56Arguments(Arguments) {}

58boolisSpirvOp() const{returnBuiltinName.rfind("__spirv_", 0) == 0; }

59};

61structNativeBuiltin {

62StringRef Name;

63 InstructionSet::InstructionSetSet;

64uint32_t Opcode;

65};

67#define GET_NativeBuiltins_DECL

68#define GET_NativeBuiltins_IMPL

70structGroupBuiltin {

71StringRef Name;

72uint32_t Opcode;

73uint32_t GroupOperation;

74boolIsElect;

75boolIsAllOrAny;

76boolIsAllEqual;

77boolIsBallot;

78boolIsInverseBallot;

79boolIsBallotBitExtract;

80boolIsBallotFindBit;

81boolIsLogical;

82boolNoGroupOperation;

83boolHasBoolArg;

84};

86#define GET_GroupBuiltins_DECL

87#define GET_GroupBuiltins_IMPL

89structIntelSubgroupsBuiltin {

90StringRef Name;

91uint32_t Opcode;

92boolIsBlock;

93boolIsWrite;

94boolIsMedia;

95};

97#define GET_IntelSubgroupsBuiltins_DECL

98#define GET_IntelSubgroupsBuiltins_IMPL

100structAtomicFloatingBuiltin {

101StringRef Name;

102uint32_t Opcode;

103};

104

105#define GET_AtomicFloatingBuiltins_DECL

106#define GET_AtomicFloatingBuiltins_IMPL

107structGroupUniformBuiltin {

108StringRef Name;

109uint32_t Opcode;

110boolIsLogical;

111};

112

113#define GET_GroupUniformBuiltins_DECL

114#define GET_GroupUniformBuiltins_IMPL

115

116structGetBuiltin {

117StringRef Name;

118 InstructionSet::InstructionSetSet;

119 BuiltIn::BuiltInValue;

120};

121

122using namespaceBuiltIn;

123#define GET_GetBuiltins_DECL

124#define GET_GetBuiltins_IMPL

125

126structImageQueryBuiltin {

127StringRef Name;

128 InstructionSet::InstructionSetSet;

129uint32_t Component;

130};

131

132#define GET_ImageQueryBuiltins_DECL

133#define GET_ImageQueryBuiltins_IMPL

134

135structConvertBuiltin {

136StringRef Name;

137 InstructionSet::InstructionSetSet;

138boolIsDestinationSigned;

139boolIsSaturated;

140boolIsRounded;

141boolIsBfloat16;

142 FPRoundingMode::FPRoundingModeRoundingMode;

143};

144

145structVectorLoadStoreBuiltin {

146StringRef Name;

147 InstructionSet::InstructionSetSet;

148uint32_t Number;

149uint32_t ElementCount;

150boolIsRounded;

151 FPRoundingMode::FPRoundingModeRoundingMode;

152};

153

154using namespaceFPRoundingMode;

155#define GET_ConvertBuiltins_DECL

156#define GET_ConvertBuiltins_IMPL

157

158using namespaceInstructionSet;

159#define GET_VectorLoadStoreBuiltins_DECL

160#define GET_VectorLoadStoreBuiltins_IMPL

161

162#define GET_CLMemoryScope_DECL

163#define GET_CLSamplerAddressingMode_DECL

164#define GET_CLMemoryFenceFlags_DECL

165#define GET_ExtendedBuiltins_DECL

166#include "SPIRVGenTables.inc"

167}// namespace SPIRV

168

169//===----------------------------------------------------------------------===//

170// Misc functions for looking up builtins and veryfying requirements using

171// TableGen records

172//===----------------------------------------------------------------------===//

173

174namespaceSPIRV {

175/// Parses the name part of the demangled builtin call.

176std::stringlookupBuiltinNameHelper(StringRef DemangledCall,

177FPDecorationId *DecorationId) {

178conststatic std::string PassPrefix ="(anonymous namespace)::";

179 std::string BuiltinName;

180// Itanium Demangler result may have "(anonymous namespace)::" prefix

181if (DemangledCall.starts_with(PassPrefix.c_str()))

182 BuiltinName = DemangledCall.substr(PassPrefix.length());

183else

184 BuiltinName = DemangledCall;

185// Extract the builtin function name and types of arguments from the call

186// skeleton.

187 BuiltinName = BuiltinName.substr(0, BuiltinName.find('('));

188

189// Account for possible "__spirv_ocl_" prefix in SPIR-V friendly LLVM IR

190if (BuiltinName.rfind("__spirv_ocl_", 0) == 0)

191 BuiltinName = BuiltinName.substr(12);

192

193// Check if the extracted name contains type information between angle

194// brackets. If so, the builtin is an instantiated template - needs to have

195// the information after angle brackets and return type removed.

196 std::size_t Pos1 = BuiltinName.rfind('<');

197if (Pos1 != std::string::npos && BuiltinName.back() =='>') {

198 std::size_t Pos2 = BuiltinName.rfind(' ', Pos1);

199if (Pos2 == std::string::npos)

200 Pos2 = 0;

201else

202 ++Pos2;

203 BuiltinName = BuiltinName.substr(Pos2, Pos1 - Pos2);

204 BuiltinName = BuiltinName.substr(BuiltinName.find_last_of(' ') + 1);

205 }

206

207// Check if the extracted name begins with:

208// - "__spirv_ImageSampleExplicitLod"

209// - "__spirv_ImageRead"

210// - "__spirv_ImageQuerySizeLod"

211// - "__spirv_UDotKHR"

212// - "__spirv_SDotKHR"

213// - "__spirv_SUDotKHR"

214// - "__spirv_SDotAccSatKHR"

215// - "__spirv_UDotAccSatKHR"

216// - "__spirv_SUDotAccSatKHR"

217// - "__spirv_ReadClockKHR"

218// - "__spirv_SubgroupBlockReadINTEL"

219// - "__spirv_SubgroupImageBlockReadINTEL"

220// - "__spirv_SubgroupImageMediaBlockReadINTEL"

221// - "__spirv_SubgroupImageMediaBlockWriteINTEL"

222// - "__spirv_Convert"

223// - "__spirv_UConvert"

224// - "__spirv_SConvert"

225// - "__spirv_FConvert"

226// - "__spirv_SatConvert"

227// and contains return type information at the end "_R<type>".

228// If so, extract the plain builtin name without the type information.

229staticconst std::regex SpvWithR(

230"(__spirv_(ImageSampleExplicitLod|ImageRead|ImageQuerySizeLod|UDotKHR|"

232"ReadClockKHR|SubgroupBlockReadINTEL|SubgroupImageBlockReadINTEL|"

233"SubgroupImageMediaBlockReadINTEL|SubgroupImageMediaBlockWriteINTEL|"

234"Convert|"

235"UConvert|SConvert|FConvert|SatConvert).*)_R[^_]*_?(\\w+)?.*");

236 std::smatchMatch;

237if (std::regex_match(BuiltinName,Match, SpvWithR) &&Match.size() > 1) {

238 std::ssub_match SubMatch;

239if (DecorationId &&Match.size() > 3) {

240 SubMatch =Match[3];

241 *DecorationId =demangledPostfixToDecorationId(SubMatch.str());

242 }

243 SubMatch =Match[1];

244 BuiltinName = SubMatch.str();

245 }

246

247return BuiltinName;

248}

249}// namespace SPIRV

250

251/// Looks up the demangled builtin call in the SPIRVBuiltins.td records using

252/// the provided \p DemangledCall and specified \p Set.

253///

254/// The lookup follows the following algorithm, returning the first successful

255/// match:

256/// 1. Search with the plain demangled name (expecting a 1:1 match).

257/// 2. Search with the prefix before or suffix after the demangled name

258/// signyfying the type of the first argument.

259///

260/// \returns Wrapper around the demangled call and found builtin definition.

261static std::unique_ptr<const SPIRV::IncomingCall>

262lookupBuiltin(StringRef DemangledCall,

263 SPIRV::InstructionSet::InstructionSet Set,

264Register ReturnRegister,constSPIRVType *ReturnType,

265constSmallVectorImpl<Register> &Arguments) {

266 std::string BuiltinName =SPIRV::lookupBuiltinNameHelper(DemangledCall);

267

268SmallVector<StringRef, 10> BuiltinArgumentTypes;

269StringRef BuiltinArgs =

270 DemangledCall.slice(DemangledCall.find('(') + 1, DemangledCall.find(')'));

271 BuiltinArgs.split(BuiltinArgumentTypes,',', -1,false);

272

273// Look up the builtin in the defined set. Start with the plain demangled

274// name, expecting a 1:1 match in the defined builtin set.

275constSPIRV::DemangledBuiltin *Builtin;

276if ((Builtin = SPIRV::lookupBuiltin(BuiltinName, Set)))

277return std::make_unique<SPIRV::IncomingCall>(

278 BuiltinName, Builtin, ReturnRegister, ReturnType,Arguments);

279

280// If the initial look up was unsuccessful and the demangled call takes at

281// least 1 argument, add a prefix or suffix signifying the type of the first

282// argument and repeat the search.

283if (BuiltinArgumentTypes.size() >= 1) {

284char FirstArgumentType = BuiltinArgumentTypes[0][0];

285// Prefix to be added to the builtin's name for lookup.

286// For example, OpenCL "abs" taking an unsigned value has a prefix "u_".

287 std::string Prefix;

288

289switch (FirstArgumentType) {

290// Unsigned:

291case'u':

292if (Set == SPIRV::InstructionSet::OpenCL_std)

293 Prefix ="u_";

294elseif (Set == SPIRV::InstructionSet::GLSL_std_450)

295 Prefix ="u";

296break;

297// Signed:

298case'c':

299case's':

300case'i':

301case'l':

302if (Set == SPIRV::InstructionSet::OpenCL_std)

303 Prefix ="s_";

304elseif (Set == SPIRV::InstructionSet::GLSL_std_450)

305 Prefix ="s";

306break;

307// Floating-point:

308case'f':

309case'd':

310case'h':

311if (Set == SPIRV::InstructionSet::OpenCL_std ||

312 Set == SPIRV::InstructionSet::GLSL_std_450)

313 Prefix ="f";

314break;

315 }

316

317// If argument-type name prefix was added, look up the builtin again.

318if (!Prefix.empty() &&

319 (Builtin = SPIRV::lookupBuiltin(Prefix + BuiltinName, Set)))

320return std::make_unique<SPIRV::IncomingCall>(

321 BuiltinName, Builtin, ReturnRegister, ReturnType,Arguments);

322

323// If lookup with a prefix failed, find a suffix to be added to the

324// builtin's name for lookup. For example, OpenCL "group_reduce_max" taking

325// an unsigned value has a suffix "u".

326 std::string Suffix;

327

328switch (FirstArgumentType) {

329// Unsigned:

330case'u':

331 Suffix ="u";

332break;

333// Signed:

334case'c':

335case's':

336case'i':

337case'l':

338 Suffix ="s";

339break;

340// Floating-point:

341case'f':

342case'd':

343case'h':

344 Suffix ="f";

345break;

346 }

347

348// If argument-type name suffix was added, look up the builtin again.

349if (!Suffix.empty() &&

350 (Builtin = SPIRV::lookupBuiltin(BuiltinName + Suffix, Set)))

351return std::make_unique<SPIRV::IncomingCall>(

352 BuiltinName, Builtin, ReturnRegister, ReturnType,Arguments);

353 }

354

355// No builtin with such name was found in the set.

356returnnullptr;

357}

358

359staticMachineInstr *getBlockStructInstr(Register ParamReg,

360MachineRegisterInfo *MRI) {

361// We expect the following sequence of instructions:

362// %0:_(pN) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.alloca)

363// or = G_GLOBAL_VALUE @block_literal_global

364// %1:_(pN) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.bitcast), %0

365// %2:_(p4) = G_ADDRSPACE_CAST %1:_(pN)

366MachineInstr *MI =MRI->getUniqueVRegDef(ParamReg);

367assert(MI->getOpcode() == TargetOpcode::G_ADDRSPACE_CAST &&

368MI->getOperand(1).isReg());

369Register BitcastReg =MI->getOperand(1).getReg();

370MachineInstr *BitcastMI =MRI->getUniqueVRegDef(BitcastReg);

371assert(isSpvIntrinsic(*BitcastMI, Intrinsic::spv_bitcast) &&

372 BitcastMI->getOperand(2).isReg());

373Register ValueReg = BitcastMI->getOperand(2).getReg();

374MachineInstr *ValueMI =MRI->getUniqueVRegDef(ValueReg);

375return ValueMI;

376}

377

378// Return an integer constant corresponding to the given register and

379// defined in spv_track_constant.

380// TODO: maybe unify with prelegalizer pass.

381staticunsignedgetConstFromIntrinsic(Register Reg,MachineRegisterInfo *MRI) {

382MachineInstr *DefMI =MRI->getUniqueVRegDef(Reg);

383assert(isSpvIntrinsic(*DefMI, Intrinsic::spv_track_constant) &&

384DefMI->getOperand(2).isReg());

385MachineInstr *DefMI2 =MRI->getUniqueVRegDef(DefMI->getOperand(2).getReg());

386assert(DefMI2->getOpcode() == TargetOpcode::G_CONSTANT &&

387 DefMI2->getOperand(1).isCImm());

388return DefMI2->getOperand(1).getCImm()->getValue().getZExtValue();

389}

390

391// Return type of the instruction result from spv_assign_type intrinsic.

392// TODO: maybe unify with prelegalizer pass.

393staticconstType *getMachineInstrType(MachineInstr *MI) {

394MachineInstr *NextMI =MI->getNextNode();

395if (!NextMI)

396returnnullptr;

397if (isSpvIntrinsic(*NextMI, Intrinsic::spv_assign_name))

398if ((NextMI = NextMI->getNextNode()) ==nullptr)

399returnnullptr;

400Register ValueReg =MI->getOperand(0).getReg();

401if ((!isSpvIntrinsic(*NextMI, Intrinsic::spv_assign_type) &&

402 !isSpvIntrinsic(*NextMI, Intrinsic::spv_assign_ptr_type)) ||

403 NextMI->getOperand(1).getReg() != ValueReg)

404returnnullptr;

405Type *Ty =getMDOperandAsType(NextMI->getOperand(2).getMetadata(), 0);

406assert(Ty &&"Type is expected");

407return Ty;

408}

409

410staticconstType *getBlockStructType(Register ParamReg,

411MachineRegisterInfo *MRI) {

412// In principle, this information should be passed to us from Clang via

413// an elementtype attribute. However, said attribute requires that

414// the function call be an intrinsic, which is not. Instead, we rely on being

415// able to trace this to the declaration of a variable: OpenCL C specification

416// section 6.12.5 should guarantee that we can do this.

417MachineInstr *MI =getBlockStructInstr(ParamReg,MRI);

418if (MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)

419returnMI->getOperand(1).getGlobal()->getType();

420assert(isSpvIntrinsic(*MI, Intrinsic::spv_alloca) &&

421"Blocks in OpenCL C must be traceable to allocation site");

422returngetMachineInstrType(MI);

423}

424

425//===----------------------------------------------------------------------===//

426// Helper functions for building misc instructions

427//===----------------------------------------------------------------------===//

428

429/// Helper function building either a resulting scalar or vector bool register

430/// depending on the expected \p ResultType.

431///

432/// \returns Tuple of the resulting register and its type.

433static std::tuple<Register, SPIRVType *>

434buildBoolRegister(MachineIRBuilder &MIRBuilder,constSPIRVType *ResultType,

435SPIRVGlobalRegistry *GR) {

436LLT Type;

437SPIRVType *BoolType = GR->getOrCreateSPIRVBoolType(MIRBuilder);

438

439if (ResultType->getOpcode() == SPIRV::OpTypeVector) {

440unsigned VectorElements = ResultType->getOperand(2).getImm();

441 BoolType =

442 GR->getOrCreateSPIRVVectorType(BoolType, VectorElements, MIRBuilder);

443constFixedVectorType *LLVMVectorType =

444 cast<FixedVectorType>(GR->getTypeForSPIRVType(BoolType));

445Type =LLT::vector(LLVMVectorType->getElementCount(), 1);

446 }else {

447Type =LLT::scalar(1);

448 }

449

450Register ResultRegister =

451 MIRBuilder.getMRI()->createGenericVirtualRegister(Type);

452 MIRBuilder.getMRI()->setRegClass(ResultRegister, GR->getRegClass(ResultType));

453 GR->assignSPIRVTypeToVReg(BoolType, ResultRegister, MIRBuilder.getMF());

454return std::make_tuple(ResultRegister, BoolType);

455}

456

457/// Helper function for building either a vector or scalar select instruction

458/// depending on the expected \p ResultType.

459staticboolbuildSelectInst(MachineIRBuilder &MIRBuilder,

460Register ReturnRegister,Register SourceRegister,

461constSPIRVType *ReturnType,

462SPIRVGlobalRegistry *GR) {

463Register TrueConst, FalseConst;

464

465if (ReturnType->getOpcode() == SPIRV::OpTypeVector) {

466unsigned Bits = GR->getScalarOrVectorBitWidth(ReturnType);

467uint64_t AllOnes =APInt::getAllOnes(Bits).getZExtValue();

468 TrueConst = GR->getOrCreateConsIntVector(AllOnes, MIRBuilder, ReturnType);

469 FalseConst = GR->getOrCreateConsIntVector(0, MIRBuilder, ReturnType);

470 }else {

471 TrueConst = GR->buildConstantInt(1, MIRBuilder, ReturnType);

472 FalseConst = GR->buildConstantInt(0, MIRBuilder, ReturnType);

473 }

474

475return MIRBuilder.buildSelect(ReturnRegister, SourceRegister, TrueConst,

476 FalseConst);

477}

478

479/// Helper function for building a load instruction loading into the

480/// \p DestinationReg.

481staticRegister buildLoadInst(SPIRVType *BaseType,Register PtrRegister,

482MachineIRBuilder &MIRBuilder,

483SPIRVGlobalRegistry *GR,LLT LowLevelType,

484Register DestinationReg =Register(0)) {

485if (!DestinationReg.isValid())

486 DestinationReg =createVirtualRegister(BaseType, GR, MIRBuilder);

487// TODO: consider using correct address space and alignment (p0 is canonical

488// type for selection though).

489MachinePointerInfo PtrInfo =MachinePointerInfo();

490 MIRBuilder.buildLoad(DestinationReg, PtrRegister, PtrInfo,Align());

491return DestinationReg;

492}

493

494/// Helper function for building a load instruction for loading a builtin global

495/// variable of \p BuiltinValue value.

496staticRegister buildBuiltinVariableLoad(

497MachineIRBuilder &MIRBuilder,SPIRVType *VariableType,

498SPIRVGlobalRegistry *GR, SPIRV::BuiltIn::BuiltIn BuiltinValue,LLT LLType,

499Register Reg =Register(0),bool isConst =true,bool hasLinkageTy =true) {

500Register NewRegister =

501 MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::pIDRegClass);

502 MIRBuilder.getMRI()->setType(

503 NewRegister,

504LLT::pointer(storageClassToAddressSpace(SPIRV::StorageClass::Function),

505 GR->getPointerSize()));

506SPIRVType *PtrType = GR->getOrCreateSPIRVPointerType(

507 VariableType, MIRBuilder, SPIRV::StorageClass::Input);

508 GR->assignSPIRVTypeToVReg(PtrType, NewRegister, MIRBuilder.getMF());

509

510// Set up the global OpVariable with the necessary builtin decorations.

511Register Variable = GR->buildGlobalVariable(

512 NewRegister, PtrType,getLinkStringForBuiltIn(BuiltinValue),nullptr,

513 SPIRV::StorageClass::Input,nullptr,/* isConst= */ isConst,

514/* HasLinkageTy */ hasLinkageTy, SPIRV::LinkageType::Import, MIRBuilder,

515false);

516

517// Load the value from the global variable.

518Register LoadedRegister =

519buildLoadInst(VariableType, Variable, MIRBuilder, GR, LLType,Reg);

520 MIRBuilder.getMRI()->setType(LoadedRegister, LLType);

521return LoadedRegister;

522}

523

524/// Helper external function for inserting ASSIGN_TYPE instuction between \p Reg

525/// and its definition, set the new register as a destination of the definition,

526/// assign SPIRVType to both registers. If SpirvTy is provided, use it as

527/// SPIRVType in ASSIGN_TYPE, otherwise create it from \p Ty. Defined in

528/// SPIRVPreLegalizer.cpp.

529extern RegisterinsertAssignInstr(RegisterReg, Type *Ty,SPIRVType *SpirvTy,

530 SPIRVGlobalRegistry *GR,

531 MachineIRBuilder &MIB,

532 MachineRegisterInfo &MRI);

533

534// TODO: Move to TableGen.

535static SPIRV::MemorySemantics::MemorySemantics

536getSPIRVMemSemantics(std::memory_order MemOrder) {

537switch (MemOrder) {

538case std::memory_order_relaxed:

539return SPIRV::MemorySemantics::None;

540case std::memory_order_acquire:

541return SPIRV::MemorySemantics::Acquire;

542case std::memory_order_release:

543return SPIRV::MemorySemantics::Release;

544case std::memory_order_acq_rel:

545return SPIRV::MemorySemantics::AcquireRelease;

546case std::memory_order_seq_cst:

547return SPIRV::MemorySemantics::SequentiallyConsistent;

548default:

549report_fatal_error("Unknown CL memory scope");

550 }

551}

552

553static SPIRV::Scope::ScopegetSPIRVScope(SPIRV::CLMemoryScope ClScope) {

554switch (ClScope) {

555case SPIRV::CLMemoryScope::memory_scope_work_item:

556return SPIRV::Scope::Invocation;

557case SPIRV::CLMemoryScope::memory_scope_work_group:

558return SPIRV::Scope::Workgroup;

559case SPIRV::CLMemoryScope::memory_scope_device:

560return SPIRV::Scope::Device;

561case SPIRV::CLMemoryScope::memory_scope_all_svm_devices:

562return SPIRV::Scope::CrossDevice;

563case SPIRV::CLMemoryScope::memory_scope_sub_group:

564return SPIRV::Scope::Subgroup;

565 }

566report_fatal_error("Unknown CL memory scope");

567}

568

569staticRegister buildConstantIntReg32(uint64_t Val,

570MachineIRBuilder &MIRBuilder,

571SPIRVGlobalRegistry *GR) {

572return GR->buildConstantInt(Val, MIRBuilder,

573 GR->getOrCreateSPIRVIntegerType(32, MIRBuilder));

574}

575

576staticRegister buildScopeReg(Register CLScopeRegister,

577 SPIRV::Scope::Scope Scope,

578MachineIRBuilder &MIRBuilder,

579SPIRVGlobalRegistry *GR,

580MachineRegisterInfo *MRI) {

581if (CLScopeRegister.isValid()) {

582auto CLScope =

583static_cast<SPIRV::CLMemoryScope>(getIConstVal(CLScopeRegister,MRI));

584 Scope =getSPIRVScope(CLScope);

585

586if (CLScope ==static_cast<unsigned>(Scope)) {

587MRI->setRegClass(CLScopeRegister, &SPIRV::iIDRegClass);

588return CLScopeRegister;

589 }

590 }

591returnbuildConstantIntReg32(Scope, MIRBuilder, GR);

592}

593

594staticvoidsetRegClassIfNull(Register Reg,MachineRegisterInfo *MRI,

595SPIRVGlobalRegistry *GR) {

596if (MRI->getRegClassOrNull(Reg))

597return;

598SPIRVType *SpvType = GR->getSPIRVTypeForVReg(Reg);

599MRI->setRegClass(Reg,

600 SpvType ? GR->getRegClass(SpvType) : &SPIRV::iIDRegClass);

601}

602

603staticRegister buildMemSemanticsReg(Register SemanticsRegister,

604Register PtrRegister,unsigned &Semantics,

605MachineIRBuilder &MIRBuilder,

606SPIRVGlobalRegistry *GR) {

607if (SemanticsRegister.isValid()) {

608MachineRegisterInfo *MRI = MIRBuilder.getMRI();

609 std::memory_order Order =

610static_cast<std::memory_order>(getIConstVal(SemanticsRegister,MRI));

611 Semantics =

612getSPIRVMemSemantics(Order) |

613getMemSemanticsForStorageClass(GR->getPointerStorageClass(PtrRegister));

614if (static_cast<unsigned>(Order) == Semantics) {

615MRI->setRegClass(SemanticsRegister, &SPIRV::iIDRegClass);

616return SemanticsRegister;

617 }

618 }

619returnbuildConstantIntReg32(Semantics, MIRBuilder, GR);

620}

621

622staticboolbuildOpFromWrapper(MachineIRBuilder &MIRBuilder,unsigned Opcode,

623constSPIRV::IncomingCall *Call,

624Register TypeReg,

625ArrayRef<uint32_t> ImmArgs = {}) {

626auto MIB = MIRBuilder.buildInstr(Opcode);

627if (TypeReg.isValid())

628 MIB.addDef(Call->ReturnRegister).addUse(TypeReg);

629unsigned Sz = Call->Arguments.size() - ImmArgs.size();

630for (unsigned i = 0; i < Sz; ++i)

631 MIB.addUse(Call->Arguments[i]);

632for (uint32_t ImmArg : ImmArgs)

633 MIB.addImm(ImmArg);

634returntrue;

635}

636

637/// Helper function for translating atomic init to OpStore.

638staticboolbuildAtomicInitInst(constSPIRV::IncomingCall *Call,

639MachineIRBuilder &MIRBuilder) {

640if (Call->isSpirvOp())

641returnbuildOpFromWrapper(MIRBuilder, SPIRV::OpStore, Call,Register(0));

642

643assert(Call->Arguments.size() == 2 &&

644"Need 2 arguments for atomic init translation");

645 MIRBuilder.buildInstr(SPIRV::OpStore)

646 .addUse(Call->Arguments[0])

647 .addUse(Call->Arguments[1]);

648returntrue;

649}

650

651/// Helper function for building an atomic load instruction.

652staticboolbuildAtomicLoadInst(constSPIRV::IncomingCall *Call,

653MachineIRBuilder &MIRBuilder,

654SPIRVGlobalRegistry *GR) {

655Register TypeReg = GR->getSPIRVTypeID(Call->ReturnType);

656if (Call->isSpirvOp())

657returnbuildOpFromWrapper(MIRBuilder, SPIRV::OpAtomicLoad, Call, TypeReg);

658

659Register PtrRegister = Call->Arguments[0];

660// TODO: if true insert call to __translate_ocl_memory_sccope before

661// OpAtomicLoad and the function implementation. We can use Translator's

662// output for transcoding/atomic_explicit_arguments.cl as an example.

663Register ScopeRegister =

664 Call->Arguments.size() > 1

665 ? Call->Arguments[1]

666 :buildConstantIntReg32(SPIRV::Scope::Device, MIRBuilder, GR);

667Register MemSemanticsReg;

668if (Call->Arguments.size() > 2) {

669// TODO: Insert call to __translate_ocl_memory_order before OpAtomicLoad.

670 MemSemanticsReg = Call->Arguments[2];

671 }else {

672int Semantics =

673 SPIRV::MemorySemantics::SequentiallyConsistent |

674getMemSemanticsForStorageClass(GR->getPointerStorageClass(PtrRegister));

675 MemSemanticsReg =buildConstantIntReg32(Semantics, MIRBuilder, GR);

676 }

677

678 MIRBuilder.buildInstr(SPIRV::OpAtomicLoad)

679 .addDef(Call->ReturnRegister)

680 .addUse(TypeReg)

681 .addUse(PtrRegister)

682 .addUse(ScopeRegister)

683 .addUse(MemSemanticsReg);

684returntrue;

685}

686

687/// Helper function for building an atomic store instruction.

688staticboolbuildAtomicStoreInst(constSPIRV::IncomingCall *Call,

689MachineIRBuilder &MIRBuilder,

690SPIRVGlobalRegistry *GR) {

691if (Call->isSpirvOp())

692returnbuildOpFromWrapper(MIRBuilder, SPIRV::OpAtomicStore, Call,Register(0));

693

694Register ScopeRegister =

695buildConstantIntReg32(SPIRV::Scope::Device, MIRBuilder, GR);

696Register PtrRegister = Call->Arguments[0];

697int Semantics =

698 SPIRV::MemorySemantics::SequentiallyConsistent |

699getMemSemanticsForStorageClass(GR->getPointerStorageClass(PtrRegister));

700Register MemSemanticsReg =buildConstantIntReg32(Semantics, MIRBuilder, GR);

701 MIRBuilder.buildInstr(SPIRV::OpAtomicStore)

702 .addUse(PtrRegister)

703 .addUse(ScopeRegister)

704 .addUse(MemSemanticsReg)

705 .addUse(Call->Arguments[1]);

706returntrue;

707}

708

709/// Helper function for building an atomic compare-exchange instruction.

710staticboolbuildAtomicCompareExchangeInst(

711constSPIRV::IncomingCall *Call,constSPIRV::DemangledBuiltin *Builtin,

712unsigned Opcode,MachineIRBuilder &MIRBuilder,SPIRVGlobalRegistry *GR) {

713if (Call->isSpirvOp())

714returnbuildOpFromWrapper(MIRBuilder, Opcode, Call,

715 GR->getSPIRVTypeID(Call->ReturnType));

716

717bool IsCmpxchg = Call->Builtin->Name.contains("cmpxchg");

718MachineRegisterInfo *MRI = MIRBuilder.getMRI();

719

720Register ObjectPtr = Call->Arguments[0];// Pointer (volatile A *object.)

721Register ExpectedArg = Call->Arguments[1];// Comparator (C* expected).

722Register Desired = Call->Arguments[2];// Value (C Desired).

723SPIRVType *SpvDesiredTy = GR->getSPIRVTypeForVReg(Desired);

724LLT DesiredLLT =MRI->getType(Desired);

725

726assert(GR->getSPIRVTypeForVReg(ObjectPtr)->getOpcode() ==

727 SPIRV::OpTypePointer);

728unsigned ExpectedType = GR->getSPIRVTypeForVReg(ExpectedArg)->getOpcode();

729 (void)ExpectedType;

730assert(IsCmpxchg ? ExpectedType == SPIRV::OpTypeInt

731 : ExpectedType == SPIRV::OpTypePointer);

732assert(GR->isScalarOfType(Desired, SPIRV::OpTypeInt));

733

734SPIRVType *SpvObjectPtrTy = GR->getSPIRVTypeForVReg(ObjectPtr);

735assert(SpvObjectPtrTy->getOperand(2).isReg() &&"SPIRV type is expected");

736autoStorageClass =static_cast<SPIRV::StorageClass::StorageClass>(

737 SpvObjectPtrTy->getOperand(1).getImm());

738auto MemSemStorage =getMemSemanticsForStorageClass(StorageClass);

739

740Register MemSemEqualReg;

741Register MemSemUnequalReg;

742uint64_t MemSemEqual =

743 IsCmpxchg

744 ? SPIRV::MemorySemantics::None

745 : SPIRV::MemorySemantics::SequentiallyConsistent | MemSemStorage;

746uint64_t MemSemUnequal =

747 IsCmpxchg

748 ? SPIRV::MemorySemantics::None

749 : SPIRV::MemorySemantics::SequentiallyConsistent | MemSemStorage;

750if (Call->Arguments.size() >= 4) {

751assert(Call->Arguments.size() >= 5 &&

752"Need 5+ args for explicit atomic cmpxchg");

753auto MemOrdEq =

754static_cast<std::memory_order>(getIConstVal(Call->Arguments[3],MRI));

755auto MemOrdNeq =

756static_cast<std::memory_order>(getIConstVal(Call->Arguments[4],MRI));

757 MemSemEqual =getSPIRVMemSemantics(MemOrdEq) | MemSemStorage;

758 MemSemUnequal =getSPIRVMemSemantics(MemOrdNeq) | MemSemStorage;

759if (static_cast<unsigned>(MemOrdEq) == MemSemEqual)

760 MemSemEqualReg = Call->Arguments[3];

761if (static_cast<unsigned>(MemOrdNeq) == MemSemEqual)

762 MemSemUnequalReg = Call->Arguments[4];

763 }

764if (!MemSemEqualReg.isValid())

765 MemSemEqualReg =buildConstantIntReg32(MemSemEqual, MIRBuilder, GR);

766if (!MemSemUnequalReg.isValid())

767 MemSemUnequalReg =buildConstantIntReg32(MemSemUnequal, MIRBuilder, GR);

768

769Register ScopeReg;

770auto Scope = IsCmpxchg ? SPIRV::Scope::Workgroup : SPIRV::Scope::Device;

771if (Call->Arguments.size() >= 6) {

772assert(Call->Arguments.size() == 6 &&

773"Extra args for explicit atomic cmpxchg");

774auto ClScope =static_cast<SPIRV::CLMemoryScope>(

775getIConstVal(Call->Arguments[5],MRI));

776 Scope =getSPIRVScope(ClScope);

777if (ClScope ==static_cast<unsigned>(Scope))

778 ScopeReg = Call->Arguments[5];

779 }

780if (!ScopeReg.isValid())

781 ScopeReg =buildConstantIntReg32(Scope, MIRBuilder, GR);

782

783Register Expected = IsCmpxchg

784 ? ExpectedArg

785 :buildLoadInst(SpvDesiredTy, ExpectedArg, MIRBuilder,

786 GR,LLT::scalar(64));

787MRI->setType(Expected, DesiredLLT);

788Register Tmp = !IsCmpxchg ?MRI->createGenericVirtualRegister(DesiredLLT)

789 : Call->ReturnRegister;

790if (!MRI->getRegClassOrNull(Tmp))

791MRI->setRegClass(Tmp, GR->getRegClass(SpvDesiredTy));

792 GR->assignSPIRVTypeToVReg(SpvDesiredTy, Tmp, MIRBuilder.getMF());

793

794SPIRVType *IntTy = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);

795 MIRBuilder.buildInstr(Opcode)

796 .addDef(Tmp)

797 .addUse(GR->getSPIRVTypeID(IntTy))

798 .addUse(ObjectPtr)

799 .addUse(ScopeReg)

800 .addUse(MemSemEqualReg)

801 .addUse(MemSemUnequalReg)

802 .addUse(Desired)

803 .addUse(Expected);

804if (!IsCmpxchg) {

805 MIRBuilder.buildInstr(SPIRV::OpStore).addUse(ExpectedArg).addUse(Tmp);

806 MIRBuilder.buildICmp(CmpInst::ICMP_EQ, Call->ReturnRegister, Tmp,Expected);

807 }

808returntrue;

809}

810

811/// Helper function for building atomic instructions.

812staticboolbuildAtomicRMWInst(constSPIRV::IncomingCall *Call,unsigned Opcode,

813MachineIRBuilder &MIRBuilder,

814SPIRVGlobalRegistry *GR) {

815if (Call->isSpirvOp())

816returnbuildOpFromWrapper(MIRBuilder, Opcode, Call,

817 GR->getSPIRVTypeID(Call->ReturnType));

818

819MachineRegisterInfo *MRI = MIRBuilder.getMRI();

820Register ScopeRegister =

821 Call->Arguments.size() >= 4 ? Call->Arguments[3] :Register();

822

823assert(Call->Arguments.size() <= 4 &&

824"Too many args for explicit atomic RMW");

825 ScopeRegister =buildScopeReg(ScopeRegister, SPIRV::Scope::Workgroup,

826 MIRBuilder, GR,MRI);

827

828Register PtrRegister = Call->Arguments[0];

829unsigned Semantics = SPIRV::MemorySemantics::None;

830Register MemSemanticsReg =

831 Call->Arguments.size() >= 3 ? Call->Arguments[2] :Register();

832 MemSemanticsReg =buildMemSemanticsReg(MemSemanticsReg, PtrRegister,

833 Semantics, MIRBuilder, GR);

834Register ValueReg = Call->Arguments[1];

835Register ValueTypeReg = GR->getSPIRVTypeID(Call->ReturnType);

836// support cl_ext_float_atomics

837if (Call->ReturnType->getOpcode() == SPIRV::OpTypeFloat) {

838if (Opcode == SPIRV::OpAtomicIAdd) {

839 Opcode = SPIRV::OpAtomicFAddEXT;

840 }elseif (Opcode == SPIRV::OpAtomicISub) {

841// Translate OpAtomicISub applied to a floating type argument to

842// OpAtomicFAddEXT with the negative value operand

843 Opcode = SPIRV::OpAtomicFAddEXT;

844Register NegValueReg =

845MRI->createGenericVirtualRegister(MRI->getType(ValueReg));

846MRI->setRegClass(NegValueReg, GR->getRegClass(Call->ReturnType));

847 GR->assignSPIRVTypeToVReg(Call->ReturnType, NegValueReg,

848 MIRBuilder.getMF());

849 MIRBuilder.buildInstr(TargetOpcode::G_FNEG)

850 .addDef(NegValueReg)

851 .addUse(ValueReg);

852insertAssignInstr(NegValueReg,nullptr, Call->ReturnType, GR, MIRBuilder,

853 MIRBuilder.getMF().getRegInfo());

854 ValueReg = NegValueReg;

855 }

856 }

857 MIRBuilder.buildInstr(Opcode)

858 .addDef(Call->ReturnRegister)

859 .addUse(ValueTypeReg)

860 .addUse(PtrRegister)

861 .addUse(ScopeRegister)

862 .addUse(MemSemanticsReg)

863 .addUse(ValueReg);

864returntrue;

865}

866

867/// Helper function for building an atomic floating-type instruction.

868staticboolbuildAtomicFloatingRMWInst(constSPIRV::IncomingCall *Call,

869unsigned Opcode,

870MachineIRBuilder &MIRBuilder,

871SPIRVGlobalRegistry *GR) {

872assert(Call->Arguments.size() == 4 &&

873"Wrong number of atomic floating-type builtin");

874Register PtrReg = Call->Arguments[0];

875Register ScopeReg = Call->Arguments[1];

876Register MemSemanticsReg = Call->Arguments[2];

877Register ValueReg = Call->Arguments[3];

878 MIRBuilder.buildInstr(Opcode)

879 .addDef(Call->ReturnRegister)

880 .addUse(GR->getSPIRVTypeID(Call->ReturnType))

881 .addUse(PtrReg)

882 .addUse(ScopeReg)

883 .addUse(MemSemanticsReg)

884 .addUse(ValueReg);

885returntrue;

886}

887

888/// Helper function for building atomic flag instructions (e.g.

889/// OpAtomicFlagTestAndSet).

890staticboolbuildAtomicFlagInst(constSPIRV::IncomingCall *Call,

891unsigned Opcode,MachineIRBuilder &MIRBuilder,

892SPIRVGlobalRegistry *GR) {

893bool IsSet = Opcode == SPIRV::OpAtomicFlagTestAndSet;

894Register TypeReg = GR->getSPIRVTypeID(Call->ReturnType);

895if (Call->isSpirvOp())

896returnbuildOpFromWrapper(MIRBuilder, Opcode, Call,

897 IsSet ? TypeReg :Register(0));

898

899MachineRegisterInfo *MRI = MIRBuilder.getMRI();

900Register PtrRegister = Call->Arguments[0];

901unsigned Semantics = SPIRV::MemorySemantics::SequentiallyConsistent;

902Register MemSemanticsReg =

903 Call->Arguments.size() >= 2 ? Call->Arguments[1] :Register();

904 MemSemanticsReg =buildMemSemanticsReg(MemSemanticsReg, PtrRegister,

905 Semantics, MIRBuilder, GR);

906

907assert((Opcode != SPIRV::OpAtomicFlagClear ||

908 (Semantics != SPIRV::MemorySemantics::Acquire &&

909 Semantics != SPIRV::MemorySemantics::AcquireRelease)) &&

910"Invalid memory order argument!");

911

912Register ScopeRegister =

913 Call->Arguments.size() >= 3 ? Call->Arguments[2] :Register();

914 ScopeRegister =

915buildScopeReg(ScopeRegister, SPIRV::Scope::Device, MIRBuilder, GR,MRI);

916

917auto MIB = MIRBuilder.buildInstr(Opcode);

918if (IsSet)

919 MIB.addDef(Call->ReturnRegister).addUse(TypeReg);

920

921 MIB.addUse(PtrRegister).addUse(ScopeRegister).addUse(MemSemanticsReg);

922returntrue;

923}

924

925/// Helper function for building barriers, i.e., memory/control ordering

926/// operations.

927staticboolbuildBarrierInst(constSPIRV::IncomingCall *Call,unsigned Opcode,

928MachineIRBuilder &MIRBuilder,

929SPIRVGlobalRegistry *GR) {

930constSPIRV::DemangledBuiltin *Builtin = Call->Builtin;

931constauto *ST =

932static_cast<constSPIRVSubtarget *>(&MIRBuilder.getMF().getSubtarget());

933if ((Opcode == SPIRV::OpControlBarrierArriveINTEL ||

934 Opcode == SPIRV::OpControlBarrierWaitINTEL) &&

935 !ST->canUseExtension(SPIRV::Extension::SPV_INTEL_split_barrier)) {

936 std::string DiagMsg = std::string(Builtin->Name) +

937": the builtin requires the following SPIR-V "

938"extension: SPV_INTEL_split_barrier";

939report_fatal_error(DiagMsg.c_str(),false);

940 }

941

942if (Call->isSpirvOp())

943returnbuildOpFromWrapper(MIRBuilder, Opcode, Call,Register(0));

944

945MachineRegisterInfo *MRI = MIRBuilder.getMRI();

946unsigned MemFlags =getIConstVal(Call->Arguments[0],MRI);

947unsigned MemSemantics = SPIRV::MemorySemantics::None;

948

949if (MemFlags & SPIRV::CLK_LOCAL_MEM_FENCE)

950 MemSemantics |= SPIRV::MemorySemantics::WorkgroupMemory;

951

952if (MemFlags & SPIRV::CLK_GLOBAL_MEM_FENCE)

953 MemSemantics |= SPIRV::MemorySemantics::CrossWorkgroupMemory;

954

955if (MemFlags & SPIRV::CLK_IMAGE_MEM_FENCE)

956 MemSemantics |= SPIRV::MemorySemantics::ImageMemory;

957

958if (Opcode == SPIRV::OpMemoryBarrier)

959 MemSemantics =getSPIRVMemSemantics(static_cast<std::memory_order>(

960getIConstVal(Call->Arguments[1],MRI))) |

961 MemSemantics;

962elseif (Opcode == SPIRV::OpControlBarrierArriveINTEL)

963 MemSemantics |= SPIRV::MemorySemantics::Release;

964elseif (Opcode == SPIRV::OpControlBarrierWaitINTEL)

965 MemSemantics |= SPIRV::MemorySemantics::Acquire;

966else

967 MemSemantics |= SPIRV::MemorySemantics::SequentiallyConsistent;

968

969Register MemSemanticsReg =

970 MemFlags == MemSemantics

971 ? Call->Arguments[0]

972 :buildConstantIntReg32(MemSemantics, MIRBuilder, GR);

973Register ScopeReg;

974 SPIRV::Scope::Scope Scope = SPIRV::Scope::Workgroup;

975 SPIRV::Scope::Scope MemScope = Scope;

976if (Call->Arguments.size() >= 2) {

977assert(

978 ((Opcode != SPIRV::OpMemoryBarrier && Call->Arguments.size() == 2) ||

979 (Opcode == SPIRV::OpMemoryBarrier && Call->Arguments.size() == 3)) &&

980"Extra args for explicitly scoped barrier");

981Register ScopeArg = (Opcode == SPIRV::OpMemoryBarrier) ? Call->Arguments[2]

982 : Call->Arguments[1];

983 SPIRV::CLMemoryScope CLScope =

984static_cast<SPIRV::CLMemoryScope>(getIConstVal(ScopeArg,MRI));

985 MemScope =getSPIRVScope(CLScope);

986if (!(MemFlags & SPIRV::CLK_LOCAL_MEM_FENCE) ||

987 (Opcode == SPIRV::OpMemoryBarrier))

988 Scope = MemScope;

989if (CLScope ==static_cast<unsigned>(Scope))

990 ScopeReg = Call->Arguments[1];

991 }

992

993if (!ScopeReg.isValid())

994 ScopeReg =buildConstantIntReg32(Scope, MIRBuilder, GR);

995

996auto MIB = MIRBuilder.buildInstr(Opcode).addUse(ScopeReg);

997if (Opcode != SPIRV::OpMemoryBarrier)

998 MIB.addUse(buildConstantIntReg32(MemScope, MIRBuilder, GR));

999 MIB.addUse(MemSemanticsReg);

1000returntrue;

1001}

1002

1003staticunsignedgetNumComponentsForDim(SPIRV::Dim::Dim dim) {

1004switch (dim) {

1005case SPIRV::Dim::DIM_1D:

1006case SPIRV::Dim::DIM_Buffer:

1007return 1;

1008case SPIRV::Dim::DIM_2D:

1009case SPIRV::Dim::DIM_Cube:

1010case SPIRV::Dim::DIM_Rect:

1011return 2;

1012case SPIRV::Dim::DIM_3D:

1013return 3;

1014default:

1015report_fatal_error("Cannot get num components for given Dim");

1016 }

1017}

1018

1019/// Helper function for obtaining the number of size components.

1020staticunsignedgetNumSizeComponents(SPIRVType *imgType) {

1021assert(imgType->getOpcode() == SPIRV::OpTypeImage);

1022auto dim =static_cast<SPIRV::Dim::Dim>(imgType->getOperand(2).getImm());

1023unsigned numComps =getNumComponentsForDim(dim);

1024bool arrayed = imgType->getOperand(4).getImm() == 1;

1025return arrayed ? numComps + 1 : numComps;

1026}

1027

1028//===----------------------------------------------------------------------===//

1029// Implementation functions for each builtin group

1030//===----------------------------------------------------------------------===//

1031

1032staticboolgenerateExtInst(constSPIRV::IncomingCall *Call,

1033MachineIRBuilder &MIRBuilder,

1034SPIRVGlobalRegistry *GR) {

1035// Lookup the extended instruction number in the TableGen records.

1036constSPIRV::DemangledBuiltin *Builtin = Call->Builtin;

1037uint32_t Number =

1038 SPIRV::lookupExtendedBuiltin(Builtin->Name, Builtin->Set)->Number;

1039

1040// Build extended instruction.

1041auto MIB =

1042 MIRBuilder.buildInstr(SPIRV::OpExtInst)

1043 .addDef(Call->ReturnRegister)

1044 .addUse(GR->getSPIRVTypeID(Call->ReturnType))

1045 .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::OpenCL_std))

1046 .addImm(Number);

1047

1048for (autoArgument : Call->Arguments)

1049 MIB.addUse(Argument);

1050returntrue;

1051}

1052

1053staticboolgenerateRelationalInst(constSPIRV::IncomingCall *Call,

1054MachineIRBuilder &MIRBuilder,

1055SPIRVGlobalRegistry *GR) {

1056// Lookup the instruction opcode in the TableGen records.

1057constSPIRV::DemangledBuiltin *Builtin = Call->Builtin;

1058unsigned Opcode =

1059 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;

1060

1061Register CompareRegister;

1062SPIRVType *RelationType;

1063 std::tie(CompareRegister, RelationType) =

1064buildBoolRegister(MIRBuilder, Call->ReturnType, GR);

1065

1066// Build relational instruction.

1067auto MIB = MIRBuilder.buildInstr(Opcode)

1068 .addDef(CompareRegister)

1069 .addUse(GR->getSPIRVTypeID(RelationType));

1070

1071for (autoArgument : Call->Arguments)

1072 MIB.addUse(Argument);

1073

1074// Build select instruction.

1075returnbuildSelectInst(MIRBuilder, Call->ReturnRegister, CompareRegister,

1076 Call->ReturnType, GR);

1077}

1078

1079staticboolgenerateGroupInst(constSPIRV::IncomingCall *Call,

1080MachineIRBuilder &MIRBuilder,

1081SPIRVGlobalRegistry *GR) {

1082constSPIRV::DemangledBuiltin *Builtin = Call->Builtin;

1083constSPIRV::GroupBuiltin *GroupBuiltin =

1084 SPIRV::lookupGroupBuiltin(Builtin->Name);

1085

1086MachineRegisterInfo *MRI = MIRBuilder.getMRI();

1087if (Call->isSpirvOp()) {

1088if (GroupBuiltin->NoGroupOperation)

1089returnbuildOpFromWrapper(MIRBuilder, GroupBuiltin->Opcode, Call,

1090 GR->getSPIRVTypeID(Call->ReturnType));

1091

1092// Group Operation is a literal

1093Register GroupOpReg = Call->Arguments[1];

1094constMachineInstr *MI =getDefInstrMaybeConstant(GroupOpReg,MRI);

1095if (!MI ||MI->getOpcode() != TargetOpcode::G_CONSTANT)

1096report_fatal_error(

1097"Group Operation parameter must be an integer constant");

1098uint64_t GrpOp =MI->getOperand(1).getCImm()->getValue().getZExtValue();

1099Register ScopeReg = Call->Arguments[0];

1100auto MIB = MIRBuilder.buildInstr(GroupBuiltin->Opcode)

1101 .addDef(Call->ReturnRegister)

1102 .addUse(GR->getSPIRVTypeID(Call->ReturnType))

1103 .addUse(ScopeReg)

1104 .addImm(GrpOp);

1105for (unsigned i = 2; i < Call->Arguments.size(); ++i)

1106 MIB.addUse(Call->Arguments[i]);

1107returntrue;

1108 }

1109

1110Register Arg0;

1111if (GroupBuiltin->HasBoolArg) {

1112SPIRVType *BoolType = GR->getOrCreateSPIRVBoolType(MIRBuilder);

1113Register BoolReg = Call->Arguments[0];

1114SPIRVType *BoolRegType = GR->getSPIRVTypeForVReg(BoolReg);

1115if (!BoolRegType)

1116report_fatal_error("Can't find a register's type definition");

1117MachineInstr *ArgInstruction =getDefInstrMaybeConstant(BoolReg,MRI);

1118if (ArgInstruction->getOpcode() == TargetOpcode::G_CONSTANT) {

1119if (BoolRegType->getOpcode() != SPIRV::OpTypeBool)

1120 Arg0 = GR->buildConstantInt(getIConstVal(BoolReg,MRI), MIRBuilder,

1121 BoolType);

1122 }else {

1123if (BoolRegType->getOpcode() == SPIRV::OpTypeInt) {

1124 Arg0 =MRI->createGenericVirtualRegister(LLT::scalar(1));

1125MRI->setRegClass(Arg0, &SPIRV::iIDRegClass);

1126 GR->assignSPIRVTypeToVReg(BoolType, Arg0, MIRBuilder.getMF());

1127 MIRBuilder.buildICmp(CmpInst::ICMP_NE, Arg0, BoolReg,

1128 GR->buildConstantInt(0, MIRBuilder, BoolRegType));

1129insertAssignInstr(Arg0,nullptr, BoolType, GR, MIRBuilder,

1130 MIRBuilder.getMF().getRegInfo());

1131 }elseif (BoolRegType->getOpcode() != SPIRV::OpTypeBool) {

1132report_fatal_error("Expect a boolean argument");

1133 }

1134// if BoolReg is a boolean register, we don't need to do anything

1135 }

1136 }

1137

1138Register GroupResultRegister = Call->ReturnRegister;

1139SPIRVType *GroupResultType = Call->ReturnType;

1140

1141// TODO: maybe we need to check whether the result type is already boolean

1142// and in this case do not insert select instruction.

1143constbool HasBoolReturnTy =

1144 GroupBuiltin->IsElect || GroupBuiltin->IsAllOrAny ||

1145 GroupBuiltin->IsAllEqual || GroupBuiltin->IsLogical ||

1146 GroupBuiltin->IsInverseBallot || GroupBuiltin->IsBallotBitExtract;

1147

1148if (HasBoolReturnTy)

1149 std::tie(GroupResultRegister, GroupResultType) =

1150buildBoolRegister(MIRBuilder, Call->ReturnType, GR);

1151

1152auto Scope = Builtin->Name.starts_with("sub_group") ? SPIRV::Scope::Subgroup

1153 : SPIRV::Scope::Workgroup;

1154Register ScopeRegister =buildConstantIntReg32(Scope, MIRBuilder, GR);

1155

1156Register VecReg;

1157if (GroupBuiltin->Opcode == SPIRV::OpGroupBroadcast &&

1158 Call->Arguments.size() > 2) {

1159// For OpGroupBroadcast "LocalId must be an integer datatype. It must be a

1160// scalar, a vector with 2 components, or a vector with 3 components.",

1161// meaning that we must create a vector from the function arguments if

1162// it's a work_group_broadcast(val, local_id_x, local_id_y) or

1163// work_group_broadcast(val, local_id_x, local_id_y, local_id_z) call.

1164Register ElemReg = Call->Arguments[1];

1165SPIRVType *ElemType = GR->getSPIRVTypeForVReg(ElemReg);

1166if (!ElemType || ElemType->getOpcode() != SPIRV::OpTypeInt)

1167report_fatal_error("Expect an integer <LocalId> argument");

1168unsigned VecLen = Call->Arguments.size() - 1;

1169 VecReg =MRI->createGenericVirtualRegister(

1170LLT::fixed_vector(VecLen,MRI->getType(ElemReg)));

1171MRI->setRegClass(VecReg, &SPIRV::vIDRegClass);

1172SPIRVType *VecType =

1173 GR->getOrCreateSPIRVVectorType(ElemType, VecLen, MIRBuilder);

1174 GR->assignSPIRVTypeToVReg(VecType, VecReg, MIRBuilder.getMF());

1175auto MIB =

1176 MIRBuilder.buildInstr(TargetOpcode::G_BUILD_VECTOR).addDef(VecReg);

1177for (unsigned i = 1; i < Call->Arguments.size(); i++) {

1178 MIB.addUse(Call->Arguments[i]);

1179setRegClassIfNull(Call->Arguments[i],MRI, GR);

1180 }

1181insertAssignInstr(VecReg,nullptr, VecType, GR, MIRBuilder,

1182 MIRBuilder.getMF().getRegInfo());

1183 }

1184

1185// Build work/sub group instruction.

1186auto MIB = MIRBuilder.buildInstr(GroupBuiltin->Opcode)

1187 .addDef(GroupResultRegister)

1188 .addUse(GR->getSPIRVTypeID(GroupResultType))

1189 .addUse(ScopeRegister);

1190

1191if (!GroupBuiltin->NoGroupOperation)

1192 MIB.addImm(GroupBuiltin->GroupOperation);

1193if (Call->Arguments.size() > 0) {

1194 MIB.addUse(Arg0.isValid() ? Arg0 : Call->Arguments[0]);

1195setRegClassIfNull(Call->Arguments[0],MRI, GR);

1196if (VecReg.isValid())

1197 MIB.addUse(VecReg);

1198else

1199for (unsigned i = 1; i < Call->Arguments.size(); i++)

1200 MIB.addUse(Call->Arguments[i]);

1201 }

1202

1203// Build select instruction.

1204if (HasBoolReturnTy)

1205buildSelectInst(MIRBuilder, Call->ReturnRegister, GroupResultRegister,

1206 Call->ReturnType, GR);

1207returntrue;

1208}

1209

1210staticboolgenerateIntelSubgroupsInst(constSPIRV::IncomingCall *Call,

1211MachineIRBuilder &MIRBuilder,

1212SPIRVGlobalRegistry *GR) {

1213constSPIRV::DemangledBuiltin *Builtin = Call->Builtin;

1214MachineFunction &MF = MIRBuilder.getMF();

1215constauto *ST =static_cast<constSPIRVSubtarget *>(&MF.getSubtarget());

1216constSPIRV::IntelSubgroupsBuiltin *IntelSubgroups =

1217 SPIRV::lookupIntelSubgroupsBuiltin(Builtin->Name);

1218

1219if (IntelSubgroups->IsMedia &&

1220 !ST->canUseExtension(SPIRV::Extension::SPV_INTEL_media_block_io)) {

1221 std::string DiagMsg = std::string(Builtin->Name) +

1222": the builtin requires the following SPIR-V "

1223"extension: SPV_INTEL_media_block_io";

1224report_fatal_error(DiagMsg.c_str(),false);

1225 }elseif (!IntelSubgroups->IsMedia &&

1226 !ST->canUseExtension(SPIRV::Extension::SPV_INTEL_subgroups)) {

1227 std::string DiagMsg = std::string(Builtin->Name) +

1228": the builtin requires the following SPIR-V "

1229"extension: SPV_INTEL_subgroups";

1230report_fatal_error(DiagMsg.c_str(),false);

1231 }

1232

1233uint32_t OpCode = IntelSubgroups->Opcode;

1234if (Call->isSpirvOp()) {

1235bool IsSet = OpCode != SPIRV::OpSubgroupBlockWriteINTEL &&

1236 OpCode != SPIRV::OpSubgroupImageBlockWriteINTEL &&

1237 OpCode != SPIRV::OpSubgroupImageMediaBlockWriteINTEL;

1238returnbuildOpFromWrapper(MIRBuilder, OpCode, Call,

1239 IsSet ? GR->getSPIRVTypeID(Call->ReturnType)

1240 :Register(0));

1241 }

1242

1243if (IntelSubgroups->IsBlock) {

1244// Minimal number or arguments set in TableGen records is 1

1245if (SPIRVType *Arg0Type = GR->getSPIRVTypeForVReg(Call->Arguments[0])) {

1246if (Arg0Type->getOpcode() == SPIRV::OpTypeImage) {

1247// TODO: add required validation from the specification:

1248// "'Image' must be an object whose type is OpTypeImage with a 'Sampled'

1249// operand of 0 or 2. If the 'Sampled' operand is 2, then some

1250// dimensions require a capability."

1251switch (OpCode) {

1252case SPIRV::OpSubgroupBlockReadINTEL:

1253 OpCode = SPIRV::OpSubgroupImageBlockReadINTEL;

1254break;

1255case SPIRV::OpSubgroupBlockWriteINTEL:

1256 OpCode = SPIRV::OpSubgroupImageBlockWriteINTEL;

1257break;

1258 }

1259 }

1260 }

1261 }

1262

1263// TODO: opaque pointers types should be eventually resolved in such a way

1264// that validation of block read is enabled with respect to the following

1265// specification requirement:

1266// "'Result Type' may be a scalar or vector type, and its component type must

1267// be equal to the type pointed to by 'Ptr'."

1268// For example, function parameter type should not be default i8 pointer, but

1269// depend on the result type of the instruction where it is used as a pointer

1270// argument of OpSubgroupBlockReadINTEL

1271

1272// Build Intel subgroups instruction

1273MachineInstrBuilder MIB =

1274 IntelSubgroups->IsWrite

1275 ? MIRBuilder.buildInstr(OpCode)

1276 : MIRBuilder.buildInstr(OpCode)

1277 .addDef(Call->ReturnRegister)

1278 .addUse(GR->getSPIRVTypeID(Call->ReturnType));

1279for (size_t i = 0; i < Call->Arguments.size(); ++i)

1280 MIB.addUse(Call->Arguments[i]);

1281returntrue;

1282}

1283

1284staticboolgenerateGroupUniformInst(constSPIRV::IncomingCall *Call,

1285MachineIRBuilder &MIRBuilder,

1286SPIRVGlobalRegistry *GR) {

1287constSPIRV::DemangledBuiltin *Builtin = Call->Builtin;

1288MachineFunction &MF = MIRBuilder.getMF();

1289constauto *ST =static_cast<constSPIRVSubtarget *>(&MF.getSubtarget());

1290if (!ST->canUseExtension(

1291 SPIRV::Extension::SPV_KHR_uniform_group_instructions)) {

1292 std::string DiagMsg = std::string(Builtin->Name) +

1293": the builtin requires the following SPIR-V "

1294"extension: SPV_KHR_uniform_group_instructions";

1295report_fatal_error(DiagMsg.c_str(),false);

1296 }

1297constSPIRV::GroupUniformBuiltin *GroupUniform =

1298 SPIRV::lookupGroupUniformBuiltin(Builtin->Name);

1299MachineRegisterInfo *MRI = MIRBuilder.getMRI();

1300

1301Register GroupResultReg = Call->ReturnRegister;

1302Register ScopeReg = Call->Arguments[0];

1303Register ValueReg = Call->Arguments[2];

1304

1305// Group Operation

1306Register ConstGroupOpReg = Call->Arguments[1];

1307constMachineInstr *Const =getDefInstrMaybeConstant(ConstGroupOpReg,MRI);

1308if (!Const || Const->getOpcode() != TargetOpcode::G_CONSTANT)

1309report_fatal_error(

1310"expect a constant group operation for a uniform group instruction",

1311false);

1312constMachineOperand &ConstOperand = Const->getOperand(1);

1313if (!ConstOperand.isCImm())

1314report_fatal_error("uniform group instructions: group operation must be an "

1315"integer constant",

1316false);

1317

1318auto MIB = MIRBuilder.buildInstr(GroupUniform->Opcode)

1319 .addDef(GroupResultReg)

1320 .addUse(GR->getSPIRVTypeID(Call->ReturnType))

1321 .addUse(ScopeReg);

1322addNumImm(ConstOperand.getCImm()->getValue(), MIB);

1323 MIB.addUse(ValueReg);

1324

1325returntrue;

1326}

1327

1328staticboolgenerateKernelClockInst(constSPIRV::IncomingCall *Call,

1329MachineIRBuilder &MIRBuilder,

1330SPIRVGlobalRegistry *GR) {

1331constSPIRV::DemangledBuiltin *Builtin = Call->Builtin;

1332MachineFunction &MF = MIRBuilder.getMF();

1333constauto *ST =static_cast<constSPIRVSubtarget *>(&MF.getSubtarget());

1334if (!ST->canUseExtension(SPIRV::Extension::SPV_KHR_shader_clock)) {

1335 std::string DiagMsg = std::string(Builtin->Name) +

1336": the builtin requires the following SPIR-V "

1337"extension: SPV_KHR_shader_clock";

1338report_fatal_error(DiagMsg.c_str(),false);

1339 }

1340

1341Register ResultReg = Call->ReturnRegister;

1342

1343// Deduce the `Scope` operand from the builtin function name.

1344 SPIRV::Scope::Scope ScopeArg =

1345StringSwitch<SPIRV::Scope::Scope>(Builtin->Name)

1346 .EndsWith("device", SPIRV::Scope::Scope::Device)

1347 .EndsWith("work_group", SPIRV::Scope::Scope::Workgroup)

1348 .EndsWith("sub_group", SPIRV::Scope::Scope::Subgroup);

1349Register ScopeReg =buildConstantIntReg32(ScopeArg, MIRBuilder, GR);

1350

1351 MIRBuilder.buildInstr(SPIRV::OpReadClockKHR)

1352 .addDef(ResultReg)

1353 .addUse(GR->getSPIRVTypeID(Call->ReturnType))

1354 .addUse(ScopeReg);

1355

1356returntrue;

1357}

1358

1359// These queries ask for a single size_t result for a given dimension index, e.g

1360// size_t get_global_id(uint dimindex). In SPIR-V, the builtins corresonding to

1361// these values are all vec3 types, so we need to extract the correct index or

1362// return defaultVal (0 or 1 depending on the query). We also handle extending

1363// or tuncating in case size_t does not match the expected result type's

1364// bitwidth.

1365//

1366// For a constant index >= 3 we generate:

1367// %res = OpConstant %SizeT 0

1368//

1369// For other indices we generate:

1370// %g = OpVariable %ptr_V3_SizeT Input

1371// OpDecorate %g BuiltIn XXX

1372// OpDecorate %g LinkageAttributes "__spirv_BuiltInXXX"

1373// OpDecorate %g Constant

1374// %loadedVec = OpLoad %V3_SizeT %g

1375//

1376// Then, if the index is constant < 3, we generate:

1377// %res = OpCompositeExtract %SizeT %loadedVec idx

1378// If the index is dynamic, we generate:

1379// %tmp = OpVectorExtractDynamic %SizeT %loadedVec %idx

1380// %cmp = OpULessThan %bool %idx %const_3

1381// %res = OpSelect %SizeT %cmp %tmp %const_0

1382//

1383// If the bitwidth of %res does not match the expected return type, we add an

1384// extend or truncate.

1385staticboolgenWorkgroupQuery(constSPIRV::IncomingCall *Call,

1386MachineIRBuilder &MIRBuilder,

1387SPIRVGlobalRegistry *GR,

1388 SPIRV::BuiltIn::BuiltIn BuiltinValue,

1389uint64_t DefaultValue) {

1390Register IndexRegister = Call->Arguments[0];

1391constunsigned ResultWidth = Call->ReturnType->getOperand(1).getImm();

1392constunsigned PointerSize = GR->getPointerSize();

1393constSPIRVType *PointerSizeType =

1394 GR->getOrCreateSPIRVIntegerType(PointerSize, MIRBuilder);

1395MachineRegisterInfo *MRI = MIRBuilder.getMRI();

1396auto IndexInstruction =getDefInstrMaybeConstant(IndexRegister,MRI);

1397

1398// Set up the final register to do truncation or extension on at the end.

1399Register ToTruncate = Call->ReturnRegister;

1400

1401// If the index is constant, we can statically determine if it is in range.

1402bool IsConstantIndex =

1403 IndexInstruction->getOpcode() == TargetOpcode::G_CONSTANT;

1404

1405// If it's out of range (max dimension is 3), we can just return the constant

1406// default value (0 or 1 depending on which query function).

1407if (IsConstantIndex &&getIConstVal(IndexRegister,MRI) >= 3) {

1408Register DefaultReg = Call->ReturnRegister;

1409if (PointerSize != ResultWidth) {

1410 DefaultReg =MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));

1411MRI->setRegClass(DefaultReg, &SPIRV::iIDRegClass);

1412 GR->assignSPIRVTypeToVReg(PointerSizeType, DefaultReg,

1413 MIRBuilder.getMF());

1414 ToTruncate = DefaultReg;

1415 }

1416auto NewRegister =

1417 GR->buildConstantInt(DefaultValue, MIRBuilder, PointerSizeType);

1418 MIRBuilder.buildCopy(DefaultReg, NewRegister);

1419 }else {// If it could be in range, we need to load from the given builtin.

1420auto Vec3Ty =

1421 GR->getOrCreateSPIRVVectorType(PointerSizeType, 3, MIRBuilder);

1422Register LoadedVector =

1423buildBuiltinVariableLoad(MIRBuilder, Vec3Ty, GR, BuiltinValue,

1424LLT::fixed_vector(3, PointerSize));

1425// Set up the vreg to extract the result to (possibly a new temporary one).

1426Register Extracted = Call->ReturnRegister;

1427if (!IsConstantIndex || PointerSize != ResultWidth) {

1428 Extracted =MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));

1429MRI->setRegClass(Extracted, &SPIRV::iIDRegClass);

1430 GR->assignSPIRVTypeToVReg(PointerSizeType, Extracted, MIRBuilder.getMF());

1431 }

1432// Use Intrinsic::spv_extractelt so dynamic vs static extraction is

1433// handled later: extr = spv_extractelt LoadedVector, IndexRegister.

1434MachineInstrBuilder ExtractInst = MIRBuilder.buildIntrinsic(

1435 Intrinsic::spv_extractelt,ArrayRef<Register>{Extracted},true,false);

1436 ExtractInst.addUse(LoadedVector).addUse(IndexRegister);

1437

1438// If the index is dynamic, need check if it's < 3, and then use a select.

1439if (!IsConstantIndex) {

1440insertAssignInstr(Extracted,nullptr, PointerSizeType, GR, MIRBuilder,

1441 *MRI);

1442

1443auto IndexType = GR->getSPIRVTypeForVReg(IndexRegister);

1444auto BoolType = GR->getOrCreateSPIRVBoolType(MIRBuilder);

1445

1446Register CompareRegister =

1447MRI->createGenericVirtualRegister(LLT::scalar(1));

1448MRI->setRegClass(CompareRegister, &SPIRV::iIDRegClass);

1449 GR->assignSPIRVTypeToVReg(BoolType, CompareRegister, MIRBuilder.getMF());

1450

1451// Use G_ICMP to check if idxVReg < 3.

1452 MIRBuilder.buildICmp(CmpInst::ICMP_ULT, CompareRegister, IndexRegister,

1453 GR->buildConstantInt(3, MIRBuilder, IndexType));

1454

1455// Get constant for the default value (0 or 1 depending on which

1456// function).

1457Register DefaultRegister =

1458 GR->buildConstantInt(DefaultValue, MIRBuilder, PointerSizeType);

1459

1460// Get a register for the selection result (possibly a new temporary one).

1461Register SelectionResult = Call->ReturnRegister;

1462if (PointerSize != ResultWidth) {

1463 SelectionResult =

1464MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));

1465MRI->setRegClass(SelectionResult, &SPIRV::iIDRegClass);

1466 GR->assignSPIRVTypeToVReg(PointerSizeType, SelectionResult,

1467 MIRBuilder.getMF());

1468 }

1469// Create the final G_SELECT to return the extracted value or the default.

1470 MIRBuilder.buildSelect(SelectionResult, CompareRegister, Extracted,

1471 DefaultRegister);

1472 ToTruncate = SelectionResult;

1473 }else {

1474 ToTruncate = Extracted;

1475 }

1476 }

1477// Alter the result's bitwidth if it does not match the SizeT value extracted.

1478if (PointerSize != ResultWidth)

1479 MIRBuilder.buildZExtOrTrunc(Call->ReturnRegister, ToTruncate);

1480returntrue;

1481}

1482

1483staticboolgenerateBuiltinVar(constSPIRV::IncomingCall *Call,

1484MachineIRBuilder &MIRBuilder,

1485SPIRVGlobalRegistry *GR) {

1486// Lookup the builtin variable record.

1487constSPIRV::DemangledBuiltin *Builtin = Call->Builtin;

1488 SPIRV::BuiltIn::BuiltInValue =

1489 SPIRV::lookupGetBuiltin(Builtin->Name, Builtin->Set)->Value;

1490

1491if (Value == SPIRV::BuiltIn::GlobalInvocationId)

1492returngenWorkgroupQuery(Call, MIRBuilder, GR,Value, 0);

1493

1494// Build a load instruction for the builtin variable.

1495unsignedBitWidth = GR->getScalarOrVectorBitWidth(Call->ReturnType);

1496LLT LLType;

1497if (Call->ReturnType->getOpcode() == SPIRV::OpTypeVector)

1498 LLType =

1499LLT::fixed_vector(Call->ReturnType->getOperand(2).getImm(),BitWidth);

1500else

1501 LLType =LLT::scalar(BitWidth);

1502

1503returnbuildBuiltinVariableLoad(MIRBuilder, Call->ReturnType, GR,Value,

1504 LLType, Call->ReturnRegister);

1505}

1506

1507staticboolgenerateAtomicInst(constSPIRV::IncomingCall *Call,

1508MachineIRBuilder &MIRBuilder,

1509SPIRVGlobalRegistry *GR) {

1510// Lookup the instruction opcode in the TableGen records.

1511constSPIRV::DemangledBuiltin *Builtin = Call->Builtin;

1512unsigned Opcode =

1513 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;

1514

1515switch (Opcode) {

1516case SPIRV::OpStore:

1517returnbuildAtomicInitInst(Call, MIRBuilder);

1518case SPIRV::OpAtomicLoad:

1519returnbuildAtomicLoadInst(Call, MIRBuilder, GR);

1520case SPIRV::OpAtomicStore:

1521returnbuildAtomicStoreInst(Call, MIRBuilder, GR);

1522case SPIRV::OpAtomicCompareExchange:

1523case SPIRV::OpAtomicCompareExchangeWeak:

1524returnbuildAtomicCompareExchangeInst(Call, Builtin, Opcode, MIRBuilder,

1525 GR);

1526case SPIRV::OpAtomicIAdd:

1527case SPIRV::OpAtomicISub:

1528case SPIRV::OpAtomicOr:

1529case SPIRV::OpAtomicXor:

1530case SPIRV::OpAtomicAnd:

1531case SPIRV::OpAtomicExchange:

1532returnbuildAtomicRMWInst(Call, Opcode, MIRBuilder, GR);

1533case SPIRV::OpMemoryBarrier:

1534returnbuildBarrierInst(Call, SPIRV::OpMemoryBarrier, MIRBuilder, GR);

1535case SPIRV::OpAtomicFlagTestAndSet:

1536case SPIRV::OpAtomicFlagClear:

1537returnbuildAtomicFlagInst(Call, Opcode, MIRBuilder, GR);

1538default:

1539if (Call->isSpirvOp())

1540returnbuildOpFromWrapper(MIRBuilder, Opcode, Call,

1541 GR->getSPIRVTypeID(Call->ReturnType));

1542returnfalse;

1543 }

1544}

1545

1546staticboolgenerateAtomicFloatingInst(constSPIRV::IncomingCall *Call,

1547MachineIRBuilder &MIRBuilder,

1548SPIRVGlobalRegistry *GR) {

1549// Lookup the instruction opcode in the TableGen records.

1550constSPIRV::DemangledBuiltin *Builtin = Call->Builtin;

1551unsigned Opcode = SPIRV::lookupAtomicFloatingBuiltin(Builtin->Name)->Opcode;

1552

1553switch (Opcode) {

1554case SPIRV::OpAtomicFAddEXT:

1555case SPIRV::OpAtomicFMinEXT:

1556case SPIRV::OpAtomicFMaxEXT:

1557returnbuildAtomicFloatingRMWInst(Call, Opcode, MIRBuilder, GR);

1558default:

1559returnfalse;

1560 }

1561}

1562

1563staticboolgenerateBarrierInst(constSPIRV::IncomingCall *Call,

1564MachineIRBuilder &MIRBuilder,

1565SPIRVGlobalRegistry *GR) {

1566// Lookup the instruction opcode in the TableGen records.

1567constSPIRV::DemangledBuiltin *Builtin = Call->Builtin;

1568unsigned Opcode =

1569 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;

1570

1571returnbuildBarrierInst(Call, Opcode, MIRBuilder, GR);

1572}

1573

1574staticboolgenerateCastToPtrInst(constSPIRV::IncomingCall *Call,

1575MachineIRBuilder &MIRBuilder) {

1576 MIRBuilder.buildInstr(TargetOpcode::G_ADDRSPACE_CAST)

1577 .addDef(Call->ReturnRegister)

1578 .addUse(Call->Arguments[0]);

1579returntrue;

1580}

1581

1582staticboolgenerateDotOrFMulInst(constSPIRV::IncomingCall *Call,

1583MachineIRBuilder &MIRBuilder,

1584SPIRVGlobalRegistry *GR) {

1585if (Call->isSpirvOp())

1586returnbuildOpFromWrapper(MIRBuilder, SPIRV::OpDot, Call,

1587 GR->getSPIRVTypeID(Call->ReturnType));

1588unsigned Opcode = GR->getSPIRVTypeForVReg(Call->Arguments[0])->getOpcode();

1589bool IsVec = Opcode == SPIRV::OpTypeVector;

1590// Use OpDot only in case of vector args and OpFMul in case of scalar args.

1591 MIRBuilder.buildInstr(IsVec ? SPIRV::OpDot : SPIRV::OpFMulS)

1592 .addDef(Call->ReturnRegister)

1593 .addUse(GR->getSPIRVTypeID(Call->ReturnType))

1594 .addUse(Call->Arguments[0])

1595 .addUse(Call->Arguments[1]);

1596returntrue;

1597}

1598

1599staticboolgenerateWaveInst(constSPIRV::IncomingCall *Call,

1600MachineIRBuilder &MIRBuilder,

1601SPIRVGlobalRegistry *GR) {

1602constSPIRV::DemangledBuiltin *Builtin = Call->Builtin;

1603 SPIRV::BuiltIn::BuiltInValue =

1604 SPIRV::lookupGetBuiltin(Builtin->Name, Builtin->Set)->Value;

1605

1606// For now, we only support a single Wave intrinsic with a single return type.

1607assert(Call->ReturnType->getOpcode() == SPIRV::OpTypeInt);

1608LLT LLType =LLT::scalar(GR->getScalarOrVectorBitWidth(Call->ReturnType));

1609

1610returnbuildBuiltinVariableLoad(

1611 MIRBuilder, Call->ReturnType, GR,Value, LLType, Call->ReturnRegister,

1612/* isConst= */false,/* hasLinkageTy= */false);

1613}

1614

1615// We expect a builtin

1616// Name(ptr sret([RetType]) %result, Type %operand1, Type %operand1)

1617// where %result is a pointer to where the result of the builtin execution

1618// is to be stored, and generate the following instructions:

1619// Res = Opcode RetType Operand1 Operand1

1620// OpStore RetVariable Res

1621staticboolgenerateICarryBorrowInst(constSPIRV::IncomingCall *Call,

1622MachineIRBuilder &MIRBuilder,

1623SPIRVGlobalRegistry *GR) {

1624constSPIRV::DemangledBuiltin *Builtin = Call->Builtin;

1625unsigned Opcode =

1626 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;

1627

1628Register SRetReg = Call->Arguments[0];

1629SPIRVType *PtrRetType = GR->getSPIRVTypeForVReg(SRetReg);

1630SPIRVType *RetType = GR->getPointeeType(PtrRetType);

1631if (!RetType)

1632report_fatal_error("The first parameter must be a pointer");

1633if (RetType->getOpcode() != SPIRV::OpTypeStruct)

1634report_fatal_error("Expected struct type result for the arithmetic with "

1635"overflow builtins");

1636

1637SPIRVType *OpType1 = GR->getSPIRVTypeForVReg(Call->Arguments[1]);

1638SPIRVType *OpType2 = GR->getSPIRVTypeForVReg(Call->Arguments[2]);

1639if (!OpType1 || !OpType2 || OpType1 != OpType2)

1640report_fatal_error("Operands must have the same type");

1641if (OpType1->getOpcode() == SPIRV::OpTypeVector)

1642switch (Opcode) {

1643case SPIRV::OpIAddCarryS:

1644 Opcode = SPIRV::OpIAddCarryV;

1645break;

1646case SPIRV::OpISubBorrowS:

1647 Opcode = SPIRV::OpISubBorrowV;

1648break;

1649 }

1650

1651MachineRegisterInfo *MRI = MIRBuilder.getMRI();

1652Register ResReg =MRI->createVirtualRegister(&SPIRV::iIDRegClass);

1653if (constTargetRegisterClass *DstRC =

1654MRI->getRegClassOrNull(Call->Arguments[1])) {

1655MRI->setRegClass(ResReg, DstRC);

1656MRI->setType(ResReg,MRI->getType(Call->Arguments[1]));

1657 }else {

1658MRI->setType(ResReg,LLT::scalar(64));

1659 }

1660 GR->assignSPIRVTypeToVReg(RetType, ResReg, MIRBuilder.getMF());

1661 MIRBuilder.buildInstr(Opcode)

1662 .addDef(ResReg)

1663 .addUse(GR->getSPIRVTypeID(RetType))

1664 .addUse(Call->Arguments[1])

1665 .addUse(Call->Arguments[2]);

1666 MIRBuilder.buildInstr(SPIRV::OpStore).addUse(SRetReg).addUse(ResReg);

1667returntrue;

1668}

1669

1670staticboolgenerateGetQueryInst(constSPIRV::IncomingCall *Call,

1671MachineIRBuilder &MIRBuilder,

1672SPIRVGlobalRegistry *GR) {

1673// Lookup the builtin record.

1674 SPIRV::BuiltIn::BuiltInValue =

1675 SPIRV::lookupGetBuiltin(Call->Builtin->Name, Call->Builtin->Set)->Value;

1676uint64_t IsDefault = (Value == SPIRV::BuiltIn::GlobalSize ||

1677Value == SPIRV::BuiltIn::WorkgroupSize ||

1678Value == SPIRV::BuiltIn::EnqueuedWorkgroupSize);

1679returngenWorkgroupQuery(Call, MIRBuilder, GR,Value, IsDefault ? 1 : 0);

1680}

1681

1682staticboolgenerateImageSizeQueryInst(constSPIRV::IncomingCall *Call,

1683MachineIRBuilder &MIRBuilder,

1684SPIRVGlobalRegistry *GR) {

1685// Lookup the image size query component number in the TableGen records.

1686constSPIRV::DemangledBuiltin *Builtin = Call->Builtin;

1687uint32_t Component =

1688 SPIRV::lookupImageQueryBuiltin(Builtin->Name, Builtin->Set)->Component;

1689// Query result may either be a vector or a scalar. If return type is not a

1690// vector, expect only a single size component. Otherwise get the number of

1691// expected components.

1692SPIRVType *RetTy = Call->ReturnType;

1693unsigned NumExpectedRetComponents =RetTy->getOpcode() == SPIRV::OpTypeVector

1694 ?RetTy->getOperand(2).getImm()

1695 : 1;

1696// Get the actual number of query result/size components.

1697SPIRVType *ImgType = GR->getSPIRVTypeForVReg(Call->Arguments[0]);

1698unsigned NumActualRetComponents =getNumSizeComponents(ImgType);

1699Register QueryResult = Call->ReturnRegister;

1700SPIRVType *QueryResultType = Call->ReturnType;

1701if (NumExpectedRetComponents != NumActualRetComponents) {

1702 QueryResult = MIRBuilder.getMRI()->createGenericVirtualRegister(

1703LLT::fixed_vector(NumActualRetComponents, 32));

1704 MIRBuilder.getMRI()->setRegClass(QueryResult, &SPIRV::vIDRegClass);

1705SPIRVType *IntTy = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);

1706 QueryResultType = GR->getOrCreateSPIRVVectorType(

1707 IntTy, NumActualRetComponents, MIRBuilder);

1708 GR->assignSPIRVTypeToVReg(QueryResultType, QueryResult, MIRBuilder.getMF());

1709 }

1710bool IsDimBuf = ImgType->getOperand(2).getImm() == SPIRV::Dim::DIM_Buffer;

1711unsigned Opcode =

1712 IsDimBuf ? SPIRV::OpImageQuerySize : SPIRV::OpImageQuerySizeLod;

1713auto MIB = MIRBuilder.buildInstr(Opcode)

1714 .addDef(QueryResult)

1715 .addUse(GR->getSPIRVTypeID(QueryResultType))

1716 .addUse(Call->Arguments[0]);

1717if (!IsDimBuf)

1718 MIB.addUse(buildConstantIntReg32(0, MIRBuilder, GR));// Lod id.

1719if (NumExpectedRetComponents == NumActualRetComponents)

1720returntrue;

1721if (NumExpectedRetComponents == 1) {

1722// Only 1 component is expected, build OpCompositeExtract instruction.

1723unsigned ExtractedComposite =

1724 Component == 3 ? NumActualRetComponents - 1 : Component;

1725assert(ExtractedComposite < NumActualRetComponents &&

1726"Invalid composite index!");

1727Register TypeReg = GR->getSPIRVTypeID(Call->ReturnType);

1728SPIRVType *NewType =nullptr;

1729if (QueryResultType->getOpcode() == SPIRV::OpTypeVector) {

1730Register NewTypeReg = QueryResultType->getOperand(1).getReg();

1731if (TypeReg != NewTypeReg &&

1732 (NewType = GR->getSPIRVTypeForVReg(NewTypeReg)) !=nullptr)

1733 TypeReg = NewTypeReg;

1734 }

1735 MIRBuilder.buildInstr(SPIRV::OpCompositeExtract)

1736 .addDef(Call->ReturnRegister)

1737 .addUse(TypeReg)

1738 .addUse(QueryResult)

1739 .addImm(ExtractedComposite);

1740if (NewType !=nullptr)

1741insertAssignInstr(Call->ReturnRegister,nullptr, NewType, GR, MIRBuilder,

1742 MIRBuilder.getMF().getRegInfo());

1743 }else {

1744// More than 1 component is expected, fill a new vector.

1745auto MIB = MIRBuilder.buildInstr(SPIRV::OpVectorShuffle)

1746 .addDef(Call->ReturnRegister)

1747 .addUse(GR->getSPIRVTypeID(Call->ReturnType))

1748 .addUse(QueryResult)

1749 .addUse(QueryResult);

1750for (unsigned i = 0; i < NumExpectedRetComponents; ++i)

1751 MIB.addImm(i < NumActualRetComponents ? i : 0xffffffff);

1752 }

1753returntrue;

1754}

1755

1756staticboolgenerateImageMiscQueryInst(constSPIRV::IncomingCall *Call,

1757MachineIRBuilder &MIRBuilder,

1758SPIRVGlobalRegistry *GR) {

1759assert(Call->ReturnType->getOpcode() == SPIRV::OpTypeInt &&

1760"Image samples query result must be of int type!");

1761

1762// Lookup the instruction opcode in the TableGen records.

1763constSPIRV::DemangledBuiltin *Builtin = Call->Builtin;

1764unsigned Opcode =

1765 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;

1766

1767Register Image = Call->Arguments[0];

1768 SPIRV::Dim::Dim ImageDimensionality =static_cast<SPIRV::Dim::Dim>(

1769 GR->getSPIRVTypeForVReg(Image)->getOperand(2).getImm());

1770 (void)ImageDimensionality;

1771

1772switch (Opcode) {

1773case SPIRV::OpImageQuerySamples:

1774assert(ImageDimensionality == SPIRV::Dim::DIM_2D &&

1775"Image must be of 2D dimensionality");

1776break;

1777case SPIRV::OpImageQueryLevels:

1778assert((ImageDimensionality == SPIRV::Dim::DIM_1D ||

1779 ImageDimensionality == SPIRV::Dim::DIM_2D ||

1780 ImageDimensionality == SPIRV::Dim::DIM_3D ||

1781 ImageDimensionality == SPIRV::Dim::DIM_Cube) &&

1782"Image must be of 1D/2D/3D/Cube dimensionality");

1783break;

1784 }

1785

1786 MIRBuilder.buildInstr(Opcode)

1787 .addDef(Call->ReturnRegister)

1788 .addUse(GR->getSPIRVTypeID(Call->ReturnType))

1789 .addUse(Image);

1790returntrue;

1791}

1792

1793// TODO: Move to TableGen.

1794static SPIRV::SamplerAddressingMode::SamplerAddressingMode

1795getSamplerAddressingModeFromBitmask(unsigned Bitmask) {

1796switch (Bitmask & SPIRV::CLK_ADDRESS_MODE_MASK) {

1797case SPIRV::CLK_ADDRESS_CLAMP:

1798return SPIRV::SamplerAddressingMode::Clamp;

1799case SPIRV::CLK_ADDRESS_CLAMP_TO_EDGE:

1800return SPIRV::SamplerAddressingMode::ClampToEdge;

1801case SPIRV::CLK_ADDRESS_REPEAT:

1802return SPIRV::SamplerAddressingMode::Repeat;

1803case SPIRV::CLK_ADDRESS_MIRRORED_REPEAT:

1804return SPIRV::SamplerAddressingMode::RepeatMirrored;

1805case SPIRV::CLK_ADDRESS_NONE:

1806return SPIRV::SamplerAddressingMode::None;

1807default:

1808report_fatal_error("Unknown CL address mode");

1809 }

1810}

1811

1812staticunsignedgetSamplerParamFromBitmask(unsigned Bitmask) {

1813return (Bitmask & SPIRV::CLK_NORMALIZED_COORDS_TRUE) ? 1 : 0;

1814}

1815

1816static SPIRV::SamplerFilterMode::SamplerFilterMode

1817getSamplerFilterModeFromBitmask(unsigned Bitmask) {

1818if (Bitmask & SPIRV::CLK_FILTER_LINEAR)

1819return SPIRV::SamplerFilterMode::Linear;

1820if (Bitmask & SPIRV::CLK_FILTER_NEAREST)

1821return SPIRV::SamplerFilterMode::Nearest;

1822return SPIRV::SamplerFilterMode::Nearest;

1823}

1824

1825staticboolgenerateReadImageInst(constStringRef DemangledCall,

1826constSPIRV::IncomingCall *Call,

1827MachineIRBuilder &MIRBuilder,

1828SPIRVGlobalRegistry *GR) {

1829Register Image = Call->Arguments[0];

1830MachineRegisterInfo *MRI = MIRBuilder.getMRI();

1831bool HasOclSampler = DemangledCall.contains_insensitive("ocl_sampler");

1832bool HasMsaa = DemangledCall.contains_insensitive("msaa");

1833if (HasOclSampler) {

1834Register Sampler = Call->Arguments[1];

1835

1836if (!GR->isScalarOfType(Sampler, SPIRV::OpTypeSampler) &&

1837getDefInstrMaybeConstant(Sampler,MRI)->getOperand(1).isCImm()) {

1838uint64_t SamplerMask =getIConstVal(Sampler,MRI);

1839 Sampler = GR->buildConstantSampler(

1840Register(),getSamplerAddressingModeFromBitmask(SamplerMask),

1841getSamplerParamFromBitmask(SamplerMask),

1842getSamplerFilterModeFromBitmask(SamplerMask), MIRBuilder,

1843 GR->getSPIRVTypeForVReg(Sampler));

1844 }

1845SPIRVType *ImageType = GR->getSPIRVTypeForVReg(Image);

1846SPIRVType *SampledImageType =

1847 GR->getOrCreateOpTypeSampledImage(ImageType, MIRBuilder);

1848Register SampledImage =MRI->createVirtualRegister(&SPIRV::iIDRegClass);

1849

1850 MIRBuilder.buildInstr(SPIRV::OpSampledImage)

1851 .addDef(SampledImage)

1852 .addUse(GR->getSPIRVTypeID(SampledImageType))

1853 .addUse(Image)

1854 .addUse(Sampler);

1855

1856Register Lod = GR->buildConstantFP(APFloat::getZero(APFloat::IEEEsingle()),

1857 MIRBuilder);

1858

1859if (Call->ReturnType->getOpcode() != SPIRV::OpTypeVector) {

1860SPIRVType *TempType =

1861 GR->getOrCreateSPIRVVectorType(Call->ReturnType, 4, MIRBuilder);

1862Register TempRegister =

1863MRI->createGenericVirtualRegister(GR->getRegType(TempType));

1864MRI->setRegClass(TempRegister, GR->getRegClass(TempType));

1865 GR->assignSPIRVTypeToVReg(TempType, TempRegister, MIRBuilder.getMF());

1866 MIRBuilder.buildInstr(SPIRV::OpImageSampleExplicitLod)

1867 .addDef(TempRegister)

1868 .addUse(GR->getSPIRVTypeID(TempType))

1869 .addUse(SampledImage)

1870 .addUse(Call->Arguments[2])// Coordinate.

1871 .addImm(SPIRV::ImageOperand::Lod)

1872 .addUse(Lod);

1873 MIRBuilder.buildInstr(SPIRV::OpCompositeExtract)

1874 .addDef(Call->ReturnRegister)

1875 .addUse(GR->getSPIRVTypeID(Call->ReturnType))

1876 .addUse(TempRegister)

1877 .addImm(0);

1878 }else {

1879 MIRBuilder.buildInstr(SPIRV::OpImageSampleExplicitLod)

1880 .addDef(Call->ReturnRegister)

1881 .addUse(GR->getSPIRVTypeID(Call->ReturnType))

1882 .addUse(SampledImage)

1883 .addUse(Call->Arguments[2])// Coordinate.

1884 .addImm(SPIRV::ImageOperand::Lod)

1885 .addUse(Lod);

1886 }

1887 }elseif (HasMsaa) {

1888 MIRBuilder.buildInstr(SPIRV::OpImageRead)

1889 .addDef(Call->ReturnRegister)

1890 .addUse(GR->getSPIRVTypeID(Call->ReturnType))

1891 .addUse(Image)

1892 .addUse(Call->Arguments[1])// Coordinate.

1893 .addImm(SPIRV::ImageOperand::Sample)

1894 .addUse(Call->Arguments[2]);

1895 }else {

1896 MIRBuilder.buildInstr(SPIRV::OpImageRead)

1897 .addDef(Call->ReturnRegister)

1898 .addUse(GR->getSPIRVTypeID(Call->ReturnType))

1899 .addUse(Image)

1900 .addUse(Call->Arguments[1]);// Coordinate.

1901 }

1902returntrue;

1903}

1904

1905staticboolgenerateWriteImageInst(constSPIRV::IncomingCall *Call,

1906MachineIRBuilder &MIRBuilder,

1907SPIRVGlobalRegistry *GR) {

1908 MIRBuilder.buildInstr(SPIRV::OpImageWrite)

1909 .addUse(Call->Arguments[0])// Image.

1910 .addUse(Call->Arguments[1])// Coordinate.

1911 .addUse(Call->Arguments[2]);// Texel.

1912returntrue;

1913}

1914

1915staticboolgenerateSampleImageInst(constStringRef DemangledCall,

1916constSPIRV::IncomingCall *Call,

1917MachineIRBuilder &MIRBuilder,

1918SPIRVGlobalRegistry *GR) {

1919MachineRegisterInfo *MRI = MIRBuilder.getMRI();

1920if (Call->Builtin->Name.contains_insensitive(

1921"__translate_sampler_initializer")) {

1922// Build sampler literal.

1923uint64_t Bitmask =getIConstVal(Call->Arguments[0],MRI);

1924Register Sampler = GR->buildConstantSampler(

1925 Call->ReturnRegister,getSamplerAddressingModeFromBitmask(Bitmask),

1926getSamplerParamFromBitmask(Bitmask),

1927getSamplerFilterModeFromBitmask(Bitmask), MIRBuilder, Call->ReturnType);

1928return Sampler.isValid();

1929 }elseif (Call->Builtin->Name.contains_insensitive("__spirv_SampledImage")) {

1930// Create OpSampledImage.

1931Register Image = Call->Arguments[0];

1932SPIRVType *ImageType = GR->getSPIRVTypeForVReg(Image);

1933SPIRVType *SampledImageType =

1934 GR->getOrCreateOpTypeSampledImage(ImageType, MIRBuilder);

1935Register SampledImage =

1936 Call->ReturnRegister.isValid()

1937 ? Call->ReturnRegister

1938 :MRI->createVirtualRegister(&SPIRV::iIDRegClass);

1939 MIRBuilder.buildInstr(SPIRV::OpSampledImage)

1940 .addDef(SampledImage)

1941 .addUse(GR->getSPIRVTypeID(SampledImageType))

1942 .addUse(Image)

1943 .addUse(Call->Arguments[1]);// Sampler.

1944returntrue;

1945 }elseif (Call->Builtin->Name.contains_insensitive(

1946"__spirv_ImageSampleExplicitLod")) {

1947// Sample an image using an explicit level of detail.

1948 std::string ReturnType = DemangledCall.str();

1949if (DemangledCall.contains("_R")) {

1950 ReturnType = ReturnType.substr(ReturnType.find("_R") + 2);

1951 ReturnType = ReturnType.substr(0, ReturnType.find('('));

1952 }

1953SPIRVType *Type =

1954 Call->ReturnType

1955 ? Call->ReturnType

1956 : GR->getOrCreateSPIRVTypeByName(ReturnType, MIRBuilder);

1957if (!Type) {

1958 std::string DiagMsg =

1959"Unable to recognize SPIRV type name: " + ReturnType;

1960report_fatal_error(DiagMsg.c_str());

1961 }

1962 MIRBuilder.buildInstr(SPIRV::OpImageSampleExplicitLod)

1963 .addDef(Call->ReturnRegister)

1964 .addUse(GR->getSPIRVTypeID(Type))

1965 .addUse(Call->Arguments[0])// Image.

1966 .addUse(Call->Arguments[1])// Coordinate.

1967 .addImm(SPIRV::ImageOperand::Lod)

1968 .addUse(Call->Arguments[3]);

1969returntrue;

1970 }

1971returnfalse;

1972}

1973

1974staticboolgenerateSelectInst(constSPIRV::IncomingCall *Call,

1975MachineIRBuilder &MIRBuilder) {

1976 MIRBuilder.buildSelect(Call->ReturnRegister, Call->Arguments[0],

1977 Call->Arguments[1], Call->Arguments[2]);

1978returntrue;

1979}

1980

1981staticboolgenerateConstructInst(constSPIRV::IncomingCall *Call,

1982MachineIRBuilder &MIRBuilder,

1983SPIRVGlobalRegistry *GR) {

1984returnbuildOpFromWrapper(MIRBuilder, SPIRV::OpCompositeConstruct, Call,

1985 GR->getSPIRVTypeID(Call->ReturnType));

1986}

1987

1988staticboolgenerateCoopMatrInst(constSPIRV::IncomingCall *Call,

1989MachineIRBuilder &MIRBuilder,

1990SPIRVGlobalRegistry *GR) {

1991constSPIRV::DemangledBuiltin *Builtin = Call->Builtin;

1992unsigned Opcode =

1993 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;

1994bool IsSet = Opcode != SPIRV::OpCooperativeMatrixStoreKHR &&

1995 Opcode != SPIRV::OpCooperativeMatrixStoreCheckedINTEL &&

1996 Opcode != SPIRV::OpCooperativeMatrixPrefetchINTEL;

1997unsigned ArgSz = Call->Arguments.size();

1998unsigned LiteralIdx = 0;

1999switch (Opcode) {

2000// Memory operand is optional and is literal.

2001case SPIRV::OpCooperativeMatrixLoadKHR:

2002 LiteralIdx = ArgSz > 3 ? 3 : 0;

2003break;

2004case SPIRV::OpCooperativeMatrixStoreKHR:

2005 LiteralIdx = ArgSz > 4 ? 4 : 0;

2006break;

2007case SPIRV::OpCooperativeMatrixLoadCheckedINTEL:

2008 LiteralIdx = ArgSz > 7 ? 7 : 0;

2009break;

2010case SPIRV::OpCooperativeMatrixStoreCheckedINTEL:

2011 LiteralIdx = ArgSz > 8 ? 8 : 0;

2012break;

2013// Cooperative Matrix Operands operand is optional and is literal.

2014case SPIRV::OpCooperativeMatrixMulAddKHR:

2015 LiteralIdx = ArgSz > 3 ? 3 : 0;

2016break;

2017 };

2018

2019SmallVector<uint32_t, 1> ImmArgs;

2020MachineRegisterInfo *MRI = MIRBuilder.getMRI();

2021if (Opcode == SPIRV::OpCooperativeMatrixPrefetchINTEL) {

2022constuint32_t CacheLevel =getConstFromIntrinsic(Call->Arguments[3],MRI);

2023auto MIB = MIRBuilder.buildInstr(SPIRV::OpCooperativeMatrixPrefetchINTEL)

2024 .addUse(Call->Arguments[0])// pointer

2025 .addUse(Call->Arguments[1])// rows

2026 .addUse(Call->Arguments[2])// columns

2027 .addImm(CacheLevel)// cache level

2028 .addUse(Call->Arguments[4]);// memory layout

2029if (ArgSz > 5)

2030 MIB.addUse(Call->Arguments[5]);// stride

2031if (ArgSz > 6) {

2032constuint32_t MemOp =getConstFromIntrinsic(Call->Arguments[6],MRI);

2033 MIB.addImm(MemOp);// memory operand

2034 }

2035returntrue;

2036 }

2037if (LiteralIdx > 0)

2038 ImmArgs.push_back(getConstFromIntrinsic(Call->Arguments[LiteralIdx],MRI));

2039Register TypeReg = GR->getSPIRVTypeID(Call->ReturnType);

2040if (Opcode == SPIRV::OpCooperativeMatrixLengthKHR) {

2041SPIRVType *CoopMatrType = GR->getSPIRVTypeForVReg(Call->Arguments[0]);

2042if (!CoopMatrType)

2043report_fatal_error("Can't find a register's type definition");

2044 MIRBuilder.buildInstr(Opcode)

2045 .addDef(Call->ReturnRegister)

2046 .addUse(TypeReg)

2047 .addUse(CoopMatrType->getOperand(0).getReg());

2048returntrue;

2049 }

2050returnbuildOpFromWrapper(MIRBuilder, Opcode, Call,

2051 IsSet ? TypeReg :Register(0), ImmArgs);

2052}

2053

2054staticboolgenerateSpecConstantInst(constSPIRV::IncomingCall *Call,

2055MachineIRBuilder &MIRBuilder,

2056SPIRVGlobalRegistry *GR) {

2057// Lookup the instruction opcode in the TableGen records.

2058constSPIRV::DemangledBuiltin *Builtin = Call->Builtin;

2059unsigned Opcode =

2060 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;

2061constMachineRegisterInfo *MRI = MIRBuilder.getMRI();

2062

2063switch (Opcode) {

2064case SPIRV::OpSpecConstant: {

2065// Build the SpecID decoration.

2066unsigned SpecId =

2067static_cast<unsigned>(getIConstVal(Call->Arguments[0],MRI));

2068buildOpDecorate(Call->ReturnRegister, MIRBuilder, SPIRV::Decoration::SpecId,

2069 {SpecId});

2070// Determine the constant MI.

2071Register ConstRegister = Call->Arguments[1];

2072constMachineInstr *Const =getDefInstrMaybeConstant(ConstRegister,MRI);

2073assert(Const &&

2074 (Const->getOpcode() == TargetOpcode::G_CONSTANT ||

2075 Const->getOpcode() == TargetOpcode::G_FCONSTANT) &&

2076"Argument should be either an int or floating-point constant");

2077// Determine the opcode and built the OpSpec MI.

2078constMachineOperand &ConstOperand = Const->getOperand(1);

2079if (Call->ReturnType->getOpcode() == SPIRV::OpTypeBool) {

2080assert(ConstOperand.isCImm() &&"Int constant operand is expected");

2081 Opcode = ConstOperand.getCImm()->getValue().getZExtValue()

2082 ? SPIRV::OpSpecConstantTrue

2083 : SPIRV::OpSpecConstantFalse;

2084 }

2085auto MIB = MIRBuilder.buildInstr(Opcode)

2086 .addDef(Call->ReturnRegister)

2087 .addUse(GR->getSPIRVTypeID(Call->ReturnType));

2088

2089if (Call->ReturnType->getOpcode() != SPIRV::OpTypeBool) {

2090if (Const->getOpcode() == TargetOpcode::G_CONSTANT)

2091addNumImm(ConstOperand.getCImm()->getValue(), MIB);

2092else

2093addNumImm(ConstOperand.getFPImm()->getValueAPF().bitcastToAPInt(), MIB);

2094 }

2095returntrue;

2096 }

2097case SPIRV::OpSpecConstantComposite: {

2098auto MIB = MIRBuilder.buildInstr(Opcode)

2099 .addDef(Call->ReturnRegister)

2100 .addUse(GR->getSPIRVTypeID(Call->ReturnType));

2101for (unsigned i = 0; i < Call->Arguments.size(); i++)

2102 MIB.addUse(Call->Arguments[i]);

2103returntrue;

2104 }

2105default:

2106returnfalse;

2107 }

2108}

2109

2110staticboolbuildNDRange(constSPIRV::IncomingCall *Call,

2111MachineIRBuilder &MIRBuilder,

2112SPIRVGlobalRegistry *GR) {

2113MachineRegisterInfo *MRI = MIRBuilder.getMRI();

2114SPIRVType *PtrType = GR->getSPIRVTypeForVReg(Call->Arguments[0]);

2115assert(PtrType->getOpcode() == SPIRV::OpTypePointer &&

2116 PtrType->getOperand(2).isReg());

2117Register TypeReg = PtrType->getOperand(2).getReg();

2118SPIRVType *StructType = GR->getSPIRVTypeForVReg(TypeReg);

2119MachineFunction &MF = MIRBuilder.getMF();

2120Register TmpReg =MRI->createVirtualRegister(&SPIRV::iIDRegClass);

2121 GR->assignSPIRVTypeToVReg(StructType, TmpReg, MF);

2122// Skip the first arg, it's the destination pointer. OpBuildNDRange takes

2123// three other arguments, so pass zero constant on absence.

2124unsigned NumArgs = Call->Arguments.size();

2125assert(NumArgs >= 2);

2126Register GlobalWorkSize = Call->Arguments[NumArgs < 4 ? 1 : 2];

2127Register LocalWorkSize =

2128 NumArgs == 2 ?Register(0) : Call->Arguments[NumArgs < 4 ? 2 : 3];

2129Register GlobalWorkOffset = NumArgs <= 3 ?Register(0) : Call->Arguments[1];

2130if (NumArgs < 4) {

2131Register Const;

2132SPIRVType *SpvTy = GR->getSPIRVTypeForVReg(GlobalWorkSize);

2133if (SpvTy->getOpcode() == SPIRV::OpTypePointer) {

2134MachineInstr *DefInstr =MRI->getUniqueVRegDef(GlobalWorkSize);

2135assert(DefInstr &&isSpvIntrinsic(*DefInstr, Intrinsic::spv_gep) &&

2136 DefInstr->getOperand(3).isReg());

2137Register GWSPtr = DefInstr->getOperand(3).getReg();

2138// TODO: Maybe simplify generation of the type of the fields.

2139unsignedSize = Call->Builtin->Name =="ndrange_3D" ? 3 : 2;

2140unsignedBitWidth = GR->getPointerSize() == 64 ? 64 : 32;

2141Type *BaseTy =IntegerType::get(MF.getFunction().getContext(),BitWidth);

2142Type *FieldTy =ArrayType::get(BaseTy,Size);

2143SPIRVType *SpvFieldTy = GR->getOrCreateSPIRVType(FieldTy, MIRBuilder);

2144 GlobalWorkSize =MRI->createVirtualRegister(&SPIRV::iIDRegClass);

2145 GR->assignSPIRVTypeToVReg(SpvFieldTy, GlobalWorkSize, MF);

2146 MIRBuilder.buildInstr(SPIRV::OpLoad)

2147 .addDef(GlobalWorkSize)

2148 .addUse(GR->getSPIRVTypeID(SpvFieldTy))

2149 .addUse(GWSPtr);

2150constSPIRVSubtarget &ST =

2151 cast<SPIRVSubtarget>(MIRBuilder.getMF().getSubtarget());

2152 Const = GR->getOrCreateConstIntArray(0,Size, *MIRBuilder.getInsertPt(),

2153 SpvFieldTy, *ST.getInstrInfo());

2154 }else {

2155 Const = GR->buildConstantInt(0, MIRBuilder, SpvTy);

2156 }

2157if (!LocalWorkSize.isValid())

2158 LocalWorkSize = Const;

2159if (!GlobalWorkOffset.isValid())

2160 GlobalWorkOffset = Const;

2161 }

2162assert(LocalWorkSize.isValid() && GlobalWorkOffset.isValid());

2163 MIRBuilder.buildInstr(SPIRV::OpBuildNDRange)

2164 .addDef(TmpReg)

2165 .addUse(TypeReg)

2166 .addUse(GlobalWorkSize)

2167 .addUse(LocalWorkSize)

2168 .addUse(GlobalWorkOffset);

2169return MIRBuilder.buildInstr(SPIRV::OpStore)

2170 .addUse(Call->Arguments[0])

2171 .addUse(TmpReg);

2172}

2173

2174// TODO: maybe move to the global register.

2175staticSPIRVType *

2176getOrCreateSPIRVDeviceEventPointer(MachineIRBuilder &MIRBuilder,

2177SPIRVGlobalRegistry *GR) {

2178LLVMContext &Context = MIRBuilder.getMF().getFunction().getContext();

2179Type *OpaqueType =StructType::getTypeByName(Context,"spirv.DeviceEvent");

2180if (!OpaqueType)

2181 OpaqueType =StructType::getTypeByName(Context,"opencl.clk_event_t");

2182if (!OpaqueType)

2183 OpaqueType =StructType::create(Context,"spirv.DeviceEvent");

2184unsigned SC0 =storageClassToAddressSpace(SPIRV::StorageClass::Function);

2185unsigned SC1 =storageClassToAddressSpace(SPIRV::StorageClass::Generic);

2186Type *PtrType =PointerType::get(PointerType::get(OpaqueType, SC0), SC1);

2187return GR->getOrCreateSPIRVType(PtrType, MIRBuilder);

2188}

2189

2190staticboolbuildEnqueueKernel(constSPIRV::IncomingCall *Call,

2191MachineIRBuilder &MIRBuilder,

2192SPIRVGlobalRegistry *GR) {

2193MachineRegisterInfo *MRI = MIRBuilder.getMRI();

2194constDataLayout &DL = MIRBuilder.getDataLayout();

2195bool IsSpirvOp = Call->isSpirvOp();

2196bool HasEvents = Call->Builtin->Name.contains("events") || IsSpirvOp;

2197constSPIRVType *Int32Ty = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);

2198

2199// Make vararg instructions before OpEnqueueKernel.

2200// Local sizes arguments: Sizes of block invoke arguments. Clang generates

2201// local size operands as an array, so we need to unpack them.

2202SmallVector<Register, 16> LocalSizes;

2203if (Call->Builtin->Name.contains("_varargs") || IsSpirvOp) {

2204constunsigned LocalSizeArrayIdx = HasEvents ? 9 : 6;

2205Register GepReg = Call->Arguments[LocalSizeArrayIdx];

2206MachineInstr *GepMI =MRI->getUniqueVRegDef(GepReg);

2207assert(isSpvIntrinsic(*GepMI, Intrinsic::spv_gep) &&

2208 GepMI->getOperand(3).isReg());

2209Register ArrayReg = GepMI->getOperand(3).getReg();

2210MachineInstr *ArrayMI =MRI->getUniqueVRegDef(ArrayReg);

2211constType *LocalSizeTy =getMachineInstrType(ArrayMI);

2212assert(LocalSizeTy &&"Local size type is expected");

2213constuint64_t LocalSizeNum =

2214 cast<ArrayType>(LocalSizeTy)->getNumElements();

2215unsigned SC =storageClassToAddressSpace(SPIRV::StorageClass::Generic);

2216constLLT LLType =LLT::pointer(SC, GR->getPointerSize());

2217constSPIRVType *PointerSizeTy = GR->getOrCreateSPIRVPointerType(

2218 Int32Ty, MIRBuilder, SPIRV::StorageClass::Function);

2219for (unsignedI = 0;I < LocalSizeNum; ++I) {

2220Register Reg =MRI->createVirtualRegister(&SPIRV::pIDRegClass);

2221MRI->setType(Reg, LLType);

2222 GR->assignSPIRVTypeToVReg(PointerSizeTy,Reg, MIRBuilder.getMF());

2223auto GEPInst = MIRBuilder.buildIntrinsic(

2224 Intrinsic::spv_gep,ArrayRef<Register>{Reg},true,false);

2225 GEPInst

2226 .addImm(GepMI->getOperand(2).getImm())// In bound.

2227 .addUse(ArrayMI->getOperand(0).getReg())// Alloca.

2228 .addUse(buildConstantIntReg32(0, MIRBuilder, GR))// Indices.

2229 .addUse(buildConstantIntReg32(I, MIRBuilder, GR));

2230 LocalSizes.push_back(Reg);

2231 }

2232 }

2233

2234// SPIRV OpEnqueueKernel instruction has 10+ arguments.

2235auto MIB = MIRBuilder.buildInstr(SPIRV::OpEnqueueKernel)

2236 .addDef(Call->ReturnRegister)

2237 .addUse(GR->getSPIRVTypeID(Int32Ty));

2238

2239// Copy all arguments before block invoke function pointer.

2240constunsigned BlockFIdx = HasEvents ? 6 : 3;

2241for (unsigned i = 0; i < BlockFIdx; i++)

2242 MIB.addUse(Call->Arguments[i]);

2243

2244// If there are no event arguments in the original call, add dummy ones.

2245if (!HasEvents) {

2246 MIB.addUse(buildConstantIntReg32(0, MIRBuilder, GR));// Dummy num events.

2247Register NullPtr = GR->getOrCreateConstNullPtr(

2248 MIRBuilder,getOrCreateSPIRVDeviceEventPointer(MIRBuilder, GR));

2249 MIB.addUse(NullPtr);// Dummy wait events.

2250 MIB.addUse(NullPtr);// Dummy ret event.

2251 }

2252

2253MachineInstr *BlockMI =getBlockStructInstr(Call->Arguments[BlockFIdx],MRI);

2254assert(BlockMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE);

2255// Invoke: Pointer to invoke function.

2256 MIB.addGlobalAddress(BlockMI->getOperand(1).getGlobal());

2257

2258Register BlockLiteralReg = Call->Arguments[BlockFIdx + 1];

2259// Param: Pointer to block literal.

2260 MIB.addUse(BlockLiteralReg);

2261

2262Type *PType =const_cast<Type *>(getBlockStructType(BlockLiteralReg,MRI));

2263// TODO: these numbers should be obtained from block literal structure.

2264// Param Size: Size of block literal structure.

2265 MIB.addUse(buildConstantIntReg32(DL.getTypeStoreSize(PType), MIRBuilder, GR));

2266// Param Aligment: Aligment of block literal structure.

2267 MIB.addUse(buildConstantIntReg32(DL.getPrefTypeAlign(PType).value(),

2268 MIRBuilder, GR));

2269

2270for (unsigned i = 0; i < LocalSizes.size(); i++)

2271 MIB.addUse(LocalSizes[i]);

2272returntrue;

2273}

2274

2275staticboolgenerateEnqueueInst(constSPIRV::IncomingCall *Call,

2276MachineIRBuilder &MIRBuilder,

2277SPIRVGlobalRegistry *GR) {

2278// Lookup the instruction opcode in the TableGen records.

2279constSPIRV::DemangledBuiltin *Builtin = Call->Builtin;

2280unsigned Opcode =

2281 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;

2282

2283switch (Opcode) {

2284case SPIRV::OpRetainEvent:

2285case SPIRV::OpReleaseEvent:

2286return MIRBuilder.buildInstr(Opcode).addUse(Call->Arguments[0]);

2287case SPIRV::OpCreateUserEvent:

2288case SPIRV::OpGetDefaultQueue:

2289return MIRBuilder.buildInstr(Opcode)

2290 .addDef(Call->ReturnRegister)

2291 .addUse(GR->getSPIRVTypeID(Call->ReturnType));

2292case SPIRV::OpIsValidEvent:

2293return MIRBuilder.buildInstr(Opcode)

2294 .addDef(Call->ReturnRegister)

2295 .addUse(GR->getSPIRVTypeID(Call->ReturnType))

2296 .addUse(Call->Arguments[0]);

2297case SPIRV::OpSetUserEventStatus:

2298return MIRBuilder.buildInstr(Opcode)

2299 .addUse(Call->Arguments[0])

2300 .addUse(Call->Arguments[1]);

2301case SPIRV::OpCaptureEventProfilingInfo:

2302return MIRBuilder.buildInstr(Opcode)

2303 .addUse(Call->Arguments[0])

2304 .addUse(Call->Arguments[1])

2305 .addUse(Call->Arguments[2]);

2306case SPIRV::OpBuildNDRange:

2307returnbuildNDRange(Call, MIRBuilder, GR);

2308case SPIRV::OpEnqueueKernel:

2309returnbuildEnqueueKernel(Call, MIRBuilder, GR);

2310default:

2311returnfalse;

2312 }

2313}

2314

2315staticboolgenerateAsyncCopy(constSPIRV::IncomingCall *Call,

2316MachineIRBuilder &MIRBuilder,

2317SPIRVGlobalRegistry *GR) {

2318// Lookup the instruction opcode in the TableGen records.

2319constSPIRV::DemangledBuiltin *Builtin = Call->Builtin;

2320unsigned Opcode =

2321 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;

2322

2323bool IsSet = Opcode == SPIRV::OpGroupAsyncCopy;

2324Register TypeReg = GR->getSPIRVTypeID(Call->ReturnType);

2325if (Call->isSpirvOp())

2326returnbuildOpFromWrapper(MIRBuilder, Opcode, Call,

2327 IsSet ? TypeReg :Register(0));

2328

2329auto Scope =buildConstantIntReg32(SPIRV::Scope::Workgroup, MIRBuilder, GR);

2330

2331switch (Opcode) {

2332case SPIRV::OpGroupAsyncCopy: {

2333SPIRVType *NewType =

2334 Call->ReturnType->getOpcode() == SPIRV::OpTypeEvent

2335 ? nullptr

2336 : GR->getOrCreateSPIRVTypeByName("spirv.Event", MIRBuilder);

2337Register TypeReg = GR->getSPIRVTypeID(NewType ? NewType : Call->ReturnType);

2338unsigned NumArgs = Call->Arguments.size();

2339Register EventReg = Call->Arguments[NumArgs - 1];

2340bool Res = MIRBuilder.buildInstr(Opcode)

2341 .addDef(Call->ReturnRegister)

2342 .addUse(TypeReg)

2343 .addUse(Scope)

2344 .addUse(Call->Arguments[0])

2345 .addUse(Call->Arguments[1])

2346 .addUse(Call->Arguments[2])

2347 .addUse(Call->Arguments.size() > 4

2348 ? Call->Arguments[3]

2349 :buildConstantIntReg32(1, MIRBuilder, GR))

2350 .addUse(EventReg);

2351if (NewType !=nullptr)

2352insertAssignInstr(Call->ReturnRegister,nullptr, NewType, GR, MIRBuilder,

2353 MIRBuilder.getMF().getRegInfo());

2354return Res;

2355 }

2356case SPIRV::OpGroupWaitEvents:

2357return MIRBuilder.buildInstr(Opcode)

2358 .addUse(Scope)

2359 .addUse(Call->Arguments[0])

2360 .addUse(Call->Arguments[1]);

2361default:

2362returnfalse;

2363 }

2364}

2365

2366staticboolgenerateConvertInst(constStringRef DemangledCall,

2367constSPIRV::IncomingCall *Call,

2368MachineIRBuilder &MIRBuilder,

2369SPIRVGlobalRegistry *GR) {

2370// Lookup the conversion builtin in the TableGen records.

2371constSPIRV::ConvertBuiltin *Builtin =

2372 SPIRV::lookupConvertBuiltin(Call->Builtin->Name, Call->Builtin->Set);

2373

2374if (!Builtin && Call->isSpirvOp()) {

2375constSPIRV::DemangledBuiltin *Builtin = Call->Builtin;

2376unsigned Opcode =

2377 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;

2378returnbuildOpFromWrapper(MIRBuilder, Opcode, Call,

2379 GR->getSPIRVTypeID(Call->ReturnType));

2380 }

2381

2382if (Builtin->IsSaturated)

2383buildOpDecorate(Call->ReturnRegister, MIRBuilder,

2384 SPIRV::Decoration::SaturatedConversion, {});

2385if (Builtin->IsRounded)

2386buildOpDecorate(Call->ReturnRegister, MIRBuilder,

2387 SPIRV::Decoration::FPRoundingMode,

2388 {(unsigned)Builtin->RoundingMode});

2389

2390 std::string NeedExtMsg;// no errors if empty

2391bool IsRightComponentsNumber =true;// check if input/output accepts vectors

2392unsigned Opcode = SPIRV::OpNop;

2393if (GR->isScalarOrVectorOfType(Call->Arguments[0], SPIRV::OpTypeInt)) {

2394// Int -> ...

2395if (GR->isScalarOrVectorOfType(Call->ReturnRegister, SPIRV::OpTypeInt)) {

2396// Int -> Int

2397if (Builtin->IsSaturated)

2398 Opcode = Builtin->IsDestinationSigned ? SPIRV::OpSatConvertUToS

2399 : SPIRV::OpSatConvertSToU;

2400else

2401 Opcode = Builtin->IsDestinationSigned ? SPIRV::OpUConvert

2402 : SPIRV::OpSConvert;

2403 }elseif (GR->isScalarOrVectorOfType(Call->ReturnRegister,

2404 SPIRV::OpTypeFloat)) {

2405// Int -> Float

2406if (Builtin->IsBfloat16) {

2407constauto *ST =static_cast<constSPIRVSubtarget *>(

2408 &MIRBuilder.getMF().getSubtarget());

2409if (!ST->canUseExtension(

2410 SPIRV::Extension::SPV_INTEL_bfloat16_conversion))

2411 NeedExtMsg ="SPV_INTEL_bfloat16_conversion";

2412 IsRightComponentsNumber =

2413 GR->getScalarOrVectorComponentCount(Call->Arguments[0]) ==

2414 GR->getScalarOrVectorComponentCount(Call->ReturnRegister);

2415 Opcode = SPIRV::OpConvertBF16ToFINTEL;

2416 }else {

2417bool IsSourceSigned =

2418 DemangledCall[DemangledCall.find_first_of('(') + 1] !='u';

2419 Opcode = IsSourceSigned ? SPIRV::OpConvertSToF : SPIRV::OpConvertUToF;

2420 }

2421 }

2422 }elseif (GR->isScalarOrVectorOfType(Call->Arguments[0],

2423 SPIRV::OpTypeFloat)) {

2424// Float -> ...

2425if (GR->isScalarOrVectorOfType(Call->ReturnRegister, SPIRV::OpTypeInt)) {

2426// Float -> Int

2427if (Builtin->IsBfloat16) {

2428constauto *ST =static_cast<constSPIRVSubtarget *>(

2429 &MIRBuilder.getMF().getSubtarget());

2430if (!ST->canUseExtension(

2431 SPIRV::Extension::SPV_INTEL_bfloat16_conversion))

2432 NeedExtMsg ="SPV_INTEL_bfloat16_conversion";

2433 IsRightComponentsNumber =

2434 GR->getScalarOrVectorComponentCount(Call->Arguments[0]) ==

2435 GR->getScalarOrVectorComponentCount(Call->ReturnRegister);

2436 Opcode = SPIRV::OpConvertFToBF16INTEL;

2437 }else {

2438 Opcode = Builtin->IsDestinationSigned ? SPIRV::OpConvertFToS

2439 : SPIRV::OpConvertFToU;

2440 }

2441 }elseif (GR->isScalarOrVectorOfType(Call->ReturnRegister,

2442 SPIRV::OpTypeFloat)) {

2443// Float -> Float

2444 Opcode = SPIRV::OpFConvert;

2445 }

2446 }

2447

2448if (!NeedExtMsg.empty()) {

2449 std::string DiagMsg = std::string(Builtin->Name) +

2450": the builtin requires the following SPIR-V "

2451"extension: " +

2452 NeedExtMsg;

2453report_fatal_error(DiagMsg.c_str(),false);

2454 }

2455if (!IsRightComponentsNumber) {

2456 std::string DiagMsg =

2457 std::string(Builtin->Name) +

2458": result and argument must have the same number of components";

2459report_fatal_error(DiagMsg.c_str(),false);

2460 }

2461assert(Opcode != SPIRV::OpNop &&

2462"Conversion between the types not implemented!");

2463

2464 MIRBuilder.buildInstr(Opcode)

2465 .addDef(Call->ReturnRegister)

2466 .addUse(GR->getSPIRVTypeID(Call->ReturnType))

2467 .addUse(Call->Arguments[0]);

2468returntrue;

2469}

2470

2471staticboolgenerateVectorLoadStoreInst(constSPIRV::IncomingCall *Call,

2472MachineIRBuilder &MIRBuilder,

2473SPIRVGlobalRegistry *GR) {

2474// Lookup the vector load/store builtin in the TableGen records.

2475constSPIRV::VectorLoadStoreBuiltin *Builtin =

2476 SPIRV::lookupVectorLoadStoreBuiltin(Call->Builtin->Name,

2477 Call->Builtin->Set);

2478// Build extended instruction.

2479auto MIB =

2480 MIRBuilder.buildInstr(SPIRV::OpExtInst)

2481 .addDef(Call->ReturnRegister)

2482 .addUse(GR->getSPIRVTypeID(Call->ReturnType))

2483 .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::OpenCL_std))

2484 .addImm(Builtin->Number);

2485for (autoArgument : Call->Arguments)

2486 MIB.addUse(Argument);

2487if (Builtin->Name.contains("load") && Builtin->ElementCount > 1)

2488 MIB.addImm(Builtin->ElementCount);

2489

2490// Rounding mode should be passed as a last argument in the MI for builtins

2491// like "vstorea_halfn_r".

2492if (Builtin->IsRounded)

2493 MIB.addImm(static_cast<uint32_t>(Builtin->RoundingMode));

2494returntrue;

2495}

2496

2497staticboolgenerateLoadStoreInst(constSPIRV::IncomingCall *Call,

2498MachineIRBuilder &MIRBuilder,

2499SPIRVGlobalRegistry *GR) {

2500// Lookup the instruction opcode in the TableGen records.

2501constSPIRV::DemangledBuiltin *Builtin = Call->Builtin;

2502unsigned Opcode =

2503 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;

2504bool IsLoad = Opcode == SPIRV::OpLoad;

2505// Build the instruction.

2506auto MIB = MIRBuilder.buildInstr(Opcode);

2507if (IsLoad) {

2508 MIB.addDef(Call->ReturnRegister);

2509 MIB.addUse(GR->getSPIRVTypeID(Call->ReturnType));

2510 }

2511// Add a pointer to the value to load/store.

2512 MIB.addUse(Call->Arguments[0]);

2513MachineRegisterInfo *MRI = MIRBuilder.getMRI();

2514// Add a value to store.

2515if (!IsLoad)

2516 MIB.addUse(Call->Arguments[1]);

2517// Add optional memory attributes and an alignment.

2518unsigned NumArgs = Call->Arguments.size();

2519if ((IsLoad && NumArgs >= 2) || NumArgs >= 3)

2520 MIB.addImm(getConstFromIntrinsic(Call->Arguments[IsLoad ? 1 : 2],MRI));

2521if ((IsLoad && NumArgs >= 3) || NumArgs >= 4)

2522 MIB.addImm(getConstFromIntrinsic(Call->Arguments[IsLoad ? 2 : 3],MRI));

2523returntrue;

2524}

2525

2526namespaceSPIRV {

2527// Try to find a builtin function attributes by a demangled function name and

2528// return a tuple <builtin group, op code, ext instruction number>, or a special

2529// tuple value <-1, 0, 0> if the builtin function is not found.

2530// Not all builtin functions are supported, only those with a ready-to-use op

2531// code or instruction number defined in TableGen.

2532// TODO: consider a major rework of mapping demangled calls into a builtin

2533// functions to unify search and decrease number of individual cases.

2534std::tuple<int, unsigned, unsigned>

2535mapBuiltinToOpcode(constStringRef DemangledCall,

2536 SPIRV::InstructionSet::InstructionSet Set) {

2537Register Reg;

2538SmallVector<Register> Args;

2539 std::unique_ptr<const IncomingCall> Call =

2540lookupBuiltin(DemangledCall, Set,Reg,nullptr, Args);

2541if (!Call)

2542return std::make_tuple(-1, 0, 0);

2543

2544switch (Call->Builtin->Group) {

2545case SPIRV::Relational:

2546case SPIRV::Atomic:

2547case SPIRV::Barrier:

2548case SPIRV::CastToPtr:

2549case SPIRV::ImageMiscQuery:

2550case SPIRV::SpecConstant:

2551case SPIRV::Enqueue:

2552case SPIRV::AsyncCopy:

2553case SPIRV::LoadStore:

2554case SPIRV::CoopMatr:

2555if (constauto *R =

2556 SPIRV::lookupNativeBuiltin(Call->Builtin->Name, Call->Builtin->Set))

2557return std::make_tuple(Call->Builtin->Group, R->Opcode, 0);

2558break;

2559case SPIRV::Extended:

2560if (constauto *R = SPIRV::lookupExtendedBuiltin(Call->Builtin->Name,

2561 Call->Builtin->Set))

2562return std::make_tuple(Call->Builtin->Group, 0, R->Number);

2563break;

2564case SPIRV::VectorLoadStore:

2565if (constauto *R = SPIRV::lookupVectorLoadStoreBuiltin(Call->Builtin->Name,

2566 Call->Builtin->Set))

2567return std::make_tuple(SPIRV::Extended, 0, R->Number);

2568break;

2569case SPIRV::Group:

2570if (constauto *R = SPIRV::lookupGroupBuiltin(Call->Builtin->Name))

2571return std::make_tuple(Call->Builtin->Group, R->Opcode, 0);

2572break;

2573case SPIRV::AtomicFloating:

2574if (constauto *R = SPIRV::lookupAtomicFloatingBuiltin(Call->Builtin->Name))

2575return std::make_tuple(Call->Builtin->Group, R->Opcode, 0);

2576break;

2577case SPIRV::IntelSubgroups:

2578if (constauto *R = SPIRV::lookupIntelSubgroupsBuiltin(Call->Builtin->Name))

2579return std::make_tuple(Call->Builtin->Group, R->Opcode, 0);

2580break;

2581case SPIRV::GroupUniform:

2582if (constauto *R = SPIRV::lookupGroupUniformBuiltin(Call->Builtin->Name))

2583return std::make_tuple(Call->Builtin->Group, R->Opcode, 0);

2584break;

2585case SPIRV::WriteImage:

2586return std::make_tuple(Call->Builtin->Group, SPIRV::OpImageWrite, 0);

2587case SPIRV::Select:

2588return std::make_tuple(Call->Builtin->Group, TargetOpcode::G_SELECT, 0);

2589case SPIRV::Construct:

2590return std::make_tuple(Call->Builtin->Group, SPIRV::OpCompositeConstruct,

2591 0);

2592case SPIRV::KernelClock:

2593return std::make_tuple(Call->Builtin->Group, SPIRV::OpReadClockKHR, 0);

2594default:

2595return std::make_tuple(-1, 0, 0);

2596 }

2597return std::make_tuple(-1, 0, 0);

2598}

2599

2600std::optional<bool>lowerBuiltin(constStringRef DemangledCall,

2601 SPIRV::InstructionSet::InstructionSet Set,

2602MachineIRBuilder &MIRBuilder,

2603constRegister OrigRet,constType *OrigRetTy,

2604constSmallVectorImpl<Register> &Args,

2605SPIRVGlobalRegistry *GR) {

2606LLVM_DEBUG(dbgs() <<"Lowering builtin call: " << DemangledCall <<"\n");

2607

2608// Lookup the builtin in the TableGen records.

2609SPIRVType *SpvType = GR->getSPIRVTypeForVReg(OrigRet);

2610assert(SpvType &&"Inconsistent return register: expected valid type info");

2611 std::unique_ptr<const IncomingCall> Call =

2612lookupBuiltin(DemangledCall, Set, OrigRet, SpvType, Args);

2613

2614if (!Call) {

2615LLVM_DEBUG(dbgs() <<"Builtin record was not found!\n");

2616return std::nullopt;

2617 }

2618

2619// TODO: check if the provided args meet the builtin requirments.

2620assert(Args.size() >= Call->Builtin->MinNumArgs &&

2621"Too few arguments to generate the builtin");

2622if (Call->Builtin->MaxNumArgs && Args.size() > Call->Builtin->MaxNumArgs)

2623LLVM_DEBUG(dbgs() <<"More arguments provided than required!\n");

2624

2625// Match the builtin with implementation based on the grouping.

2626switch (Call->Builtin->Group) {

2627case SPIRV::Extended:

2628returngenerateExtInst(Call.get(), MIRBuilder, GR);

2629case SPIRV::Relational:

2630returngenerateRelationalInst(Call.get(), MIRBuilder, GR);

2631case SPIRV::Group:

2632returngenerateGroupInst(Call.get(), MIRBuilder, GR);

2633case SPIRV::Variable:

2634returngenerateBuiltinVar(Call.get(), MIRBuilder, GR);

2635case SPIRV::Atomic:

2636returngenerateAtomicInst(Call.get(), MIRBuilder, GR);

2637case SPIRV::AtomicFloating:

2638returngenerateAtomicFloatingInst(Call.get(), MIRBuilder, GR);

2639case SPIRV::Barrier:

2640returngenerateBarrierInst(Call.get(), MIRBuilder, GR);

2641case SPIRV::CastToPtr:

2642returngenerateCastToPtrInst(Call.get(), MIRBuilder);

2643case SPIRV::Dot:

2644returngenerateDotOrFMulInst(Call.get(), MIRBuilder, GR);

2645case SPIRV::Wave:

2646returngenerateWaveInst(Call.get(), MIRBuilder, GR);

2647case SPIRV::ICarryBorrow:

2648returngenerateICarryBorrowInst(Call.get(), MIRBuilder, GR);

2649case SPIRV::GetQuery:

2650returngenerateGetQueryInst(Call.get(), MIRBuilder, GR);

2651case SPIRV::ImageSizeQuery:

2652returngenerateImageSizeQueryInst(Call.get(), MIRBuilder, GR);

2653case SPIRV::ImageMiscQuery:

2654returngenerateImageMiscQueryInst(Call.get(), MIRBuilder, GR);

2655case SPIRV::ReadImage:

2656returngenerateReadImageInst(DemangledCall, Call.get(), MIRBuilder, GR);

2657case SPIRV::WriteImage:

2658returngenerateWriteImageInst(Call.get(), MIRBuilder, GR);

2659case SPIRV::SampleImage:

2660returngenerateSampleImageInst(DemangledCall, Call.get(), MIRBuilder, GR);

2661case SPIRV::Select:

2662returngenerateSelectInst(Call.get(), MIRBuilder);

2663case SPIRV::Construct:

2664returngenerateConstructInst(Call.get(), MIRBuilder, GR);

2665case SPIRV::SpecConstant:

2666returngenerateSpecConstantInst(Call.get(), MIRBuilder, GR);

2667case SPIRV::Enqueue:

2668returngenerateEnqueueInst(Call.get(), MIRBuilder, GR);

2669case SPIRV::AsyncCopy:

2670returngenerateAsyncCopy(Call.get(), MIRBuilder, GR);

2671case SPIRV::Convert:

2672returngenerateConvertInst(DemangledCall, Call.get(), MIRBuilder, GR);

2673case SPIRV::VectorLoadStore:

2674returngenerateVectorLoadStoreInst(Call.get(), MIRBuilder, GR);

2675case SPIRV::LoadStore:

2676returngenerateLoadStoreInst(Call.get(), MIRBuilder, GR);

2677case SPIRV::IntelSubgroups:

2678returngenerateIntelSubgroupsInst(Call.get(), MIRBuilder, GR);

2679case SPIRV::GroupUniform:

2680returngenerateGroupUniformInst(Call.get(), MIRBuilder, GR);

2681case SPIRV::KernelClock:

2682returngenerateKernelClockInst(Call.get(), MIRBuilder, GR);

2683case SPIRV::CoopMatr:

2684returngenerateCoopMatrInst(Call.get(), MIRBuilder, GR);

2685 }

2686returnfalse;

2687}

2688

2689Type *parseBuiltinCallArgumentType(StringRef TypeStr,LLVMContext &Ctx) {

2690// Parse strings representing OpenCL builtin types.

2691if (hasBuiltinTypePrefix(TypeStr)) {

2692// OpenCL builtin types in demangled call strings have the following format:

2693// e.g. ocl_image2d_ro

2694 [[maybe_unused]]bool IsOCLBuiltinType = TypeStr.consume_front("ocl_");

2695assert(IsOCLBuiltinType &&"Invalid OpenCL builtin prefix");

2696

2697// Check if this is pointer to a builtin type and not just pointer

2698// representing a builtin type. In case it is a pointer to builtin type,

2699// this will require additional handling in the method calling

2700// parseBuiltinCallArgumentBaseType(...) as this function only retrieves the

2701// base types.

2702if (TypeStr.ends_with("*"))

2703 TypeStr = TypeStr.slice(0, TypeStr.find_first_of(" *"));

2704

2705returnparseBuiltinTypeNameToTargetExtType("opencl." + TypeStr.str() +"_t",

2706 Ctx);

2707 }

2708

2709// Parse type name in either "typeN" or "type vector[N]" format, where

2710// N is the number of elements of the vector.

2711Type *BaseType;

2712unsigned VecElts = 0;

2713

2714BaseType =parseBasicTypeName(TypeStr, Ctx);

2715if (!BaseType)

2716// Unable to recognize SPIRV type name.

2717returnnullptr;

2718

2719// Handle "typeN*" or "type vector[N]*".

2720 TypeStr.consume_back("*");

2721

2722if (TypeStr.consume_front(" vector["))

2723 TypeStr = TypeStr.substr(0, TypeStr.find(']'));

2724

2725 TypeStr.getAsInteger(10, VecElts);

2726if (VecElts > 0)

2727BaseType =VectorType::get(

2728BaseType->isVoidTy() ?Type::getInt8Ty(Ctx) :BaseType, VecElts,false);

2729

2730returnBaseType;

2731}

2732

2733boolparseBuiltinTypeStr(SmallVector<StringRef, 10> &BuiltinArgsTypeStrs,

2734constStringRef DemangledCall,LLVMContext &Ctx) {

2735auto Pos1 = DemangledCall.find('(');

2736if (Pos1 ==StringRef::npos)

2737returnfalse;

2738auto Pos2 = DemangledCall.find(')');

2739if (Pos2 ==StringRef::npos || Pos1 > Pos2)

2740returnfalse;

2741 DemangledCall.slice(Pos1 + 1, Pos2)

2742 .split(BuiltinArgsTypeStrs,',', -1,false);

2743returntrue;

2744}

2745

2746Type *parseBuiltinCallArgumentBaseType(constStringRef DemangledCall,

2747unsigned ArgIdx,LLVMContext &Ctx) {

2748SmallVector<StringRef, 10> BuiltinArgsTypeStrs;

2749parseBuiltinTypeStr(BuiltinArgsTypeStrs, DemangledCall, Ctx);

2750if (ArgIdx >= BuiltinArgsTypeStrs.size())

2751returnnullptr;

2752StringRef TypeStr = BuiltinArgsTypeStrs[ArgIdx].trim();

2753returnparseBuiltinCallArgumentType(TypeStr, Ctx);

2754}

2755

2756structBuiltinType {

2757StringRef Name;

2758uint32_t Opcode;

2759};

2760

2761#define GET_BuiltinTypes_DECL

2762#define GET_BuiltinTypes_IMPL

2763

2764structOpenCLType {

2765StringRef Name;

2766StringRef SpirvTypeLiteral;

2767};

2768

2769#define GET_OpenCLTypes_DECL

2770#define GET_OpenCLTypes_IMPL

2771

2772#include "SPIRVGenTables.inc"

2773}// namespace SPIRV

2774

2775//===----------------------------------------------------------------------===//

2776// Misc functions for parsing builtin types.

2777//===----------------------------------------------------------------------===//

2778

2779staticType *parseTypeString(constStringRef Name,LLVMContext &Context) {

2780if (Name.starts_with("void"))

2781returnType::getVoidTy(Context);

2782elseif (Name.starts_with("int") ||Name.starts_with("uint"))

2783returnType::getInt32Ty(Context);

2784elseif (Name.starts_with("float"))

2785returnType::getFloatTy(Context);

2786elseif (Name.starts_with("half"))

2787returnType::getHalfTy(Context);

2788report_fatal_error("Unable to recognize type!");

2789}

2790

2791//===----------------------------------------------------------------------===//

2792// Implementation functions for builtin types.

2793//===----------------------------------------------------------------------===//

2794

2795staticSPIRVType *getNonParameterizedType(constTargetExtType *ExtensionType,

2796constSPIRV::BuiltinType *TypeRecord,

2797MachineIRBuilder &MIRBuilder,

2798SPIRVGlobalRegistry *GR) {

2799unsigned Opcode = TypeRecord->Opcode;

2800// Create or get an existing type from GlobalRegistry.

2801return GR->getOrCreateOpTypeByOpcode(ExtensionType, MIRBuilder, Opcode);

2802}

2803

2804staticSPIRVType *getSamplerType(MachineIRBuilder &MIRBuilder,

2805SPIRVGlobalRegistry *GR) {

2806// Create or get an existing type from GlobalRegistry.

2807return GR->getOrCreateOpTypeSampler(MIRBuilder);

2808}

2809

2810staticSPIRVType *getPipeType(constTargetExtType *ExtensionType,

2811MachineIRBuilder &MIRBuilder,

2812SPIRVGlobalRegistry *GR) {

2813assert(ExtensionType->getNumIntParameters() == 1 &&

2814"Invalid number of parameters for SPIR-V pipe builtin!");

2815// Create or get an existing type from GlobalRegistry.

2816return GR->getOrCreateOpTypePipe(MIRBuilder,

2817 SPIRV::AccessQualifier::AccessQualifier(

2818 ExtensionType->getIntParameter(0)));

2819}

2820

2821staticSPIRVType *getCoopMatrType(constTargetExtType *ExtensionType,

2822MachineIRBuilder &MIRBuilder,

2823SPIRVGlobalRegistry *GR) {

2824assert(ExtensionType->getNumIntParameters() == 4 &&

2825"Invalid number of parameters for SPIR-V coop matrices builtin!");

2826assert(ExtensionType->getNumTypeParameters() == 1 &&

2827"SPIR-V coop matrices builtin type must have a type parameter!");

2828constSPIRVType *ElemType =

2829 GR->getOrCreateSPIRVType(ExtensionType->getTypeParameter(0), MIRBuilder);

2830// Create or get an existing type from GlobalRegistry.

2831return GR->getOrCreateOpTypeCoopMatr(

2832 MIRBuilder, ExtensionType, ElemType, ExtensionType->getIntParameter(0),

2833 ExtensionType->getIntParameter(1), ExtensionType->getIntParameter(2),

2834 ExtensionType->getIntParameter(3));

2835}

2836

2837staticSPIRVType *

2838getImageType(constTargetExtType *ExtensionType,

2839const SPIRV::AccessQualifier::AccessQualifier Qualifier,

2840MachineIRBuilder &MIRBuilder,SPIRVGlobalRegistry *GR) {

2841assert(ExtensionType->getNumTypeParameters() == 1 &&

2842"SPIR-V image builtin type must have sampled type parameter!");

2843constSPIRVType *SampledType =

2844 GR->getOrCreateSPIRVType(ExtensionType->getTypeParameter(0), MIRBuilder);

2845assert((ExtensionType->getNumIntParameters() == 7 ||

2846 ExtensionType->getNumIntParameters() == 6) &&

2847"Invalid number of parameters for SPIR-V image builtin!");

2848

2849 SPIRV::AccessQualifier::AccessQualifier accessQualifier =

2850 SPIRV::AccessQualifier::None;

2851if (ExtensionType->getNumIntParameters() == 7) {

2852 accessQualifier = Qualifier == SPIRV::AccessQualifier::WriteOnly

2853 ? SPIRV::AccessQualifier::WriteOnly

2854 : SPIRV::AccessQualifier::AccessQualifier(

2855 ExtensionType->getIntParameter(6));

2856 }

2857

2858// Create or get an existing type from GlobalRegistry.

2859return GR->getOrCreateOpTypeImage(

2860 MIRBuilder, SampledType,

2861 SPIRV::Dim::Dim(ExtensionType->getIntParameter(0)),

2862 ExtensionType->getIntParameter(1), ExtensionType->getIntParameter(2),

2863 ExtensionType->getIntParameter(3), ExtensionType->getIntParameter(4),

2864 SPIRV::ImageFormat::ImageFormat(ExtensionType->getIntParameter(5)),

2865 accessQualifier);

2866}

2867

2868staticSPIRVType *getSampledImageType(constTargetExtType *OpaqueType,

2869MachineIRBuilder &MIRBuilder,

2870SPIRVGlobalRegistry *GR) {

2871SPIRVType *OpaqueImageType =getImageType(

2872 OpaqueType, SPIRV::AccessQualifier::ReadOnly, MIRBuilder, GR);

2873// Create or get an existing type from GlobalRegistry.

2874return GR->getOrCreateOpTypeSampledImage(OpaqueImageType, MIRBuilder);

2875}

2876

2877namespaceSPIRV {

2878TargetExtType *parseBuiltinTypeNameToTargetExtType(std::string TypeName,

2879LLVMContext &Context) {

2880StringRef NameWithParameters = TypeName;

2881

2882// Pointers-to-opaque-structs representing OpenCL types are first translated

2883// to equivalent SPIR-V types. OpenCL builtin type names should have the

2884// following format: e.g. %opencl.event_t

2885if (NameWithParameters.starts_with("opencl.")) {

2886constSPIRV::OpenCLType *OCLTypeRecord =

2887 SPIRV::lookupOpenCLType(NameWithParameters);

2888if (!OCLTypeRecord)

2889report_fatal_error("Missing TableGen record for OpenCL type: " +

2890 NameWithParameters);

2891 NameWithParameters = OCLTypeRecord->SpirvTypeLiteral;

2892// Continue with the SPIR-V builtin type...

2893 }

2894

2895// Names of the opaque structs representing a SPIR-V builtins without

2896// parameters should have the following format: e.g. %spirv.Event

2897assert(NameWithParameters.starts_with("spirv.") &&

2898"Unknown builtin opaque type!");

2899

2900// Parameterized SPIR-V builtins names follow this format:

2901// e.g. %spirv.Image._void_1_0_0_0_0_0_0, %spirv.Pipe._0

2902if (!NameWithParameters.contains('_'))

2903returnTargetExtType::get(Context, NameWithParameters);

2904

2905SmallVector<StringRef> Parameters;

2906unsigned BaseNameLength = NameWithParameters.find('_') - 1;

2907 SplitString(NameWithParameters.substr(BaseNameLength + 1), Parameters,"_");

2908

2909SmallVector<Type *, 1> TypeParameters;

2910bool HasTypeParameter = !isDigit(Parameters[0][0]);

2911if (HasTypeParameter)

2912 TypeParameters.push_back(parseTypeString(Parameters[0], Context));

2913SmallVector<unsigned> IntParameters;

2914for (unsigned i = HasTypeParameter ? 1 : 0; i < Parameters.size(); i++) {

2915unsigned IntParameter = 0;

2916bool ValidLiteral = !Parameters[i].getAsInteger(10, IntParameter);

2917 (void)ValidLiteral;

2918assert(ValidLiteral &&

2919"Invalid format of SPIR-V builtin parameter literal!");

2920 IntParameters.push_back(IntParameter);

2921 }

2922returnTargetExtType::get(Context,

2923 NameWithParameters.substr(0, BaseNameLength),

2924 TypeParameters, IntParameters);

2925}

2926

2927SPIRVType *lowerBuiltinType(constType *OpaqueType,

2928 SPIRV::AccessQualifier::AccessQualifier AccessQual,

2929MachineIRBuilder &MIRBuilder,

2930SPIRVGlobalRegistry *GR) {

2931// In LLVM IR, SPIR-V and OpenCL builtin types are represented as either

2932// target(...) target extension types or pointers-to-opaque-structs. The

2933// approach relying on structs is deprecated and works only in the non-opaque

2934// pointer mode (-opaque-pointers=0).

2935// In order to maintain compatibility with LLVM IR generated by older versions

2936// of Clang and LLVM/SPIR-V Translator, the pointers-to-opaque-structs are

2937// "translated" to target extension types. This translation is temporary and

2938// will be removed in the future release of LLVM.

2939constTargetExtType *BuiltinType = dyn_cast<TargetExtType>(OpaqueType);

2940if (!BuiltinType)

2941BuiltinType =parseBuiltinTypeNameToTargetExtType(

2942 OpaqueType->getStructName().str(), MIRBuilder.getContext());

2943

2944unsigned NumStartingVRegs = MIRBuilder.getMRI()->getNumVirtRegs();

2945

2946constStringRef Name =BuiltinType->getName();

2947LLVM_DEBUG(dbgs() <<"Lowering builtin type: " <<Name <<"\n");

2948

2949// Lookup the demangled builtin type in the TableGen records.

2950constSPIRV::BuiltinType *TypeRecord = SPIRV::lookupBuiltinType(Name);

2951if (!TypeRecord)

2952report_fatal_error("Missing TableGen record for builtin type: " +Name);

2953

2954// "Lower" the BuiltinType into TargetType. The following get<...>Type methods

2955// use the implementation details from TableGen records or TargetExtType

2956// parameters to either create a new OpType<...> machine instruction or get an

2957// existing equivalent SPIRVType from GlobalRegistry.

2958SPIRVType *TargetType;

2959switch (TypeRecord->Opcode) {

2960case SPIRV::OpTypeImage:

2961 TargetType =getImageType(BuiltinType, AccessQual, MIRBuilder, GR);

2962break;

2963case SPIRV::OpTypePipe:

2964 TargetType =getPipeType(BuiltinType, MIRBuilder, GR);

2965break;

2966case SPIRV::OpTypeDeviceEvent:

2967 TargetType = GR->getOrCreateOpTypeDeviceEvent(MIRBuilder);

2968break;

2969case SPIRV::OpTypeSampler:

2970 TargetType =getSamplerType(MIRBuilder, GR);

2971break;

2972case SPIRV::OpTypeSampledImage:

2973 TargetType =getSampledImageType(BuiltinType, MIRBuilder, GR);

2974break;

2975case SPIRV::OpTypeCooperativeMatrixKHR:

2976 TargetType =getCoopMatrType(BuiltinType, MIRBuilder, GR);

2977break;

2978default:

2979 TargetType =

2980getNonParameterizedType(BuiltinType, TypeRecord, MIRBuilder, GR);

2981break;

2982 }

2983

2984// Emit OpName instruction if a new OpType<...> instruction was added

2985// (equivalent type was not found in GlobalRegistry).

2986if (NumStartingVRegs < MIRBuilder.getMRI()->getNumVirtRegs())

2987buildOpName(GR->getSPIRVTypeID(TargetType),Name, MIRBuilder);

2988

2989return TargetType;

2990}

2991}// namespace SPIRV

2992}// namespace llvm

MRI

unsigned const MachineRegisterInfo * MRI

Definition:AArch64AdvSIMDScalarPass.cpp:105

DefMI

MachineInstrBuilder MachineInstrBuilder & DefMI

Definition:AArch64ExpandPseudoInsts.cpp:113

Arguments

AMDGPU Lower Kernel Arguments

Definition:AMDGPULowerKernelArguments.cpp:504

MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL

Definition:ARMSLSHardening.cpp:73

RetTy

return RetTy

Definition:DeadArgumentElimination.cpp:361

LLVM_DEBUG

#define LLVM_DEBUG(...)

Definition:Debug.h:106

Name

std::string Name

Definition:ELFObjHandler.cpp:77

Size

uint64_t Size

Definition:ELFObjHandler.cpp:81

IRTranslator LLVM IR MI

Definition:IRTranslator.cpp:112

#define I(x, y, z)

Definition:MD5.cpp:58

Reg

unsigned Reg

Definition:MachineSink.cpp:2028

isDigit

static bool isDigit(const char C)

Definition:RustDemangle.cpp:170

assert

assert(ImpDefSCC.getReg()==AMDGPU::SCC &&ImpDefSCC.isDef())

SPIRVBuiltins.h

SPIRV

spirv structurize SPIRV

Definition:SPIRVStructurizer.cpp:1243

A given derived pointer can have multiple base pointers through phi/selects.

Definition:SafepointIRVerifier.cpp:316

StringExtras.h

This file contains some functions that are useful when dealing with strings.

ValueTracking.h

BaseTy

BaseType

llvm::APFloat::bitcastToAPInt

APInt bitcastToAPInt() const

Definition:APFloat.h:1351

llvm::APFloat::getZero

static APFloat getZero(const fltSemantics &Sem, bool Negative=false)

Factory for Positive and Negative Zero.

Definition:APFloat.h:1081

llvm::APInt::getAllOnes

static APInt getAllOnes(unsigned numBits)

Return an APInt of a specified width with all bits set.

Definition:APInt.h:234

llvm::APInt::getZExtValue

uint64_t getZExtValue() const

Get zero extended value.

Definition:APInt.h:1520

llvm::Argument

This class represents an incoming formal argument to a Function.

Definition:Argument.h:31

llvm::ArrayRef

ArrayRef - Represent a constant reference to an array (0 or more elements consecutively in memory),...

Definition:ArrayRef.h:41

llvm::ArrayType::get

static ArrayType * get(Type *ElementType, uint64_t NumElements)

This static method is the primary way to construct an ArrayType.

llvm::CmpInst::ICMP_ULT

@ ICMP_ULT

unsigned less than

Definition:InstrTypes.h:698

llvm::CmpInst::ICMP_EQ

@ ICMP_EQ

equal

Definition:InstrTypes.h:694

llvm::CmpInst::ICMP_NE

@ ICMP_NE

not equal

Definition:InstrTypes.h:695

llvm::ConstantFP::getValueAPF

const APFloat & getValueAPF() const

Definition:Constants.h:314

llvm::ConstantInt::getValue

const APInt & getValue() const

Return the constant as an APInt value reference.

Definition:Constants.h:148

llvm::DataLayout

A parsed version of the target data layout string in and methods for querying it.

Definition:DataLayout.h:63

llvm::Expected

Tagged union holding either a T or a Error.

Definition:Error.h:481

llvm::FixedVectorType

Class to represent fixed width SIMD vectors.

Definition:DerivedTypes.h:563

llvm::Function::getContext

LLVMContext & getContext() const

getContext - Return a reference to the LLVMContext associated with this function.

Definition:Function.cpp:369

llvm::IntegerType::get

static IntegerType * get(LLVMContext &C, unsigned NumBits)

This static method is the primary way of constructing an IntegerType.

Definition:Type.cpp:311

llvm::LLT

Definition:LowLevelType.h:39

llvm::LLT::vector

static constexpr LLT vector(ElementCount EC, unsigned ScalarSizeInBits)

Get a low-level vector of some number of elements and element width.

Definition:LowLevelType.h:64

llvm::LLT::scalar

static constexpr LLT scalar(unsigned SizeInBits)

Get a low-level scalar or aggregate "bag of bits".

Definition:LowLevelType.h:42

llvm::LLT::pointer

static constexpr LLT pointer(unsigned AddressSpace, unsigned SizeInBits)

Get a low-level pointer in the given address space.

Definition:LowLevelType.h:57

llvm::LLT::fixed_vector

static constexpr LLT fixed_vector(unsigned NumElements, unsigned ScalarSizeInBits)

Get a low-level fixed-width vector of some number of elements and element width.

Definition:LowLevelType.h:100

llvm::LLVMContext

This is an important class for using LLVM in a threaded context.

Definition:LLVMContext.h:67

llvm::MachineFunction

Definition:MachineFunction.h:267

llvm::MachineFunction::getSubtarget

const TargetSubtargetInfo & getSubtarget() const

getSubtarget - Return the subtarget for which this machine code is being compiled.

Definition:MachineFunction.h:733

llvm::MachineFunction::getRegInfo

MachineRegisterInfo & getRegInfo()

getRegInfo - Return information about the registers currently in use.

Definition:MachineFunction.h:743

llvm::MachineFunction::getFunction

Function & getFunction()

Return the LLVM function that this machine code represents.

Definition:MachineFunction.h:704

llvm::MachineIRBuilder

Helper class to build MachineInstr.

Definition:MachineIRBuilder.h:235

llvm::MachineIRBuilder::getContext

LLVMContext & getContext() const

Definition:MachineIRBuilder.h:301

llvm::MachineIRBuilder::buildSelect

MachineInstrBuilder buildSelect(const DstOp &Res, const SrcOp &Tst, const SrcOp &Op0, const SrcOp &Op1, std::optional< unsigned > Flags=std::nullopt)

Build and insert a Res = G_SELECT Tst, Op0, Op1.

Definition:MachineIRBuilder.cpp:950

llvm::MachineIRBuilder::buildICmp

MachineInstrBuilder buildICmp(CmpInst::Predicate Pred, const DstOp &Res, const SrcOp &Op0, const SrcOp &Op1, std::optional< unsigned > Flags=std::nullopt)

Build and insert a Res = G_ICMP Pred, Op0, Op1.

Definition:MachineIRBuilder.cpp:920

llvm::MachineIRBuilder::getInsertPt

MachineBasicBlock::iterator getInsertPt()

Current insertion point for new instructions.

Definition:MachineIRBuilder.h:333

llvm::MachineIRBuilder::buildIntrinsic

MachineInstrBuilder buildIntrinsic(Intrinsic::ID ID, ArrayRef< Register > Res, bool HasSideEffects, bool isConvergent)

Build and insert a G_INTRINSIC instruction.

Definition:MachineIRBuilder.cpp:870

llvm::MachineIRBuilder::buildLoad

MachineInstrBuilder buildLoad(const DstOp &Res, const SrcOp &Addr, MachineMemOperand &MMO)

Build and insert Res = G_LOAD Addr, MMO.

Definition:MachineIRBuilder.h:959

llvm::MachineIRBuilder::buildZExtOrTrunc

MachineInstrBuilder buildZExtOrTrunc(const DstOp &Res, const SrcOp &Op)

Build and insert Res = G_ZEXT Op, Res = G_TRUNC Op, or Res = COPY Op depending on the differing sizes...

Definition:MachineIRBuilder.cpp:576

llvm::MachineIRBuilder::buildInstr

MachineInstrBuilder buildInstr(unsigned Opcode)

Build and insert <empty> = Opcode <empty>.

Definition:MachineIRBuilder.h:417

llvm::MachineIRBuilder::getMF

MachineFunction & getMF()

Getter for the function we currently build.

Definition:MachineIRBuilder.h:287

llvm::MachineIRBuilder::getMRI

MachineRegisterInfo * getMRI()

Getter for MRI.

Definition:MachineIRBuilder.h:309

llvm::MachineIRBuilder::buildCopy

MachineInstrBuilder buildCopy(const DstOp &Res, const SrcOp &Op)

Build and insert Res = COPY Op.

Definition:MachineIRBuilder.cpp:312

llvm::MachineIRBuilder::getDataLayout

const DataLayout & getDataLayout() const

Definition:MachineIRBuilder.h:297

llvm::MachineInstrBuilder

Definition:MachineInstrBuilder.h:71

llvm::MachineInstrBuilder::addImm

const MachineInstrBuilder & addImm(int64_t Val) const

Add a new immediate operand.

Definition:MachineInstrBuilder.h:133

llvm::MachineInstrBuilder::addUse

const MachineInstrBuilder & addUse(Register RegNo, unsigned Flags=0, unsigned SubReg=0) const

Add a virtual register use operand.

Definition:MachineInstrBuilder.h:125

llvm::MachineInstrBuilder::addDef

const MachineInstrBuilder & addDef(Register RegNo, unsigned Flags=0, unsigned SubReg=0) const

Add a virtual register definition operand.

Definition:MachineInstrBuilder.h:118

llvm::MachineInstr

Representation of each machine instruction.

Definition:MachineInstr.h:71

llvm::MachineInstr::getOpcode

unsigned getOpcode() const

Returns the opcode of this MachineInstr.

Definition:MachineInstr.h:577

llvm::MachineInstr::getOperand

const MachineOperand & getOperand(unsigned i) const

Definition:MachineInstr.h:587

llvm::MachineOperand

MachineOperand class - Representation of each machine instruction operand.

Definition:MachineOperand.h:48

llvm::MachineOperand::getGlobal

const GlobalValue * getGlobal() const

Definition:MachineOperand.h:582

llvm::MachineOperand::getCImm

const ConstantInt * getCImm() const

Definition:MachineOperand.h:561

llvm::MachineOperand::isCImm

bool isCImm() const

isCImm - Test if this is a MO_CImmediate operand.

Definition:MachineOperand.h:333

llvm::MachineOperand::getImm

int64_t getImm() const

Definition:MachineOperand.h:556

llvm::MachineOperand::isReg

bool isReg() const

isReg - Tests if this is a MO_Register operand.

Definition:MachineOperand.h:329

llvm::MachineOperand::getMetadata

const MDNode * getMetadata() const

Definition:MachineOperand.h:676

llvm::MachineOperand::getReg

getReg - Returns the register number.

Definition:MachineOperand.h:369

llvm::MachineOperand::getFPImm

const ConstantFP * getFPImm() const

Definition:MachineOperand.h:566

llvm::MachineRegisterInfo

MachineRegisterInfo - Keep track of information for virtual and physical registers,...

Definition:MachineRegisterInfo.h:51

llvm::MachineRegisterInfo::createVirtualRegister

createVirtualRegister - Create and return a new virtual register in the function with the specified r...

Definition:MachineRegisterInfo.cpp:156

llvm::MachineRegisterInfo::setType

void setType(Register VReg, LLT Ty)

Set the low-level type of VReg to Ty.

Definition:MachineRegisterInfo.cpp:187

llvm::MachineRegisterInfo::setRegClass

void setRegClass(Register Reg, const TargetRegisterClass *RC)

setRegClass - Set the register class of the specified virtual register.

Definition:MachineRegisterInfo.cpp:58

llvm::MachineRegisterInfo::createGenericVirtualRegister

Create and return a new generic virtual register with low-level type Ty.

Definition:MachineRegisterInfo.cpp:193

llvm::MachineRegisterInfo::getNumVirtRegs

unsigned getNumVirtRegs() const

getNumVirtRegs - Return the number of virtual registers created.

Definition:MachineRegisterInfo.h:802

llvm::PointerType::get

static PointerType * get(Type *ElementType, unsigned AddressSpace)

This constructs a pointer to an object of the specified type in a numbered address space.

llvm::Register

Wrapper class representing virtual and physical registers.

Definition:Register.h:19

llvm::Register::isValid

constexpr bool isValid() const

Definition:Register.h:115

llvm::SPIRVGlobalRegistry

Definition:SPIRVGlobalRegistry.h:30

llvm::SPIRVGlobalRegistry::getOrCreateOpTypePipe

SPIRVType * getOrCreateOpTypePipe(MachineIRBuilder &MIRBuilder, SPIRV::AccessQualifier::AccessQualifier AccQual)

Definition:SPIRVGlobalRegistry.cpp:1339

llvm::SPIRVGlobalRegistry::getSPIRVTypeForVReg

SPIRVType * getSPIRVTypeForVReg(Register VReg, const MachineFunction *MF=nullptr) const

Definition:SPIRVGlobalRegistry.cpp:1106

llvm::SPIRVGlobalRegistry::assignSPIRVTypeToVReg

void assignSPIRVTypeToVReg(SPIRVType *Type, Register VReg, const MachineFunction &MF)

Definition:SPIRVGlobalRegistry.cpp:82

llvm::SPIRVGlobalRegistry::getOrCreateSPIRVBoolType

SPIRVType * getOrCreateSPIRVBoolType(MachineIRBuilder &MIRBuilder)

Definition:SPIRVGlobalRegistry.cpp:1530

llvm::SPIRVGlobalRegistry::getOrCreateConsIntVector

Definition:SPIRVGlobalRegistry.cpp:609

llvm::SPIRVGlobalRegistry::getTypeForSPIRVType

const Type * getTypeForSPIRVType(const SPIRVType *Ty) const

Definition:SPIRVGlobalRegistry.h:351

llvm::SPIRVGlobalRegistry::buildConstantSampler

Register buildConstantSampler(Register Res, unsigned AddrMode, unsigned Param, unsigned FilerMode, MachineIRBuilder &MIRBuilder, SPIRVType *SpvType)

Definition:SPIRVGlobalRegistry.cpp:649

llvm::SPIRVGlobalRegistry::getScalarOrVectorComponentCount

unsigned getScalarOrVectorComponentCount(Register VReg) const

Definition:SPIRVGlobalRegistry.cpp:1186

llvm::SPIRVGlobalRegistry::getOrCreateOpTypeImage

SPIRVType * getOrCreateOpTypeImage(MachineIRBuilder &MIRBuilder, SPIRVType *SampledType, SPIRV::Dim::Dim Dim, uint32_t Depth, uint32_t Arrayed, uint32_t Multisampled, uint32_t Sampled, SPIRV::ImageFormat::ImageFormat ImageFormat, SPIRV::AccessQualifier::AccessQualifier AccQual)

Definition:SPIRVGlobalRegistry.cpp:1302

llvm::SPIRVGlobalRegistry::getPointerSize

unsigned getPointerSize() const

Definition:SPIRVGlobalRegistry.h:445

llvm::SPIRVGlobalRegistry::getOrCreateOpTypeByOpcode

SPIRVType * getOrCreateOpTypeByOpcode(const Type *Ty, MachineIRBuilder &MIRBuilder, unsigned Opcode)

Definition:SPIRVGlobalRegistry.cpp:1398

llvm::SPIRVGlobalRegistry::buildConstantFP

Definition:SPIRVGlobalRegistry.cpp:406

llvm::SPIRVGlobalRegistry::getPointeeType

SPIRVType * getPointeeType(SPIRVType *PtrType)

Definition:SPIRVGlobalRegistry.cpp:1258

llvm::SPIRVGlobalRegistry::getSPIRVTypeID

Definition:SPIRVGlobalRegistry.cpp:969

llvm::SPIRVGlobalRegistry::getOrCreateSPIRVType

SPIRVType * getOrCreateSPIRVType(const Type *Type, MachineIRBuilder &MIRBuilder, SPIRV::AccessQualifier::AccessQualifier AQ=SPIRV::AccessQualifier::ReadWrite, bool EmitIR=true)

Definition:SPIRVGlobalRegistry.cpp:1125

llvm::SPIRVGlobalRegistry::isScalarOfType

bool isScalarOfType(Register VReg, unsigned TypeOpcode) const

Definition:SPIRVGlobalRegistry.cpp:1164

llvm::SPIRVGlobalRegistry::buildGlobalVariable

Register buildGlobalVariable(Register Reg, SPIRVType *BaseType, StringRef Name, const GlobalValue *GV, SPIRV::StorageClass::StorageClass Storage, const MachineInstr *Init, bool IsConst, bool HasLinkageTy, SPIRV::LinkageType::LinkageType LinkageType, MachineIRBuilder &MIRBuilder, bool IsInstSelector)

Definition:SPIRVGlobalRegistry.cpp:675

llvm::SPIRVGlobalRegistry::getOrCreateOpTypeSampledImage

SPIRVType * getOrCreateOpTypeSampledImage(SPIRVType *ImageType, MachineIRBuilder &MIRBuilder)

Definition:SPIRVGlobalRegistry.cpp:1362

llvm::SPIRVGlobalRegistry::getOrCreateSPIRVTypeByName

SPIRVType * getOrCreateSPIRVTypeByName(StringRef TypeStr, MachineIRBuilder &MIRBuilder, SPIRV::StorageClass::StorageClass SC=SPIRV::StorageClass::Function, SPIRV::AccessQualifier::AccessQualifier AQ=SPIRV::AccessQualifier::ReadWrite)

Definition:SPIRVGlobalRegistry.cpp:1419

llvm::SPIRVGlobalRegistry::getRegClass

const TargetRegisterClass * getRegClass(SPIRVType *SpvType) const

Definition:SPIRVGlobalRegistry.cpp:1654

llvm::SPIRVGlobalRegistry::isScalarOrVectorOfType

bool isScalarOrVectorOfType(Register VReg, unsigned TypeOpcode) const

Definition:SPIRVGlobalRegistry.cpp:1171

llvm::SPIRVGlobalRegistry::getOrCreateConstIntArray

Register getOrCreateConstIntArray(uint64_t Val, size_t Num, MachineInstr &I, SPIRVType *SpvType, const SPIRVInstrInfo &TII)

Definition:SPIRVGlobalRegistry.cpp:543

llvm::SPIRVGlobalRegistry::getOrCreateOpTypeDeviceEvent

SPIRVType * getOrCreateOpTypeDeviceEvent(MachineIRBuilder &MIRBuilder)

Definition:SPIRVGlobalRegistry.cpp:1352

llvm::SPIRVGlobalRegistry::getOrCreateSPIRVPointerType

SPIRVType * getOrCreateSPIRVPointerType(SPIRVType *BaseType, MachineIRBuilder &MIRBuilder, SPIRV::StorageClass::StorageClass SClass=SPIRV::StorageClass::Function)

Definition:SPIRVGlobalRegistry.cpp:1594

llvm::SPIRVGlobalRegistry::getOrCreateOpTypeCoopMatr

SPIRVType * getOrCreateOpTypeCoopMatr(MachineIRBuilder &MIRBuilder, const TargetExtType *ExtensionType, const SPIRVType *ElemType, uint32_t Scope, uint32_t Rows, uint32_t Columns, uint32_t Use)

Definition:SPIRVGlobalRegistry.cpp:1377

llvm::SPIRVGlobalRegistry::getOrCreateSPIRVVectorType

SPIRVType * getOrCreateSPIRVVectorType(SPIRVType *BaseType, unsigned NumElements, MachineIRBuilder &MIRBuilder)

Definition:SPIRVGlobalRegistry.cpp:1550

llvm::SPIRVGlobalRegistry::getOrCreateSPIRVIntegerType

SPIRVType * getOrCreateSPIRVIntegerType(unsigned BitWidth, MachineIRBuilder &MIRBuilder)

Definition:SPIRVGlobalRegistry.cpp:1466

llvm::SPIRVGlobalRegistry::getRegType

LLT getRegType(SPIRVType *SpvType) const

Definition:SPIRVGlobalRegistry.cpp:1680

llvm::SPIRVGlobalRegistry::getPointerStorageClass

SPIRV::StorageClass::StorageClass getPointerStorageClass(Register VReg) const

Definition:SPIRVGlobalRegistry.cpp:1289

llvm::SPIRVGlobalRegistry::getOrCreateOpTypeSampler

SPIRVType * getOrCreateOpTypeSampler(MachineIRBuilder &MIRBuilder)

Definition:SPIRVGlobalRegistry.cpp:1330

llvm::SPIRVGlobalRegistry::getOrCreateConstNullPtr

Definition:SPIRVGlobalRegistry.cpp:626

llvm::SPIRVGlobalRegistry::getScalarOrVectorBitWidth

unsigned getScalarOrVectorBitWidth(const SPIRVType *Type) const

Definition:SPIRVGlobalRegistry.cpp:1218

llvm::SPIRVGlobalRegistry::buildConstantInt

Register buildConstantInt(uint64_t Val, MachineIRBuilder &MIRBuilder, SPIRVType *SpvType, bool EmitIR=true, bool ZeroAsNull=true)

Definition:SPIRVGlobalRegistry.cpp:359

llvm::SPIRVSubtarget

Definition:SPIRVSubtarget.h:38

llvm::SmallVectorBase::size

size_t size() const

Definition:SmallVector.h:78

llvm::SmallVectorImpl

This class consists of common code factored out of the SmallVector class to reduce code duplication b...

Definition:SmallVector.h:573

llvm::SmallVectorTemplateBase::push_back

void push_back(const T &Elt)

Definition:SmallVector.h:413

llvm::SmallVector

This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.

Definition:SmallVector.h:1196

llvm::StringRef

StringRef - Represent a constant reference to a string, i.e.

Definition:StringRef.h:51

llvm::StringRef::split

std::pair< StringRef, StringRef > split(char Separator) const

Split into two substrings around the first occurrence of a separator character.

Definition:StringRef.h:700

llvm::StringRef::consume_back

bool consume_back(StringRef Suffix)

Returns true if this StringRef has the given suffix and removes that suffix.

Definition:StringRef.h:655

llvm::StringRef::getAsInteger

bool getAsInteger(unsigned Radix, T &Result) const

Parse the current string as an integer of the specified radix.

Definition:StringRef.h:470

llvm::StringRef::str

std::string str() const

str - Get the contents as an std::string.

Definition:StringRef.h:229

llvm::StringRef::substr

constexpr StringRef substr(size_t Start, size_t N=npos) const

Return a reference to the substring from [Start, Start + N).

Definition:StringRef.h:571

llvm::StringRef::starts_with

bool starts_with(StringRef Prefix) const

Check if this string starts with the given Prefix.

Definition:StringRef.h:265

llvm::StringRef::contains_insensitive

bool contains_insensitive(StringRef Other) const

Return true if the given string is a substring of *this, and false otherwise.

Definition:StringRef.h:436

llvm::StringRef::slice

StringRef slice(size_t Start, size_t End) const

Return a reference to the substring from [Start, End).

Definition:StringRef.h:684

llvm::StringRef::contains

bool contains(StringRef Other) const

Return true if the given string is a substring of *this, and false otherwise.

Definition:StringRef.h:424

llvm::StringRef::consume_front

bool consume_front(StringRef Prefix)

Returns true if this StringRef has the given prefix and removes that prefix.

Definition:StringRef.h:635

llvm::StringRef::find_first_of

size_t find_first_of(char C, size_t From=0) const

Find the first character in the string that is C, or npos if not found.

Definition:StringRef.h:377

llvm::StringRef::rfind

size_t rfind(char C, size_t From=npos) const

Search for the last character C in the string.

Definition:StringRef.h:347

llvm::StringRef::find

size_t find(char C, size_t From=0) const

Search for the first character C in the string.

Definition:StringRef.h:297

llvm::StringRef::ends_with

bool ends_with(StringRef Suffix) const

Check if this string ends with the given Suffix.

Definition:StringRef.h:277

llvm::StringRef::npos

static constexpr size_t npos

Definition:StringRef.h:53

llvm::StringSwitch

A switch()-like statement whose cases are string literals.

Definition:StringSwitch.h:44

llvm::StringSwitch::EndsWith

StringSwitch & EndsWith(StringLiteral S, T Value)

Definition:StringSwitch.h:76

llvm::StructType

Class to represent struct types.

Definition:DerivedTypes.h:218

llvm::StructType::getTypeByName

static StructType * getTypeByName(LLVMContext &C, StringRef Name)

Return the type with the specified name, or null if there is none by that name.

Definition:Type.cpp:731

llvm::StructType::create

static StructType * create(LLVMContext &Context, StringRef Name)

This creates an identified struct.

Definition:Type.cpp:612

llvm::TargetExtType

Class to represent target extensions types, which are generally unintrospectable from target-independ...

Definition:DerivedTypes.h:744

llvm::TargetExtType::getNumIntParameters

unsigned getNumIntParameters() const

Definition:DerivedTypes.h:802

llvm::TargetExtType::get

static TargetExtType * get(LLVMContext &Context, StringRef Name, ArrayRef< Type * > Types={}, ArrayRef< unsigned > Ints={})

Return a target extension type having the specified name and optional type and integer parameters.

Definition:Type.cpp:895

llvm::TargetExtType::getTypeParameter

Type * getTypeParameter(unsigned i) const

Definition:DerivedTypes.h:792

llvm::TargetExtType::getNumTypeParameters

unsigned getNumTypeParameters() const

Definition:DerivedTypes.h:793

llvm::TargetExtType::getIntParameter

unsigned getIntParameter(unsigned i) const

Definition:DerivedTypes.h:801

llvm::TargetRegisterClass

Definition:TargetRegisterInfo.h:44

llvm::Type

The instances of the Type class are immutable: once they are created, they are never changed.

Definition:Type.h:45

llvm::Type::getHalfTy

static Type * getHalfTy(LLVMContext &C)

llvm::Type::getStructName

StringRef getStructName() const

llvm::Type::getVoidTy

static Type * getVoidTy(LLVMContext &C)

llvm::Type::getInt8Ty

static IntegerType * getInt8Ty(LLVMContext &C)

llvm::Type::getInt32Ty

static IntegerType * getInt32Ty(LLVMContext &C)

llvm::Type::getFloatTy

static Type * getFloatTy(LLVMContext &C)

llvm::Value

LLVM Value Representation.

Definition:Value.h:74

llvm::Value::Value

Value(Type *Ty, unsigned scid)

Definition:Value.cpp:53

llvm::VectorType::get

static VectorType * get(Type *ElementType, ElementCount EC)

This static method is the primary way to construct an VectorType.

llvm::ilist_node_with_parent::getNextNode

NodeTy * getNextNode()

Get the next node, or nullptr for the list tail.

Definition:ilist_node.h:353

LLVMTypeRef LLVMVectorType(LLVMTypeRef ElementType, unsigned ElementCount)

Create a vector type that contains a defined type and has a specific number of elements.

Definition:Core.cpp:884

llvm::SPIRV::lookupBuiltinNameHelper

std::string lookupBuiltinNameHelper(StringRef DemangledCall, FPDecorationId *DecorationId)

Parses the name part of the demangled builtin call.

Definition:SPIRVBuiltins.cpp:176

llvm::SPIRV::parseBuiltinCallArgumentType

Type * parseBuiltinCallArgumentType(StringRef TypeStr, LLVMContext &Ctx)

Definition:SPIRVBuiltins.cpp:2689

llvm::SPIRV::parseBuiltinTypeStr

bool parseBuiltinTypeStr(SmallVector< StringRef, 10 > &BuiltinArgsTypeStrs, const StringRef DemangledCall, LLVMContext &Ctx)

Definition:SPIRVBuiltins.cpp:2733

llvm::SPIRV::mapBuiltinToOpcode

std::tuple< int, unsigned, unsigned > mapBuiltinToOpcode(const StringRef DemangledCall, SPIRV::InstructionSet::InstructionSet Set)

Helper function for finding a builtin function attributes by a demangled function name.

Definition:SPIRVBuiltins.cpp:2535

llvm::SPIRV::parseBuiltinCallArgumentBaseType

Type * parseBuiltinCallArgumentBaseType(const StringRef DemangledCall, unsigned ArgIdx, LLVMContext &Ctx)

Parses the provided ArgIdx argument base type in the DemangledCall skeleton.

Definition:SPIRVBuiltins.cpp:2746

llvm::SPIRV::parseBuiltinTypeNameToTargetExtType

TargetExtType * parseBuiltinTypeNameToTargetExtType(std::string TypeName, LLVMContext &Context)

Translates a string representing a SPIR-V or OpenCL builtin type to a TargetExtType that can be furth...

Definition:SPIRVBuiltins.cpp:2878

llvm::SPIRV::lowerBuiltin

std::optional< bool > lowerBuiltin(const StringRef DemangledCall, SPIRV::InstructionSet::InstructionSet Set, MachineIRBuilder &MIRBuilder, const Register OrigRet, const Type *OrigRetTy, const SmallVectorImpl< Register > &Args, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:2600

llvm::SPIRV::lowerBuiltinType

SPIRVType * lowerBuiltinType(const Type *OpaqueType, SPIRV::AccessQualifier::AccessQualifier AccessQual, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:2927

llvm::XCOFF::StorageClass

StorageClass

Definition:XCOFF.h:170

llvm

This is an optimization pass for GlobalISel generic memory operations.

Definition:AddressRanges.h:18

llvm::buildOpName

void buildOpName(Register Target, const StringRef &Name, MachineIRBuilder &MIRBuilder)

Definition:SPIRVUtils.cpp:103

llvm::generateGetQueryInst

static bool generateGetQueryInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:1670

llvm::generateLoadStoreInst

static bool generateLoadStoreInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:2497

llvm::generateConstructInst

static bool generateConstructInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:1981

llvm::buildAtomicFlagInst

static bool buildAtomicFlagInst(const SPIRV::IncomingCall *Call, unsigned Opcode, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Helper function for building atomic flag instructions (e.g.

Definition:SPIRVBuiltins.cpp:890

llvm::buildBuiltinVariableLoad

static Register buildBuiltinVariableLoad(MachineIRBuilder &MIRBuilder, SPIRVType *VariableType, SPIRVGlobalRegistry *GR, SPIRV::BuiltIn::BuiltIn BuiltinValue, LLT LLType, Register Reg=Register(0), bool isConst=true, bool hasLinkageTy=true)

Helper function for building a load instruction for loading a builtin global variable of BuiltinValue...

Definition:SPIRVBuiltins.cpp:496

llvm::generateImageSizeQueryInst

static bool generateImageSizeQueryInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:1682

llvm::getSamplerFilterModeFromBitmask

static SPIRV::SamplerFilterMode::SamplerFilterMode getSamplerFilterModeFromBitmask(unsigned Bitmask)

Definition:SPIRVBuiltins.cpp:1817

llvm::buildAtomicStoreInst

static bool buildAtomicStoreInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Helper function for building an atomic store instruction.

Definition:SPIRVBuiltins.cpp:688

llvm::addNumImm

void addNumImm(const APInt &Imm, MachineInstrBuilder &MIB)

Definition:SPIRVUtils.cpp:83

llvm::getBlockStructType

static const Type * getBlockStructType(Register ParamReg, MachineRegisterInfo *MRI)

Definition:SPIRVBuiltins.cpp:410

llvm::generateGroupInst

static bool generateGroupInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:1079

llvm::demangledPostfixToDecorationId

FPDecorationId demangledPostfixToDecorationId(const std::string &S)

Definition:SPIRVUtils.h:411

llvm::DiagnosticPredicateTy::Match

@ Match

llvm::getNumComponentsForDim

static unsigned getNumComponentsForDim(SPIRV::Dim::Dim dim)

Definition:SPIRVBuiltins.cpp:1003

llvm::insertAssignInstr

Register insertAssignInstr(Register Reg, Type *Ty, SPIRVType *SpirvTy, SPIRVGlobalRegistry *GR, MachineIRBuilder &MIB, MachineRegisterInfo &MRI)

Helper external function for inserting ASSIGN_TYPE instuction between Reg and its definition,...

Definition:SPIRVPreLegalizer.cpp:429

llvm::generateICarryBorrowInst

static bool generateICarryBorrowInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:1621

llvm::buildScopeReg

static Register buildScopeReg(Register CLScopeRegister, SPIRV::Scope::Scope Scope, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR, MachineRegisterInfo *MRI)

Definition:SPIRVBuiltins.cpp:576

llvm::FPDecorationId

FPDecorationId

Definition:SPIRVUtils.h:409

llvm::buildBoolRegister

static std::tuple< Register, SPIRVType * > buildBoolRegister(MachineIRBuilder &MIRBuilder, const SPIRVType *ResultType, SPIRVGlobalRegistry *GR)

Helper function building either a resulting scalar or vector bool register depending on the expected ...

Definition:SPIRVBuiltins.cpp:434

llvm::getNumSizeComponents

static unsigned getNumSizeComponents(SPIRVType *imgType)

Helper function for obtaining the number of size components.

Definition:SPIRVBuiltins.cpp:1020

llvm::getIConstVal

uint64_t getIConstVal(Register ConstReg, const MachineRegisterInfo *MRI)

Definition:SPIRVUtils.cpp:326

llvm::buildConstantIntReg32

static Register buildConstantIntReg32(uint64_t Val, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:569

llvm::getSampledImageType

static SPIRVType * getSampledImageType(const TargetExtType *OpaqueType, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:2868

llvm::getMemSemanticsForStorageClass

SPIRV::MemorySemantics::MemorySemantics getMemSemanticsForStorageClass(SPIRV::StorageClass::StorageClass SC)

Definition:SPIRVUtils.cpp:245

llvm::storageClassToAddressSpace

constexpr unsigned storageClassToAddressSpace(SPIRV::StorageClass::StorageClass SC)

Definition:SPIRVUtils.h:163

llvm::generateDotOrFMulInst

static bool generateDotOrFMulInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:1582

llvm::generateSampleImageInst

static bool generateSampleImageInst(const StringRef DemangledCall, const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:1915

llvm::generateBarrierInst

static bool generateBarrierInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:1563

llvm::getCoopMatrType

static SPIRVType * getCoopMatrType(const TargetExtType *ExtensionType, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:2821

llvm::generateKernelClockInst

static bool generateKernelClockInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:1328

llvm::setRegClassIfNull

static void setRegClassIfNull(Register Reg, MachineRegisterInfo *MRI, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:594

llvm::generateGroupUniformInst

static bool generateGroupUniformInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:1284

llvm::generateWaveInst

static bool generateWaveInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:1599

llvm::getImageType

static SPIRVType * getImageType(const TargetExtType *ExtensionType, const SPIRV::AccessQualifier::AccessQualifier Qualifier, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:2838

llvm::buildOpDecorate

void buildOpDecorate(Register Reg, MachineIRBuilder &MIRBuilder, SPIRV::Decoration::Decoration Dec, const std::vector< uint32_t > &DecArgs, StringRef StrImm)

Definition:SPIRVUtils.cpp:130

llvm::createVirtualRegister

Register createVirtualRegister(SPIRVType *SpvType, SPIRVGlobalRegistry *GR, MachineRegisterInfo *MRI, const MachineFunction &MF)

Definition:SPIRVUtils.cpp:748

llvm::buildBarrierInst

static bool buildBarrierInst(const SPIRV::IncomingCall *Call, unsigned Opcode, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Helper function for building barriers, i.e., memory/control ordering operations.

Definition:SPIRVBuiltins.cpp:927

llvm::generateAsyncCopy

static bool generateAsyncCopy(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:2315

llvm::getSPIRVScope

static SPIRV::Scope::Scope getSPIRVScope(SPIRV::CLMemoryScope ClScope)

Definition:SPIRVBuiltins.cpp:553

llvm::getSamplerType

static SPIRVType * getSamplerType(MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:2804

llvm::dbgs

raw_ostream & dbgs()

dbgs() - This returns a reference to a raw_ostream for debugging messages.

Definition:Debug.cpp:163

llvm::buildLoadInst

static Register buildLoadInst(SPIRVType *BaseType, Register PtrRegister, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR, LLT LowLevelType, Register DestinationReg=Register(0))

Helper function for building a load instruction loading into the DestinationReg.

Definition:SPIRVBuiltins.cpp:481

llvm::generateEnqueueInst

static bool generateEnqueueInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:2275

llvm::report_fatal_error

void report_fatal_error(Error Err, bool gen_crash_diag=true)

Report a serious error, calling any installed error handler.

Definition:Error.cpp:167

llvm::buildSelectInst

static bool buildSelectInst(MachineIRBuilder &MIRBuilder, Register ReturnRegister, Register SourceRegister, const SPIRVType *ReturnType, SPIRVGlobalRegistry *GR)

Helper function for building either a vector or scalar select instruction depending on the expected R...

Definition:SPIRVBuiltins.cpp:459

llvm::getMachineInstrType

static const Type * getMachineInstrType(MachineInstr *MI)

Definition:SPIRVBuiltins.cpp:393

llvm::getSamplerAddressingModeFromBitmask

static SPIRV::SamplerAddressingMode::SamplerAddressingMode getSamplerAddressingModeFromBitmask(unsigned Bitmask)

Definition:SPIRVBuiltins.cpp:1795

llvm::generateAtomicInst

static bool generateAtomicInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:1507

llvm::generateConvertInst

static bool generateConvertInst(const StringRef DemangledCall, const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:2366

llvm::buildMemSemanticsReg

static Register buildMemSemanticsReg(Register SemanticsRegister, Register PtrRegister, unsigned &Semantics, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:603

llvm::getConstFromIntrinsic

static unsigned getConstFromIntrinsic(Register Reg, MachineRegisterInfo *MRI)

Definition:SPIRVBuiltins.cpp:381

llvm::generateImageMiscQueryInst

static bool generateImageMiscQueryInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:1756

llvm::generateSelectInst

static bool generateSelectInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder)

Definition:SPIRVBuiltins.cpp:1974

llvm::buildAtomicLoadInst

static bool buildAtomicLoadInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Helper function for building an atomic load instruction.

Definition:SPIRVBuiltins.cpp:652

llvm::generateIntelSubgroupsInst

static bool generateIntelSubgroupsInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:1210

llvm::generateSpecConstantInst

static bool generateSpecConstantInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:2054

llvm::getOrCreateSPIRVDeviceEventPointer

static SPIRVType * getOrCreateSPIRVDeviceEventPointer(MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:2176

llvm::parseBasicTypeName

Type * parseBasicTypeName(StringRef &TypeName, LLVMContext &Ctx)

Definition:SPIRVUtils.cpp:459

llvm::generateVectorLoadStoreInst

static bool generateVectorLoadStoreInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:2471

llvm::genWorkgroupQuery

static bool genWorkgroupQuery(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR, SPIRV::BuiltIn::BuiltIn BuiltinValue, uint64_t DefaultValue)

Definition:SPIRVBuiltins.cpp:1385

llvm::generateCoopMatrInst

static bool generateCoopMatrInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:1988

llvm::lookupBuiltin

static std::unique_ptr< const SPIRV::IncomingCall > lookupBuiltin(StringRef DemangledCall, SPIRV::InstructionSet::InstructionSet Set, Register ReturnRegister, const SPIRVType *ReturnType, const SmallVectorImpl< Register > &Arguments)

Looks up the demangled builtin call in the SPIRVBuiltins.td records using the provided DemangledCall ...

Definition:SPIRVBuiltins.cpp:262

llvm::buildAtomicFloatingRMWInst

static bool buildAtomicFloatingRMWInst(const SPIRV::IncomingCall *Call, unsigned Opcode, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Helper function for building an atomic floating-type instruction.

Definition:SPIRVBuiltins.cpp:868

llvm::getDefInstrMaybeConstant

MachineInstr * getDefInstrMaybeConstant(Register &ConstReg, const MachineRegisterInfo *MRI)

Definition:SPIRVUtils.cpp:307

llvm::BitWidth

constexpr unsigned BitWidth

Definition:BitmaskEnum.h:217

llvm::SPIRVType

const MachineInstr SPIRVType

Definition:SPIRVGlobalRegistry.h:28

llvm::generateReadImageInst

static bool generateReadImageInst(const StringRef DemangledCall, const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:1825

llvm::hasBuiltinTypePrefix

bool hasBuiltinTypePrefix(StringRef Name)

Definition:SPIRVUtils.cpp:429

llvm::buildEnqueueKernel

static bool buildEnqueueKernel(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:2190

llvm::getMDOperandAsType

Type * getMDOperandAsType(const MDNode *N, unsigned I)

Definition:SPIRVUtils.cpp:338

llvm::buildAtomicRMWInst

static bool buildAtomicRMWInst(const SPIRV::IncomingCall *Call, unsigned Opcode, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Helper function for building atomic instructions.

Definition:SPIRVBuiltins.cpp:812

llvm::getSPIRVMemSemantics

static SPIRV::MemorySemantics::MemorySemantics getSPIRVMemSemantics(std::memory_order MemOrder)

Definition:SPIRVBuiltins.cpp:536

llvm::generateRelationalInst

static bool generateRelationalInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:1053

llvm::buildAtomicInitInst

static bool buildAtomicInitInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder)

Helper function for translating atomic init to OpStore.

Definition:SPIRVBuiltins.cpp:638

llvm::generateWriteImageInst

static bool generateWriteImageInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:1905

llvm::getPipeType

static SPIRVType * getPipeType(const TargetExtType *ExtensionType, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:2810

llvm::parseTypeString

static Type * parseTypeString(const StringRef Name, LLVMContext &Context)

Definition:SPIRVBuiltins.cpp:2779

llvm::isSpvIntrinsic

bool isSpvIntrinsic(const MachineInstr &MI, Intrinsic::ID IntrinsicID)

Definition:SPIRVUtils.cpp:332

llvm::generateCastToPtrInst

static bool generateCastToPtrInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder)

Definition:SPIRVBuiltins.cpp:1574

llvm::generateAtomicFloatingInst

static bool generateAtomicFloatingInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:1546

llvm::generateExtInst

static bool generateExtInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:1032

llvm::buildNDRange

static bool buildNDRange(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:2110

llvm::getNonParameterizedType

static SPIRVType * getNonParameterizedType(const TargetExtType *ExtensionType, const SPIRV::BuiltinType *TypeRecord, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:2795

llvm::getBlockStructInstr

static MachineInstr * getBlockStructInstr(Register ParamReg, MachineRegisterInfo *MRI)

Definition:SPIRVBuiltins.cpp:359

llvm::buildOpFromWrapper

static bool buildOpFromWrapper(MachineIRBuilder &MIRBuilder, unsigned Opcode, const SPIRV::IncomingCall *Call, Register TypeReg, ArrayRef< uint32_t > ImmArgs={})

Definition:SPIRVBuiltins.cpp:622

llvm::fltNanEncoding::AllOnes

@ AllOnes

llvm::getSamplerParamFromBitmask

static unsigned getSamplerParamFromBitmask(unsigned Bitmask)

Definition:SPIRVBuiltins.cpp:1812

llvm::buildAtomicCompareExchangeInst

static bool buildAtomicCompareExchangeInst(const SPIRV::IncomingCall *Call, const SPIRV::DemangledBuiltin *Builtin, unsigned Opcode, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Helper function for building an atomic compare-exchange instruction.

Definition:SPIRVBuiltins.cpp:710

llvm::IntegerStyle::Number

@ Number

llvm::getLinkStringForBuiltIn

std::string getLinkStringForBuiltIn(SPIRV::BuiltIn::BuiltIn BuiltInValue)

Definition:SPIRVBaseInfo.cpp:175

llvm::generateBuiltinVar

static bool generateBuiltinVar(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)

Definition:SPIRVBuiltins.cpp:1483

llvm::APFloatBase::IEEEsingle

static const fltSemantics & IEEEsingle() LLVM_READNONE

Definition:APFloat.cpp:257

llvm::Align

This struct is a compact representation of a valid (non-zero power of two) alignment.

Definition:Alignment.h:39

llvm::MachinePointerInfo

This class contains a discriminated union of information about pointers in memory operands,...

Definition:MachineMemOperand.h:41

llvm::MemOp

Definition:TargetLowering.h:115

llvm::SPIRV::AtomicFloatingBuiltin

Definition:SPIRVBuiltins.cpp:100

llvm::SPIRV::AtomicFloatingBuiltin::Name

StringRef Name

Definition:SPIRVBuiltins.cpp:101

llvm::SPIRV::AtomicFloatingBuiltin::Opcode

uint32_t Opcode

Definition:SPIRVBuiltins.cpp:102

llvm::SPIRV::BuiltinType

Definition:SPIRVBuiltins.cpp:2756

llvm::SPIRV::BuiltinType::Name

StringRef Name

Definition:SPIRVBuiltins.cpp:2757

llvm::SPIRV::BuiltinType::Opcode

uint32_t Opcode

Definition:SPIRVBuiltins.cpp:2758

llvm::SPIRV::ConvertBuiltin

Definition:SPIRVBuiltins.cpp:135

llvm::SPIRV::ConvertBuiltin::IsSaturated

bool IsSaturated

Definition:SPIRVBuiltins.cpp:139

llvm::SPIRV::ConvertBuiltin::RoundingMode

FPRoundingMode::FPRoundingMode RoundingMode

Definition:SPIRVBuiltins.cpp:142

llvm::SPIRV::ConvertBuiltin::IsDestinationSigned

bool IsDestinationSigned

Definition:SPIRVBuiltins.cpp:138

llvm::SPIRV::ConvertBuiltin::IsRounded

bool IsRounded

Definition:SPIRVBuiltins.cpp:140

llvm::SPIRV::ConvertBuiltin::Name

StringRef Name

Definition:SPIRVBuiltins.cpp:136

llvm::SPIRV::ConvertBuiltin::IsBfloat16

bool IsBfloat16

Definition:SPIRVBuiltins.cpp:141

llvm::SPIRV::ConvertBuiltin::Set

InstructionSet::InstructionSet Set

Definition:SPIRVBuiltins.cpp:137

llvm::SPIRV::DemangledBuiltin

Definition:SPIRVBuiltins.cpp:32

llvm::SPIRV::DemangledBuiltin::Set

InstructionSet::InstructionSet Set

Definition:SPIRVBuiltins.cpp:34

llvm::SPIRV::DemangledBuiltin::MaxNumArgs

uint8_t MaxNumArgs

Definition:SPIRVBuiltins.cpp:37

llvm::SPIRV::DemangledBuiltin::MinNumArgs

uint8_t MinNumArgs

Definition:SPIRVBuiltins.cpp:36

llvm::SPIRV::DemangledBuiltin::Name

StringRef Name

Definition:SPIRVBuiltins.cpp:33

llvm::SPIRV::DemangledBuiltin::Group

BuiltinGroup Group

Definition:SPIRVBuiltins.cpp:35

llvm::SPIRV::GetBuiltin

Definition:SPIRVBuiltins.cpp:116

llvm::SPIRV::GetBuiltin::Set

InstructionSet::InstructionSet Set

Definition:SPIRVBuiltins.cpp:118

llvm::SPIRV::GetBuiltin::Name

StringRef Name

Definition:SPIRVBuiltins.cpp:117

llvm::SPIRV::GetBuiltin::Value

BuiltIn::BuiltIn Value

Definition:SPIRVBuiltins.cpp:119

llvm::SPIRV::GroupBuiltin

Definition:SPIRVBuiltins.cpp:70

llvm::SPIRV::GroupBuiltin::IsBallotFindBit

bool IsBallotFindBit

Definition:SPIRVBuiltins.cpp:80

llvm::SPIRV::GroupBuiltin::IsBallot

bool IsBallot

Definition:SPIRVBuiltins.cpp:77

llvm::SPIRV::GroupBuiltin::Name

StringRef Name

Definition:SPIRVBuiltins.cpp:71

llvm::SPIRV::GroupBuiltin::HasBoolArg

bool HasBoolArg

Definition:SPIRVBuiltins.cpp:83

llvm::SPIRV::GroupBuiltin::IsAllEqual

bool IsAllEqual

Definition:SPIRVBuiltins.cpp:76

llvm::SPIRV::GroupBuiltin::IsAllOrAny

bool IsAllOrAny

Definition:SPIRVBuiltins.cpp:75

llvm::SPIRV::GroupBuiltin::IsBallotBitExtract

bool IsBallotBitExtract

Definition:SPIRVBuiltins.cpp:79

llvm::SPIRV::GroupBuiltin::NoGroupOperation

bool NoGroupOperation

Definition:SPIRVBuiltins.cpp:82

llvm::SPIRV::GroupBuiltin::IsInverseBallot

bool IsInverseBallot

Definition:SPIRVBuiltins.cpp:78

llvm::SPIRV::GroupBuiltin::IsLogical

bool IsLogical

Definition:SPIRVBuiltins.cpp:81

llvm::SPIRV::GroupBuiltin::Opcode

uint32_t Opcode

Definition:SPIRVBuiltins.cpp:72

llvm::SPIRV::GroupBuiltin::IsElect

bool IsElect

Definition:SPIRVBuiltins.cpp:74

llvm::SPIRV::GroupBuiltin::GroupOperation

uint32_t GroupOperation

Definition:SPIRVBuiltins.cpp:73

llvm::SPIRV::GroupUniformBuiltin

Definition:SPIRVBuiltins.cpp:107

llvm::SPIRV::GroupUniformBuiltin::Opcode

uint32_t Opcode

Definition:SPIRVBuiltins.cpp:109

llvm::SPIRV::GroupUniformBuiltin::Name

StringRef Name

Definition:SPIRVBuiltins.cpp:108

llvm::SPIRV::GroupUniformBuiltin::IsLogical

bool IsLogical

Definition:SPIRVBuiltins.cpp:110

llvm::SPIRV::ImageQueryBuiltin

Definition:SPIRVBuiltins.cpp:126

llvm::SPIRV::ImageQueryBuiltin::Name

StringRef Name

Definition:SPIRVBuiltins.cpp:127

llvm::SPIRV::ImageQueryBuiltin::Component

uint32_t Component

Definition:SPIRVBuiltins.cpp:129

llvm::SPIRV::ImageQueryBuiltin::Set

InstructionSet::InstructionSet Set

Definition:SPIRVBuiltins.cpp:128

llvm::SPIRV::IncomingCall

Definition:SPIRVBuiltins.cpp:43

llvm::SPIRV::IncomingCall::Arguments

const SmallVectorImpl< Register > & Arguments

Definition:SPIRVBuiltins.cpp:49

llvm::SPIRV::IncomingCall::isSpirvOp

bool isSpirvOp() const

Definition:SPIRVBuiltins.cpp:58

llvm::SPIRV::IncomingCall::BuiltinName

const std::string BuiltinName

Definition:SPIRVBuiltins.cpp:44

llvm::SPIRV::IncomingCall::ReturnType

const SPIRVType * ReturnType

Definition:SPIRVBuiltins.cpp:48

llvm::SPIRV::IncomingCall::ReturnRegister

const Register ReturnRegister

Definition:SPIRVBuiltins.cpp:47

llvm::SPIRV::IncomingCall::Builtin

const DemangledBuiltin * Builtin

Definition:SPIRVBuiltins.cpp:45

llvm::SPIRV::IncomingCall::IncomingCall

IncomingCall(const std::string BuiltinName, const DemangledBuiltin *Builtin, const Register ReturnRegister, const SPIRVType *ReturnType, const SmallVectorImpl< Register > &Arguments)

Definition:SPIRVBuiltins.cpp:51

llvm::SPIRV::IntelSubgroupsBuiltin

Definition:SPIRVBuiltins.cpp:89

llvm::SPIRV::IntelSubgroupsBuiltin::Opcode

uint32_t Opcode

Definition:SPIRVBuiltins.cpp:91

llvm::SPIRV::IntelSubgroupsBuiltin::IsWrite

bool IsWrite

Definition:SPIRVBuiltins.cpp:93

llvm::SPIRV::IntelSubgroupsBuiltin::IsBlock

bool IsBlock

Definition:SPIRVBuiltins.cpp:92