<table border="1" cellspacing="0" cellpadding="8">
    <tr>
        <th>Issue</th>
        <td>
            <a href=http://email.email.llvm.org/c/eJytVtuS2jgQ_Rrz0jWULxjMgx-4LBmqMpupSbL7SMl2Y2sRltcSMMzXb0viZpJNpiqhXGBduvv06dYRmSyO6SP3wpnnzz1_4r6XwLZwaLnmdQmPy2fYYFujUMDqAnK5bbgwKweuK9AVwsePfz1Bw_INK7EAXsPLp9kTxH2_H4CS4IWJcbLHVnFZe9HELUVDP3jIsmCYhEkBk6c55IKR29M-CAa0yzfWldaNIjsvXNBTUthd1iccNHhhBcr6U4P1jHDtNNKcEPvtQ9PKfzDX0Mr8wSEJQz-m7yjOhkniY4hxEMf-aDBKgiTAPAhY6Of0kozGgzjzwnG_y0ktNc8pP10xbdPWx4bnTMDs63xC6WtNLCnQ0lJEUIDBgbUNLOfEL-yU42bL6qMzYU0jyIGmZJXZwZWJAQKJAcskuSI_NpZ1xK2Dzx-eX1wp1C5T-O8OazKSrCAnC6Vli4qCmQJxDUyRSYGvwFq03t2KItysBdlQJBv_LtcMBScYFJpA0cNMVY6wbk_hTtkaOCaVjvUjUihro1FpyJlCUzu3Y-ifHjtcrUohMyZWK9hLys5kuSy-kBlVfU050e_Ezq4MYrQ0XeapW3Al12uFurvAOqPWmXGC_SdVFbzR1EUHcA3lcLoeN5yrSh7gUCHx3l7JN1Vc2Jj0ei13iZpo1HzL31zFqM537l3B1rLr7Iaxy85Coqq9cKThINuNO2DB3FWQynbebVLRRJcXzSETMt8si9f-K5hk7XDOt3Y4pYAtssIue9F92j-J5ibCOZQtL7qm7wVwMzze4Tn-Itzif_BGvxHvzfDtDszl04X99kMnr-9z8m5q7twdjD4o22ReGF6aNjzJ9k4UdK6totCxZbYtbzvqcCLIEmXckYPP1NU3lTDbRGdb_N1ta6PajvDZN4fuivfmXMdTCu_F8ytOo2h4qjI7C5aZvPg5n17Kd71l5tga04rt8WrQYskVnVQjhFbt6qu9VYZ4ql1gk1NHUeKpOK0Q_d_F6rTm7OG2Y0fn83156Qrf3xYq6bMpiFSKZwKNeBg1UXKLZ0mx4nySDlKzaHFb-C8VqzfKFrPYszrHXpFGxTgasx7b6Uq26R90wbz1dq1If3CFmuvy7tY0R0apnRHORTzwA79XpTgajMOCJaN1PBgHwwizcRgNBxldllGAQ9YTjG4NlRIf1HM1HsC6oHfipvfrCHga-qF5In8Uj-NBf0xAWIh5hEkc-ci8gY9bxkXf-OnLtuy1qXWZ7UpFi4IaQV0XGbFe1ogWMCHUXAtMz_esu8Id_Ret_-Z2Jq0_X8dr2kLGha0H_Zn58PwVMvpLhHXRs5mkNo3_AD8H7Rg>54010</a>
        </td>
    </tr>

    <tr>
        <th>Summary</th>
        <td>
            warp id computation patterns not leveraged to use SGPRs afterwards in AMDGPU backend
        </td>
    </tr>

    <tr>
      <th>Labels</th>
      <td>
            new issue
      </td>
    </tr>

    <tr>
      <th>Assignees</th>
      <td>
      </td>
    </tr>

    <tr>
      <th>Reporter</th>
      <td>
          Epliz
      </td>
    </tr>
</table>

<pre>
    Hi,

