MLIR 24.0.0git
MPIToLLVM.cpp
Go to the documentation of this file.
1//===- MPIToLLVM.cpp - MPI to LLVM dialect conversion ---------------------===//
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//
10// Copyright (C) by Argonne National Laboratory
11// See COPYRIGHT in top-level directory
12// of MPICH source repository.
13//
14
15#include "mlir/Conversion/MPIToLLVM/MPIToLLVM.h"
16#include "mlir/Conversion/ConvertToLLVM/ToLLVMInterface.h"
17#include "mlir/Conversion/LLVMCommon/Pattern.h"
18#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h"
19#include "mlir/Dialect/DLTI/DLTI.h"
20#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
21#include "mlir/Dialect/LLVMIR/LLVMTypes.h"
22#include "mlir/Dialect/MPI/IR/MPI.h"
23#include "mlir/Dialect/MPI/IR/Utils.h"
24#include "mlir/Transforms/DialectConversion.h"
25#include <memory>
26
27using namespace mlir;
28
29namespace {
30
31template <typename Op, typename... Args>
32static Op getOrDefineGlobal(ModuleOp &moduleOp, const Location loc,
33 ConversionPatternRewriter &rewriter, StringRef name,
34 Args &&...args) {
35 Op ret;
36 if (!(ret = moduleOp.lookupSymbol<Op>(name))) {
37 ConversionPatternRewriter::InsertionGuard guard(rewriter);
38 rewriter.setInsertionPointToStart(moduleOp.getBody());
39 ret = Op::create(rewriter, loc, std::forward<Args>(args)...);
40 }
41 return ret;
42}
43
44static LLVM::LLVMFuncOp getOrDefineFunction(ModuleOp &moduleOp,
45 const Location loc,
46 ConversionPatternRewriter &rewriter,
47 StringRef name,
48 LLVM::LLVMFunctionType type) {
49 return getOrDefineGlobal<LLVM::LLVMFuncOp>(
50 moduleOp, loc, rewriter, name, name, type, LLVM::Linkage::External);
51}
52
53std::pair<Value, Value> getRawPtrAndSize(const Location loc,
54 ConversionPatternRewriter &rewriter,
55 Value memRef, int64_t rank,
56 Type elType) {
57 Type ptrType = LLVM::LLVMPointerType::get(rewriter.getContext());
58 Type i32Type = rewriter.getI32Type();
59 auto descriptorType = cast<LLVM::LLVMStructType>(memRef.getType());
60 // The offset and the sizes of a memref descriptor have the converted index
61 // type, which is not necessarily `i64`. Take it from the descriptor itself.
62 auto indexType = cast<IntegerType>(descriptorType.getBody()[2]);
63
64 Value dataPtr =
65 LLVM::ExtractValueOp::create(rewriter, loc, ptrType, memRef, 1);
66 Value offset =
67 LLVM::ExtractValueOp::create(rewriter, loc, indexType, memRef, 2);
68 Value resPtr =
69 LLVM::GEPOp::create(rewriter, loc, ptrType, elType, dataPtr, offset);
70 Value size = LLVM::ConstantOp::create(rewriter, loc, i32Type,
71 rewriter.getI32IntegerAttr(1));
72 if (descriptorType.getBody().size() > 3) {
73 for (int64_t i = 0; i < rank; ++i) {
74 Value dim = LLVM::ExtractValueOp::create(rewriter, loc, memRef,
75 ArrayRef<int64_t>{3, i});
76 // The MPI interface counts elements in an `i32`, so adjust the
77 // index-typed extent to that width. Extents are non-negative, hence the
78 // zero extension.
79 if (indexType.getWidth() > 32)
80 dim = LLVM::TruncOp::create(rewriter, loc, i32Type, dim);
81 else if (indexType.getWidth() < 32)
82 dim = LLVM::ZExtOp::create(rewriter, loc, i32Type, dim);
83 size = LLVM::MulOp::create(rewriter, loc, i32Type, dim, size);
84 }
85 }
86 return {resPtr, size};
87}
88
89/// When lowering the mpi dialect to functions calls certain details
90/// differ between various MPI implementations. This class will provide
91/// these in a generic way, depending on the MPI implementation that got
92/// selected by the DLTI attribute on the module.
93class MPIImplTraits {
94 ModuleOp &moduleOp;
95
96public:
97 /// Instantiate a new MPIImplTraits object according to the DLTI attribute
98 /// on the given module. Default to MPICH if no attribute is present or
99 /// the value is unknown.
100 static std::unique_ptr<MPIImplTraits> get(ModuleOp &moduleOp);
101
102 explicit MPIImplTraits(ModuleOp &moduleOp) : moduleOp(moduleOp) {}
103
104 virtual ~MPIImplTraits() = default;
105
106 ModuleOp &getModuleOp() { return moduleOp; }
107
108 /// Gets or creates MPI_COMM_WORLD as a Value.
109 /// Different MPI implementations have different communicator types.
110 /// Using i64 as a portable, intermediate type.
111 /// Appropriate cast needs to take place before calling MPI functions.
112 virtual Value getCommWorld(Location loc,
113 ConversionPatternRewriter &rewriter) = 0;
114
115 /// Type converter provides i64 type for communicator type.
116 /// Converts to native type, which might be ptr or int or whatever.
117 virtual Value castComm(Location loc, ConversionPatternRewriter &rewriter,
118 Value comm) = 0;
119
120 /// Get the MPI_STATUS_IGNORE value (typically a pointer type).
121 virtual intptr_t getStatusIgnore() = 0;
122
123 /// Get the MPI_IN_PLACE value (void *).
124 virtual void *getInPlace() = 0;
125
126 /// Gets or creates an MPI datatype as a value which corresponds to the given
127 /// type.
128 virtual Value getDataType(Location loc, ConversionPatternRewriter &rewriter,
129 Type type) = 0;
130
131 /// Gets or creates an MPI_Op value which corresponds to the given
132 /// enum value.
133 virtual Value getMPIOp(Location loc, ConversionPatternRewriter &rewriter,
134 mpi::MPI_ReductionOpEnum opAttr) = 0;
135};
136
137//===----------------------------------------------------------------------===//
138// Implementation details for MPICH ABI compatible MPI implementations
139//===----------------------------------------------------------------------===//
140
141class MPICHImplTraits : public MPIImplTraits {
142 static constexpr int MPI_FLOAT = 0x4c00040a;
143 static constexpr int MPI_DOUBLE = 0x4c00080b;
144 static constexpr int MPI_INT8_T = 0x4c000137;
145 static constexpr int MPI_INT16_T = 0x4c000238;
146 static constexpr int MPI_INT32_T = 0x4c000439;
147 static constexpr int MPI_INT64_T = 0x4c00083a;
148 static constexpr int MPI_UINT8_T = 0x4c00013b;
149 static constexpr int MPI_UINT16_T = 0x4c00023c;
150 static constexpr int MPI_UINT32_T = 0x4c00043d;
151 static constexpr int MPI_UINT64_T = 0x4c00083e;
152 static constexpr int MPI_MAX = 0x58000001;
153 static constexpr int MPI_MIN = 0x58000002;
154 static constexpr int MPI_SUM = 0x58000003;
155 static constexpr int MPI_PROD = 0x58000004;
156 static constexpr int MPI_LAND = 0x58000005;
157 static constexpr int MPI_BAND = 0x58000006;
158 static constexpr int MPI_LOR = 0x58000007;
159 static constexpr int MPI_BOR = 0x58000008;
160 static constexpr int MPI_LXOR = 0x58000009;
161 static constexpr int MPI_BXOR = 0x5800000a;
162 static constexpr int MPI_MINLOC = 0x5800000b;
163 static constexpr int MPI_MAXLOC = 0x5800000c;
164 static constexpr int MPI_REPLACE = 0x5800000d;
165 static constexpr int MPI_NO_OP = 0x5800000e;
166
167public:
168 using MPIImplTraits::MPIImplTraits;
169
170 ~MPICHImplTraits() override = default;
171
172 Value getCommWorld(const Location loc,
173 ConversionPatternRewriter &rewriter) override {
174 static constexpr int MPI_COMM_WORLD = 0x44000000;
175 return LLVM::ConstantOp::create(rewriter, loc, rewriter.getI64Type(),
176 MPI_COMM_WORLD);
177 }
178
179 Value castComm(const Location loc, ConversionPatternRewriter &rewriter,
180 Value comm) override {
181 return LLVM::TruncOp::create(rewriter, loc, rewriter.getI32Type(), comm);
182 }
183
184 intptr_t getStatusIgnore() override { return 1; }
185
186 void *getInPlace() override { return reinterpret_cast<void *>(-1); }
187
188 Value getDataType(const Location loc, ConversionPatternRewriter &rewriter,
189 Type type) override {
190 int32_t mtype = 0;
191 if (type.isF32())
192 mtype = MPI_FLOAT;
193 else if (type.isF64())
194 mtype = MPI_DOUBLE;
195 else if (type.isInteger(64) && !type.isUnsignedInteger())
196 mtype = MPI_INT64_T;
197 else if (type.isInteger(64))
198 mtype = MPI_UINT64_T;
199 else if (type.isInteger(32) && !type.isUnsignedInteger())
200 mtype = MPI_INT32_T;
201 else if (type.isInteger(32))
202 mtype = MPI_UINT32_T;
203 else if (type.isInteger(16) && !type.isUnsignedInteger())
204 mtype = MPI_INT16_T;
205 else if (type.isInteger(16))
206 mtype = MPI_UINT16_T;
207 else if (type.isInteger(8) && !type.isUnsignedInteger())
208 mtype = MPI_INT8_T;
209 else if (type.isInteger(8))
210 mtype = MPI_UINT8_T;
211 else
212 assert(false && "unsupported type");
213 return LLVM::ConstantOp::create(rewriter, loc, rewriter.getI32Type(),
214 mtype);
215 }
216
217 Value getMPIOp(const Location loc, ConversionPatternRewriter &rewriter,
218 mpi::MPI_ReductionOpEnum opAttr) override {
219 int32_t op = MPI_NO_OP;
220 switch (opAttr) {
221 case mpi::MPI_ReductionOpEnum::MPI_OP_NULL:
222 op = MPI_NO_OP;
223 break;
224 case mpi::MPI_ReductionOpEnum::MPI_MAX:
225 op = MPI_MAX;
226 break;
227 case mpi::MPI_ReductionOpEnum::MPI_MIN:
228 op = MPI_MIN;
229 break;
230 case mpi::MPI_ReductionOpEnum::MPI_SUM:
231 op = MPI_SUM;
232 break;
233 case mpi::MPI_ReductionOpEnum::MPI_PROD:
234 op = MPI_PROD;
235 break;
236 case mpi::MPI_ReductionOpEnum::MPI_LAND:
237 op = MPI_LAND;
238 break;
239 case mpi::MPI_ReductionOpEnum::MPI_BAND:
240 op = MPI_BAND;
241 break;
242 case mpi::MPI_ReductionOpEnum::MPI_LOR:
243 op = MPI_LOR;
244 break;
245 case mpi::MPI_ReductionOpEnum::MPI_BOR:
246 op = MPI_BOR;
247 break;
248 case mpi::MPI_ReductionOpEnum::MPI_LXOR:
249 op = MPI_LXOR;
250 break;
251 case mpi::MPI_ReductionOpEnum::MPI_BXOR:
252 op = MPI_BXOR;
253 break;
254 case mpi::MPI_ReductionOpEnum::MPI_MINLOC:
255 op = MPI_MINLOC;
256 break;
257 case mpi::MPI_ReductionOpEnum::MPI_MAXLOC:
258 op = MPI_MAXLOC;
259 break;
260 case mpi::MPI_ReductionOpEnum::MPI_REPLACE:
261 op = MPI_REPLACE;
262 break;
263 }
264 return LLVM::ConstantOp::create(rewriter, loc, rewriter.getI32Type(), op);
265 }
266};
267
268//===----------------------------------------------------------------------===//
269// Implementation details for OpenMPI
270//===----------------------------------------------------------------------===//
271class OMPIImplTraits : public MPIImplTraits {
272 LLVM::GlobalOp getOrDefineExternalStruct(const Location loc,
273 ConversionPatternRewriter &rewriter,
274 StringRef name,
275 LLVM::LLVMStructType type) {
276
277 return getOrDefineGlobal<LLVM::GlobalOp>(
278 getModuleOp(), loc, rewriter, name, type, /*isConstant=*/false,
279 LLVM::Linkage::External, name,
280 /*value=*/Attribute(), /*alignment=*/0, 0);
281 }
282
283public:
284 using MPIImplTraits::MPIImplTraits;
285
286 ~OMPIImplTraits() override = default;
287
288 Value getCommWorld(const Location loc,
289 ConversionPatternRewriter &rewriter) override {
290 auto *context = rewriter.getContext();
291 // get external opaque struct pointer type
292 auto commStructT =
293 LLVM::LLVMStructType::getOpaque("ompi_communicator_t", context);
294 StringRef name = "ompi_mpi_comm_world";
295
296 // make sure global op definition exists
297 getOrDefineExternalStruct(loc, rewriter, name, commStructT);
298
299 // get address of symbol
300 auto comm = LLVM::AddressOfOp::create(rewriter, loc,
301 LLVM::LLVMPointerType::get(context),
302 SymbolRefAttr::get(context, name));
303 return LLVM::PtrToIntOp::create(rewriter, loc, rewriter.getI64Type(), comm);
304 }
305
306 Value castComm(const Location loc, ConversionPatternRewriter &rewriter,
307 Value comm) override {
308 return LLVM::IntToPtrOp::create(
309 rewriter, loc, LLVM::LLVMPointerType::get(rewriter.getContext()), comm);
310 }
311
312 intptr_t getStatusIgnore() override { return 0; }
313
314 void *getInPlace() override { return reinterpret_cast<void *>(1); }
315
316 Value getDataType(const Location loc, ConversionPatternRewriter &rewriter,
317 Type type) override {
318 StringRef mtype;
319 if (type.isF32())
320 mtype = "ompi_mpi_float";
321 else if (type.isF64())
322 mtype = "ompi_mpi_double";
323 else if (type.isInteger(64) && !type.isUnsignedInteger())
324 mtype = "ompi_mpi_int64_t";
325 else if (type.isInteger(64))
326 mtype = "ompi_mpi_uint64_t";
327 else if (type.isInteger(32) && !type.isUnsignedInteger())
328 mtype = "ompi_mpi_int32_t";
329 else if (type.isInteger(32))
330 mtype = "ompi_mpi_uint32_t";
331 else if (type.isInteger(16) && !type.isUnsignedInteger())
332 mtype = "ompi_mpi_int16_t";
333 else if (type.isInteger(16))
334 mtype = "ompi_mpi_uint16_t";
335 else if (type.isInteger(8) && !type.isUnsignedInteger())
336 mtype = "ompi_mpi_int8_t";
337 else if (type.isInteger(8))
338 mtype = "ompi_mpi_uint8_t";
339 else
340 assert(false && "unsupported type");
341
342 auto *context = rewriter.getContext();
343 // get external opaque struct pointer type
344 auto typeStructT =
345 LLVM::LLVMStructType::getOpaque("ompi_predefined_datatype_t", context);
346 // make sure global op definition exists
347 getOrDefineExternalStruct(loc, rewriter, mtype, typeStructT);
348 // get address of symbol
349 return LLVM::AddressOfOp::create(rewriter, loc,
350 LLVM::LLVMPointerType::get(context),
351 SymbolRefAttr::get(context, mtype));
352 }
353
354 Value getMPIOp(const Location loc, ConversionPatternRewriter &rewriter,
355 mpi::MPI_ReductionOpEnum opAttr) override {
356 StringRef op;
357 switch (opAttr) {
358 case mpi::MPI_ReductionOpEnum::MPI_OP_NULL:
359 op = "ompi_mpi_no_op";
360 break;
361 case mpi::MPI_ReductionOpEnum::MPI_MAX:
362 op = "ompi_mpi_max";
363 break;
364 case mpi::MPI_ReductionOpEnum::MPI_MIN:
365 op = "ompi_mpi_min";
366 break;
367 case mpi::MPI_ReductionOpEnum::MPI_SUM:
368 op = "ompi_mpi_sum";
369 break;
370 case mpi::MPI_ReductionOpEnum::MPI_PROD:
371 op = "ompi_mpi_prod";
372 break;
373 case mpi::MPI_ReductionOpEnum::MPI_LAND:
374 op = "ompi_mpi_land";
375 break;
376 case mpi::MPI_ReductionOpEnum::MPI_BAND:
377 op = "ompi_mpi_band";
378 break;
379 case mpi::MPI_ReductionOpEnum::MPI_LOR:
380 op = "ompi_mpi_lor";
381 break;
382 case mpi::MPI_ReductionOpEnum::MPI_BOR:
383 op = "ompi_mpi_bor";
384 break;
385 case mpi::MPI_ReductionOpEnum::MPI_LXOR:
386 op = "ompi_mpi_lxor";
387 break;
388 case mpi::MPI_ReductionOpEnum::MPI_BXOR:
389 op = "ompi_mpi_bxor";
390 break;
391 case mpi::MPI_ReductionOpEnum::MPI_MINLOC:
392 op = "ompi_mpi_minloc";
393 break;
394 case mpi::MPI_ReductionOpEnum::MPI_MAXLOC:
395 op = "ompi_mpi_maxloc";
396 break;
397 case mpi::MPI_ReductionOpEnum::MPI_REPLACE:
398 op = "ompi_mpi_replace";
399 break;
400 }
401 auto *context = rewriter.getContext();
402 // get external opaque struct pointer type
403 auto opStructT =
404 LLVM::LLVMStructType::getOpaque("ompi_predefined_op_t", context);
405 // make sure global op definition exists
406 getOrDefineExternalStruct(loc, rewriter, op, opStructT);
407 // get address of symbol
408 return LLVM::AddressOfOp::create(rewriter, loc,
409 LLVM::LLVMPointerType::get(context),
410 SymbolRefAttr::get(context, op));
411 }
412};
413
414std::unique_ptr<MPIImplTraits> MPIImplTraits::get(ModuleOp &moduleOp) {
415 auto attr = dlti::query(moduleOp, {"MPI:Implementation"}, false);
416 if (failed(attr))
417 return std::make_unique<MPICHImplTraits>(moduleOp);
418 auto strAttr = dyn_cast<StringAttr>(attr.value());
419 if (strAttr && strAttr.getValue() == "OpenMPI")
420 return std::make_unique<OMPIImplTraits>(moduleOp);
421 if (!strAttr || strAttr.getValue() != "MPICH")
422 moduleOp.emitWarning() << "Unknown \"MPI:Implementation\" value in DLTI ("
423 << (strAttr ? strAttr.getValue() : "<NULL>")
424 << "), defaulting to MPICH";
425 return std::make_unique<MPICHImplTraits>(moduleOp);
426}
427
428//===----------------------------------------------------------------------===//
429// InitOpLowering
430//===----------------------------------------------------------------------===//
431
432struct InitOpLowering : public ConvertOpToLLVMPattern<mpi::InitOp> {
433 using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern;
434
435 LogicalResult
436 matchAndRewrite(mpi::InitOp op, OpAdaptor adaptor,
437 ConversionPatternRewriter &rewriter) const override {
438 Location loc = op.getLoc();
439
440 // ptrType `!llvm.ptr`
441 Type ptrType = LLVM::LLVMPointerType::get(rewriter.getContext());
442
443 // instantiate nullptr `%nullptr = llvm.mlir.zero : !llvm.ptr`
444 auto nullPtrOp = LLVM::ZeroOp::create(rewriter, loc, ptrType);
445 Value llvmnull = nullPtrOp.getRes();
446
447 // grab a reference to the global module op:
448 auto moduleOp = op->getParentOfType<ModuleOp>();
449
450 // LLVM Function type representing `i32 MPI_Init(ptr, ptr)`
451 auto initFuncType =
452 LLVM::LLVMFunctionType::get(rewriter.getI32Type(), {ptrType, ptrType});
453 // get or create function declaration:
454 LLVM::LLVMFuncOp initDecl =
455 getOrDefineFunction(moduleOp, loc, rewriter, "MPI_Init", initFuncType);
456
457 // replace init with function call
458 rewriter.replaceOpWithNewOp<LLVM::CallOp>(op, initDecl,
459 ValueRange{llvmnull, llvmnull});
460
461 return success();
462 }
463};
464
465//===----------------------------------------------------------------------===//
466// FinalizeOpLowering
467//===----------------------------------------------------------------------===//
468
469struct FinalizeOpLowering : public ConvertOpToLLVMPattern<mpi::FinalizeOp> {
470 using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern;
471
472 LogicalResult
473 matchAndRewrite(mpi::FinalizeOp op, OpAdaptor adaptor,
474 ConversionPatternRewriter &rewriter) const override {
475 // get loc
476 Location loc = op.getLoc();
477
478 // grab a reference to the global module op:
479 auto moduleOp = op->getParentOfType<ModuleOp>();
480
481 // LLVM Function type representing `i32 MPI_Finalize()`
482 auto initFuncType = LLVM::LLVMFunctionType::get(rewriter.getI32Type(), {});
483 // get or create function declaration:
484 LLVM::LLVMFuncOp initDecl = getOrDefineFunction(
485 moduleOp, loc, rewriter, "MPI_Finalize", initFuncType);
486
487 // replace init with function call
488 rewriter.replaceOpWithNewOp<LLVM::CallOp>(op, initDecl, ValueRange{});
489
490 return success();
491 }
492};
493
494//===----------------------------------------------------------------------===//
495// CommWorldOpLowering
496//===----------------------------------------------------------------------===//
497
498struct CommWorldOpLowering : public ConvertOpToLLVMPattern<mpi::CommWorldOp> {
499 using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern;
500
501 LogicalResult
502 matchAndRewrite(mpi::CommWorldOp op, OpAdaptor adaptor,
503 ConversionPatternRewriter &rewriter) const override {
504 // grab a reference to the global module op:
505 auto moduleOp = op->getParentOfType<ModuleOp>();
506 auto mpiTraits = MPIImplTraits::get(moduleOp);
507 // get MPI_COMM_WORLD
508 rewriter.replaceOp(op, mpiTraits->getCommWorld(op.getLoc(), rewriter));
509
510 return success();
511 }
512};
513
514//===----------------------------------------------------------------------===//
515// CommSplitOpLowering
516//===----------------------------------------------------------------------===//
517
518struct CommSplitOpLowering : public ConvertOpToLLVMPattern<mpi::CommSplitOp> {
519 using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern;
520
521 LogicalResult
522 matchAndRewrite(mpi::CommSplitOp op, OpAdaptor adaptor,
523 ConversionPatternRewriter &rewriter) const override {
524 // grab a reference to the global module op:
525 auto moduleOp = op->getParentOfType<ModuleOp>();
526 auto mpiTraits = MPIImplTraits::get(moduleOp);
527 Type i32 = rewriter.getI32Type();
528 Type ptrType = LLVM::LLVMPointerType::get(op->getContext());
529 Location loc = op.getLoc();
530
531 // get communicator
532 Value comm = mpiTraits->castComm(loc, rewriter, adaptor.getComm());
533 auto one = LLVM::ConstantOp::create(rewriter, loc, i32, 1);
534 auto outPtr =
535 LLVM::AllocaOp::create(rewriter, loc, ptrType, comm.getType(), one);
536
537 // int MPI_Comm_split(MPI_Comm comm, int color, int key, MPI_Comm * newcomm)
538 auto funcType =
539 LLVM::LLVMFunctionType::get(i32, {comm.getType(), i32, i32, ptrType});
540 // get or create function declaration:
541 LLVM::LLVMFuncOp funcDecl = getOrDefineFunction(moduleOp, loc, rewriter,
542 "MPI_Comm_split", funcType);
543
544 auto callOp =
545 LLVM::CallOp::create(rewriter, loc, funcDecl,
546 ValueRange{comm, adaptor.getColor(),
547 adaptor.getKey(), outPtr.getRes()});
548
549 // load the communicator into a register
550 Value res = LLVM::LoadOp::create(rewriter, loc, i32, outPtr.getResult());
551 res = LLVM::SExtOp::create(rewriter, loc, rewriter.getI64Type(), res);
552
553 // if retval is checked, replace uses of retval with the results from the
554 // call op
555 SmallVector<Value> replacements;
556 if (op.getRetval())
557 replacements.push_back(callOp.getResult());
558
559 // replace op
560 replacements.push_back(res);
561 rewriter.replaceOp(op, replacements);
562
563 return success();
564 }
565};
566
567//===----------------------------------------------------------------------===//
568// CommRankOpLowering
569//===----------------------------------------------------------------------===//
570
571struct CommRankOpLowering : public ConvertOpToLLVMPattern<mpi::CommRankOp> {
572 using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern;
573
574 LogicalResult
575 matchAndRewrite(mpi::CommRankOp op, OpAdaptor adaptor,
576 ConversionPatternRewriter &rewriter) const override {
577 // get some helper vars
578 Location loc = op.getLoc();
579 MLIRContext *context = rewriter.getContext();
580 Type i32 = rewriter.getI32Type();
581
582 // ptrType `!llvm.ptr`
583 Type ptrType = LLVM::LLVMPointerType::get(context);
584
585 // grab a reference to the global module op:
586 auto moduleOp = op->getParentOfType<ModuleOp>();
587
588 auto mpiTraits = MPIImplTraits::get(moduleOp);
589 // get communicator
590 Value comm = mpiTraits->castComm(loc, rewriter, adaptor.getComm());
591
592 // LLVM Function type representing `i32 MPI_Comm_rank(ptr, ptr)`
593 auto rankFuncType =
594 LLVM::LLVMFunctionType::get(i32, {comm.getType(), ptrType});
595 // get or create function declaration:
596 LLVM::LLVMFuncOp initDecl = getOrDefineFunction(
597 moduleOp, loc, rewriter, "MPI_Comm_rank", rankFuncType);
598
599 // replace with function call
600 auto one = LLVM::ConstantOp::create(rewriter, loc, i32, 1);
601 auto rankptr = LLVM::AllocaOp::create(rewriter, loc, ptrType, i32, one);
602 auto callOp = LLVM::CallOp::create(rewriter, loc, initDecl,
603 ValueRange{comm, rankptr.getRes()});
604
605 // load the rank into a register
606 auto loadedRank =
607 LLVM::LoadOp::create(rewriter, loc, i32, rankptr.getResult());
608
609 // if retval is checked, replace uses of retval with the results from the
610 // call op
611 SmallVector<Value> replacements;
612 if (op.getRetval())
613 replacements.push_back(callOp.getResult());
614
615 // replace all uses, then erase op
616 replacements.push_back(loadedRank.getRes());
617 rewriter.replaceOp(op, replacements);
618
619 return success();
620 }
621};
622
623//===----------------------------------------------------------------------===//
624// CommSizeOpLowering
625//===----------------------------------------------------------------------===//
626
627static Value createOrFoldCommSize(ConversionPatternRewriter &rewriter,
628 Location loc, Value commOrg,
629 Value commAdapt) {
630 auto i32 = rewriter.getI32Type();
631 auto nRanksOp = mpi::CommSizeOp::create(rewriter, loc, i32, commOrg);
632 if (succeeded(FoldToDLTIConst(nRanksOp, "MPI:comm_world_size", rewriter)))
633 return nRanksOp.getSize();
634 rewriter.eraseOp(nRanksOp);
635 return mpi::CommSizeOp::create(rewriter, loc, i32, commAdapt).getSize();
636}
637
638struct CommSizeOpLowering : public ConvertOpToLLVMPattern<mpi::CommSizeOp> {
639 using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern;
640
641 LogicalResult
642 matchAndRewrite(mpi::CommSizeOp op, OpAdaptor adaptor,
643 ConversionPatternRewriter &rewriter) const override {
644 // get some helper vars
645 Location loc = op.getLoc();
646 MLIRContext *context = rewriter.getContext();
647 Type i32 = rewriter.getI32Type();
648
649 // ptrType `!llvm.ptr`
650 Type ptrType = LLVM::LLVMPointerType::get(context);
651
652 // grab a reference to the global module op:
653 auto moduleOp = op->getParentOfType<ModuleOp>();
654
655 auto mpiTraits = MPIImplTraits::get(moduleOp);
656 // get communicator
657 Value comm = mpiTraits->castComm(loc, rewriter, adaptor.getComm());
658
659 // LLVM Function type representing `i32 MPI_Comm_rank(ptr, ptr)`
660 auto SizeFuncType =
661 LLVM::LLVMFunctionType::get(i32, {comm.getType(), ptrType});
662 // get or create function declaration:
663 LLVM::LLVMFuncOp initDecl = getOrDefineFunction(
664 moduleOp, loc, rewriter, "MPI_Comm_size", SizeFuncType);
665
666 // replace with function call
667 auto one = LLVM::ConstantOp::create(rewriter, loc, i32, 1);
668 auto sizeptr = LLVM::AllocaOp::create(rewriter, loc, ptrType, i32, one);
669 auto callOp = LLVM::CallOp::create(rewriter, loc, initDecl,
670 ValueRange{comm, sizeptr.getRes()});
671
672 // load the Size into a register
673 auto loadedSize =
674 LLVM::LoadOp::create(rewriter, loc, i32, sizeptr.getResult());
675
676 // if retval is checked, replace uses of retval with the results from the
677 // call op
678 SmallVector<Value> replacements;
679 if (op.getRetval())
680 replacements.push_back(callOp.getResult());
681
682 // replace all uses, then erase op
683 replacements.push_back(loadedSize.getRes());
684 rewriter.replaceOp(op, replacements);
685
686 return success();
687 }
688};
689
690//===----------------------------------------------------------------------===//
691// SendOpLowering
692//===----------------------------------------------------------------------===//
693
694struct SendOpLowering : public ConvertOpToLLVMPattern<mpi::SendOp> {
695 using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern;
696
697 LogicalResult
698 matchAndRewrite(mpi::SendOp op, OpAdaptor adaptor,
699 ConversionPatternRewriter &rewriter) const override {
700 // get some helper vars
701 Location loc = op.getLoc();
702 MLIRContext *context = rewriter.getContext();
703 Type i32 = rewriter.getI32Type();
704 Type elemType = op.getRef().getType().getElementType();
705 int64_t rank = op.getRef().getType().getRank();
706
707 // ptrType `!llvm.ptr`
708 Type ptrType = LLVM::LLVMPointerType::get(context);
709
710 // grab a reference to the global module op:
711 auto moduleOp = op->getParentOfType<ModuleOp>();
712
713 // get MPI_COMM_WORLD, dataType and pointer
714 auto [dataPtr, size] =
715 getRawPtrAndSize(loc, rewriter, adaptor.getRef(), rank, elemType);
716 auto mpiTraits = MPIImplTraits::get(moduleOp);
717 Value dataType = mpiTraits->getDataType(loc, rewriter, elemType);
718 Value comm = mpiTraits->castComm(loc, rewriter, adaptor.getComm());
719
720 // LLVM Function type representing `i32 MPI_send(data, count, datatype, dst,
721 // tag, comm)`
722 auto funcType = LLVM::LLVMFunctionType::get(
723 i32, {ptrType, i32, dataType.getType(), i32, i32, comm.getType()});
724 // get or create function declaration:
725 LLVM::LLVMFuncOp funcDecl =
726 getOrDefineFunction(moduleOp, loc, rewriter, "MPI_Send", funcType);
727
728 // replace op with function call
729 auto funcCall = LLVM::CallOp::create(rewriter, loc, funcDecl,
730 ValueRange{dataPtr, size, dataType,
731 adaptor.getDest(),
732 adaptor.getTag(), comm});
733 if (op.getRetval())
734 rewriter.replaceOp(op, funcCall.getResult());
735 else
736 rewriter.eraseOp(op);
737
738 return success();
739 }
740};
741
742//===----------------------------------------------------------------------===//
743// RecvOpLowering
744//===----------------------------------------------------------------------===//
745
746struct RecvOpLowering : public ConvertOpToLLVMPattern<mpi::RecvOp> {
747 using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern;
748
749 LogicalResult
750 matchAndRewrite(mpi::RecvOp op, OpAdaptor adaptor,
751 ConversionPatternRewriter &rewriter) const override {
752 // get some helper vars
753 Location loc = op.getLoc();
754 MLIRContext *context = rewriter.getContext();
755 Type i32 = rewriter.getI32Type();
756 Type i64 = rewriter.getI64Type();
757 Type elemType = op.getRef().getType().getElementType();
758 int64_t rank = op.getRef().getType().getRank();
759
760 // ptrType `!llvm.ptr`
761 Type ptrType = LLVM::LLVMPointerType::get(context);
762
763 // grab a reference to the global module op:
764 auto moduleOp = op->getParentOfType<ModuleOp>();
765
766 // get MPI_COMM_WORLD, dataType, status_ignore and pointer
767 auto [dataPtr, size] =
768 getRawPtrAndSize(loc, rewriter, adaptor.getRef(), rank, elemType);
769 auto mpiTraits = MPIImplTraits::get(moduleOp);
770 Value dataType = mpiTraits->getDataType(loc, rewriter, elemType);
771 Value comm = mpiTraits->castComm(loc, rewriter, adaptor.getComm());
772 Value statusIgnore = LLVM::ConstantOp::create(rewriter, loc, i64,
773 mpiTraits->getStatusIgnore());
774 statusIgnore =
775 LLVM::IntToPtrOp::create(rewriter, loc, ptrType, statusIgnore);
776
777 // LLVM Function type representing `i32 MPI_Recv(data, count, datatype, dst,
778 // tag, comm)`
779 auto funcType =
780 LLVM::LLVMFunctionType::get(i32, {ptrType, i32, dataType.getType(), i32,
781 i32, comm.getType(), ptrType});
782 // get or create function declaration:
783 LLVM::LLVMFuncOp funcDecl =
784 getOrDefineFunction(moduleOp, loc, rewriter, "MPI_Recv", funcType);
785
786 // replace op with function call
787 auto funcCall = LLVM::CallOp::create(
788 rewriter, loc, funcDecl,
789 ValueRange{dataPtr, size, dataType, adaptor.getSource(),
790 adaptor.getTag(), comm, statusIgnore});
791 if (op.getRetval())
792 rewriter.replaceOp(op, funcCall.getResult());
793 else
794 rewriter.eraseOp(op);
795
796 return success();
797 }
798};
799
800//===----------------------------------------------------------------------===//
801// AllGatherOpLowering
802//===----------------------------------------------------------------------===//
803
804struct AllGatherOpLowering : public ConvertOpToLLVMPattern<mpi::AllGatherOp> {
805 using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern;
806
807 LogicalResult
808 matchAndRewrite(mpi::AllGatherOp op, OpAdaptor adaptor,
809 ConversionPatternRewriter &rewriter) const override {
810 Location loc = op.getLoc();
811 MLIRContext *context = rewriter.getContext();
812 Type sElemType = op.getSendbuf().getType().getElementType();
813 Type rElemType = op.getRecvbuf().getType().getElementType();
814 int64_t sRank = op.getSendbuf().getType().getRank();
815 int64_t rRank = op.getRecvbuf().getType().getRank();
816 auto [sendPtr, sendSize] =
817 getRawPtrAndSize(loc, rewriter, adaptor.getSendbuf(), sRank, sElemType);
818 auto [recvPtr, recvSize] =
819 getRawPtrAndSize(loc, rewriter, adaptor.getRecvbuf(), rRank, rElemType);
820
821 auto moduleOp = op->getParentOfType<ModuleOp>();
822 auto mpiTraits = MPIImplTraits::get(moduleOp);
823 Value sDataType = mpiTraits->getDataType(loc, rewriter, sElemType);
824 Value rDataType = mpiTraits->getDataType(loc, rewriter, rElemType);
825 Value comm = mpiTraits->castComm(loc, rewriter, adaptor.getComm());
826
827 Type ptrType = LLVM::LLVMPointerType::get(context);
828 Type i32 = rewriter.getI32Type();
829 // int MPI_Allgather(
830 // const void* buffer_send, int count_send, MPI_Datatype datatype_send,
831 // void* buffer_recv, int count_recv, MPI_Datatype datatype_recv,
832 // MPI_Comm communicator);
833 auto funcType = LLVM::LLVMFunctionType::get(
834 i32, {ptrType, i32, sDataType.getType(), ptrType, i32,
835 rDataType.getType(), comm.getType()});
836 // get or create function declaration:
837 LLVM::LLVMFuncOp funcDecl =
838 getOrDefineFunction(moduleOp, loc, rewriter, "MPI_Allgather", funcType);
839
840 // count_recv is the number of elements received from each rank, not total
841 Value nRanks =
842 createOrFoldCommSize(rewriter, loc, op.getComm(), adaptor.getComm());
843 Value recvCountPerRank =
844 LLVM::UDivOp::create(rewriter, loc, i32, recvSize, nRanks);
845
846 // replace op with function call
847 auto funcCall =
848 LLVM::CallOp::create(rewriter, loc, funcDecl,
849 ValueRange{sendPtr, sendSize, sDataType, recvPtr,
850 recvCountPerRank, rDataType, comm});
851
852 if (op.getRetval())
853 rewriter.replaceOp(op, funcCall.getResult());
854 else
855 rewriter.eraseOp(op);
856
857 return success();
858 }
859};
860
861//===----------------------------------------------------------------------===//
862// AllReduceOpLowering
863//===----------------------------------------------------------------------===//
864
865struct AllReduceOpLowering : public ConvertOpToLLVMPattern<mpi::AllReduceOp> {
866 using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern;
867
868 LogicalResult
869 matchAndRewrite(mpi::AllReduceOp op, OpAdaptor adaptor,
870 ConversionPatternRewriter &rewriter) const override {
871 Location loc = op.getLoc();
872 MLIRContext *context = rewriter.getContext();
873 Type i32 = rewriter.getI32Type();
874 Type i64 = rewriter.getI64Type();
875 Type elemType = op.getSendbuf().getType().getElementType();
876 int64_t sRank = op.getSendbuf().getType().getRank();
877 int64_t rRank = op.getRecvbuf().getType().getRank();
878
879 // ptrType `!llvm.ptr`
880 Type ptrType = LLVM::LLVMPointerType::get(context);
881 auto moduleOp = op->getParentOfType<ModuleOp>();
882 auto mpiTraits = MPIImplTraits::get(moduleOp);
883 auto [sendPtr, sendSize] =
884 getRawPtrAndSize(loc, rewriter, adaptor.getSendbuf(), sRank, elemType);
885 auto [recvPtr, recvSize] =
886 getRawPtrAndSize(loc, rewriter, adaptor.getRecvbuf(), rRank, elemType);
887
888 // If input and output are the same, request in-place operation.
889 if (adaptor.getSendbuf() == adaptor.getRecvbuf()) {
890 sendPtr = LLVM::ConstantOp::create(
891 rewriter, loc, i64,
892 reinterpret_cast<int64_t>(mpiTraits->getInPlace()));
893 sendPtr = LLVM::IntToPtrOp::create(rewriter, loc, ptrType, sendPtr);
894 }
895
896 Value dataType = mpiTraits->getDataType(loc, rewriter, elemType);
897 Value mpiOp = mpiTraits->getMPIOp(loc, rewriter, op.getOp());
898 Value commWorld = mpiTraits->castComm(loc, rewriter, adaptor.getComm());
899
900 // 'int MPI_Allreduce(const void *sendbuf, void *recvbuf, int count,
901 // MPI_Datatype datatype, MPI_Op op, MPI_Comm comm)'
902 auto funcType = LLVM::LLVMFunctionType::get(
903 i32, {ptrType, ptrType, i32, dataType.getType(), mpiOp.getType(),
904 commWorld.getType()});
905 // get or create function declaration:
906 LLVM::LLVMFuncOp funcDecl =
907 getOrDefineFunction(moduleOp, loc, rewriter, "MPI_Allreduce", funcType);
908
909 // replace op with function call
910 auto funcCall = LLVM::CallOp::create(
911 rewriter, loc, funcDecl,
912 ValueRange{sendPtr, recvPtr, sendSize, dataType, mpiOp, commWorld});
913
914 if (op.getRetval())
915 rewriter.replaceOp(op, funcCall.getResult());
916 else
917 rewriter.eraseOp(op);
918
919 return success();
920 }
921};
922
923//===----------------------------------------------------------------------===//
924// ReduceScatterBlockOpLowering
925//===----------------------------------------------------------------------===//
926
927struct ReduceScatterBlockOpLowering
928 : public ConvertOpToLLVMPattern<mpi::ReduceScatterBlockOp> {
929 using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern;
930
931 LogicalResult
932 matchAndRewrite(mpi::ReduceScatterBlockOp op, OpAdaptor adaptor,
933 ConversionPatternRewriter &rewriter) const override {
934 Location loc = op.getLoc();
935 MLIRContext *context = rewriter.getContext();
936 Type i32 = rewriter.getI32Type();
937 Type i64 = rewriter.getI64Type();
938 Type elemType = op.getSendbuf().getType().getElementType();
939 int64_t sRank = op.getSendbuf().getType().getRank();
940 int64_t rRank = op.getRecvbuf().getType().getRank();
941
942 // ptrType `!llvm.ptr`
943 Type ptrType = LLVM::LLVMPointerType::get(context);
944 auto moduleOp = op->getParentOfType<ModuleOp>();
945 auto mpiTraits = MPIImplTraits::get(moduleOp);
946 auto [sendPtr, sendSize] =
947 getRawPtrAndSize(loc, rewriter, adaptor.getSendbuf(), sRank, elemType);
948 auto [recvPtr, recvSize] =
949 getRawPtrAndSize(loc, rewriter, adaptor.getRecvbuf(), rRank, elemType);
950
951 // If input and output are the same, request in-place operation.
952 if (adaptor.getSendbuf() == adaptor.getRecvbuf()) {
953 sendPtr = LLVM::ConstantOp::create(
954 rewriter, loc, i64,
955 reinterpret_cast<int64_t>(mpiTraits->getInPlace()));
956 sendPtr = LLVM::IntToPtrOp::create(rewriter, loc, ptrType, sendPtr);
957 }
958
959 Value dataType = mpiTraits->getDataType(loc, rewriter, elemType);
960 Value mpiOp = mpiTraits->getMPIOp(loc, rewriter, op.getOp());
961 Value comm = mpiTraits->castComm(loc, rewriter, adaptor.getComm());
962
963 Value nRanks =
964 createOrFoldCommSize(rewriter, loc, op.getComm(), adaptor.getComm());
965 Value totalExpected =
966 LLVM::MulOp::create(rewriter, loc, i32, recvSize, nRanks);
967 Value sizeIsValid = LLVM::ICmpOp::create(
968 rewriter, loc, LLVM::ICmpPredicate::eq, sendSize, totalExpected);
969 cf::AssertOp::create(rewriter, loc, sizeIsValid,
970 "Send buffer's size must be the receive buffer's size "
971 "times the number of ranks");
972
973 // 'int MPI_Reduce_scatter_block(const void *sendbuf, void *recvbuf,
974 // int recvcount, MPI_Datatype datatype, MPI_Op op, MPI_Comm comm)'
975 auto funcType = LLVM::LLVMFunctionType::get(
976 i32, {ptrType, ptrType, i32, dataType.getType(), mpiOp.getType(),
977 comm.getType()});
978 // get or create function declaration:
979 LLVM::LLVMFuncOp funcDecl = getOrDefineFunction(
980 moduleOp, loc, rewriter, "MPI_Reduce_scatter_block", funcType);
981
982 // replace op with function call
983 auto funcCall = LLVM::CallOp::create(
984 rewriter, loc, funcDecl,
985 ValueRange{sendPtr, recvPtr, recvSize, dataType, mpiOp, comm});
986
987 if (op.getRetval())
988 rewriter.replaceOp(op, funcCall.getResult());
989 else
990 rewriter.eraseOp(op);
991
992 return success();
993 }
994};
995
996//===----------------------------------------------------------------------===//
997// ConvertToLLVMPatternInterface implementation
998//===----------------------------------------------------------------------===//
999
1000/// Implement the interface to convert Func to LLVM.
1001struct FuncToLLVMDialectInterface : public ConvertToLLVMPatternInterface {
1002 FuncToLLVMDialectInterface(Dialect *dialect)
1003 : ConvertToLLVMPatternInterface(dialect) {}
1004
1005 /// Hook for derived dialect interface to provide conversion patterns
1006 /// and mark dialect legal for the conversion target.
1007 void populateConvertToLLVMConversionPatterns(
1008 ConversionTarget &target, LLVMTypeConverter &typeConverter,
1009 RewritePatternSet &patterns) const final {
1010 mpi::populateMPIToLLVMConversionPatterns(typeConverter, patterns);
1011 }
1012};
1013} // namespace
1014
1015//===----------------------------------------------------------------------===//
1016// Pattern Population
1017//===----------------------------------------------------------------------===//
1018
1019void mpi::populateMPIToLLVMConversionPatterns(LLVMTypeConverter &converter,
1020 RewritePatternSet &patterns) {
1021 // Using i64 as a portable, intermediate type for !mpi.comm.
1022 // It would be nicer to somehow get the right type directly, but TLDI is not
1023 // available here.
1024 converter.addConversion([](mpi::CommType type) {
1025 return IntegerType::get(type.getContext(), 64);
1026 });
1027 patterns.add<CommRankOpLowering, CommSizeOpLowering, CommSplitOpLowering,
1028 CommWorldOpLowering, FinalizeOpLowering, InitOpLowering,
1029 SendOpLowering, RecvOpLowering, AllGatherOpLowering,
1030 AllReduceOpLowering, ReduceScatterBlockOpLowering>(converter);
1031}
1032
1033void mpi::registerConvertMPIToLLVMInterface(DialectRegistry &registry) {
1034 registry.addExtension(+[](MLIRContext *ctx, mpi::MPIDialect *dialect) {
1035 dialect->addInterfaces<FuncToLLVMDialectInterface>();
1036 });
1037}
success
return success()
int64_t
mlir::ConvertOpToLLVMPattern::ConvertOpToLLVMPattern
ConvertOpToLLVMPattern(const LLVMTypeConverter &typeConverter, PatternBenefit benefit=1)
Definition Pattern.h:239
mlir::DialectRegistry
The DialectRegistry maps a dialect namespace to a constructor for the matching dialect.
Definition DialectRegistry.h:150
mlir::DialectRegistry::addExtension
bool addExtension(TypeID extensionID, std::unique_ptr< DialectExtensionBase > extension)
Add the given extension to the registry.
Definition DialectRegistry.h:255
mlir::LLVMTypeConverter
Conversion from types to the LLVM IR dialect.
Definition TypeConverter.h:35
mlir::Location
This class defines the main interface for locations in MLIR and acts as a non-nullable wrapper around...
Definition Location.h:76
mlir::MLIRContext
MLIRContext is the top-level object for a collection of MLIR operations.
Definition MLIRContext.h:63
mlir::Op
This provides public APIs that all operations should have.
Definition OpDefinition.h:1711
mlir::RewritePatternSet::add
RewritePatternSet & add(ConstructorArg &&arg, ConstructorArgs &&...args)
Add an instance of each of the pattern types 'Ts' to the pattern list with the given arguments.
Definition PatternMatch.h:861
mlir::Type
Instances of the Type class are uniqued, have an immutable identifier and an optional mutable compone...
Definition Types.h:74
mlir::Type::isF64
bool isF64() const
Definition Types.cpp:41
mlir::Type::isF32
bool isF32() const
Definition Types.cpp:40
mlir::Type::isUnsignedInteger
bool isUnsignedInteger() const
Return true if this is an unsigned integer type (with the specified width).
Definition Types.cpp:90
mlir::Type::isInteger
bool isInteger() const
Return true if this is an integer type (with the specified width).
Definition Types.cpp:58
mlir::Value
This class represents an instance of an SSA value in the MLIR system, representing a computable value...
Definition Value.h:96
mlir::Value::getType
Type getType() const
Return the type of this value.
Definition Value.h:105
mlir::dlti::query
FailureOr< Attribute > query(Operation *op, ArrayRef< DataLayoutEntryKey > keys, bool emitError=false)
Perform a DLTI-query at op, recursively querying each key of keys on query interface-implementing att...
Definition DLTI.cpp:537
mlir::mpi::FoldToDLTIConst
LogicalResult FoldToDLTIConst(OpT op, const char *key, mlir::PatternRewriter &b)
Definition Utils.h:19
mlir::mpi::populateMPIToLLVMConversionPatterns
void populateMPIToLLVMConversionPatterns(LLVMTypeConverter &converter, RewritePatternSet &patterns)
Definition MPIToLLVM.cpp:1019
mlir::mpi::registerConvertMPIToLLVMInterface
void registerConvertMPIToLLVMInterface(DialectRegistry &registry)
Definition MPIToLLVM.cpp:1033
mlir::remark::failed
detail::InFlightRemark failed(Location loc, RemarkOpts opts)
Report an optimization remark that failed.
Definition Remarks.h:732
mlir
Include the generated interface declarations.
Definition ABIRewriteContext.h:32
mlir::getOrDefineFunction
LLVM::LLVMFuncOp getOrDefineFunction(Operation *moduleOp, Location loc, OpBuilder &b, StringRef name, LLVM::LLVMFunctionType type)
Note that these functions don't take a SymbolTable because GPU module lowerings can have name collisi...
Definition GPUOpsLowering.cpp:27
mlir::get
auto get(MLIRContext *context, Ts &&...params)
Helper method that injects context only if needed, this helps unify some of the attribute constructio...
Definition BytecodeImplementation.h:461
Generated on for MLIR by doxygen 1.14.0