[all-commits] [llvm/llvm-project] 44e631: [mlir][transforms] Revamp the implementation of ma...

[Nicolas Vasilache via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 44e6318ceacdc00d4f9b0fbb2814d6dc03e27f7d
      https://github.com/llvm/llvm-project/commit/44e6318ceacdc00d4f9b0fbb2814d6dc03e27f7d
  Author: Nicolas Vasilache <nicolasvasilache at users.noreply.github.com>
  Date:   2023-07-26 (Wed, 26 Jul 2023)

  Changed paths:
    M mlir/include/mlir/Dialect/GPU/TransformOps/GPUDeviceMappingAttr.td
    M mlir/include/mlir/Dialect/GPU/TransformOps/GPUTransformOps.h
    M mlir/include/mlir/Dialect/GPU/TransformOps/GPUTransformOps.td
    M mlir/include/mlir/Dialect/GPU/TransformOps/Utils.h
    M mlir/include/mlir/Dialect/SCF/IR/DeviceMappingInterface.td
    M mlir/include/mlir/Dialect/Utils/StaticValueUtils.h
    M mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
    M mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp
    M mlir/lib/Dialect/GPU/TransformOps/Utils.cpp
    M mlir/lib/Dialect/Linalg/TransformOps/GPUHeuristics.cpp
    M mlir/test/Dialect/GPU/transform-gpu-failing.mlir
    M mlir/test/Dialect/GPU/transform-gpu.mlir
    M mlir/test/Dialect/Linalg/transform-op-gpu-map-copy-to-threads.mlir

  Log Message:
  -----------
  [mlir][transforms] Revamp the implementation of mapping loops to GPUs

This revision significantly simplifies the specification and implementation of mapping loops to GPU ids.

Each type of mapping (block, warpgroup, warp, thread) now comes with 2 mapping modes:
  1. a 3-D "grid-like" mode, subject to alignment considerations on threadIdx.x, on which predication
     may occur on a per-dimension 3-D sub-rectangle basis.
  2. a n-D linearized mode, on which predication may only occur on a linear basis.

In the process, better size and alignment requirement inference are introduced along with improved runtime verification messages.

The `warp_dims` attribute was deemed confusing and is removed from the transform in favor of better size inference.

Differential Revision: https://reviews.llvm.org/D155941