[all-commits] [llvm/llvm-project] cdf7ca: [MLIR][ROCDL] Add conversion for gpu.lane_id to ROCDL

[SJW via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: cdf7ca6db76b09d58e7c5ac1d8156034dd29c3c7
      https://github.com/llvm/llvm-project/commit/cdf7ca6db76b09d58e7c5ac1d8156034dd29c3c7
  Author: SJW <swaters at amd.com>
  Date:   2023-07-26 (Wed, 26 Jul 2023)

  Changed paths:
    M mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
    M mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
    M mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl.mlir
    M mlir/test/Target/LLVMIR/rocdl.mlir

  Log Message:
  -----------
  [MLIR][ROCDL] Add conversion for gpu.lane_id to ROCDL

Creates rocdl.lane_id op with llvm conversion to:

  __device__ static unsigned int __lane_id() {
      return  __builtin_amdgcn_mbcnt_hi(
                 -1, __builtin_amdgcn_mbcnt_lo(-1, 0));
  }

Reviewed By: krzysz00

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