I am writing HIP kernels and compiling with the LLVM packaged in ROCM 5.0.1 so (HIP version: 5.0.13601-bb16828d AMD clang version 14.0.0 (https://github.com/RadeonOpenCompute/llvm-project roc-5.0.1 22051 235b6880e2e515507478181ec11a20c1ec87945b).

I noticed that the typical CUDA patterns to compute a warp ID, used in many CUDA applications, is not leveraged to put the warp id in SGPRs and subsequent loads/stores using it as index are not using scalar operations.

I believe this is a very frequent pattern in CUDA.

Here is a test case:
```
__global__ void warpIdTest(float* warp_scales, float* lane_offsets, float* a, float* res, int N) {
  // test kernel to show whether the warp ID/ lane ID patterns get optimized to use
  // SGPRs for the warp ID

  // doesn't work with 1D indexing
  int tid = blockIdx.x * blockDim.x + threadIdx.x;
  // doesn't work with 1D indexing with 2D grid
  //int tid = blockIdx.x * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x;
  // doesn't workd with 1D indexing with 3D grid
  //int tid = blockIdx.x * blockDim.x * blockDim.y * blockDim.z
  //          + threadIdx.z * blockDim.y * blockDim.x
  //          + threadIdx.y * blockDim.x + threadIdx.x;

  // wid is the "warp ID" and could be put in a SGPR
  int wid = tid / warpSize;
  int lid = tid % warpSize;
  if (tid < N) {
    // warp_scales[wid] could be loaded with a scalar load
   // the fma should have a scalar register as operand
    res[tid] = lane_offsets[lid] + warp_scales[wid] * a[tid];
  }
}
```
Would it be possible to get some optimization for that?

Thanks in advance
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJytVclu4zgQ_Rr5UoihxfJy0MHLuGOgMx2ku2eOBiWWZI5p0SPSdpyv7yLpTe4twEwgOKLIevVq4atc8WP2KIJ4GoSzIBz73wWwDRwaYURdwePiGdbY1Cg1sJpDoTZbIe3OQZgVmBXCx49_PcGWFWtWIQdRw8un6ROk3bAbgVYQxEMLssdGC1UHydhvJf0wesjzqD-MhxzGTzMoJCPY0zmIenQqtNYrY7aa7IJ4Tk9Fbnd5l3jQ4oVxVPWnLdZT4rUzSN-k3G8eto36BwsDjSoePJM4DlP6TdK8PxyGGGMapWk46A2G0TDCIopYHBb0MhyMemkexKNuOye1MqKg-MyKGRe2OW5FwSRMv87GFL4xlCUNRrkUERVgcGDNFhYzyi_stM_NhtVHb8K2W0kAhoLV9oTQ1gdIpAy4TBIU4ThfDkg4gM8fnl98KfQu1_jvDmsyUowTyFwb1aAmZ7ZAwgDTZMLxFViDDt3vaOLNGlBb8uT838WaoxREg1wTKXqYrcoRyubk7hStpWNDaVk_IrlyNga1gYJptLXzJ_rh6XHL5bKSKmdyuYS9ouhslAv-hcyo6iXFRP_H7uvSMkaXpst36hZcqrLUaNobrLVqvJkg2n9SVSEYTLx3AN9QnqfvcZtzvVIHOKyQ8t5ck2-rOHc-6fVa7goNpdGIjXjzFaM638H7gpWqDXaTsctJrlDXQTwwcFDN2l-waOYrSGU7n7ahGEpXkMwgl6pYL_hr9xVssG45Exu3nJDDBhl320FyH_ZvvPkP8QyqRvC26XsJ3CyPd3yO_5Eu_wnf5H_ke7N8uyNz-WvTfvslyOv7QN6dmju4g9UH7ZosiONL08Yn2d5JTvfaKQpdW-ba8rajDqcEuURZOAL4TF19Uwl7TLaOpT88VlrV9gmffnfprnxv7nU6IfdBOrvytIqGpyqzs2DZjxec8-2leMsNs9fWmq7YHq8GDVZC0021QujUrr7aO2VIJ8Y7tjG1FCWdyNMOpf-HXL3WnBFuO3Zwvt-Xl7bw_e2okj7bgiitRS7RiodVE602eJYUJ84n6SA1S-a3hf-yYvVau2LyPasL7PAs4aNkxDpGGInZeWj4eeSxLsL13agh4TrPlpKOkDF34DSZPzx_hZzmO9a8s2tk9ouBbIfv3Qy2F1DrnZXhedoLo7CzyljZZzT78xHnvB8N-6P-iLZGLI94Uca9fkcymkE6o-xSB9d4AAdB75TpjsjiMLZPEg5SMuyOyrTHYiwSHKZJiCzohbhhQnYtj65qqk6TOUr5rtK0Kakp9HWTUQWqGtG5I3y2MyvVZH_QfH7rOMeZI_4NL1rYkQ">