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

    <tr>
        <th>Summary</th>
        <td>
            [bug][opt][AMDGPU]At the -O2 optimization level, the division computation result for uint64_t type is incorrect
        </td>
    </tr>

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

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

    <tr>
      <th>Reporter</th>
      <td>
          0oyyo0
      </td>
    </tr>
</table>

<pre>
    I tested the following code on gfx906 and found that the results were not as expected.
The core of this code is to check whether the input string can be safely converted to a uint64_t integer.
When compiled using the following command, the calculation result is incorrect. The cause of the error is located in the incorrect calculation result of the line "a2 / IntegerType{10}".
`hipcc demo.cpp -o demo -std=c++17 -O2`

<img width="1151" height="305" alt="Image" src="https://github.com/user-attachments/assets/dc1689e2-9826-455d-adf1-dc851cb5d6b4" />


When DivExpand is disabled by adding the "disable-idiv-expansion" option, the calculation result is correct.
`hipcc demo.cpp -o demo -std=c++17 -O2 -mllvm -amdgpu-codegenprepare-disable-idiv-expansion=true`

<img width="1348" height="561" alt="Image" src="https://github.com/user-attachments/assets/8d64022c-1e13-44ca-a13f-761fadc79d31" />


Please advise: What is the specific role of the expandDivRem32 function in the amdgpu-codegenprepare pass? Why does using it lead to incorrect computation results?


