[all-commits] [llvm/llvm-project] 0a6283: [mlir][gpu][transforms] Add support for mapping to...

[Nicolas Vasilache via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 0a628369698101b9db7eb16cee1f8ba33a23db4c
      https://github.com/llvm/llvm-project/commit/0a628369698101b9db7eb16cee1f8ba33a23db4c
  Author: Nicolas Vasilache <Nico.Vasilache at amd.com>
  Date:   2025-07-07 (Mon, 07 Jul 2025)

  Changed paths:
    M mlir/include/mlir/Dialect/GPU/IR/GPUDeviceMappingAttr.td
    M mlir/include/mlir/Dialect/GPU/TransformOps/Utils.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/test/Dialect/GPU/transform-gpu.mlir

  Log Message:
  -----------
  [mlir][gpu][transforms] Add support for mapping to lanes (#146912)

This revision adds a new attribute for mapping `scf.forall` to linear
lane ids.

Example:
```
    // %arg2 and %arg3 map to lanes [0, 6) and are turned into epxressions
    // involving threadIdx.x/y by the map_nested_forall_to_threads
    // transformation. This results in a if (linear_thread_id < 6) conditional.
    scf.forall (%arg2, %arg3) in (2, 3) {
       ...
    } {mapping = [#gpu.lane<linear_dim_0>, #gpu.lane<linear_dim_1>]}
 ```

---------

Co-authored-by: Oleksandr "Alex" Zinenko <git at ozinenko.com>



To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications