<table border="1" cellspacing="0" cellpadding="8">
    <tr>
        <th>Issue</th>
        <td>
            <a href=https://github.com/llvm/llvm-project/issues/65720>65720</a>
        </td>
    </tr>

    <tr>
        <th>Summary</th>
        <td>
            [LICM] Loop-invariant load instruction is not move to preheader of loop
        </td>
    </tr>

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

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

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

<pre>
    **Base llvm amd-stg-open, execute this command:**
opt -filter-print-funcs=_ZN6hipcub10DeviceSpmv11CsrMVKernelIiEEvNS0_10SpmvParamsIT_iEE -S -passes="licm" benchmark_device_spmv-hip-amdgcn-amd-amdhsa-gfx906.ll (rename benchmark_device_spmv-hip-amdgcn-amd-amdhsa-gfx906.log to benchmark_device_spmv-hip-amdgcn-amd-amdhsa-gfx906.ll)

**IR output:**
define protected amdgpu_kernel void @_ZN6hipcub10DeviceSpmv11CsrMVKernelIiEEvNS0_10SpmvParamsIT_iEE(ptr addrspace(4) nocapture readonly byref(%"struct.hipcub::DeviceSpmv::SpmvParams") align 8 %0) local_unnamed_addr #4 comdat align 2 {

  .......
.lr.ph: ; preds = %21
  %30 = add nuw nsw i64 %22, %5
  %31 = sext i32 %2 to i64
  %32 = getelementptr inbounds i32, ptr %.sroa.2.0.copyload, i64 %31
  %33 = addrspacecast ptr %32 to ptr addrspace(1)
  %34 = getelementptr inbounds i32, ptr %.sroa.6.0.copyload, i64 %31
  %35 = addrspacecast ptr %34 to ptr addrspace(1)
  br label %36

._crit_edge.loopexit:                             ; preds = %60
  br label %._crit_edge

._crit_edge: ; preds = %._crit_edge.loopexit, %21
  ret void

36: ; preds = %60, %.lr.ph
  %37 = phi i64 [ 0, %.lr.ph ], [ %62, %60 ]
  %38 = mul nuw nsw i64 %37, %28
  %39 = add nuw nsw i64 %30, %38
  %40 = trunc i64 %39 to i32
  **%41 = load i32, ptr addrspace(1) %33, align 4, !tbaa !19**
  %42 = icmp sgt i32 %41, %40
  br i1 %42, label %43, label %60

43: ; preds = %36
  %sext39 = shl i64 %39, 32
  %44 = ashr exact i64 %sext39, 32
  %45 = getelementptr inbounds i32, ptr %.sroa.0.0.copyload, i64 %44
  %46 = addrspacecast ptr %45 to ptr addrspace(1)
  %47 = load i32, ptr addrspace(1) %46, align 4, !tbaa !19
  %48 = mul nsw i32 %47, %.sroa.932.0.copyload
  %49 = getelementptr inbounds i32, ptr %.sroa.4.0.copyload, i64 %44
  %50 = addrspacecast ptr %49 to ptr addrspace(1)
  %51 = load i32, ptr addrspace(1) %50, align 4, !tbaa !19
  %52 = sext i32 %51 to i64
  %53 = getelementptr inbounds i32, ptr %.sroa.5.0.copyload, i64 %52
  %54 = addrspacecast ptr %53 to ptr addrspace(1)
  %55 = load i32, ptr addrspace(1) %54, align 4, !tbaa !19
  %56 = mul nsw i32 %48, %55
  %57 = atomicrmw add ptr addrspace(3) @_ZZN6hipcub10DeviceSpmv11CsrMVKernelIiEEvNS0_10SpmvParamsIT_iEEE7partial, i32 %56 syncscope("agent-one-as") monotonic, align 4
  fence syncscope("workgroup") release
  tail call void @llvm.amdgcn.s.barrier()
  fence syncscope("workgroup") acquire
  br i1 %6, label %58, label %60

58: ; preds = %43
  %59 = load i32, ptr addrspace(3) @_ZZN6hipcub10DeviceSpmv11CsrMVKernelIiEEvNS0_10SpmvParamsIT_iEEE7partial, align 4, !tbaa !19
  store i32 %59, ptr addrspace(1) %35, align 4, !tbaa !19
  br label %60

60: ; preds = %58, %43, %36
  %61 = add nuw nsw i64 %37, 1
  %62 = and i64 %61, 4294967295
  %63 = icmp ugt i64 %29, %62
  br i1 %63, label %36, label %._crit_edge.loopexit, !llvm.loop !22
}

**question:**
**The instruction(%41 = load i32, ptr addrspace(1) %33, align 4, !tbaa !19) is loop-invariant, but it is not optimized to move to block .lr.ph**