```
#include <hip/hip_runtime.h>
#include <iostream>
#include <limits>
#include <string>
#include <type_traits>

template <typename IntegerType>
struct string_to_integer_check {
  const char* d_str;
  std::size_t n;

 __device__ bool operator()(int)
  {
    if (d_str[0] == '-' && std::is_unsigned_v<IntegerType>) { return false; }

    auto iter = d_str + static_cast<int>((d_str[0] == '-' || d_str[0] == '+'));
    auto const iter_end = d_str + n;
    if (iter == iter_end) { return false; }

 // printf("hhh\n");

    auto const sign = d_str[0] == '-' ? static_cast<IntegerType>(-1) : IntegerType{1};
    auto const bound_val =
 sign > 0 ? std::numeric_limits<IntegerType>::max() : std::numeric_limits<IntegerType>::min();

    IntegerType value = 0;
 while (iter != iter_end) {
      auto const chr = *iter++;
      if (chr < '0' || chr > '9') { return false; }

      auto const digit       = static_cast<IntegerType>(chr - '0');

      // auto const bound_check = (bound_val - sign * digit) / IntegerType{10} * sign;
      IntegerType const a1 = sign * digit;
      IntegerType const a2 = bound_val - a1;
 IntegerType const a3 = a2 / IntegerType{10};

      IntegerType const bound_check = a3 * sign;
      // printf("value: %llu, bound_check: %llu, digit: %llu,bound_val: %llu, sign: %llu\n", value,bound_check,digit, bound_val, sign);

      if (value > bound_check) {
        return false;
      }

      value = value * IntegerType{10} + digit;
    }

 return true;
  }
};

template<typename IntegerType>
__global__ void check_string_kernel(const char* d_str, std::size_t str_len, bool* result)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
 if (idx == 0) {
        string_to_integer_check<IntegerType> checker{d_str, str_len};
        result[0] = checker(0);
    }
}

int main()
{
 std::string int_str("9223372036854775806");
    
    char* d_str;
 bool* d_result;
    hipMalloc(&d_str, int_str.size());
 hipMalloc(&d_result, sizeof(bool));
    
    hipMemcpy(d_str, int_str.c_str(), int_str.size(), hipMemcpyHostToDevice);
    
    dim3 blockSize(256);
    dim3 gridSize(1);

 check_string_kernel<uint64_t><<<gridSize, blockSize>>>(d_str, int_str.size(), d_result);
    
    hipDeviceSynchronize();
    
    bool h_result;
 hipMemcpy(&h_result, d_result, sizeof(bool), hipMemcpyDeviceToHost);
    
 if (!h_result) {
        std::cout << "Got 'false' when check if \"9223372036854775806\" is a uint64_t. Expect 'true'." << std::endl;
 } else {
        std::cout << "check passed" << std::endl;
    }
    
 hipFree(d_str);
    hipFree(d_result);
    
    return 0;
}
```
</pre>
<img width="1" height="1" alt="" src="http://email.email.llvm.org/o/eJy0WMtu4zjWfhpmcyBDoizZXmThS9J_LRrd-LsGvRRo8tjiFCUKJOUk9fQDkpLlW1LTiwECRCbP5Tt3ksxaeWwRn0mxIcXuifWu1uY51R8fOn3aa_Hx_A0cWocCXI1w0ErpN9kegWuBoFs4Ht5XaQmsFXDQfevJmAu0Bm2vnIU3NAitdsAs4HuH3KGYkXT9vUbg2iDoA7ha2ihSWnAaeI38B7zV6Go0QZpsu96BdSYoZy3sESw7oPoArtsTmgBRA4Netq6cVw5k6_CIxuv6u8YWuG46qVBAb72QW3uahrWC0G3Y4EzxXjEndTsY4pHJlmtjkLsZBPSstwN8BDRGG0-kNGcejGwH4APPI5kDr5ItAqGUUSD0Fb5F4N8_OiSLTZaSxY5Q6u0gZVrLjnMQ2OgZ7zpIdPiGxDpB8h0ndEPoJltA8gclZep50jXJt7I5wpsUria5F5ZlRUYohRrlsXZxLU8Lv8TU8Ptbw47oV6zhcaV2rrMkXxP6SujrUbq638-4bgh97S2ahDnHeN1g6yyhr8xaDB-CZ-VyhTRZLWmZzItCJEwcskTwZZHxfSHK_dzr8WLzlwHyELSdPL28dz69pAUhLdv7CO4_gAkxRpFQOuwkUshTgp7BSt16obpz4euruI5R_ccehqRR6tRAwhpx7PrEp_AR285gxwwmn6DKd870-FV08vnyNjpFmf1vorMU5TyllCcZZnkyn3OWsCw_JIsyOzDBFyuRZw-i86dCZhGYOEmLJF_D377wffnWCLZDLg-Sg9FqqpAQx508_T82OYVD3_IQh6FQHroQOmYtyV_h7_oDhEY7FK90oJCFir8oMN10vbsMrmc9I_YOj3_pmtBctlz1AoHk21p2hL7WsqtM3zrZ4KweTL0ik9o6g6x5tKdkI726-53Ysx7tuI8OK2fYxJiuHTadYu6837IGrxpCILTO9Hxsh5XT1dDrqtg4yWJD0jX4zmgd8JoZQtcgKusMyeNWSOY1yddW_sTKQRs3_F5VCTxJjlUFe60V6A4Nc9oQuiR0RehSts5_eDGjJgB5AEKXUUexSUmxA5-W-Q4IXSSELoDQktBy0ixt1bdh_ojqRPLtjZV05aWDQdebFg5M-TTbgG-GEScAsN4ngEPjdUUDgdANWJ8FvOLMOh-31gWBy18BXGzJYgufkIS6XwQPrEYvDgiinz2OCltxg6WdiKOPRryebOT5tbWxqqEzsnWHYAut65oUW9_lzpDuUXkHT4g-szx_vfHZbTCWSRYw5uvb8eQhPnLH3h8HqhNTQVm6HpG8QDroG_Kg7Rs0kldjCd3qDkQNe4_5FyD8Q17ZDrl76aMLSjgx1WPwUjrY8lZLhVO0aPYgWoOgK6t5HXOR0LWnjsNi8s-YA5Fs6wOQXqReXH7xy6uYbL-ugSv1Qh6li6sBxi-C6vUlI4hbB8GYc3dBHbpMMHM5xTkZQux7jccR4D8-zwQiT33pmsuIRHUsi1Zci_2ahQaWS1QsG1geUOeB-ouD17VH7iXcesRLfGDbg_INSeezmdBCqd4fUS6EXW8Mhl8snQ28Jox6zytje9jGFD_zRRV0O8RpO_nrLOQuG2LmjqXycgX2piDgNmUnP1yn7lR4wxe96y9Dumxugz9JGnSFQ9WwPWxexG8crF_O1ao6Kr1nqqrgpKWI95BqGLM_0LSofNU8GKveazcj1TpTKWyjd7XytPFUEofnxehsHUjxHvNWaf7jm3ifvQdfhJ872YSfG3C1QSbC9oUj0vU4WaKQ2MfuQvLJaeGuL0Sr0ZDF5sK2aMtlr49hDgZNU-XMTJfp1aCcIhLC4W1u2Lkxn90xOTFe9WTrIgZfMStK83xB07xcFvPFolim5eXwG50B8PDMMwZBVAPqM1Mtu9-ZUpoHNeXZ6kH5zAd0PP2Muu54xtj68vmJ-hBao9e4egTQs2PDu4_zmeRCHz-bvPoMxnaS8H_auu96F85sj1QJ2eQxkf6KAmhRXtEFgqORYtjPrmr_UQ3k2_GWHYbsNv5NIrYX-jzBSxw3Xzp2OwXmsb-ihX99tLw2up04b0nDobW-DvKluwkt64tofR25Cz9H_d-19_ed3liAhGaT6Pv6G1Kb695BdJm_u_6mnR_BsVXSBbyFx4owULzUYvtJ5ocNf-Oa3jxm8BJeWLy80A7pYhbublHZGQG2Qg34fXdFZfG_BBtx-VsZil9IPhf96KNadq8GcUqE1VUJnvc-z4Kh0w-HtKGZjNe6J_Gci1W-Yk_4nC2KeZqv5qvVU_28QI50ifNDmmO5F4uUlliWyzndLw-M0cOTfKYpLdJlmqWrLMtXszTnHOeCl6LIylLsyTzFhkk187f9mTbHJ2ltj89ZkZWL8kmxPSobHtEobfENwq5vTcXuyTx7pmTfHy2Zp0paZycxTjoVXt_2_ZEUO1JsdOfix_r33W9__osUu3V8T0v-iI8ZjfwZr7gKT6jGhw0hT9L61ftLMBy0mV7F_PC7est66o16_uLtwGMd_iWd0f9G7gh9DRZaQl8HF5ye6X8CAAD__5YCGVQ">