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// AffineToStandard 16//===----------------------------------------------------------------------===// 17 18def ConvertAffineToStandard : Pass<"lower-affine"> { 19 let summary = "Lower Affine operations to a combination of Standard and SCF " 20 "operations"; 21 let description = [{ 22 23 Convert operations from the affine dialect into operations from the SCF and 24 standard dialects. 25 26 `affine.for` operations are converted to `scf.for` operations that are free 27 of certain structural restrictions (on their bounds and step). `affine.if` 28 is similarly converted to the `scf.if` operation. `affine.apply` operations 29 are converted into sequences of primitive arithmetic operations from the 30 standard dialect that have the same effect, using operands of the `index` 31 type. Consequently, named maps and sets thare are no longer in use may be 32 removed from the module. 33 34 For example, `%r = affine.apply affine_map<(d0, d1)[s0] -> (d0 + 2*d1 + 35 s0)>(%d0, %d1)[%s0]` 36 can be converted into: 37 38 ```mlir 39 %d0 = <...> 40 %d1 = <...> 41 %s0 = <...> 42 %0 = arith.constant 2 : index 43 %1 = arith.muli %0, %d1 44 %2 = arith.addi %d0, %1 45 %r = arith.addi %2, %s0 46 ``` 47 48 #### Input invariant 49 50 - no `Tensor` types; 51 52 These restrictions may be lifted in the future. 53 54 #### Output IR 55 56 Functions with `affine.for` and `affine.if` operations eliminated. These 57 functions may contain operations from the Standard dialect in addition to 58 those already present before the pass. 59 60 #### Invariants 61 62 - Functions without a body are not modified. 63 - The semantics of the other functions is preserved. 64 - Individual operations other than those mentioned above are not modified 65 if they do not depend on the loop iterator value or on the result of 66 `affine.apply`. 67 }]; 68 let constructor = "mlir::createLowerAffinePass()"; 69 let dependentDialects = [ 70 "memref::MemRefDialect", 71 "scf::SCFDialect", 72 "vector::VectorDialect" 73 ]; 74} 75 76//===----------------------------------------------------------------------===// 77// AMDGPUToROCDL 78//===----------------------------------------------------------------------===// 79 80def ConvertAMDGPUToROCDL : Pass<"convert-amdgpu-to-rocdl"> { 81 let summary = "Convert AMDGPU dialect to ROCDL dialect"; 82 let description = [{ 83 This pass converts supported AMDGPU ops to ROCDL dialect intrinsics. 84 }]; 85 let constructor = "mlir::createConvertAMDGPUToROCDLPass()"; 86 let dependentDialects = [ 87 "LLVM::LLVMDialect", 88 "ROCDL::ROCDLDialect", 89 ]; 90 let options = [Option<"chipset", "chipset", "std::string", 91 /*default=*/"\"gfx000\"", 92 "Chipset that these operations will run on">]; 93} 94 95//===----------------------------------------------------------------------===// 96// ArithmeticToLLVM 97//===----------------------------------------------------------------------===// 98 99def ConvertArithmeticToLLVM : Pass<"convert-arith-to-llvm"> { 100 let summary = "Convert Arithmetic dialect to LLVM dialect"; 101 let description = [{ 102 This pass converts supported Arithmetic ops to LLVM dialect instructions. 103 }]; 104 let constructor = "mlir::arith::createConvertArithmeticToLLVMPass()"; 105 let dependentDialects = ["LLVM::LLVMDialect"]; 106 let options = [ 107 Option<"indexBitwidth", "index-bitwidth", "unsigned", 108 /*default=kDeriveIndexBitwidthFromDataLayout*/"0", 109 "Bitwidth of the index type, 0 to use size of machine word">, 110 ]; 111} 112 113//===----------------------------------------------------------------------===// 114// ArithmeticToSPIRV 115//===----------------------------------------------------------------------===// 116 117def ConvertArithmeticToSPIRV : Pass<"convert-arith-to-spirv", "ModuleOp"> { 118 let summary = "Convert Arithmetic dialect to SPIR-V dialect"; 119 let constructor = "mlir::arith::createConvertArithmeticToSPIRVPass()"; 120 let dependentDialects = ["spirv::SPIRVDialect"]; 121 let options = [ 122 Option<"emulateNon32BitScalarTypes", "emulate-non-32-bit-scalar-types", 123 "bool", /*default=*/"true", 124 "Emulate non-32-bit scalar types with 32-bit ones if " 125 "missing native support"> 126 ]; 127} 128 129//===----------------------------------------------------------------------===// 130// ArmNeon2dToIntr 131//===----------------------------------------------------------------------===// 132 133def ConvertArmNeon2dToIntr : Pass<"arm-neon-2d-to-intr"> { 134 let summary = "Convert Arm NEON structured ops to intrinsics"; 135 let constructor = "mlir::createConvertArmNeon2dToIntrPass()"; 136 let dependentDialects = ["arm_neon::ArmNeonDialect", "vector::VectorDialect"]; 137} 138 139//===----------------------------------------------------------------------===// 140// AsyncToLLVM 141//===----------------------------------------------------------------------===// 142 143def ConvertAsyncToLLVM : Pass<"convert-async-to-llvm", "ModuleOp"> { 144 let summary = "Convert the operations from the async dialect into the LLVM " 145 "dialect"; 146 let description = [{ 147 Convert `async.execute` operations to LLVM coroutines and use async runtime 148 API to execute them. 149 }]; 150 let constructor = "mlir::createConvertAsyncToLLVMPass()"; 151 let dependentDialects = [ 152 "arith::ArithmeticDialect", 153 "LLVM::LLVMDialect", 154 ]; 155} 156 157//===----------------------------------------------------------------------===// 158// BufferizationToMemRef 159//===----------------------------------------------------------------------===// 160 161def ConvertBufferizationToMemRef : Pass<"convert-bufferization-to-memref"> { 162 let summary = "Convert operations from the Bufferization dialect to the " 163 "MemRef dialect"; 164 let description = [{ 165 166 This pass converts bufferization operations into memref operations. 167 168 In the current state, this pass only transforms a `bufferization.clone` 169 operation into `memref.alloc` and `memref.copy` operations. This conversion 170 is needed, since some clone operations could remain after applying several 171 transformation processes. Currently, only `canonicalize` transforms clone 172 operations or even eliminates them. This can lead to errors if any clone op 173 survived after all conversion passes (starting from the bufferization 174 dialect) are performed. 175 176 See: 177 https://llvm.discourse.group/t/bufferization-error-related-to-memref-clone/4665 178 179 To avoid these errors, this pass can be performed as a last clean-up pass to 180 transform remaining operations and to proceed in other dialects (memref 181 e.g.). 182 183 Note that this pass only transforms the operation without any further 184 analyses. This pass does not consider any memory analysis or optimization 185 and hence does not resolve any memory leaks. 186 187 }]; 188 let constructor = "mlir::createBufferizationToMemRefPass()"; 189 let dependentDialects = ["arith::ArithmeticDialect", "memref::MemRefDialect"]; 190} 191 192//===----------------------------------------------------------------------===// 193// ComplexToLLVM 194//===----------------------------------------------------------------------===// 195 196def ConvertComplexToLLVM : Pass<"convert-complex-to-llvm"> { 197 let summary = "Convert Complex dialect to LLVM dialect"; 198 let constructor = "mlir::createConvertComplexToLLVMPass()"; 199 let dependentDialects = ["LLVM::LLVMDialect"]; 200} 201 202//===----------------------------------------------------------------------===// 203// ComplexToLibm 204//===----------------------------------------------------------------------===// 205 206def ConvertComplexToLibm : Pass<"convert-complex-to-libm", "ModuleOp"> { 207 let summary = "Convert Complex dialect to libm calls"; 208 let description = [{ 209 This pass converts supported Complex ops to libm calls. 210 }]; 211 let constructor = "mlir::createConvertComplexToLibmPass()"; 212 let dependentDialects = [ 213 "func::FuncDialect", 214 ]; 215} 216 217//===----------------------------------------------------------------------===// 218// ComplexToStandard 219//===----------------------------------------------------------------------===// 220 221def ConvertComplexToStandard : Pass<"convert-complex-to-standard"> { 222 let summary = "Convert Complex dialect to standard dialect"; 223 let constructor = "mlir::createConvertComplexToStandardPass()"; 224 let dependentDialects = ["math::MathDialect"]; 225} 226 227//===----------------------------------------------------------------------===// 228// ControlFlowToLLVM 229//===----------------------------------------------------------------------===// 230 231def ConvertControlFlowToLLVM : Pass<"convert-cf-to-llvm", "ModuleOp"> { 232 let summary = "Convert ControlFlow operations to the LLVM dialect"; 233 let description = [{ 234 Convert ControlFlow operations into LLVM IR dialect operations. 235 236 If other operations are present and their results are required by the LLVM 237 IR dialect operations, the pass will fail. Any LLVM IR operations or types 238 already present in the IR will be kept as is. 239 }]; 240 let constructor = "mlir::cf::createConvertControlFlowToLLVMPass()"; 241 let dependentDialects = ["LLVM::LLVMDialect"]; 242 let options = [ 243 Option<"indexBitwidth", "index-bitwidth", "unsigned", 244 /*default=kDeriveIndexBitwidthFromDataLayout*/"0", 245 "Bitwidth of the index type, 0 to use size of machine word">, 246 ]; 247} 248 249//===----------------------------------------------------------------------===// 250// ControlFlowToSPIRV 251//===----------------------------------------------------------------------===// 252 253def ConvertControlFlowToSPIRV : Pass<"convert-cf-to-spirv", "ModuleOp"> { 254 let summary = "Convert ControlFlow dialect to SPIR-V dialect"; 255 let constructor = "mlir::createConvertControlFlowToSPIRVPass()"; 256 let dependentDialects = ["spirv::SPIRVDialect"]; 257 let options = [ 258 Option<"emulateNon32BitScalarTypes", "emulate-non-32-bit-scalar-types", 259 "bool", /*default=*/"true", 260 "Emulate non-32-bit scalar types with 32-bit ones if " 261 "missing native support"> 262 ]; 263} 264 265//===----------------------------------------------------------------------===// 266// FuncToLLVM 267//===----------------------------------------------------------------------===// 268 269def ConvertFuncToLLVM : Pass<"convert-func-to-llvm", "ModuleOp"> { 270 let summary = "Convert from the Func dialect to the LLVM dialect"; 271 let description = [{ 272 Convert Func dialect operations into the LLVM IR dialect operations. 273 274 #### Input invariant 275 276 - no `tensor` types; 277 - all `vector` are one-dimensional; 278 - all blocks are reachable by following the successors of the first basic 279 block; 280 281 If other operations are present and their results are required by the LLVM 282 IR dialect operations, the pass will fail. Any LLVM IR operations or types 283 already present in the IR will be kept as is. 284 285 #### Output IR 286 287 Functions converted to LLVM IR. Function arguments types are converted 288 one-to-one. Function results are converted one-to-one and, in case more than 289 1 value is returned, packed into an LLVM IR struct type. Function calls and 290 returns are updated accordingly. Block argument types are updated to use 291 LLVM IR types. 292 }]; 293 let constructor = "mlir::createConvertFuncToLLVMPass()"; 294 let dependentDialects = ["LLVM::LLVMDialect"]; 295 let options = [ 296 Option<"useBarePtrCallConv", "use-bare-ptr-memref-call-conv", "bool", 297 /*default=*/"false", 298 "Replace FuncOp's MemRef arguments with bare pointers to the MemRef " 299 "element types">, 300 Option<"indexBitwidth", "index-bitwidth", "unsigned", 301 /*default=kDeriveIndexBitwidthFromDataLayout*/"0", 302 "Bitwidth of the index type, 0 to use size of machine word">, 303 Option<"dataLayout", "data-layout", "std::string", 304 /*default=*/"\"\"", 305 "String description (LLVM format) of the data layout that is " 306 "expected on the produced module"> 307 ]; 308} 309 310//===----------------------------------------------------------------------===// 311// FuncToSPIRV 312//===----------------------------------------------------------------------===// 313 314def ConvertFuncToSPIRV : Pass<"convert-func-to-spirv", "ModuleOp"> { 315 let summary = "Convert Func dialect to SPIR-V dialect"; 316 let constructor = "mlir::createConvertFuncToSPIRVPass()"; 317 let dependentDialects = ["spirv::SPIRVDialect"]; 318 let options = [ 319 Option<"emulateNon32BitScalarTypes", "emulate-non-32-bit-scalar-types", 320 "bool", /*default=*/"true", 321 "Emulate non-32-bit scalar types with 32-bit ones if " 322 "missing native support"> 323 ]; 324} 325 326//===----------------------------------------------------------------------===// 327// GPUCommon 328//===----------------------------------------------------------------------===// 329 330def GpuToLLVMConversionPass : Pass<"gpu-to-llvm", "ModuleOp"> { 331 let summary = "Convert GPU dialect to LLVM dialect with GPU runtime calls"; 332 let constructor = "mlir::createGpuToLLVMConversionPass()"; 333 let dependentDialects = ["LLVM::LLVMDialect"]; 334} 335 336def LowerHostCodeToLLVM : Pass<"lower-host-to-llvm", "ModuleOp"> { 337 let summary = "Lowers the host module code and `gpu.launch_func` to LLVM"; 338 let constructor = "mlir::createLowerHostCodeToLLVMPass()"; 339 let dependentDialects = ["LLVM::LLVMDialect"]; 340} 341 342//===----------------------------------------------------------------------===// 343// GPUToNVVM 344//===----------------------------------------------------------------------===// 345 346def ConvertGpuOpsToNVVMOps : Pass<"convert-gpu-to-nvvm", "gpu::GPUModuleOp"> { 347 let summary = "Generate NVVM operations for gpu operations"; 348 let constructor = "mlir::createLowerGpuOpsToNVVMOpsPass()"; 349 let dependentDialects = [ 350 "cf::ControlFlowDialect", 351 "memref::MemRefDialect", 352 "NVVM::NVVMDialect", 353 ]; 354 let options = [ 355 Option<"indexBitwidth", "index-bitwidth", "unsigned", 356 /*default=kDeriveIndexBitwidthFromDataLayout*/"0", 357 "Bitwidth of the index type, 0 to use size of machine word"> 358 ]; 359} 360 361//===----------------------------------------------------------------------===// 362// GPUToROCDL 363//===----------------------------------------------------------------------===// 364 365def ConvertGpuOpsToROCDLOps : Pass<"convert-gpu-to-rocdl", "gpu::GPUModuleOp"> { 366 let summary = "Generate ROCDL operations for gpu operations"; 367 let constructor = "mlir::createLowerGpuOpsToROCDLOpsPass()"; 368 let dependentDialects = ["ROCDL::ROCDLDialect"]; 369 let options = [ 370 Option<"chipset", "chipset", "std::string", 371 /*default=*/"\"gfx000\"", 372 "Chipset that these operations will run on">, 373 Option<"indexBitwidth", "index-bitwidth", "unsigned", 374 /*default=kDeriveIndexBitwidthFromDataLayout*/"0", 375 "Bitwidth of the index type, 0 to use size of machine word">, 376 Option<"runtime", "runtime", "::mlir::gpu::amd::Runtime", 377 "::mlir::gpu::amd::Runtime::Unknown", 378 "Runtime code will be run on (default is Unknown, can also use HIP or OpenCl)", 379 [{::llvm::cl::values( 380 clEnumValN(::mlir::gpu::amd::Runtime::Unknown, "unknown", "Unknown (default)"), 381 clEnumValN(::mlir::gpu::amd::Runtime::HIP, "HIP", "HIP"), 382 clEnumValN(::mlir::gpu::amd::Runtime::OpenCL, "OpenCL", "OpenCL") 383 )}]> 384 ]; 385} 386 387//===----------------------------------------------------------------------===// 388// GPUToSPIRV 389//===----------------------------------------------------------------------===// 390 391def ConvertGPUToSPIRV : Pass<"convert-gpu-to-spirv", "ModuleOp"> { 392 let summary = "Convert GPU dialect to SPIR-V dialect"; 393 let description = [{ 394 This pass converts supported GPU device ops to SPIR-V ops. It does not 395 handle GPU host ops. 396 397 A `gpu.func` op can have parameters to pass in resources. But in SPIR-V 398 entry functions cannot take parameters; they use descriptors to access 399 resources. By default, parameters to a `gpu.func` op will be converted to 400 global variables. These global variables will be assigned sequential binding 401 numbers following their order in the original `gpu.func` op, starting from 402 0, in set 0. One can attach `spv.interface_var_abi` to those parameters 403 to control the set and binding if wanted. 404 }]; 405 let constructor = "mlir::createConvertGPUToSPIRVPass()"; 406 let dependentDialects = ["spirv::SPIRVDialect"]; 407} 408 409//===----------------------------------------------------------------------===// 410// GPUToVulkan 411//===----------------------------------------------------------------------===// 412 413def ConvertGpuLaunchFuncToVulkanLaunchFunc 414 : Pass<"convert-gpu-launch-to-vulkan-launch", "ModuleOp"> { 415 let summary = "Convert gpu.launch_func to vulkanLaunch external call"; 416 let description = [{ 417 This pass is only intended for the mlir-vulkan-runner. 418 }]; 419 let constructor = "mlir::createConvertGpuLaunchFuncToVulkanLaunchFuncPass()"; 420 let dependentDialects = ["spirv::SPIRVDialect"]; 421} 422 423def ConvertVulkanLaunchFuncToVulkanCalls 424 : Pass<"launch-func-to-vulkan", "ModuleOp"> { 425 let summary = "Convert vulkanLaunch external call to Vulkan runtime external " 426 "calls"; 427 let description = [{ 428 This pass is only intended for the mlir-vulkan-runner. 429 }]; 430 let constructor = "mlir::createConvertVulkanLaunchFuncToVulkanCallsPass()"; 431 let dependentDialects = ["LLVM::LLVMDialect"]; 432} 433 434//===----------------------------------------------------------------------===// 435// LinalgToLLVM 436//===----------------------------------------------------------------------===// 437 438def ConvertLinalgToLLVM : Pass<"convert-linalg-to-llvm", "ModuleOp"> { 439 let summary = "Convert the operations from the linalg dialect into the LLVM " 440 "dialect"; 441 let constructor = "mlir::createConvertLinalgToLLVMPass()"; 442 let dependentDialects = ["scf::SCFDialect", "LLVM::LLVMDialect"]; 443} 444 445//===----------------------------------------------------------------------===// 446// LinalgToStandard 447//===----------------------------------------------------------------------===// 448 449def ConvertLinalgToStandard : Pass<"convert-linalg-to-std", "ModuleOp"> { 450 let summary = "Convert the operations from the linalg dialect into the " 451 "Standard dialect"; 452 let constructor = "mlir::createConvertLinalgToStandardPass()"; 453 let dependentDialects = ["func::FuncDialect", "memref::MemRefDialect"]; 454} 455 456//===----------------------------------------------------------------------===// 457// LinalgToSPIRV 458//===----------------------------------------------------------------------===// 459 460def ConvertLinalgToSPIRV : Pass<"convert-linalg-to-spirv", "ModuleOp"> { 461 let summary = "Convert Linalg dialect to SPIR-V dialect"; 462 let description = [{ 463 This pass converts supported Linalg ops to SPIR-V ops. It's quite 464 experimental and are expected to migrate to other proper conversions. 465 }]; 466 let constructor = "mlir::createLinalgToSPIRVPass()"; 467 let dependentDialects = ["spirv::SPIRVDialect"]; 468} 469 470//===----------------------------------------------------------------------===// 471// MathToLibm 472//===----------------------------------------------------------------------===// 473 474def ConvertMathToLibm : Pass<"convert-math-to-libm", "ModuleOp"> { 475 let summary = "Convert Math dialect to libm calls"; 476 let description = [{ 477 This pass converts supported Math ops to libm calls. 478 }]; 479 let constructor = "mlir::createConvertMathToLibmPass()"; 480 let dependentDialects = [ 481 "arith::ArithmeticDialect", 482 "func::FuncDialect", 483 "vector::VectorDialect", 484 ]; 485} 486 487//===----------------------------------------------------------------------===// 488// MathToLLVM 489//===----------------------------------------------------------------------===// 490 491def ConvertMathToLLVM : Pass<"convert-math-to-llvm"> { 492 let summary = "Convert Math dialect to LLVM dialect"; 493 let description = [{ 494 This pass converts supported Math ops to LLVM dialect intrinsics. 495 }]; 496 let constructor = "mlir::createConvertMathToLLVMPass()"; 497 let dependentDialects = ["LLVM::LLVMDialect"]; 498} 499 500//===----------------------------------------------------------------------===// 501// MathToSPIRV 502//===----------------------------------------------------------------------===// 503 504def ConvertMathToSPIRV : Pass<"convert-math-to-spirv", "ModuleOp"> { 505 let summary = "Convert Math dialect to SPIR-V dialect"; 506 let constructor = "mlir::createConvertMathToSPIRVPass()"; 507 let dependentDialects = ["spirv::SPIRVDialect"]; 508} 509 510//===----------------------------------------------------------------------===// 511// MemRefToLLVM 512//===----------------------------------------------------------------------===// 513 514def ConvertMemRefToLLVM : Pass<"convert-memref-to-llvm", "ModuleOp"> { 515 let summary = "Convert operations from the MemRef dialect to the LLVM " 516 "dialect"; 517 let constructor = "mlir::createMemRefToLLVMPass()"; 518 let dependentDialects = ["LLVM::LLVMDialect"]; 519 let options = [ 520 Option<"useAlignedAlloc", "use-aligned-alloc", "bool", /*default=*/"false", 521 "Use aligned_alloc in place of malloc for heap allocations">, 522 Option<"indexBitwidth", "index-bitwidth", "unsigned", 523 /*default=kDeriveIndexBitwidthFromDataLayout*/"0", 524 "Bitwidth of the index type, 0 to use size of machine word">, 525 Option<"useGenericFunctions", "use-generic-functions", 526 "bool", 527 /*default=*/"false", 528 "Use generic allocation and deallocation functions instead of the " 529 "classic 'malloc', 'aligned_alloc' and 'free' functions"> 530 ]; 531} 532 533//===----------------------------------------------------------------------===// 534// MemRefToSPIRV 535//===----------------------------------------------------------------------===// 536 537def ConvertMemRefToSPIRV : Pass<"convert-memref-to-spirv", "ModuleOp"> { 538 let summary = "Convert MemRef dialect to SPIR-V dialect"; 539 let constructor = "mlir::createConvertMemRefToSPIRVPass()"; 540 let dependentDialects = ["spirv::SPIRVDialect"]; 541 let options = [ 542 Option<"boolNumBits", "bool-num-bits", 543 "int", /*default=*/"8", 544 "The number of bits to store a boolean value"> 545 ]; 546} 547 548//===----------------------------------------------------------------------===// 549// NVGPUToNVVM 550//===----------------------------------------------------------------------===// 551 552def ConvertNVGPUToNVVM : Pass<"convert-nvgpu-to-nvvm"> { 553 let summary = "Convert NVGPU dialect to NVVM dialect"; 554 let description = [{ 555 This pass converts supported NVGPU ops to NVVM dialect intrinsics. 556 }]; 557 let constructor = "mlir::createConvertNVGPUToNVVMPass()"; 558 let dependentDialects = [ 559 "NVVM::NVVMDialect", 560 ]; 561} 562 563 564//===----------------------------------------------------------------------===// 565// OpenACCToSCF 566//===----------------------------------------------------------------------===// 567 568def ConvertOpenACCToSCF : Pass<"convert-openacc-to-scf", "ModuleOp"> { 569 let summary = "Convert the OpenACC ops to OpenACC with SCF dialect"; 570 let constructor = "mlir::createConvertOpenACCToSCFPass()"; 571 let dependentDialects = ["scf::SCFDialect", "acc::OpenACCDialect"]; 572} 573 574//===----------------------------------------------------------------------===// 575// OpenACCToLLVM 576//===----------------------------------------------------------------------===// 577 578def ConvertOpenACCToLLVM : Pass<"convert-openacc-to-llvm", "ModuleOp"> { 579 let summary = "Convert the OpenACC ops to LLVM dialect"; 580 let constructor = "mlir::createConvertOpenACCToLLVMPass()"; 581 let dependentDialects = ["LLVM::LLVMDialect"]; 582} 583 584//===----------------------------------------------------------------------===// 585// OpenMPToLLVM 586//===----------------------------------------------------------------------===// 587 588def ConvertOpenMPToLLVM : Pass<"convert-openmp-to-llvm", "ModuleOp"> { 589 let summary = "Convert the OpenMP ops to OpenMP ops with LLVM dialect"; 590 let constructor = "mlir::createConvertOpenMPToLLVMPass()"; 591 let dependentDialects = ["LLVM::LLVMDialect"]; 592} 593 594//===----------------------------------------------------------------------===// 595// PDLToPDLInterp 596//===----------------------------------------------------------------------===// 597 598def ConvertPDLToPDLInterp : Pass<"convert-pdl-to-pdl-interp", "ModuleOp"> { 599 let summary = "Convert PDL ops to PDL interpreter ops"; 600 let constructor = "mlir::createPDLToPDLInterpPass()"; 601 let dependentDialects = ["pdl_interp::PDLInterpDialect"]; 602} 603 604//===----------------------------------------------------------------------===// 605// ReconcileUnrealizedCasts 606//===----------------------------------------------------------------------===// 607 608def ReconcileUnrealizedCasts : Pass<"reconcile-unrealized-casts"> { 609 let summary = "Simplify and eliminate unrealized conversion casts"; 610 let description = [{ 611 Eliminate `unrealized_conversion_cast` operations, commonly introduced by 612 partial dialect conversions, that transitively convert a value to another 613 value of the same type, that is: 614 615 ``` 616 %0 = "producer.op"() : () -> !type.A 617 %1 = unrealized_conversion_cast %0 : !type.A to !type.B 618 %2 = unrealized_conversion_cast %1 : !type.B to !type.A 619 "consumer.op"(%2) : (!type.A) -> () 620 ``` 621 622 Such situations appear when the consumer operation is converted by one pass 623 and the producer operation is converted by another pass, each of which 624 produces an unrealized cast. This pass can be used to clean up the IR. 625 }]; 626 let constructor = "mlir::createReconcileUnrealizedCastsPass()"; 627} 628 629//===----------------------------------------------------------------------===// 630// SCFToControlFlow 631//===----------------------------------------------------------------------===// 632 633def SCFToControlFlow : Pass<"convert-scf-to-cf"> { 634 let summary = "Convert SCF dialect to ControlFlow dialect, replacing structured" 635 " control flow with a CFG"; 636 let constructor = "mlir::createConvertSCFToCFPass()"; 637 let dependentDialects = ["cf::ControlFlowDialect"]; 638} 639 640//===----------------------------------------------------------------------===// 641// SCFToOpenMP 642//===----------------------------------------------------------------------===// 643 644def ConvertSCFToOpenMP : Pass<"convert-scf-to-openmp", "ModuleOp"> { 645 let summary = "Convert SCF parallel loop to OpenMP parallel + workshare " 646 "constructs."; 647 let constructor = "mlir::createConvertSCFToOpenMPPass()"; 648 let dependentDialects = ["omp::OpenMPDialect", "LLVM::LLVMDialect", 649 "memref::MemRefDialect"]; 650} 651 652//===----------------------------------------------------------------------===// 653// SCFToSPIRV 654//===----------------------------------------------------------------------===// 655 656def SCFToSPIRV : Pass<"convert-scf-to-spirv", "ModuleOp"> { 657 let summary = "Convert SCF dialect to SPIR-V dialect."; 658 let description = [{ 659 This pass converts SCF ops into SPIR-V structured control flow ops. 660 SPIR-V structured control flow ops does not support yielding values. 661 So for SCF ops yielding values, SPIR-V variables are created for 662 holding the values and load/store operations are emitted for updating 663 them. 664 }]; 665 let constructor = "mlir::createConvertSCFToSPIRVPass()"; 666 let dependentDialects = ["spirv::SPIRVDialect"]; 667} 668 669//===----------------------------------------------------------------------===// 670// SCFToGPU 671//===----------------------------------------------------------------------===// 672 673def ConvertAffineForToGPU 674 : InterfacePass<"convert-affine-for-to-gpu", "FunctionOpInterface"> { 675 let summary = "Convert top-level AffineFor Ops to GPU kernels"; 676 let constructor = "mlir::createAffineForToGPUPass()"; 677 let dependentDialects = ["gpu::GPUDialect"]; 678 let options = [ 679 Option<"numBlockDims", "gpu-block-dims", "unsigned", /*default=*/"1u", 680 "Number of GPU block dimensions for mapping">, 681 Option<"numThreadDims", "gpu-thread-dims", "unsigned", /*default=*/"1u", 682 "Number of GPU thread dimensions for mapping"> 683 ]; 684} 685 686def ConvertParallelLoopToGpu : Pass<"convert-parallel-loops-to-gpu"> { 687 let summary = "Convert mapped scf.parallel ops to gpu launch operations"; 688 let constructor = "mlir::createParallelLoopToGpuPass()"; 689 let dependentDialects = ["AffineDialect", "gpu::GPUDialect"]; 690} 691 692//===----------------------------------------------------------------------===// 693// ShapeToStandard 694//===----------------------------------------------------------------------===// 695 696def ConvertShapeToStandard : Pass<"convert-shape-to-std", "ModuleOp"> { 697 let summary = "Convert operations from the shape dialect into the standard " 698 "dialect"; 699 let constructor = "mlir::createConvertShapeToStandardPass()"; 700 let dependentDialects = [ 701 "scf::SCFDialect", 702 ]; 703} 704 705def ConvertShapeConstraints : Pass<"convert-shape-constraints"> { 706 let summary = "Convert shape constraint operations to the standard dialect"; 707 let description = [{ 708 This pass eliminates shape constraints from the program, converting them to 709 eager (side-effecting) error handling code. 710 711 This pass is separate from the regular convert-shape-to-standard, despite 712 converting between the same dialects, because converting shape constraints 713 can happen at a different part of the program than general shape 714 computation lowering. 715 }]; 716 let constructor = "mlir::createConvertShapeConstraintsPass()"; 717 let dependentDialects = ["cf::ControlFlowDialect", "scf::SCFDialect"]; 718} 719 720//===----------------------------------------------------------------------===// 721// SPIRVToLLVM 722//===----------------------------------------------------------------------===// 723 724def ConvertSPIRVToLLVM : Pass<"convert-spirv-to-llvm", "ModuleOp"> { 725 let summary = "Convert SPIR-V dialect to LLVM dialect"; 726 let description = [{ 727 See https://mlir.llvm.org/docs/SPIRVToLLVMDialectConversion/ 728 for more details. 729 }]; 730 let constructor = "mlir::createConvertSPIRVToLLVMPass()"; 731 let dependentDialects = ["LLVM::LLVMDialect"]; 732} 733 734//===----------------------------------------------------------------------===// 735// TensorToLinalg 736//===----------------------------------------------------------------------===// 737 738def ConvertTensorToLinalg : Pass<"convert-tensor-to-linalg", "ModuleOp"> { 739 let summary = "Convert some Tensor dialect ops to Linalg dialect"; 740 let constructor = "mlir::createConvertTensorToLinalgPass()"; 741 let dependentDialects = [ 742 "arith::ArithmeticDialect", 743 "linalg::LinalgDialect", 744 ]; 745} 746 747 748//===----------------------------------------------------------------------===// 749// TensorToSPIRV 750//===----------------------------------------------------------------------===// 751 752def ConvertTensorToSPIRV : Pass<"convert-tensor-to-spirv", "ModuleOp"> { 753 let summary = "Convert Tensor dialect to SPIR-V dialect"; 754 let constructor = "mlir::createConvertTensorToSPIRVPass()"; 755 let dependentDialects = ["spirv::SPIRVDialect"]; 756 let options = [ 757 Option<"emulateNon32BitScalarTypes", "emulate-non-32-bit-scalar-types", 758 "bool", /*default=*/"true", 759 "Emulate non-32-bit scalar types with 32-bit ones if " 760 "missing native support"> 761 ]; 762} 763 764//===----------------------------------------------------------------------===// 765// TosaToArith 766//===----------------------------------------------------------------------===// 767 768def TosaToArith : Pass<"tosa-to-arith"> { 769 let summary = "Lower TOSA to the Arith dialect"; 770 let dependentDialects = [ 771 "arith::ArithmeticDialect", 772 ]; 773 let description = [{ 774 Pass that converts TOSA operations to the equivalent operations using the 775 operations in the Arith dialect. The ApplyScale operator is optionally 776 included as it is often preserved until the final invocation. 777 }]; 778 779 let options = [ 780 Option<"includeApplyRescale", "include-apply-rescale", 781 "bool", /*default=*/"false", 782 "Whether to include the lowering for tosa.apply_rescale to arith">, 783 Option<"use32Bit", "use-32-bit", 784 "bool", /*default=*/"false", 785 "Whether to prioritze lowering to 32-bit operations"> 786 ]; 787 788 let constructor = "tosa::createTosaToArith()"; 789} 790 791//===----------------------------------------------------------------------===// 792// TosaToLinalg 793//===----------------------------------------------------------------------===// 794 795def TosaToLinalg 796 : InterfacePass<"tosa-to-linalg", "FunctionOpInterface"> { 797 let summary = "Lower TOSA to LinAlg on tensors"; 798 let description = [{ 799 Pass that converts TOSA operations to the equivalent operations using the 800 tensor operations in LinAlg. 801 }]; 802 803 let constructor = "tosa::createTosaToLinalg()"; 804} 805 806//===----------------------------------------------------------------------===// 807// TosaToLinalgNamed 808//===----------------------------------------------------------------------===// 809 810def TosaToLinalgNamed 811 : InterfacePass<"tosa-to-linalg-named", "FunctionOpInterface"> { 812 let summary = "Lower TOSA to LinAlg named operations"; 813 let description = [{ 814 Pass that converts TOSA operations to the equivalent operations using the 815 Linalg named operations. 816 }]; 817 818 let constructor = "tosa::createTosaToLinalgNamed()"; 819} 820 821//===----------------------------------------------------------------------===// 822// TosaToSCF 823//===----------------------------------------------------------------------===// 824 825def TosaToSCF : Pass<"tosa-to-scf"> { 826 let summary = "Lower TOSA to the SCF dialect"; 827 let dependentDialects = ["tensor::TensorDialect, scf::SCFDialect"]; 828 let description = [{ 829 Pass that converts TOSA's control flow operations to the equivalent SCF 830 operations. 831 }]; 832 833 let constructor = "tosa::createTosaToSCF()"; 834} 835 836//===----------------------------------------------------------------------===// 837// TosaToTensor 838//===----------------------------------------------------------------------===// 839 840def TosaToTensor : Pass<"tosa-to-tensor"> { 841 let summary = "Lower TOSA to the Tensor dialect"; 842 let dependentDialects = [ 843 "tensor::TensorDialect", 844 ]; 845 let description = [{ 846 Pass that converts TOSA operations to the equivalent operations using the 847 operations in the Tensor dialect. 848 }]; 849 850 let constructor = "tosa::createTosaToTensor()"; 851} 852 853//===----------------------------------------------------------------------===// 854// VectorToGPU 855//===----------------------------------------------------------------------===// 856 857def ConvertVectorToGPU : Pass<"convert-vector-to-gpu"> { 858 let summary = "Lower the operations from the vector dialect into the GPU " 859 "dialect"; 860 let constructor = "mlir::createConvertVectorToGPUPass()"; 861 let dependentDialects = [ 862 "memref::MemRefDialect", "gpu::GPUDialect", "AffineDialect", 863 "vector::VectorDialect", "nvgpu::NVGPUDialect" 864 ]; 865 866 let options = [ 867 Option<"useNvGpu", "use-nvgpu", "bool", /*default=*/"false", 868 "convert to NvGPU ops instead of GPU dialect ops"> 869 ]; 870} 871 872//===----------------------------------------------------------------------===// 873// VectorToSCF 874//===----------------------------------------------------------------------===// 875 876def ConvertVectorToSCF : Pass<"convert-vector-to-scf"> { 877 let summary = "Lower the operations from the vector dialect into the SCF " 878 "dialect"; 879 let constructor = "mlir::createConvertVectorToSCFPass()"; 880 let dependentDialects = [ 881 "AffineDialect", 882 "memref::MemRefDialect", 883 "scf::SCFDialect" 884 ]; 885 let options = [ 886 Option<"fullUnroll", "full-unroll", "bool", /*default=*/"false", 887 "Perform full unrolling when converting vector transfers to SCF">, 888 Option<"targetRank", "target-rank", "unsigned", /*default=*/"1", 889 "Target vector rank to which transfer ops should be lowered">, 890 Option<"lowerPermutationMaps", "lower-permutation-maps", "bool", 891 /*default=*/"false", "Replace permutation maps with vector " 892 "transposes/broadcasts before lowering transfer ops">, 893 Option<"lowerTensors", "lower-tensors", "bool", /*default=*/"false", 894 "Lower transfer ops that operate on tensors"> 895 ]; 896} 897 898//===----------------------------------------------------------------------===// 899// VectorToLLVM 900//===----------------------------------------------------------------------===// 901 902def ConvertVectorToLLVM : Pass<"convert-vector-to-llvm", "ModuleOp"> { 903 let summary = "Lower the operations from the vector dialect into the LLVM " 904 "dialect"; 905 let description = [{ 906 907 Convert operations from the vector dialect into the LLVM IR dialect 908 operations. The lowering pass provides several options to control 909 the kinds of optimizations that are allowed. It also provides options 910 that enable the use of one or more architectural-specific dialects 911 (AMX, X86Vector, ArmNeon, ArmSVE, etc.) in combination with the 912 architectural-neutral vector dialect lowering. 913 914 }]; 915 let constructor = "mlir::createConvertVectorToLLVMPass()"; 916 // Override explicitly in C++ to allow conditional dialect dependence. 917 // let dependentDialects; 918 let options = [ 919 Option<"reassociateFPReductions", "reassociate-fp-reductions", 920 "bool", /*default=*/"false", 921 "Allows llvm to reassociate floating-point reductions for speed">, 922 Option<"force32BitVectorIndices", "force-32bit-vector-indices", 923 "bool", /*default=*/"true", 924 "Allows compiler to assume vector indices fit in 32-bit if that " 925 "yields faster code">, 926 Option<"amx", "enable-amx", 927 "bool", /*default=*/"false", 928 "Enables the use of AMX dialect while lowering the vector " 929 "dialect.">, 930 Option<"armNeon", "enable-arm-neon", 931 "bool", /*default=*/"false", 932 "Enables the use of ArmNeon dialect while lowering the vector " 933 "dialect.">, 934 Option<"armSVE", "enable-arm-sve", 935 "bool", /*default=*/"false", 936 "Enables the use of ArmSVE dialect while lowering the vector " 937 "dialect.">, 938 Option<"x86Vector", "enable-x86vector", 939 "bool", /*default=*/"false", 940 "Enables the use of X86Vector dialect while lowering the vector " 941 "dialect."> 942 ]; 943} 944 945//===----------------------------------------------------------------------===// 946// VectorToSPIRV 947//===----------------------------------------------------------------------===// 948 949def ConvertVectorToSPIRV : Pass<"convert-vector-to-spirv", "ModuleOp"> { 950 let summary = "Convert Vector dialect to SPIR-V dialect"; 951 let constructor = "mlir::createConvertVectorToSPIRVPass()"; 952 let dependentDialects = ["spirv::SPIRVDialect"]; 953} 954 955#endif // MLIR_CONVERSION_PASSES 956