[benchmark_device_spmv-hip-amdgcn-amd-amdhsa-gfx906.log](https://github.com/llvm/llvm-project/files/12554567/benchmark_device_spmv-hip-amdgcn-amd-amdhsa-gfx906.log)
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJy0WE1v47wR_jX0hbAg8UOWDj4k6wQI-r4viu6ih14MmhrLbCRSS1JO0l9fkLJi2XGyTrY1ECsShzOjZ575oIVzqtYAS8RvEV_NRO93xi5b-6xERrJ8tjHVyxKRG0RuboUD3DT7Fou2mjtfz00HGpFvGJ5B9h6w3ymHpWlboStEb4ZtKF2h9MZ0Hs-3qvFg551V2s-3vZYO0dX6X3_lO9XJfpOlK9grCd-7dp9l35z9859_A6uheVB3d_u_vqfrLA1rfxdWtO7hx1rd3eH5dzzvhHMQdCFCGiVbRAjegJa7VtjHdRWVrl3X7uc71c1FW9VSh0v42zkxr7fPZZonTYMRKSxo0cKX9psae_M1y4iUA1CH7wjdwz-w6X3X-zMwK9gqDbizxoP0UIWA1F2_foxo4b1RFUYs_T1kESk6b7GoKus6IQGRgiFSYm2k6HxvAVsQldHNC968WNgiUiDCESHO2176ZLAcPKc3R-vD_dEWIiQoFY2qNS4wIjwN942Roln3OoSiWgcfMCKUBXJVwh_ECUaL2ylqGCfDZ7hNGpt0O0RvMKK3uLNQOYzoKhgh2bgDEU7T-FhUFdb9E9buCaucRTES6I0I51PxLIo7ePZYURLlQtxVzqZSJErV4KGBFrQPYCq9Mb2uXNgXNIdniPDEWSMSkqSJNN1LY0QVFg8-0BNX6ejqEBQpnB-10OjFecyyV2YNCthn3cqvcYt_4Bb7lVsbixuxgSZK59OAJmtplV9DVUPSGNPBswq5gD_6vIl1nl4yNFH9jsWLxLno0cCRI6cs-JiEU8U0v6gvTw-7D2SdYLqIMt1ODZDzW3wqixFfxQf8NmoauZqncWWiqYia2r455zddjL4XU_nyvXygowd0Ks-G_PG21_JVsowpQclRbihhnA3pE-g05ds5OQa2h9Uh2dlgOfMbIcI1K6cVcfBjSDkl2w67-jU5WXZwmk2JoLJhS1h7JQWjJ7cjc4ZvRi9GcGRsdCEUhQN8btccwQhqp1hwNuShcDuL4VlIP8oOGt7K88_mbXo5b9m0SLH8_bxl_IpywhZXB5PlHwZzonNC10C9QxRHqg5vV9KTcjnZXX4WKHYFUDz9AKjyCqD49azn6ZVAcfKmEfHsbSfi9LOI8MuI8CkhOXsfEU6vQYRfjwi7FpH8InWKsYtP2zgfqCu8aZW07VOsd-ce0OhBGKV-a5a6W3TCeiWaiOYhVDl2L1o6aTqIsxMRNWg_NxrmYpyLWqONN1rJKQCHd9iClnCu48nYx9qavjsosNCAcDDu8UI1WIrmOCSGgT4ZBtPEJRthrQIbdZWfMiTkz15ZeFNh85OKyosPCiwvLhZYRqdRK3_Jmv9HzH7FPeeNhdfIlh82NX6NwumocopSnl5EiY8kH3rYeWPKs3d7eiys02kuH-qK0NUok8cWykjJynxBymka5fTYdPv6tY2RcpxGyFtKnHZZekqR9yesLHI1PAs35KAYLVZvT04_e3BeGX12cBr-_7EDrPRwTAky8eDyPxxLSqwcDm7Old4Lq4SOL7DpPVY-rGnjsem8atV_oAqlsjV7iEfHxshHfBgFp34P3_z2a6fSOCgWO-87FwG5R-S-Vn7XbxJpWkTuA7CHy7yz5t8gPSL3W9WAQ-Q-CxWY5wtE7r9on5SzakmrkpZiBsssL1nGS8Ly2W5ZbATNSMFhsS2AFJLTrOIFpEVKNkxmbKaWJCU0LdMiXbAFyZNNuqAlFZCnZSoZE4il0ArVJJEextYz5VwPy5wvSDqLvHLx9w1CNDzhuBiKFl_N7DK-8qavXSiFynl31OKVb-IPI388fPsT8RX-4ySkB6oceTRGdgxmZ2EHogKLzTbSYdbbZvnpKER_Qxji-_w3AAD__4ZjE9o">