xref: /llvm-project/mlir/include/mlir/Conversion/Passes.td (revision 25ae1a266d50f24a8fffc57152d7f3c3fcb65517)
1//===-- Passes.td - Conversion pass definition file --------*- tablegen -*-===//
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#ifndef MLIR_CONVERSION_PASSES
10#define MLIR_CONVERSION_PASSES
11
12include "mlir/Pass/PassBase.td"
13
14
15//===----------------------------------------------------------------------===//
16// ToLLVM
17//===----------------------------------------------------------------------===//
18
19def ConvertToLLVMPass : Pass<"convert-to-llvm"> {
20  let summary = "Convert to LLVM via dialect interfaces found in the input IR";
21  let description = [{
22    This is a generic pass to convert to LLVM, it uses the
23    `ConvertToLLVMPatternInterface` dialect interface to delegate to dialects
24    the injection of conversion patterns.
25
26    If `dynamic` is set to `true`, the pass will look for
27    `ConvertToLLVMAttrInterface` attributes and use them to further configure
28    the conversion process. This option also uses the `DataLayoutAnalysis`
29    analysis to configure the type converter. Enabling this option incurs in
30    extra overhead.
31  }];
32
33  let constructor = "mlir::createConvertToLLVMPass()";
34  let options = [
35    ListOption<"filterDialects", "filter-dialects", "std::string",
36               "Test conversion patterns of only the specified dialects">,
37    Option<"useDynamic", "dynamic", "bool", "false",
38           "Use op conversion attributes to configure the conversion">,
39  ];
40}
41
42//===----------------------------------------------------------------------===//
43// AffineToStandard
44//===----------------------------------------------------------------------===//
45
46def ConvertAffineToStandard : Pass<"lower-affine"> {
47  let summary = "Lower Affine operations to a combination of Standard and SCF "
48                "operations";
49  let description = [{
50
51    Convert operations from the affine dialect into operations from the SCF and
52    standard dialects.
53
54    `affine.for` operations are converted to `scf.for` operations that are free
55    of certain structural restrictions (on their bounds and step). `affine.if`
56    is similarly converted to the `scf.if` operation. `affine.apply` operations
57    are converted into sequences of primitive arithmetic operations from the
58    standard dialect that have the same effect, using operands of the `index`
59    type. Consequently, named maps and sets thare are no longer in use may be
60    removed from the module.
61
62    For example, `%r = affine.apply affine_map<(d0, d1)[s0] -> (d0 + 2*d1 +
63    s0)>(%d0, %d1)[%s0]`
64    can be converted into:
65
66    ```mlir
67    %d0 = <...>
68    %d1 = <...>
69    %s0 = <...>
70    %0 = arith.constant 2 : index
71    %1 = arith.muli %0, %d1
72    %2 = arith.addi %d0, %1
73    %r = arith.addi %2, %s0
74    ```
75
76    #### Input invariant
77
78    -   no `Tensor` types;
79
80    These restrictions may be lifted in the future.
81
82    #### Output IR
83
84    Functions with `affine.for` and `affine.if` operations eliminated. These
85    functions may contain operations from the Standard dialect in addition to
86    those already present before the pass.
87
88    #### Invariants
89
90    -   Functions without a body are not modified.
91    -   The semantics of the other functions is preserved.
92    -   Individual operations other than those mentioned above are not modified
93        if they do not depend on the loop iterator value or on the result of
94        `affine.apply`.
95  }];
96  let constructor = "mlir::createLowerAffinePass()";
97  let dependentDialects = [
98    "memref::MemRefDialect",
99    "scf::SCFDialect",
100    "vector::VectorDialect"
101  ];
102}
103
104//===----------------------------------------------------------------------===//
105// AMDGPUToROCDL
106//===----------------------------------------------------------------------===//
107
108def ConvertAMDGPUToROCDL : Pass<"convert-amdgpu-to-rocdl"> {
109  let summary = "Convert AMDGPU dialect to ROCDL dialect";
110  let description = [{
111    This pass converts supported AMDGPU ops to ROCDL dialect intrinsics.
112  }];
113  let constructor = "mlir::createConvertAMDGPUToROCDLPass()";
114  let dependentDialects = [
115    "LLVM::LLVMDialect",
116    "ROCDL::ROCDLDialect",
117  ];
118  let options = [Option<"chipset", "chipset", "std::string",
119                        /*default=*/"\"gfx000\"",
120                        "Chipset that these operations will run on">];
121}
122
123//===----------------------------------------------------------------------===//
124// ArithToAMDGPU
125//===----------------------------------------------------------------------===//
126def ArithToAMDGPUConversionPass : Pass<"convert-arith-to-amdgpu"> {
127  let summary = "Convert Arith operations to AMDGPU-specific implementations";
128  let description = [{
129    Convert `arith` operations (currently extf and truncf on 8-bit floats)
130    to operations in the `amdgpu` dialect. This pass is done in two steps
131    in order to avoid running a notional arith-to-rocdl and arith-to-llvm
132    simultaniously.
133  }];
134
135  let dependentDialects = ["amdgpu::AMDGPUDialect", "vector::VectorDialect"];
136
137  let options = [
138    Option<"chipset", "chipset", "std::string",
139                        /*default=*/"\"gfx000\"",
140                        "Chipset that these operations will run on">,
141    Option<"saturateFP8Truncf", "saturate-fp8-truncf", "bool",
142           /*default=*/"false",
143           "Use saturating truncation for 8-bit float types">,
144    Option<"allowPackedF16Rtz", "allow-packed-f16-round-to-zero", "bool",
145           /*default=*/"false",
146           "Whether we should allow f32->f16 packed round-to-zero conversion">,
147  ];
148}
149
150//===----------------------------------------------------------------------===//
151// ArithToEmitC
152//===----------------------------------------------------------------------===//
153
154def ConvertArithToEmitC : Pass<"convert-arith-to-emitc"> {
155  let summary = "Convert Arith dialect to EmitC dialect";
156  let dependentDialects = ["emitc::EmitCDialect"];
157}
158
159//===----------------------------------------------------------------------===//
160// ArithToLLVM
161//===----------------------------------------------------------------------===//
162
163def ArithToLLVMConversionPass : Pass<"convert-arith-to-llvm"> {
164  let summary = "Convert Arith dialect to LLVM dialect";
165  let description = [{
166    This pass converts supported Arith ops to LLVM dialect instructions.
167  }];
168  let dependentDialects = ["LLVM::LLVMDialect"];
169  let options = [
170    Option<"indexBitwidth", "index-bitwidth", "unsigned",
171           /*default=kDeriveIndexBitwidthFromDataLayout*/"0",
172           "Bitwidth of the index type, 0 to use size of machine word">,
173  ];
174}
175
176//===----------------------------------------------------------------------===//
177// ArithToSPIRV
178//===----------------------------------------------------------------------===//
179
180def ConvertArithToSPIRV : Pass<"convert-arith-to-spirv"> {
181  let summary = "Convert Arith dialect to SPIR-V dialect";
182  let constructor = "mlir::arith::createConvertArithToSPIRVPass()";
183  let dependentDialects = ["spirv::SPIRVDialect"];
184  let options = [
185    Option<"emulateLT32BitScalarTypes", "emulate-lt-32-bit-scalar-types",
186           "bool", /*default=*/"true",
187           "Emulate narrower scalar types with 32-bit ones if not supported by "
188           "the target">,
189  ];
190}
191
192//===----------------------------------------------------------------------===//
193// ArithToArmSME
194//===----------------------------------------------------------------------===//
195
196def ArithToArmSMEConversionPass : Pass<"convert-arith-to-arm-sme"> {
197  let summary = "Convert Arith dialect to ArmSME dialect";
198  let dependentDialects = ["arm_sme::ArmSMEDialect"];
199}
200
201//===----------------------------------------------------------------------===//
202// ArmNeon2dToIntr
203//===----------------------------------------------------------------------===//
204
205def ConvertArmNeon2dToIntr : Pass<"arm-neon-2d-to-intr"> {
206  let summary = "Convert Arm NEON structured ops to intrinsics";
207  let constructor = "mlir::createConvertArmNeon2dToIntrPass()";
208  let dependentDialects = ["arm_neon::ArmNeonDialect", "vector::VectorDialect"];
209}
210
211//===----------------------------------------------------------------------===//
212// AsyncToLLVM
213//===----------------------------------------------------------------------===//
214
215def ConvertAsyncToLLVMPass : Pass<"convert-async-to-llvm", "ModuleOp"> {
216  let summary = "Convert the operations from the async dialect into the LLVM "
217                "dialect";
218  let description = [{
219    Convert `async.execute` operations to LLVM coroutines and use async runtime
220    API to execute them.
221  }];
222  let dependentDialects = [
223    "arith::ArithDialect",
224    "async::AsyncDialect",
225    "LLVM::LLVMDialect",
226    "func::FuncDialect",
227  ];
228}
229
230//===----------------------------------------------------------------------===//
231// BufferizationToMemRef
232//===----------------------------------------------------------------------===//
233
234def ConvertBufferizationToMemRef : Pass<"convert-bufferization-to-memref"> {
235  let summary = "Convert operations from the Bufferization dialect to the "
236                "MemRef dialect";
237  let description = [{
238
239    This pass converts bufferization operations into memref operations.
240
241    In the current state, this pass only transforms a `bufferization.clone`
242    operation into `memref.alloc` and `memref.copy` operations and
243    `bufferization.dealloc` operations (the same way as the
244    `-bufferization-lower-deallocations` pass). The conversion of `clone`
245    operations is needed, since some clone operations could remain after
246    applying several transformation processes. Currently, only `canonicalize`
247    transforms clone operations or even eliminates them. This can lead to errors
248    if any clone op survived after all conversion passes (starting from the
249    bufferization dialect) are performed.
250
251    See:
252    https://llvm.discourse.group/t/bufferization-error-related-to-memref-clone/4665
253
254    To avoid these errors, this pass can be performed as a last clean-up pass to
255    transform remaining operations and to proceed in other dialects (memref
256    e.g.).
257
258    Note that this pass only transforms the operation without any further
259    analyses. This pass does not consider any memory analysis or optimization
260    and hence does not resolve any memory leaks.
261
262  }];
263  let constructor = "mlir::createBufferizationToMemRefPass()";
264  let dependentDialects = [
265    "arith::ArithDialect", "memref::MemRefDialect", "scf::SCFDialect",
266    "func::FuncDialect"
267  ];
268}
269
270//===----------------------------------------------------------------------===//
271// ComplexToLLVM
272//===----------------------------------------------------------------------===//
273
274def ConvertComplexToLLVMPass : Pass<"convert-complex-to-llvm"> {
275  let summary = "Convert Complex dialect to LLVM dialect";
276  let dependentDialects = ["LLVM::LLVMDialect"];
277}
278
279//===----------------------------------------------------------------------===//
280// ComplexToLibm
281//===----------------------------------------------------------------------===//
282
283def ConvertComplexToLibm : Pass<"convert-complex-to-libm", "ModuleOp"> {
284  let summary = "Convert Complex dialect to libm calls";
285  let description = [{
286    This pass converts supported Complex ops to libm calls.
287  }];
288  let constructor = "mlir::createConvertComplexToLibmPass()";
289  let dependentDialects = [
290    "func::FuncDialect",
291  ];
292}
293
294//===----------------------------------------------------------------------===//
295// ComplexToSPIRV
296//===----------------------------------------------------------------------===//
297
298def ConvertComplexToSPIRVPass : Pass<"convert-complex-to-spirv"> {
299  let summary = "Convert Complex dialect to SPIRV dialect";
300  let dependentDialects = ["spirv::SPIRVDialect"];
301}
302
303//===----------------------------------------------------------------------===//
304// ComplexToStandard
305//===----------------------------------------------------------------------===//
306
307def ConvertComplexToStandard : Pass<"convert-complex-to-standard"> {
308  let summary = "Convert Complex dialect to standard dialect";
309  let constructor = "mlir::createConvertComplexToStandardPass()";
310  let dependentDialects = ["math::MathDialect"];
311}
312
313//===----------------------------------------------------------------------===//
314// ControlFlowToLLVM
315//===----------------------------------------------------------------------===//
316
317def ConvertControlFlowToLLVMPass : Pass<"convert-cf-to-llvm", "ModuleOp"> {
318  let summary = "Convert ControlFlow operations to the LLVM dialect";
319  let description = [{
320    Convert ControlFlow operations into LLVM IR dialect operations.
321
322    If other operations are present and their results are required by the LLVM
323    IR dialect operations, the pass will fail.  Any LLVM IR operations or types
324    already present in the IR will be kept as is.
325  }];
326  let dependentDialects = ["LLVM::LLVMDialect"];
327  let options = [
328    Option<"indexBitwidth", "index-bitwidth", "unsigned",
329           /*default=kDeriveIndexBitwidthFromDataLayout*/"0",
330           "Bitwidth of the index type, 0 to use size of machine word">
331  ];
332}
333
334//===----------------------------------------------------------------------===//
335// ControlFlowToSCF
336//===----------------------------------------------------------------------===//
337
338def LiftControlFlowToSCFPass : Pass<"lift-cf-to-scf"> {
339  let summary = "Lift ControlFlow dialect to SCF dialect";
340  let description = [{
341    Lifts ControlFlow operations to SCF dialect operations.
342
343    This pass is prefixed with "lift" instead of "convert" as it is not always
344    guaranteed to replace all ControlFlow ops.
345    If a region contains only a single kind of return-like operation, all
346    ControlFlow operations will be replaced successfully.
347    Otherwise a single ControlFlow switch branching to one block per return-like
348    operation kind remains.
349
350    This pass may need to create unreachable terminators in case of infinite
351    loops, which is only supported for 'func.func' for now. If you potentially
352    have infinite loops inside CFG regions not belonging to 'func.func',
353    consider using `transformCFGToSCF` function directly with corresponding
354    `CFGToSCFInterface::createUnreachableTerminator` implementation.
355  }];
356
357  let dependentDialects = ["scf::SCFDialect",
358                           "arith::ArithDialect",
359                           "ub::UBDialect",
360                           // TODO: This is only necessary until we have a
361                           //       ub.unreachable op.
362                           "func::FuncDialect"];
363}
364
365//===----------------------------------------------------------------------===//
366// ControlFlowToSPIRV
367//===----------------------------------------------------------------------===//
368
369def ConvertControlFlowToSPIRV : Pass<"convert-cf-to-spirv"> {
370  let summary = "Convert ControlFlow dialect to SPIR-V dialect";
371  let constructor = "mlir::createConvertControlFlowToSPIRVPass()";
372  let dependentDialects = ["spirv::SPIRVDialect"];
373  let options = [
374    Option<"emulateLT32BitScalarTypes", "emulate-lt-32-bit-scalar-types",
375           "bool", /*default=*/"true",
376           "Emulate narrower scalar types with 32-bit ones if not supported by"
377           " the target">
378  ];
379}
380
381//===----------------------------------------------------------------------===//
382// FuncToEmitC
383//===----------------------------------------------------------------------===//
384
385def ConvertFuncToEmitC : Pass<"convert-func-to-emitc", "ModuleOp"> {
386  let summary = "Convert Func dialect to EmitC dialect";
387  let dependentDialects = ["emitc::EmitCDialect"];
388}
389
390//===----------------------------------------------------------------------===//
391// FuncToLLVM
392//===----------------------------------------------------------------------===//
393
394def SetLLVMModuleDataLayoutPass : Pass<"set-llvm-module-datalayout", "ModuleOp"> {
395  let summary = "Attach a datalayout string as a module attribute";
396  let description = [{
397    Verify that the dataLayout string is a valid LLVM datalayout string and
398    attach it as an attribute `LLVMDialect::getDataLayoutAttrName()` to the
399    module, overriding the existing one.
400  }];
401  let options = [
402    Option<"dataLayout", "data-layout", "std::string",
403           /*default=*/"\"\"",
404           "String description (LLVM format) of the data layout that is "
405           "expected on the produced module">,
406  ];
407}
408
409def ConvertFuncToLLVMPass : Pass<"convert-func-to-llvm", "ModuleOp"> {
410  let summary = "Convert from the Func dialect to the LLVM dialect";
411  let description = [{
412    Convert Func dialect operations into the LLVM IR dialect operations.
413
414    #### Input invariant
415
416    -   no `tensor` types;
417    -   all `vector` are one-dimensional;
418    -   all blocks are reachable by following the successors of the first basic
419        block;
420
421    If other operations are present and their results are required by the LLVM
422    IR dialect operations, the pass will fail.  Any LLVM IR operations or types
423    already present in the IR will be kept as is.
424
425    An LLVM datalayout string can be attached as an attribute to the module on
426    which the pass anchors. Such an attribute is attached by calling the
427    set-module-datalayout pass. If present, an llvm::DataLayout object is
428    created from this attribute and used in the conversion to LLVM.
429
430    #### Output IR
431
432    Functions converted to LLVM IR. Function arguments types are converted
433    one-to-one. Function results are converted one-to-one and, in case more than
434    1 value is returned, packed into an LLVM IR struct type. Function calls and
435    returns are updated accordingly. Block argument types are updated to use
436    LLVM IR types.
437  }];
438  let dependentDialects = ["LLVM::LLVMDialect"];
439  let options = [
440    Option<"useBarePtrCallConv", "use-bare-ptr-memref-call-conv", "bool",
441           /*default=*/"false",
442           "Replace FuncOp's MemRef arguments with bare pointers to the MemRef "
443           "element types">,
444    Option<"indexBitwidth", "index-bitwidth", "unsigned",
445           /*default=kDeriveIndexBitwidthFromDataLayout*/"0",
446           "Bitwidth of the index type, 0 to use size of machine word">,
447  ];
448}
449
450//===----------------------------------------------------------------------===//
451// FuncToSPIRV
452//===----------------------------------------------------------------------===//
453
454def ConvertFuncToSPIRV : Pass<"convert-func-to-spirv"> {
455  let summary = "Convert Func dialect to SPIR-V dialect";
456  let constructor = "mlir::createConvertFuncToSPIRVPass()";
457  let dependentDialects = ["spirv::SPIRVDialect"];
458  let options = [
459    Option<"emulateLT32BitScalarTypes", "emulate-lt-32-bit-scalar-types",
460           "bool", /*default=*/"true",
461           "Emulate narrower scalar types with 32-bit ones if not supported by"
462           " the target">
463  ];
464}
465
466//===----------------------------------------------------------------------===//
467// GPUCommon
468//===----------------------------------------------------------------------===//
469
470def GpuToLLVMConversionPass : Pass<"gpu-to-llvm", "ModuleOp"> {
471  let summary = "Convert GPU dialect to LLVM dialect with GPU runtime calls";
472
473  let description = [{
474    Creates a pass to convert a GPU operations into a sequence of GPU runtime
475    calls.
476
477    This pass does not generate code to call GPU runtime APIs directly but
478    instead uses a small wrapper library that exports a stable and conveniently
479    typed ABI on top of GPU runtimes such as CUDA or ROCm (HIP).
480  }];
481
482  let options = [
483    Option<"hostBarePtrCallConv", "use-bare-pointers-for-host", "bool",
484           /*default=*/"false",
485           "Use bare pointers to pass memref arguments to host functions. "
486           "All memrefs must have static shape.">,
487    Option<"kernelBarePtrCallConv", "use-bare-pointers-for-kernels", "bool",
488           /*default=*/"false",
489             "Use bare pointers to pass memref arguments to kernels. "
490             "The kernel must use the same setting for this option."
491          >,
492    Option<"kernelIntersperseSizeCallConv", "intersperse-sizes-for-kernels", "bool",
493           /*default=*/"false",
494           "Inserts a size_t argument following each memref argument, "
495           "containing the static size in bytes of the buffer. Incompatible "
496           "arguments are rejected. This is intended for use by the Vulkan "
497           "runtime with the kernel bare pointer calling convention, to enable "
498           "dynamic binding of buffers as arguments without static type info."
499          >
500  ];
501
502  let dependentDialects = [
503    "LLVM::LLVMDialect",
504    "memref::MemRefDialect",
505  ];
506}
507
508def LowerHostCodeToLLVMPass : Pass<"lower-host-to-llvm", "ModuleOp"> {
509  let summary = "Lowers the host module code and `gpu.launch_func` to LLVM";
510
511  let description = [{
512    Creates a pass to emulate `gpu.launch_func` call in LLVM dialect and lower
513    the host module code to LLVM.
514
515    This transformation creates a sequence of global variables that are later
516    linked to the variables in the kernel module, and a series of copies to/from
517    them to emulate the memory transfer from the host or to the device sides. It
518    also converts the remaining Arithmetic, Func, and MemRef dialects into LLVM
519    dialect, emitting C wrappers.
520  }];
521
522  let dependentDialects = ["LLVM::LLVMDialect"];
523}
524
525//===----------------------------------------------------------------------===//
526// GPUToLLVMSPV
527//===----------------------------------------------------------------------===//
528
529def ConvertGpuOpsToLLVMSPVOps : Pass<"convert-gpu-to-llvm-spv", "gpu::GPUModuleOp"> {
530  let summary =
531    "Generate LLVM operations to be ingested by a SPIR-V backend for gpu operations";
532  let dependentDialects = ["LLVM::LLVMDialect"];
533  let options = [
534    Option<"use64bitIndex", "use-64bit-index",
535           "bool", /*default=*/"false",
536           "Use 64-bit integers to convert index types">,
537  ];
538}
539
540//===----------------------------------------------------------------------===//
541// GPUToNVVM
542//===----------------------------------------------------------------------===//
543
544def ConvertGpuOpsToNVVMOps : Pass<"convert-gpu-to-nvvm", "gpu::GPUModuleOp"> {
545  let summary = "Generate NVVM operations for gpu operations";
546  let dependentDialects = [
547    "cf::ControlFlowDialect",
548    "memref::MemRefDialect",
549    "NVVM::NVVMDialect",
550  ];
551  let options = [
552    Option<"indexBitwidth", "index-bitwidth", "unsigned",
553           /*default=kDeriveIndexBitwidthFromDataLayout*/"0",
554           "Bitwidth of the index type, 0 to use size of machine word">,
555    Option<"hasRedux", "has-redux", "bool", /*default=*/"false",
556           "Target gpu supports redux">,
557    Option<"useBarePtrCallConv", "use-bare-ptr-memref-call-conv", "bool",
558           /*default=*/"false",
559           "Replace memref arguments in GPU functions with bare pointers. "
560           "All memrefs must have static shape.">
561  ];
562}
563
564//===----------------------------------------------------------------------===//
565// GPUToROCDL
566//===----------------------------------------------------------------------===//
567
568def ConvertGpuOpsToROCDLOps : Pass<"convert-gpu-to-rocdl", "gpu::GPUModuleOp"> {
569  let summary = "Generate ROCDL operations for gpu operations";
570  let constructor = "mlir::createLowerGpuOpsToROCDLOpsPass()";
571  let dependentDialects = [
572    "ROCDL::ROCDLDialect",
573    "cf::ControlFlowDialect",
574    "memref::MemRefDialect",
575  ];
576  let options = [
577    Option<"chipset", "chipset", "std::string",
578           /*default=*/"\"gfx000\"",
579           "Chipset that these operations will run on">,
580    Option<"indexBitwidth", "index-bitwidth", "unsigned",
581           /*default=kDeriveIndexBitwidthFromDataLayout*/"0",
582           "Bitwidth of the index type, 0 to use size of machine word">,
583    Option<"useBarePtrCallConv", "use-bare-ptr-memref-call-conv", "bool",
584           /*default=*/"false",
585           "Replace memref arguments in GPU functions with bare pointers."
586           "All memrefs must have static shape">,
587    Option<"runtime", "runtime", "::mlir::gpu::amd::Runtime",
588          "::mlir::gpu::amd::Runtime::Unknown",
589          "Runtime code will be run on (default is Unknown, can also use HIP or OpenCl)",
590          [{::llvm::cl::values(
591            clEnumValN(::mlir::gpu::amd::Runtime::Unknown, "unknown", "Unknown (default)"),
592            clEnumValN(::mlir::gpu::amd::Runtime::HIP, "HIP", "HIP"),
593            clEnumValN(::mlir::gpu::amd::Runtime::OpenCL, "OpenCL", "OpenCL")
594          )}]>
595  ];
596}
597
598//===----------------------------------------------------------------------===//
599// GPUToSPIRV
600//===----------------------------------------------------------------------===//
601
602def ConvertGPUToSPIRV : Pass<"convert-gpu-to-spirv", "ModuleOp"> {
603  let summary = "Convert GPU dialect to SPIR-V dialect";
604  let description = [{
605    This pass converts supported GPU device ops to SPIR-V ops. It does not
606    handle GPU host ops.
607
608    A `gpu.func` op can have parameters to pass in resources. But in SPIR-V
609    entry functions cannot take parameters; they use descriptors to access
610    resources. By default, parameters to a `gpu.func` op will be converted to
611    global variables. These global variables will be assigned sequential binding
612    numbers following their order in the original `gpu.func` op, starting from
613    0, in set 0. One can attach `spirv.interface_var_abi` to those parameters
614    to control the set and binding if wanted.
615  }];
616  let constructor = "mlir::createConvertGPUToSPIRVPass()";
617  let dependentDialects = [
618    "func::FuncDialect",
619    "spirv::SPIRVDialect",
620  ];
621  let options = [
622    Option<"use64bitIndex", "use-64bit-index",
623           "bool", /*default=*/"false",
624           "Use 64-bit integers to convert index types">,
625  ];
626}
627
628//===----------------------------------------------------------------------===//
629// ConvertIndexToLLVMPass
630//===----------------------------------------------------------------------===//
631
632def ConvertIndexToLLVMPass : Pass<"convert-index-to-llvm"> {
633  let summary = "Lower the `index` dialect to the `llvm` dialect.";
634  let description = [{
635    This pass lowers Index dialect operations to LLVM dialect operations.
636    Operation conversions are 1-to-1 except for the exotic divides: `ceildivs`,
637    `ceildivu`, and `floordivs`, which expand to series of LLVM operations.
638    Importantly, the index bitwidth should be correctly set to the target
639    pointer width via `index-bitwidth`.
640  }];
641
642  let dependentDialects = ["::mlir::LLVM::LLVMDialect"];
643
644  let options = [
645    Option<"indexBitwidth", "index-bitwidth", "unsigned",
646           /*default=kDeriveIndexBitwidthFromDataLayout*/"0",
647           "Bitwidth of the index type, 0 to use size of machine word">,
648  ];
649}
650
651//===----------------------------------------------------------------------===//
652// ConvertIndexToSPIRVPass
653//===----------------------------------------------------------------------===//
654
655def ConvertIndexToSPIRVPass : Pass<"convert-index-to-spirv"> {
656  let summary = "Lower the `index` dialect to the `spirv` dialect.";
657  let description = [{
658    This pass lowers Index dialect operations to SPIR-V dialect operations.
659    Operation conversions are 1-to-1 except for the exotic divides: `ceildivs`,
660    `ceildivu`, and `floordivs`. The index bitwidth will be 32 or 64 as
661    specified by use-64bit-index.
662  }];
663
664  let dependentDialects = ["::mlir::spirv::SPIRVDialect"];
665
666  let options = [
667    Option<"use64bitIndex", "use-64bit-index",
668           "bool", /*default=*/"false",
669           "Use 64-bit integers to convert index types">
670  ];
671}
672
673//===----------------------------------------------------------------------===//
674// LinalgToStandard
675//===----------------------------------------------------------------------===//
676
677def ConvertLinalgToStandard : Pass<"convert-linalg-to-std", "ModuleOp"> {
678  let summary = "Convert the operations from the linalg dialect into the "
679                "Standard dialect";
680  let constructor = "mlir::createConvertLinalgToStandardPass()";
681  let dependentDialects = ["func::FuncDialect", "memref::MemRefDialect"];
682}
683
684//===----------------------------------------------------------------------===//
685// MathToLibm
686//===----------------------------------------------------------------------===//
687
688def ConvertMathToLibm : Pass<"convert-math-to-libm", "ModuleOp"> {
689  let summary = "Convert Math dialect to libm calls";
690  let description = [{
691    This pass converts supported Math ops to libm calls.
692  }];
693  let constructor = "mlir::createConvertMathToLibmPass()";
694  let dependentDialects = [
695    "arith::ArithDialect",
696    "func::FuncDialect",
697    "vector::VectorDialect",
698  ];
699}
700
701//===----------------------------------------------------------------------===//
702// MathToLLVM
703//===----------------------------------------------------------------------===//
704
705def ConvertMathToLLVMPass : Pass<"convert-math-to-llvm"> {
706  let summary = "Convert Math dialect to LLVM dialect";
707  let dependentDialects = ["LLVM::LLVMDialect"];
708  let options = [
709    Option<"approximateLog1p", "approximate-log1p", "bool", "true",
710           "Enable approximation of Log1p.">
711  ];
712}
713
714//===----------------------------------------------------------------------===//
715// MathToLibm
716//===----------------------------------------------------------------------===//
717
718def ConvertMathToROCDL : Pass<"convert-math-to-rocdl", "ModuleOp"> {
719  let summary = "Convert Math dialect to ROCDL library calls";
720  let description = [{
721    This pass converts supported Math ops to ROCDL library calls.
722  }];
723  let dependentDialects = [
724    "arith::ArithDialect",
725    "func::FuncDialect",
726    "ROCDL::ROCDLDialect",
727    "vector::VectorDialect",
728  ];
729}
730
731//===----------------------------------------------------------------------===//
732// MathToSPIRV
733//===----------------------------------------------------------------------===//
734
735def ConvertMathToSPIRV : Pass<"convert-math-to-spirv"> {
736  let summary = "Convert Math dialect to SPIR-V dialect";
737  let constructor = "mlir::createConvertMathToSPIRVPass()";
738  let dependentDialects = ["spirv::SPIRVDialect"];
739}
740
741//===----------------------------------------------------------------------===//
742// MathToEmitC
743//===----------------------------------------------------------------------===//
744
745def ConvertMathToEmitC  : Pass<"convert-math-to-emitc"> {
746  let summary = "Convert some Math operations to EmitC call_opaque operations";
747  let description = [{
748    This pass converts supported Math ops to `call_opaque` ops targeting libc/libm
749    functions. Unlike convert-math-to-funcs pass, converting to `call_opaque` ops
750    allows to overload the same function with different argument types.
751  }];
752  let dependentDialects = ["emitc::EmitCDialect"];
753  let options = [
754    Option<"languageTarget", "language-target", "::mlir::emitc::LanguageTarget",
755      /*default=*/"::mlir::emitc::LanguageTarget::c99", "Select the language standard target for callees (c99 or cpp11).",
756      [{::llvm::cl::values(
757            clEnumValN(::mlir::emitc::LanguageTarget::c99, "c99", "c99"),
758            clEnumValN(::mlir::emitc::LanguageTarget::cpp11, "cpp11", "cpp11")
759          )}]>
760  ];
761}
762
763//===----------------------------------------------------------------------===//
764// MathToFuncs
765//===----------------------------------------------------------------------===//
766
767def ConvertMathToFuncs : Pass<"convert-math-to-funcs", "ModuleOp"> {
768  let summary = "Convert Math operations to calls of outlined implementations.";
769  let description = [{
770    This pass converts supported Math ops to calls of compiler generated
771    functions implementing these operations in software.
772    The LLVM dialect is used for LinkonceODR linkage of the generated functions.
773  }];
774  let dependentDialects = [
775    "arith::ArithDialect",
776    "cf::ControlFlowDialect",
777    "func::FuncDialect",
778    "scf::SCFDialect",
779    "vector::VectorDialect",
780    "LLVM::LLVMDialect",
781  ];
782  let options = [
783    Option<"minWidthOfFPowIExponent", "min-width-of-fpowi-exponent", "unsigned",
784           /*default=*/"1",
785           "Convert FPowI only if the width of its exponent's integer type "
786           "is greater than or equal to this value">,
787    // Most backend targets support a native ctlz operation, so by default
788    // ctrlz conversion is disabled.
789    Option<"convertCtlz", "convert-ctlz", "bool", /*default=*/"false",
790           "Convert math.ctlz to a software implementation. Enable "
791           "for targets that do not natively support ctlz.">,
792  ];
793}
794
795//===----------------------------------------------------------------------===//
796// MemRefToEmitC
797//===----------------------------------------------------------------------===//
798
799def ConvertMemRefToEmitC : Pass<"convert-memref-to-emitc"> {
800  let summary = "Convert MemRef dialect to EmitC dialect";
801  let dependentDialects = ["emitc::EmitCDialect"];
802}
803
804//===----------------------------------------------------------------------===//
805// MemRefToLLVM
806//===----------------------------------------------------------------------===//
807
808def FinalizeMemRefToLLVMConversionPass :
809    Pass<"finalize-memref-to-llvm", "ModuleOp"> {
810  let summary = "Finalize MemRef dialect to LLVM dialect conversion";
811  let description = [{
812    Finalize the conversion of the operations from the MemRef
813    dialect to the LLVM dialect.
814    This conversion will not convert some complex MemRef
815    operations. Make sure to run `expand-strided-metadata`
816    beforehand for these.
817  }];
818  let dependentDialects = ["LLVM::LLVMDialect"];
819  let options = [
820    Option<"useAlignedAlloc", "use-aligned-alloc", "bool", /*default=*/"false",
821           "Use aligned_alloc in place of malloc for heap allocations">,
822    Option<"indexBitwidth", "index-bitwidth", "unsigned",
823           /*default=kDeriveIndexBitwidthFromDataLayout*/"0",
824           "Bitwidth of the index type, 0 to use size of machine word">,
825    Option<"useGenericFunctions", "use-generic-functions",
826           "bool",
827           /*default=*/"false",
828           "Use generic allocation and deallocation functions instead of the "
829           "classic 'malloc', 'aligned_alloc' and 'free' functions">
830  ];
831}
832
833//===----------------------------------------------------------------------===//
834// MemRefToSPIRV
835//===----------------------------------------------------------------------===//
836
837def MapMemRefStorageClass : Pass<"map-memref-spirv-storage-class"> {
838  let summary = "Map numeric MemRef memory spaces to SPIR-V storage classes";
839  let constructor = "mlir::createMapMemRefStorageClassPass()";
840  let dependentDialects = ["spirv::SPIRVDialect"];
841  let options = [
842    Option<"clientAPI", "client-api", "std::string", /*default=*/"\"vulkan\"",
843           "The client API to use for populating mappings">
844  ];
845}
846
847def ConvertMemRefToSPIRV : Pass<"convert-memref-to-spirv"> {
848  let summary = "Convert MemRef dialect to SPIR-V dialect";
849  let constructor = "mlir::createConvertMemRefToSPIRVPass()";
850  let dependentDialects = ["spirv::SPIRVDialect"];
851  let options = [
852    Option<"boolNumBits", "bool-num-bits",
853           "int", /*default=*/"8",
854           "The number of bits to store a boolean value">,
855    Option<"use64bitIndex", "use-64bit-index",
856           "bool", /*default=*/"false",
857           "Use 64-bit integers to convert index types">
858  ];
859}
860
861//===----------------------------------------------------------------------===//
862// MeshToMPI
863//===----------------------------------------------------------------------===//
864
865def ConvertMeshToMPIPass : Pass<"convert-mesh-to-mpi"> {
866  let summary = "Convert Mesh dialect to MPI dialect.";
867  let constructor = "mlir::createConvertMeshToMPIPass()";
868  let description = [{
869    This pass converts communication operations from the Mesh dialect to the
870    MPI dialect.
871    If it finds a global named "static_mpi_rank" it will use that splat value
872    instead of calling MPI_Comm_rank. This allows optimizations like constant
873    shape propagation and fusion because shard/partition sizes depend on the
874    rank.
875  }];
876  let dependentDialects = [
877    "memref::MemRefDialect",
878    "mpi::MPIDialect",
879    "scf::SCFDialect",
880    "bufferization::BufferizationDialect"
881  ];
882}
883
884//===----------------------------------------------------------------------===//
885// NVVMToLLVM
886//===----------------------------------------------------------------------===//
887
888def ConvertNVVMToLLVMPass : Pass<"convert-nvvm-to-llvm"> {
889  let summary = "Convert NVVM to PTX with Inline Assembly in LLVM dialect";
890  let description = [{
891    This pass generates PTX instructions using inline assembly for NVVM
892    operations implements `BasicPtxBuilderInterface`.
893  }];
894  let dependentDialects = [
895    "NVVM::NVVMDialect",
896  ];
897}
898
899//===----------------------------------------------------------------------===//
900// NVGPUToNVVM
901//===----------------------------------------------------------------------===//
902
903def ConvertNVGPUToNVVMPass : Pass<"convert-nvgpu-to-nvvm"> {
904  let summary = "Convert NVGPU dialect to NVVM dialect";
905  let description = [{
906    This pass converts supported NVGPU ops to NVVM dialect intrinsics.
907  }];
908
909  let dependentDialects = [
910    "NVVM::NVVMDialect",
911  ];
912}
913
914
915//===----------------------------------------------------------------------===//
916// OpenACCToSCF
917//===----------------------------------------------------------------------===//
918
919def ConvertOpenACCToSCF : Pass<"convert-openacc-to-scf", "ModuleOp"> {
920  let summary = "Convert the OpenACC ops to OpenACC with SCF dialect";
921  let constructor = "mlir::createConvertOpenACCToSCFPass()";
922  let dependentDialects = ["scf::SCFDialect", "acc::OpenACCDialect"];
923}
924
925//===----------------------------------------------------------------------===//
926// OpenMPToLLVM
927//===----------------------------------------------------------------------===//
928
929def ConvertOpenMPToLLVMPass : Pass<"convert-openmp-to-llvm", "ModuleOp"> {
930  let summary = "Convert the OpenMP ops to OpenMP ops with LLVM dialect";
931  let dependentDialects = ["LLVM::LLVMDialect"];
932}
933
934//===----------------------------------------------------------------------===//
935// PDLToPDLInterp
936//===----------------------------------------------------------------------===//
937
938def ConvertPDLToPDLInterp : Pass<"convert-pdl-to-pdl-interp", "ModuleOp"> {
939  let summary = "Convert PDL ops to PDL interpreter ops";
940  let constructor = "mlir::createPDLToPDLInterpPass()";
941  let dependentDialects = ["pdl_interp::PDLInterpDialect"];
942}
943
944//===----------------------------------------------------------------------===//
945// ReconcileUnrealizedCasts
946//===----------------------------------------------------------------------===//
947
948def ReconcileUnrealizedCasts : Pass<"reconcile-unrealized-casts"> {
949  let summary = "Simplify and eliminate unrealized conversion casts";
950  let description = [{
951    Eliminate `unrealized_conversion_cast` operations, commonly introduced by
952    partial dialect conversions, that transitively convert a value to another
953    value of the same type, that is:
954
955    ```
956    %0 = "producer.op"() : () -> !type.A
957    %1 = unrealized_conversion_cast %0 : !type.A to !type.B
958    %2 = unrealized_conversion_cast %1 : !type.B to !type.C
959    %3 = unrealized_conversion_cast %2 : !type.C to !type.A
960    "consumer.op"(%3) : (!type.A) -> ()
961    ```
962
963    Such situations appear when the consumer operation is converted by one pass
964    and the producer operation is converted by another pass, each of which
965    produces an unrealized cast. This pass can be used to clean up the IR.
966  }];
967  let constructor = "mlir::createReconcileUnrealizedCastsPass()";
968}
969
970//===----------------------------------------------------------------------===//
971// SCFToControlFlow
972//===----------------------------------------------------------------------===//
973
974def SCFToControlFlow : Pass<"convert-scf-to-cf"> {
975  let summary = "Convert SCF dialect to ControlFlow dialect, replacing structured"
976                " control flow with a CFG";
977  let constructor = "mlir::createConvertSCFToCFPass()";
978  let dependentDialects = ["cf::ControlFlowDialect"];
979}
980
981//===----------------------------------------------------------------------===//
982// SCFToOpenMP
983//===----------------------------------------------------------------------===//
984
985def ConvertSCFToOpenMPPass : Pass<"convert-scf-to-openmp", "ModuleOp"> {
986  let summary = "Convert SCF parallel loop to OpenMP parallel + workshare "
987                "constructs.";
988
989  let options = [
990    Option<"numThreads", "num-threads", "unsigned",
991           /*default=kUseOpenMPDefaultNumThreads*/"0",
992           "Number of threads to use">
993  ];
994
995  let dependentDialects = ["omp::OpenMPDialect", "LLVM::LLVMDialect",
996                           "memref::MemRefDialect"];
997}
998
999//===----------------------------------------------------------------------===//
1000// SCFToSPIRV
1001//===----------------------------------------------------------------------===//
1002
1003def SCFToSPIRV : Pass<"convert-scf-to-spirv"> {
1004  let summary = "Convert SCF dialect to SPIR-V dialect.";
1005  let description = [{
1006    Converts SCF ops into SPIR-V structured control flow ops.
1007    SPIR-V structured control flow ops do not support yielding values.
1008    So for SCF ops yielding values, SPIR-V variables are created for
1009    holding the values and load/store operations are emitted for updating
1010    them.
1011  }];
1012  let dependentDialects = ["spirv::SPIRVDialect"];
1013}
1014
1015//===----------------------------------------------------------------------===//
1016// SCFToGPU
1017//===----------------------------------------------------------------------===//
1018
1019def ConvertAffineForToGPU
1020    : InterfacePass<"convert-affine-for-to-gpu", "FunctionOpInterface"> {
1021  let summary = "Convert top-level AffineFor Ops to GPU kernels";
1022  let constructor = "mlir::createAffineForToGPUPass()";
1023  let dependentDialects = ["gpu::GPUDialect"];
1024  let options = [
1025    Option<"numBlockDims", "gpu-block-dims", "unsigned", /*default=*/"1u",
1026           "Number of GPU block dimensions for mapping">,
1027    Option<"numThreadDims", "gpu-thread-dims", "unsigned", /*default=*/"1u",
1028           "Number of GPU thread dimensions for mapping">
1029  ];
1030}
1031
1032def ConvertParallelLoopToGpu : Pass<"convert-parallel-loops-to-gpu"> {
1033  let summary = "Convert mapped scf.parallel ops to gpu launch operations";
1034  let constructor = "mlir::createParallelLoopToGpuPass()";
1035  let dependentDialects = ["affine::AffineDialect", "gpu::GPUDialect"];
1036}
1037
1038//===----------------------------------------------------------------------===//
1039// SCFToEmitC
1040//===----------------------------------------------------------------------===//
1041
1042def SCFToEmitC : Pass<"convert-scf-to-emitc"> {
1043  let summary = "Convert SCF dialect to EmitC dialect, maintaining structured"
1044                " control flow";
1045  let dependentDialects = ["emitc::EmitCDialect"];
1046}
1047
1048//===----------------------------------------------------------------------===//
1049// ShapeToStandard
1050//===----------------------------------------------------------------------===//
1051
1052def ConvertShapeToStandard : Pass<"convert-shape-to-std", "ModuleOp"> {
1053  let summary = "Convert operations from the shape dialect into the standard "
1054                "dialect";
1055  let constructor = "mlir::createConvertShapeToStandardPass()";
1056  let dependentDialects = [
1057    "scf::SCFDialect",
1058  ];
1059}
1060
1061def ConvertShapeConstraints : Pass<"convert-shape-constraints"> {
1062  let summary = "Convert shape constraint operations to the standard dialect";
1063  let description = [{
1064    This pass eliminates shape constraints from the program, converting them to
1065    eager (side-effecting) error handling code.
1066
1067    This pass is separate from the regular convert-shape-to-standard, despite
1068    converting between the same dialects, because converting shape constraints
1069    can happen at a different part of the program than general shape
1070    computation lowering.
1071  }];
1072  let constructor = "mlir::createConvertShapeConstraintsPass()";
1073  let dependentDialects = ["cf::ControlFlowDialect", "scf::SCFDialect"];
1074}
1075
1076//===----------------------------------------------------------------------===//
1077// SPIRVToLLVM
1078//===----------------------------------------------------------------------===//
1079
1080def ConvertSPIRVToLLVMPass : Pass<"convert-spirv-to-llvm", "ModuleOp"> {
1081  let summary = "Convert SPIR-V dialect to LLVM dialect";
1082  let description = [{
1083    See https://mlir.llvm.org/docs/SPIRVToLLVMDialectConversion/
1084    for more details.
1085  }];
1086  let dependentDialects = ["LLVM::LLVMDialect"];
1087
1088  let options = [
1089    Option<"clientAPI", "client-api", "::mlir::spirv::ClientAPI",
1090	   /*default=*/"::mlir::spirv::ClientAPI::Unknown",
1091	   "Derive StorageClass to address space mapping from the client API",
1092	   [{::llvm::cl::values(
1093	     clEnumValN(::mlir::spirv::ClientAPI::Unknown, "Unknown", "Unknown (default)"),
1094	     clEnumValN(::mlir::spirv::ClientAPI::Metal, "Metal", "Metal"),
1095	     clEnumValN(::mlir::spirv::ClientAPI::OpenCL, "OpenCL", "OpenCL"),
1096	     clEnumValN(::mlir::spirv::ClientAPI::Vulkan, "Vulkan", "Vulkan"),
1097	     clEnumValN(::mlir::spirv::ClientAPI::WebGPU, "WebGPU", "WebGPU")
1098	   )}]>,
1099  ];
1100}
1101
1102//===----------------------------------------------------------------------===//
1103// TensorToLinalg
1104//===----------------------------------------------------------------------===//
1105
1106def ConvertTensorToLinalg : Pass<"convert-tensor-to-linalg", "ModuleOp"> {
1107  let summary = "Convert some Tensor dialect ops to Linalg dialect";
1108  let constructor = "mlir::createConvertTensorToLinalgPass()";
1109  let dependentDialects = [
1110    "arith::ArithDialect",
1111    "linalg::LinalgDialect",
1112  ];
1113}
1114
1115//===----------------------------------------------------------------------===//
1116// TensorToSPIRV
1117//===----------------------------------------------------------------------===//
1118
1119def ConvertTensorToSPIRV : Pass<"convert-tensor-to-spirv"> {
1120  let summary = "Convert Tensor dialect to SPIR-V dialect";
1121  let constructor = "mlir::createConvertTensorToSPIRVPass()";
1122  let dependentDialects = ["spirv::SPIRVDialect"];
1123  let options = [
1124    Option<"emulateLT32BitScalarTypes", "emulate-lt-32-bit-scalar-types",
1125           "bool", /*default=*/"true",
1126           "Emulate narrower scalar types with 32-bit ones if not supported by"
1127           " the target">
1128  ];
1129}
1130
1131//===----------------------------------------------------------------------===//
1132// TosaToArith
1133//===----------------------------------------------------------------------===//
1134
1135def TosaToArith : Pass<"tosa-to-arith"> {
1136  let summary = "Lower TOSA to the Arith dialect";
1137  let dependentDialects = [
1138    "arith::ArithDialect",
1139  ];
1140  let description = [{
1141    Pass that converts TOSA operations to the equivalent operations using the
1142    operations in the Arith dialect. The ApplyScale operator is optionally
1143    included as it is often preserved until the final invocation.
1144  }];
1145
1146  let options = [
1147    Option<"includeApplyRescale", "include-apply-rescale",
1148           "bool", /*default=*/"false",
1149           "Whether to include the lowering for tosa.apply_rescale to arith">,
1150    Option<"use32Bit", "use-32-bit",
1151           "bool", /*default=*/"false",
1152           "Whether to prioritze lowering to 32-bit operations">
1153  ];
1154
1155  let constructor = "tosa::createTosaToArith()";
1156}
1157
1158//===----------------------------------------------------------------------===//
1159// TosaToLinalg
1160//===----------------------------------------------------------------------===//
1161
1162def TosaToLinalg
1163    : InterfacePass<"tosa-to-linalg", "FunctionOpInterface"> {
1164  let summary = "Lower TOSA to LinAlg on tensors";
1165  let description = [{
1166    Pass that converts TOSA operations to the equivalent operations using the
1167    tensor operations in LinAlg.
1168  }];
1169
1170  let constructor = "tosa::createTosaToLinalg()";
1171  let options = [
1172    Option<"disableTosaDecompositions", "disable-tosa-decompositions",
1173           "bool", /*default=*/"false",
1174           "Disable tosa decompositions pass">,
1175    Option<"aggressiveReduceConstant", "aggressive-reduce-constant",
1176           "bool", /*default=*/"false",
1177           "Always perform the reduce constant optimization">
1178  ];
1179}
1180
1181//===----------------------------------------------------------------------===//
1182// TosaToLinalgNamed
1183//===----------------------------------------------------------------------===//
1184
1185def TosaToLinalgNamed
1186    : InterfacePass<"tosa-to-linalg-named", "FunctionOpInterface"> {
1187  let summary = "Lower TOSA to LinAlg named operations";
1188  let description = [{
1189    Pass that converts TOSA operations to the equivalent operations using the
1190    Linalg named operations.
1191  }];
1192
1193  let options = [
1194      Option<"preferConv2DKernelLayoutHWCF", "prefer-conv2d-kernel-layout-hwcf",
1195           "bool", /*default=*/"false",
1196           "Prefer generating linalg.conv_2d_nhwc_hwcf over linalg.conv_2d_nhwc_fhwc">
1197  ];
1198
1199  let constructor = "tosa::createTosaToLinalgNamed()";
1200}
1201
1202//===----------------------------------------------------------------------===//
1203// TosaToMLProgram
1204//===----------------------------------------------------------------------===//
1205
1206def TosaToMLProgram : Pass<"tosa-to-mlprogram", "ModuleOp"> {
1207  let summary = "Lower TOSA to the MLProgram dialect";
1208  let dependentDialects = ["ml_program::MLProgramDialect"];
1209  let description = [{
1210    Pass that converts TOSA's variable operator operations to the equivalent
1211    MLProgram operations.
1212  }];
1213}
1214
1215//===----------------------------------------------------------------------===//
1216// TosaToSCF
1217//===----------------------------------------------------------------------===//
1218
1219def TosaToSCF : Pass<"tosa-to-scf"> {
1220  let summary = "Lower TOSA to the SCF dialect";
1221  let dependentDialects = ["tensor::TensorDialect, scf::SCFDialect"];
1222  let description = [{
1223    Pass that converts TOSA's control flow operations to the equivalent SCF
1224    operations.
1225  }];
1226
1227  let constructor = "tosa::createTosaToSCF()";
1228}
1229
1230//===----------------------------------------------------------------------===//
1231// TosaToTensor
1232//===----------------------------------------------------------------------===//
1233
1234def TosaToTensor : Pass<"tosa-to-tensor"> {
1235  let summary = "Lower TOSA to the Tensor dialect";
1236  let dependentDialects = [
1237    "tensor::TensorDialect",
1238  ];
1239  let description = [{
1240    Pass that converts TOSA operations to the equivalent operations using the
1241    operations in the Tensor dialect.
1242  }];
1243
1244  let constructor = "tosa::createTosaToTensor()";
1245}
1246
1247//===----------------------------------------------------------------------===//
1248// UBToLLVM
1249//===----------------------------------------------------------------------===//
1250
1251def UBToLLVMConversionPass : Pass<"convert-ub-to-llvm"> {
1252  let summary = "Convert UB dialect to LLVM dialect";
1253  let description = [{
1254    This pass converts supported UB ops to LLVM dialect instructions.
1255  }];
1256  let dependentDialects = ["LLVM::LLVMDialect"];
1257  let options = [
1258    Option<"indexBitwidth", "index-bitwidth", "unsigned",
1259           /*default=kDeriveIndexBitwidthFromDataLayout*/"0",
1260           "Bitwidth of the index type, 0 to use size of machine word">,
1261  ];
1262}
1263
1264//===----------------------------------------------------------------------===//
1265// UBToSPIRV
1266//===----------------------------------------------------------------------===//
1267
1268def UBToSPIRVConversionPass : Pass<"convert-ub-to-spirv"> {
1269  let summary = "Convert UB dialect to SPIR-V dialect";
1270  let description = [{
1271    This pass converts supported UB ops to SPIR-V dialect ops.
1272  }];
1273  let dependentDialects = ["spirv::SPIRVDialect"];
1274}
1275
1276//===----------------------------------------------------------------------===//
1277// VectorToGPU
1278//===----------------------------------------------------------------------===//
1279
1280def ConvertVectorToGPU : Pass<"convert-vector-to-gpu"> {
1281  let summary = "Lower the operations from the vector dialect into the GPU "
1282                "dialect";
1283  let constructor = "mlir::createConvertVectorToGPUPass()";
1284  let dependentDialects = [
1285    "memref::MemRefDialect", "gpu::GPUDialect", "affine::AffineDialect",
1286    "vector::VectorDialect", "nvgpu::NVGPUDialect"
1287  ];
1288
1289  let options = [
1290    Option<"useNvGpu", "use-nvgpu", "bool", /*default=*/"false",
1291      "convert to NvGPU ops instead of GPU dialect ops">
1292  ];
1293}
1294
1295//===----------------------------------------------------------------------===//
1296// VectorToSCF
1297//===----------------------------------------------------------------------===//
1298
1299def ConvertVectorToSCF : Pass<"convert-vector-to-scf"> {
1300  let summary = "Lower the operations from the vector dialect into the SCF "
1301                "dialect";
1302  let constructor = "mlir::createConvertVectorToSCFPass()";
1303  let dependentDialects = [
1304    "affine::AffineDialect",
1305    "memref::MemRefDialect",
1306    "scf::SCFDialect",
1307    "tensor::TensorDialect"
1308  ];
1309  let options = [
1310    Option<"fullUnroll", "full-unroll", "bool", /*default=*/"false",
1311           "Perform full unrolling when converting vector transfers to SCF">,
1312    Option<"targetRank", "target-rank", "unsigned", /*default=*/"1",
1313           "Target vector rank to which transfer ops should be lowered">,
1314    Option<"lowerTensors", "lower-tensors", "bool", /*default=*/"false",
1315           "Lower transfer ops that operate on tensors">,
1316    Option<"lowerScalable", "lower-scalable", "bool", /*default=*/"false",
1317           "Add scalable vector specific lowerings (that introduce loops)">
1318  ];
1319}
1320
1321//===----------------------------------------------------------------------===//
1322// VectorToArmSME
1323//===----------------------------------------------------------------------===//
1324
1325def ConvertVectorToArmSME : Pass<"convert-vector-to-arm-sme"> {
1326  let summary = "Lower the operations from the vector dialect into the ArmSME "
1327                "dialect";
1328  let constructor = "mlir::createConvertVectorToArmSMEPass()";
1329  let description = [{
1330    Pass that converts vector dialect operations into equivalent ArmSME dialect
1331    operations.
1332  }];
1333  let dependentDialects = ["arm_sme::ArmSMEDialect", "arm_sve::ArmSVEDialect"];
1334}
1335
1336//===----------------------------------------------------------------------===//
1337// ArmSMEToSCF
1338//===----------------------------------------------------------------------===//
1339
1340def ConvertArmSMEToSCF : Pass<"convert-arm-sme-to-scf"> {
1341  let summary = "Lower the operations from the ArmSME dialect into the SCF "
1342                "dialect";
1343  let constructor = "mlir::createConvertArmSMEToSCFPass()";
1344  let dependentDialects = [
1345    "scf::SCFDialect",
1346    "arith::ArithDialect",
1347    "vector::VectorDialect",
1348    "arm_sme::ArmSMEDialect"
1349  ];
1350}
1351
1352//===----------------------------------------------------------------------===//
1353// ArmSMEToLLVM
1354//===----------------------------------------------------------------------===//
1355
1356def ConvertArmSMEToLLVM : InterfacePass<"convert-arm-sme-to-llvm", "FunctionOpInterface"> {
1357  let summary = "Lower the operations from the ArmSME dialect into the LLVM "
1358                "dialect";
1359  let constructor = "mlir::createConvertArmSMEToLLVMPass()";
1360  let dependentDialects = [
1361    "arm_sme::ArmSMEDialect",
1362    "LLVM::LLVMDialect"
1363  ];
1364  let options = [
1365    Option<"dumpTileLiveRanges", "dump-tile-live-ranges",
1366           "bool", /*default=*/"false",
1367           "Dump the live ranges of SME tiles (for debugging)">
1368  ];
1369}
1370
1371//===----------------------------------------------------------------------===//
1372// VectorToLLVM
1373//===----------------------------------------------------------------------===//
1374
1375def ConvertVectorToLLVMPass : Pass<"convert-vector-to-llvm"> {
1376  let summary = "Lower the operations from the vector dialect into the LLVM "
1377                "dialect";
1378  let description = [{
1379
1380    Convert operations from the vector dialect into the LLVM IR dialect
1381    operations. The lowering pass provides several options to control
1382    the kinds of optimizations that are allowed. It also provides options
1383    that enable the use of one or more architectural-specific dialects
1384    (AMX, X86Vector, ArmNeon, ArmSVE, etc.) in combination with the
1385    architectural-neutral vector dialect lowering.
1386
1387  }];
1388  // Override explicitly in C++ to allow conditional dialect dependence.
1389  // let dependentDialects;
1390  let options = [
1391    Option<"reassociateFPReductions", "reassociate-fp-reductions",
1392           "bool", /*default=*/"false",
1393           "Allows llvm to reassociate floating-point reductions for speed">,
1394    Option<"force32BitVectorIndices", "force-32bit-vector-indices",
1395           "bool", /*default=*/"true",
1396           "Allows compiler to assume vector indices fit in 32-bit if that "
1397     "yields faster code">,
1398    Option<"amx", "enable-amx",
1399           "bool", /*default=*/"false",
1400           "Enables the use of AMX dialect while lowering the vector "
1401	   "dialect.">,
1402    Option<"armNeon", "enable-arm-neon",
1403           "bool", /*default=*/"false",
1404           "Enables the use of ArmNeon dialect while lowering the vector "
1405	   "dialect.">,
1406    Option<"armSVE", "enable-arm-sve",
1407           "bool", /*default=*/"false",
1408           "Enables the use of ArmSVE dialect while lowering the vector "
1409       "dialect.">,
1410    Option<"x86Vector", "enable-x86vector",
1411           "bool", /*default=*/"false",
1412           "Enables the use of X86Vector dialect while lowering the vector "
1413	   "dialect.">,
1414    Option<"vectorTransformsOptions", "vector-transform-options",
1415           "vector::VectorTransformsOptions",
1416           /*default=*/"vector::VectorTransformsOptions()",
1417           "Options to lower some operations like contractions and transposes.">,
1418  ];
1419}
1420
1421//===----------------------------------------------------------------------===//
1422// VectorToSPIRV
1423//===----------------------------------------------------------------------===//
1424
1425def ConvertVectorToSPIRV : Pass<"convert-vector-to-spirv"> {
1426  let summary = "Convert Vector dialect to SPIR-V dialect";
1427  let constructor = "mlir::createConvertVectorToSPIRVPass()";
1428  let dependentDialects = [
1429    "spirv::SPIRVDialect",
1430    "ub::UBDialect"
1431  ];
1432}
1433
1434//===----------------------------------------------------------------------===//
1435// VectorToXeGPU
1436//===----------------------------------------------------------------------===//
1437
1438def ConvertVectorToXeGPU : Pass<"convert-vector-to-xegpu"> {
1439  let summary = "Lower the operations from the vector dialect into the XeGPU "
1440                "dialect";
1441  let constructor = "mlir::createConvertVectorToXeGPUPass()";
1442  let dependentDialects = [
1443    "memref::MemRefDialect", "arith::ArithDialect",
1444    "vector::VectorDialect", "xegpu::XeGPUDialect"
1445  ];
1446}
1447
1448#endif // MLIR_CONVERSION_PASSES
1449