clang 23.0.0git
CIRGenAsm.cpp
Go to the documentation of this file.
1//===--- CIRGenAsm.cpp - Inline Assembly Support for CIR CodeGen ---------===//
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 contains code to emit inline assembly.
10//
11//===----------------------------------------------------------------------===//
12
14
15#include "CIRGenFunction.h"
17
18using namespace clang;
19using namespace clang::CIRGen;
20using namespace cir;
21
22static AsmFlavor inferFlavor(const CIRGenModule &cgm, const AsmStmt &s) {
23 AsmFlavor gnuAsmFlavor =
24 cgm.getCodeGenOpts().getInlineAsmDialect() == CodeGenOptions::IAD_ATT
25 ? AsmFlavor::x86_att
26 : AsmFlavor::x86_intel;
27
28 return isa<MSAsmStmt>(&s) ? AsmFlavor::x86_intel : gnuAsmFlavor;
29}
30
31static void collectClobbers(const CIRGenFunction &cgf, const AsmStmt &s,
32 std::string &constraints, bool &hasUnwindClobber,
33 bool &readOnly, bool readNone) {
34
35 hasUnwindClobber = false;
36 const CIRGenModule &cgm = cgf.getCIRGenModule();
37
38 // Clobbers
39 for (unsigned i = 0, e = s.getNumClobbers(); i != e; i++) {
40 std::string clobber = s.getClobber(i);
41 if (clobber == "memory") {
42 readOnly = readNone = false;
43 } else if (clobber == "unwind") {
44 hasUnwindClobber = true;
45 continue;
46 } else if (clobber != "cc") {
47 clobber = cgf.getTarget().getNormalizedGCCRegisterName(clobber);
48 if (cgm.getCodeGenOpts().StackClashProtector &&
49 cgf.getTarget().isSPRegName(clobber))
50 cgm.getDiags().Report(s.getAsmLoc(),
51 diag::warn_stack_clash_protection_inline_asm);
52 }
53
54 if (isa<MSAsmStmt>(&s)) {
55 if (clobber == "eax" || clobber == "edx") {
56 if (constraints.find("=&A") != std::string::npos)
57 continue;
58 std::string::size_type position1 =
59 constraints.find("={" + clobber + "}");
60 if (position1 != std::string::npos) {
61 constraints.insert(position1 + 1, "&");
62 continue;
63 }
64 std::string::size_type position2 = constraints.find("=A");
65 if (position2 != std::string::npos) {
66 constraints.insert(position2 + 1, "&");
67 continue;
68 }
69 }
70 }
71 if (!constraints.empty())
72 constraints += ',';
73
74 constraints += "~{";
75 constraints += clobber;
76 constraints += '}';
77 }
78
79 // Add machine specific clobbers
80 std::string_view machineClobbers = cgf.getTarget().getClobbers();
81 if (!machineClobbers.empty()) {
82 if (!constraints.empty())
83 constraints += ',';
84 constraints += machineClobbers;
85 }
86}
87
88static void
92
93 for (unsigned i = 0, e = s.getNumOutputs(); i != e; ++i) {
94 StringRef name;
95 if (const GCCAsmStmt *gas = dyn_cast<GCCAsmStmt>(&s))
96 name = gas->getOutputName(i);
97 TargetInfo::ConstraintInfo info(s.getOutputConstraint(i), name);
98 // `validateOutputConstraint` modifies the `info` object by setting the
99 // read/write, clobber, allows-register, and allows-memory process.
100 bool isValid = cgf.getTarget().validateOutputConstraint(info);
101 (void)isValid;
102 assert(isValid && "Failed to parse output constraint");
103 out.push_back(info);
104 }
105
106 for (unsigned i = 0, e = s.getNumInputs(); i != e; ++i) {
107 StringRef name;
108 if (const GCCAsmStmt *gas = dyn_cast<GCCAsmStmt>(&s))
109 name = gas->getInputName(i);
110 TargetInfo::ConstraintInfo info(s.getInputConstraint(i), name);
111 // `validateInputConstraint` modifies the `info` object by setting the
112 // read/write, clobber, allows-register, and allows-memory process.
113 bool isValid = cgf.getTarget().validateInputConstraint(out, info);
114 assert(isValid && "Failed to parse input constraint");
115 (void)isValid;
116 in.push_back(info);
117 }
118}
119
120std::pair<mlir::Value, mlir::Type> CIRGenFunction::emitAsmInputLValue(
121 const TargetInfo::ConstraintInfo &info, LValue inputValue,
122 QualType inputType, std::string &constraintString, SourceLocation loc) {
123
124 if (info.allowsRegister() || !info.allowsMemory()) {
125 if (hasScalarEvaluationKind(inputType))
126 return {emitLoadOfLValue(inputValue, loc).getValue(), mlir::Type()};
127
128 mlir::Type ty = convertType(inputType);
129 uint64_t size = cgm.getDataLayout().getTypeSizeInBits(ty);
130 if ((size <= 64 && llvm::isPowerOf2_64(size)) ||
131 getTargetHooks().isScalarizableAsmOperand(*this, ty)) {
132 ty = cir::IntType::get(&getMLIRContext(), size, false);
133
134 return {builder.createLoad(
135 getLoc(loc),
136 inputValue.getAddress().withElementType(builder, ty)),
137 mlir::Type()};
138 }
139 }
140
141 Address addr = inputValue.getAddress();
142 constraintString += '*';
143 return {addr.getPointer(), addr.getElementType()};
144}
145
146std::pair<mlir::Value, mlir::Type>
148 const Expr *inputExpr,
149 std::string &constraintString) {
150 mlir::Location loc = getLoc(inputExpr->getExprLoc());
151
152 // If this can't be a register or memory, i.e., has to be a constant
153 // (immediate or symbolic), try to emit it as such.
154 if (!info.allowsRegister() && !info.allowsMemory()) {
155 if (info.requiresImmediateConstant()) {
156 Expr::EvalResult evalResult;
157 inputExpr->EvaluateAsRValue(evalResult, getContext(), true);
158
159 llvm::APSInt intResult;
160 if (evalResult.Val.toIntegralConstant(intResult, inputExpr->getType(),
161 getContext()))
162 return {builder.getConstInt(loc, intResult), mlir::Type()};
163 }
164
165 Expr::EvalResult result;
166 if (inputExpr->EvaluateAsInt(result, getContext()))
167 return {builder.getConstInt(loc, result.Val.getInt()), mlir::Type()};
168 }
169
170 if (info.allowsRegister() || !info.allowsMemory())
172 return {emitScalarExpr(inputExpr), mlir::Type()};
173 if (inputExpr->getStmtClass() == Expr::CXXThisExprClass)
174 return {emitScalarExpr(inputExpr), mlir::Type()};
175 inputExpr = inputExpr->IgnoreParenNoopCasts(getContext());
176 LValue dest = emitLValue(inputExpr);
177 return emitAsmInputLValue(info, dest, inputExpr->getType(), constraintString,
178 inputExpr->getExprLoc());
179}
180
181static void emitAsmStores(CIRGenFunction &cgf, const AsmStmt &s,
182 const llvm::ArrayRef<mlir::Value> regResults,
183 const llvm::ArrayRef<mlir::Type> resultRegTypes,
184 const llvm::ArrayRef<mlir::Type> resultTruncRegTypes,
185 const llvm::ArrayRef<LValue> resultRegDests,
186 const llvm::ArrayRef<QualType> resultRegQualTys,
187 const llvm::BitVector &resultTypeRequiresCast,
188 const llvm::BitVector &resultRegIsFlagReg) {
189 CIRGenBuilderTy &builder = cgf.getBuilder();
190 CIRGenModule &cgm = cgf.cgm;
191 mlir::MLIRContext *ctx = builder.getContext();
192
193 assert(regResults.size() == resultRegTypes.size());
194 assert(regResults.size() == resultTruncRegTypes.size());
195 assert(regResults.size() == resultRegDests.size());
196
197 // ResultRegDests can be also populated by addReturnRegisterOutputs() above,
198 // in which case its size may grow.
199 assert(resultTypeRequiresCast.size() <= resultRegDests.size());
200 assert(resultRegIsFlagReg.size() <= resultRegDests.size());
201
202 for (unsigned i = 0, e = regResults.size(); i != e; ++i) {
203 mlir::Value tmp = regResults[i];
204 mlir::Type truncTy = resultTruncRegTypes[i];
205
206 if (i < resultRegIsFlagReg.size() && resultRegIsFlagReg[i])
208
209 // If the result type of the LLVM IR asm doesn't match the result type of
210 // the expression, do the conversion.
211 if (resultRegTypes[i] != truncTy) {
212
213 // Truncate the integer result to the right size, note that TruncTy can be
214 // a pointer.
215 if (mlir::isa<cir::FPTypeInterface>(truncTy)) {
216 tmp = builder.createFloatingCast(tmp, truncTy);
217 } else if (isa<cir::PointerType>(truncTy) &&
218 isa<cir::IntType>(tmp.getType())) {
219 uint64_t resSize = cgm.getDataLayout().getTypeSizeInBits(truncTy);
220 tmp = builder.createIntCast(
221 tmp, cir::IntType::get(ctx, (unsigned)resSize, false));
222 tmp = builder.createIntToPtr(tmp, truncTy);
223 } else if (isa<cir::PointerType>(tmp.getType()) &&
224 isa<cir::IntType>(truncTy)) {
225 uint64_t tmpSize = cgm.getDataLayout().getTypeSizeInBits(tmp.getType());
226 tmp = builder.createPtrToInt(
227 tmp, cir::IntType::get(ctx, (unsigned)tmpSize, false));
228 tmp = builder.createIntCast(tmp, truncTy);
229 } else if (isa<cir::IntType>(truncTy)) {
230 tmp = builder.createIntCast(tmp, truncTy);
231 } else if (isa<cir::VectorType>(truncTy)) {
233 }
234 }
235
236 LValue dest = resultRegDests[i];
237 // ResultTypeRequiresCast elements correspond to the first
238 // ResultTypeRequiresCast.size() elements of RegResults.
239 if ((i < resultTypeRequiresCast.size()) && resultTypeRequiresCast[i]) {
240 unsigned size = cgf.getContext().getTypeSize(resultRegQualTys[i]);
241 Address addr =
242 dest.getAddress().withElementType(builder, resultRegTypes[i]);
243 if (cgm.getTargetCIRGenInfo().isScalarizableAsmOperand(cgf, truncTy)) {
244 builder.createStore(cgf.getLoc(s.getAsmLoc()), tmp, addr);
245 continue;
246 }
247
248 QualType ty =
249 cgf.getContext().getIntTypeForBitwidth(size, /*Signed=*/false);
250 if (ty.isNull()) {
251 const Expr *outExpr = s.getOutputExpr(i);
252 cgm.getDiags().Report(outExpr->getExprLoc(),
253 diag::err_store_value_to_reg);
254 return;
255 }
256 dest = cgf.makeAddrLValue(addr, ty);
257 }
258
259 cgf.emitStoreThroughLValue(RValue::get(tmp), dest);
260 }
261}
262
263mlir::LogicalResult CIRGenFunction::emitAsmStmt(const AsmStmt &s) {
264 // Assemble the final asm string.
265 std::string asmString = s.generateAsmString(getContext());
266 SourceLocation srcLoc = s.getAsmLoc();
267 mlir::Location loc = getLoc(srcLoc);
268
269 // Get all the output and input constraints together.
270 SmallVector<TargetInfo::ConstraintInfo> outputConstraintInfos;
271 SmallVector<TargetInfo::ConstraintInfo> inputConstraintInfos;
272 collectInOutConstraintInfos(*this, s, outputConstraintInfos,
273 inputConstraintInfos);
274
275 bool isGCCAsmGoto = false;
276
277 std::string constraints;
278 SmallVector<LValue> resultRegDests;
279 SmallVector<QualType> resultRegQualTys;
280 SmallVector<mlir::Type> resultRegTypes;
281 SmallVector<mlir::Type> resultTruncRegTypes;
283 SmallVector<mlir::Type> argElemTypes;
287 SmallVector<mlir::Value> inOutArgs;
288 llvm::BitVector resultTypeRequiresCast;
289 llvm::BitVector resultRegIsFlagReg;
290
291 // Keep track of input constraints.
292 std::string inOutConstraints;
293 SmallVector<mlir::Type> inOutArgTypes;
294 SmallVector<mlir::Type> inOutArgElemTypes;
295
296 // Keep track of out constraints for tied input operand.
297 SmallVector<std::string> outputConstraints;
298
299 // Keep track of defined physregs.
300 llvm::SmallSet<std::string, 8> physRegOutputs;
301
302 // An inline asm can be marked readonly if it meets the following conditions:
303 // - it doesn't have any sideeffects
304 // - it doesn't clobber memory
305 // - it doesn't return a value by-reference
306 // It can be marked readnone if it doesn't have any input memory constraints
307 // in addition to meeting the conditions listed above.
308 bool readOnly = true, readNone = true;
309
310 std::string outputConstraint;
311 for (unsigned i = 0, e = s.getNumOutputs(); i != e; ++i) {
312 TargetInfo::ConstraintInfo &info = outputConstraintInfos[i];
313
314 // Simplify the output constraint.
315 outputConstraint = s.getOutputConstraint(i);
316 outputConstraint = getTarget().simplifyConstraint(
317 StringRef(outputConstraint).drop_front());
318
319 const Expr *outExpr = s.getOutputExpr(i);
320 outExpr = outExpr->IgnoreParenNoopCasts(getContext());
321
322 std::string gccReg;
323 outputConstraint = s.addVariableConstraints(
324 outputConstraint, *outExpr, getTarget(), info.earlyClobber(),
325 [&](const Stmt *unspStmt, StringRef msg) {
326 cgm.errorUnsupported(unspStmt, msg);
327 },
328 &gccReg);
329
330 // Give an error on multiple outputs to same physreg.
331 if (!gccReg.empty() && !physRegOutputs.insert(gccReg).second)
332 cgm.error(srcLoc, "multiple outputs to hard register: " + gccReg);
333
334 outputConstraints.push_back(outputConstraint);
335 LValue dest = emitLValue(outExpr);
336
337 if (!constraints.empty())
338 constraints += ',';
339
340 // If this is a register output, then make the inline a sm return it
341 // by-value. If this is a memory result, return the value by-reference.
342 QualType qty = outExpr->getType();
343 const bool isScalarOrAggregate =
345 if (!info.allowsMemory() && isScalarOrAggregate) {
346 constraints += "=" + outputConstraint;
347 resultRegQualTys.push_back(qty);
348 resultRegDests.push_back(dest);
349
350 bool isFlagReg = llvm::StringRef(outputConstraint).starts_with("{@cc");
351 resultRegIsFlagReg.push_back(isFlagReg);
352
353 mlir::Type ty = convertTypeForMem(qty);
354 const bool requiresCast =
355 info.allowsRegister() &&
356 (cgm.getTargetCIRGenInfo().isScalarizableAsmOperand(*this, ty) ||
358
359 resultTruncRegTypes.push_back(ty);
360 resultTypeRequiresCast.push_back(requiresCast);
361
362 if (requiresCast) {
363 unsigned size = getContext().getTypeSize(qty);
364 if (size == 0)
365 cgm.error(outExpr->getExprLoc(), "output size should not be zero");
366
367 ty = cir::IntType::get(&getMLIRContext(), size, false);
368 }
369
370 resultRegTypes.push_back(ty);
371 // If this output is tied to an input, and if the input is larger, then
372 // we need to set the actual result type of the inline asm node to be the
373 // same as the input type.
374 if (info.hasMatchingInput()) {
375 unsigned inputNo;
376 for (inputNo = 0; inputNo != s.getNumInputs(); ++inputNo) {
377 TargetInfo::ConstraintInfo &input = inputConstraintInfos[inputNo];
378 if (input.hasTiedOperand() && input.getTiedOperand() == i)
379 break;
380 }
381 assert(inputNo != s.getNumInputs() && "Didn't find matching input!");
382
383 QualType inputTy = s.getInputExpr(inputNo)->getType();
384 QualType outputType = outExpr->getType();
385
386 uint64_t inputSize = getContext().getTypeSize(inputTy);
387 if (getContext().getTypeSize(outputType) < inputSize) {
388 // Form the asm to return the value as a larger integer or fp type.
389 resultRegTypes.back() = convertType(inputTy);
390 }
391
392 if (mlir::Type adjTy = cgm.getTargetCIRGenInfo().adjustInlineAsmType(
393 *this, outputConstraint, resultRegTypes.back()))
394 resultRegTypes.back() = adjTy;
395 else
396 cgm.getDiags().Report(srcLoc, diag::err_asm_invalid_type_in_input)
397 << outExpr->getType() << outputConstraint;
398 }
399
400 // Update largest vector width for any vector types.
402 } else {
403 Address destAddr = dest.getAddress();
404
405 // Matrix types in memory are represented by arrays, but accessed through
406 // vector pointers, with the alignment specified on the access operation.
407 // For inline assembly, update pointer arguments to use vector pointers.
408 // Otherwise there will be a mis-match if the matrix is also an
409 // input-argument which is represented as vector.
410 if (isa<MatrixType>(outExpr->getType().getCanonicalType()))
411 destAddr =
412 destAddr.withElementType(builder, convertType(outExpr->getType()));
413
414 argTypes.push_back(destAddr.getType());
415 argElemTypes.push_back(destAddr.getElementType());
416 outArgs.push_back(destAddr.getPointer());
417 args.push_back(destAddr.getPointer());
418 constraints += "=*";
419 constraints += outputConstraint;
420 readOnly = readNone = false;
421 }
422
423 if (info.isReadWrite()) {
424 inOutConstraints += ',';
425 const Expr *inputExpr = s.getOutputExpr(i);
426
427 // argValue: mlir::Value, argElementType: mlir::Type.
428 auto [argValue, argElementType] =
429 emitAsmInputLValue(info, dest, inputExpr->getType(), inOutConstraints,
430 inputExpr->getExprLoc());
431
432 if (mlir::Type adjTy = getTargetHooks().adjustInlineAsmType(
433 *this, outputConstraint, argValue.getType()))
434 argValue = builder.createBitcast(argValue, adjTy);
435
436 // Update largest vector width for any vector types.
438
439 // Only tie earlyclobber physregs.
440 if (info.allowsRegister() && (gccReg.empty() || info.earlyClobber()))
441 inOutConstraints += llvm::utostr(i);
442 else
443 inOutConstraints += outputConstraint;
444
445 inOutArgTypes.push_back(argValue.getType());
446 inOutArgElemTypes.push_back(argElementType);
447 inOutArgs.push_back(argValue);
448 }
449
450 } // iterate over output operands
451
452 for (unsigned i = 0, e = s.getNumInputs(); i != e; ++i) {
453 TargetInfo::ConstraintInfo &info = inputConstraintInfos[i];
454 const Expr *inputExpr = s.getInputExpr(i);
455
456 if (info.allowsMemory())
457 readNone = false;
458
459 if (!constraints.empty())
460 constraints += ',';
461
462 std::string inputConstraint(s.getInputConstraint(i));
463 inputConstraint =
464 getTarget().simplifyConstraint(inputConstraint, &outputConstraintInfos);
465
466 inputConstraint = s.addVariableConstraints(
467 inputConstraint, *inputExpr->IgnoreParenNoopCasts(getContext()),
468 getTarget(), /*EarlyClobber=*/false,
469 [&](const Stmt *unspStmt, StringRef msg) {
470 cgm.errorUnsupported(unspStmt, msg);
471 });
472
473 std::string replaceConstraint(inputConstraint);
474 // argValue: mlir::Value, argElementType: mlir::Type.
475 auto [argValue, argElemType] = emitAsmInput(info, inputExpr, constraints);
476
477 // If this input argument is tied to a larger output result, extend the
478 // input to be the same size as the output. The LLVM backend wants to see
479 // the input and output of a matching constraint be the same size. Note
480 // that GCC does not define what the top bits are here. We use zext because
481 // that is usually cheaper, but LLVM IR should really get an anyext someday.
482 if (info.hasTiedOperand()) {
483 unsigned output = info.getTiedOperand();
484 QualType outputType = s.getOutputExpr(output)->getType();
485 QualType inputTy = inputExpr->getType();
486
487 if (getContext().getTypeSize(outputType) >
488 getContext().getTypeSize(inputTy)) {
489 // Use ptrtoint as appropriate so that we can do our extension.
490 if (isa<cir::PointerType>(argValue.getType()))
491 argValue = builder.createPtrToInt(argValue, uIntPtrTy);
492 mlir::Type outputTy = convertType(outputType);
493 if (isa<cir::IntType>(outputTy))
494 argValue = builder.createIntCast(argValue, outputTy);
495 else if (isa<cir::PointerType>(outputTy))
496 argValue = builder.createIntCast(argValue, uIntPtrTy);
497 else if (isa<cir::FPTypeInterface>(outputTy))
498 argValue = builder.createFloatingCast(argValue, outputTy);
499 }
500
501 // Deal with the tied operands' constraint code in adjustInlineAsmType.
502 replaceConstraint = outputConstraints[output];
503 }
504
505 if (mlir::Type adjTy = getTargetHooks().adjustInlineAsmType(
506 *this, replaceConstraint, argValue.getType()))
507 argValue = builder.createBitcast(argValue, adjTy);
508 else
509 cgm.getDiags().Report(s.getAsmLoc(), diag::err_asm_invalid_type_in_input)
510 << inputExpr->getType() << inputConstraint;
511
512 // Update largest vector width for any vector types.
514
515 argTypes.push_back(argValue.getType());
516 argElemTypes.push_back(argElemType);
517 inArgs.push_back(argValue);
518 args.push_back(argValue);
519 constraints += inputConstraint;
520 } // iterate over input operands
521
522 // Append the "input" part of inout constraints.
523 for (unsigned i = 0, e = inOutArgs.size(); i != e; ++i) {
524 args.push_back(inOutArgs[i]);
525 argTypes.push_back(inOutArgTypes[i]);
526 argElemTypes.push_back(inOutArgElemTypes[i]);
527 }
528 constraints += inOutConstraints;
529
530 bool hasUnwindClobber = false;
531 collectClobbers(*this, s, constraints, hasUnwindClobber, readOnly, readNone);
532
533 std::array<mlir::ValueRange, 3> operands = {outArgs, inArgs, inOutArgs};
534
535 mlir::Type resultType;
536
537 if (resultRegTypes.size() == 1)
538 resultType = resultRegTypes[0];
539 else if (resultRegTypes.size() > 1)
540 resultType = builder.getAnonRecordTy(resultRegTypes, /*packed=*/false,
541 /*padded=*/false);
542
543 bool hasSideEffect = s.isVolatile() || s.getNumOutputs() == 0;
544
545 std::vector<mlir::Value> regResults;
546 cir::InlineAsmOp ia = cir::InlineAsmOp::create(
547 builder, getLoc(s.getAsmLoc()), resultType, operands, asmString,
548 constraints, hasSideEffect, inferFlavor(cgm, s), mlir::ArrayAttr());
549
550 if (isGCCAsmGoto) {
552 } else if (hasUnwindClobber) {
554 } else {
556
557 mlir::Value result;
558 if (ia.getNumResults())
559 result = ia.getResult(0);
560
562
563 int i = 0;
564 for (auto typ : argElemTypes) {
565 if (typ) {
566 auto op = args[i++];
567 assert(mlir::isa<cir::PointerType>(op.getType()) &&
568 "pointer type expected");
569 assert(cast<cir::PointerType>(op.getType()).getPointee() == typ &&
570 "element type differs from pointee type!");
571
572 operandAttrs.push_back(mlir::UnitAttr::get(&getMLIRContext()));
573 } else {
574 // We need to add an attribute for every arg since later, during
575 // the lowering to LLVM IR the attributes will be assigned to the
576 // CallInsn argument by index, i.e. we can't skip null type here
577 operandAttrs.push_back(mlir::Attribute());
578 }
579 }
580 assert(args.size() == operandAttrs.size() &&
581 "The number of attributes is not even with the number of operands");
582
583 ia.setOperandAttrsAttr(builder.getArrayAttr(operandAttrs));
584
585 if (resultRegTypes.size() == 1) {
586 regResults.push_back(result);
587 } else if (resultRegTypes.size() > 1) {
588 CharUnits alignment = CharUnits::One();
589 mlir::Value dest =
590 emitAlloca("__asm_result", resultType, loc, alignment, false);
591 Address addr = Address(dest, alignment);
592 builder.createStore(loc, result, addr);
593
594 for (unsigned i = 0, e = resultRegTypes.size(); i != e; ++i) {
595 cir::PointerType typ = builder.getPointerTo(resultRegTypes[i]);
596 cir::GetMemberOp ptr = builder.createGetMember(loc, typ, dest, "", i);
597 cir::LoadOp tmp = builder.createLoad(loc, Address(ptr, alignment));
598 regResults.push_back(tmp);
599 }
600 }
601 }
602
603 emitAsmStores(*this, s, regResults, resultRegTypes, resultTruncRegTypes,
604 resultRegDests, resultRegQualTys, resultTypeRequiresCast,
605 resultRegIsFlagReg);
606
607 return mlir::success();
608}
static AsmFlavor inferFlavor(const CIRGenModule &cgm, const AsmStmt &s)
Definition CIRGenAsm.cpp:22
static void collectInOutConstraintInfos(const CIRGenFunction &cgf, const AsmStmt &s, SmallVectorImpl< TargetInfo::ConstraintInfo > &out, SmallVectorImpl< TargetInfo::ConstraintInfo > &in)
Definition CIRGenAsm.cpp:89
static void collectClobbers(const CIRGenFunction &cgf, const AsmStmt &s, std::string &constraints, bool &hasUnwindClobber, bool &readOnly, bool readNone)
Definition CIRGenAsm.cpp:31
static void emitAsmStores(CIRGenFunction &cgf, const AsmStmt &s, const llvm::ArrayRef< mlir::Value > regResults, const llvm::ArrayRef< mlir::Type > resultRegTypes, const llvm::ArrayRef< mlir::Type > resultTruncRegTypes, const llvm::ArrayRef< LValue > resultRegDests, const llvm::ArrayRef< QualType > resultRegQualTys, const llvm::BitVector &resultTypeRequiresCast, const llvm::BitVector &resultRegIsFlagReg)
__device__ __2f16 float __ockl_bool s
mlir::Value createIntToPtr(mlir::Value src, mlir::Type newTy)
mlir::Value createPtrToInt(mlir::Value src, mlir::Type newTy)
mlir::Value createIntCast(mlir::Value src, mlir::Type newTy)
llvm::TypeSize getTypeSizeInBits(mlir::Type ty) const
APSInt & getInt()
Definition APValue.h:489
bool toIntegralConstant(APSInt &Result, QualType SrcTy, const ASTContext &Ctx) const
Try to convert this value to an integral constant.
Definition APValue.cpp:963
QualType getIntTypeForBitwidth(unsigned DestWidth, unsigned Signed) const
getIntTypeForBitwidth - sets integer QualTy according to specified details: bitwidth,...
uint64_t getTypeSize(QualType T) const
Return the size of the specified (complete) type T, in bits.
AsmStmt is the base class for GCCAsmStmt and MSAsmStmt.
Definition Stmt.h:3269
mlir::Value getPointer() const
Definition Address.h:95
mlir::Type getElementType() const
Definition Address.h:122
Address withElementType(CIRGenBuilderTy &builder, mlir::Type ElemTy) const
Return address with different element type, a bitcast pointer, and the same alignment.
mlir::Type getType() const
Definition Address.h:114
mlir::Value createFloatingCast(mlir::Value v, mlir::Type destType)
cir::StoreOp createStore(mlir::Location loc, mlir::Value val, Address dst, bool isVolatile=false, mlir::IntegerAttr align={}, cir::SyncScopeKindAttr scope={}, cir::MemOrderAttr order={})
static bool hasScalarEvaluationKind(clang::QualType type)
mlir::Type convertType(clang::QualType t)
RValue emitLoadOfLValue(LValue lv, SourceLocation loc)
Given an expression that represents a value lvalue, this method emits the address of the lvalue,...
const TargetInfo & getTarget() const
LValue emitLValue(const clang::Expr *e)
Emit code to compute a designator that specifies the location of the expression.
mlir::Location getLoc(clang::SourceLocation srcLoc)
Helpers to convert Clang's SourceLocation to a MLIR Location.
const TargetCIRGenInfo & getTargetHooks() const
mlir::LogicalResult emitAsmStmt(const clang::AsmStmt &s)
std::pair< mlir::Value, mlir::Type > emitAsmInputLValue(const TargetInfo::ConstraintInfo &info, LValue inputValue, QualType inputType, std::string &constraintString, SourceLocation loc)
mlir::Type convertTypeForMem(QualType t)
mlir::Value emitAlloca(llvm::StringRef name, mlir::Type ty, mlir::Location loc, clang::CharUnits alignment, bool insertIntoFnEntryBlock, mlir::Value arraySize=nullptr)
static bool hasAggregateEvaluationKind(clang::QualType type)
mlir::Value emitScalarExpr(const clang::Expr *e, bool ignoreResultAssign=false)
Emit the computation of the specified expression of scalar type.
CIRGenBuilderTy & getBuilder()
mlir::MLIRContext & getMLIRContext()
LValue makeAddrLValue(Address addr, QualType ty, AlignmentSource source=AlignmentSource::Type)
std::pair< mlir::Value, mlir::Type > emitAsmInput(const TargetInfo::ConstraintInfo &info, const Expr *inputExpr, std::string &constraintString)
clang::ASTContext & getContext() const
void emitStoreThroughLValue(RValue src, LValue dst, bool isInit=false)
Store the specified rvalue into the specified lvalue, where both are guaranteed to the have the same ...
This class organizes the cross-function state that is used while generating CIR code.
clang::DiagnosticsEngine & getDiags() const
const cir::CIRDataLayout getDataLayout() const
const clang::CodeGenOptions & getCodeGenOpts() const
const TargetCIRGenInfo & getTargetCIRGenInfo()
Address getAddress() const
static RValue get(mlir::Value v)
Definition CIRGenValue.h:83
mlir::Value getValue() const
Return the value of this scalar value.
Definition CIRGenValue.h:57
virtual bool isScalarizableAsmOperand(CIRGenFunction &cgf, mlir::Type ty) const
Definition TargetInfo.h:108
CharUnits - This is an opaque type for sizes expressed in character units.
Definition CharUnits.h:38
static CharUnits One()
One - Construct a CharUnits quantity of one.
Definition CharUnits.h:58
DiagnosticBuilder Report(SourceLocation Loc, unsigned DiagID)
Issue the message to the client.
This represents one expression.
Definition Expr.h:112
bool EvaluateAsInt(EvalResult &Result, const ASTContext &Ctx, SideEffectsKind AllowSideEffects=SE_NoSideEffects, bool InConstantContext=false) const
EvaluateAsInt - Return true if this is a constant which we can fold and convert to an integer,...
Expr * IgnoreParenNoopCasts(const ASTContext &Ctx) LLVM_READONLY
Skip past any parentheses and casts which do not change the value (including ptr->int casts of the sa...
Definition Expr.cpp:3116
bool EvaluateAsRValue(EvalResult &Result, const ASTContext &Ctx, bool InConstantContext=false) const
EvaluateAsRValue - Return true if this is a constant which we can fold to an rvalue using any crazy t...
SourceLocation getExprLoc() const LLVM_READONLY
getExprLoc - Return the preferred location for the arrow when diagnosing a problem with a generic exp...
Definition Expr.cpp:276
QualType getType() const
Definition Expr.h:144
This represents a GCC inline-assembly statement extension.
Definition Stmt.h:3438
A (possibly-)qualified type.
Definition TypeBase.h:937
bool isNull() const
Return true if this QualType doesn't point to a type yet.
Definition TypeBase.h:1004
QualType getCanonicalType() const
Definition TypeBase.h:8344
Encodes a location in the source.
Stmt - This represents one statement.
Definition Stmt.h:86
StmtClass getStmtClass() const
Definition Stmt.h:1485
bool validateInputConstraint(MutableArrayRef< ConstraintInfo > OutputConstraints, ConstraintInfo &info) const
std::string simplifyConstraint(StringRef Constraint, SmallVectorImpl< ConstraintInfo > *OutCons=nullptr) const
StringRef getNormalizedGCCRegisterName(StringRef Name, bool ReturnCanonical=false) const
Returns the "normalized" GCC register name.
bool validateOutputConstraint(ConstraintInfo &Info) const
virtual std::string_view getClobbers() const =0
Returns a string of target-specific clobbers, in LLVM format.
virtual bool isSPRegName(StringRef) const
The JSON file list parser is used to communicate input to InstallAPI.
bool isa(CodeGen::Address addr)
Definition Address.h:330
U cast(CodeGen::Address addr)
Definition Address.h:327
static bool asmMemoryEffects()
static bool asmVectorType()
static bool asmUnwindClobber()
static bool asmLLVMAssume()
EvalResult is a struct with detailed info about an evaluated expression.
Definition Expr.h:645
APValue Val
Val - This is the value the expression can be folded to.
Definition Expr.h:647
bool hasMatchingInput() const
Return true if this output operand has a matching (tied) input operand.
bool hasTiedOperand() const
Return true if this input operand is a matching constraint that ties it to an output operand.