[bug][opt][AMDGPU]At the -O2 optimization level, the division computation result for uint64_t type is incorrect

[0oyyo0]

①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 / uint64_t{10}".
hipcc demo.cpp -o demo -std=c++17 -O2

③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

④Current analysis: In the amdgpu-codegenprepare pass, udiv is converted to fmul.
⑤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>

__device__ bool check_string(const char* d_str, size_t len) {
    auto iter = d_str;
    auto const iter_end = d_str + len;

    auto const bound_val = std::numeric_limits<uint64_t>::max();
    uint64_t value = 0;

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

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

      // auto const bound_check = (bound_val - sign * digit) / uint64_t{10} * sign;
      uint64_t const a1 = digit;
      uint64_t const a2 = bound_val - a1;
      uint64_t const bound_check = a2 / uint64_t{10};
      
      printf("bound_val: %llu, digit: %llu, value: %llu, bound_check: %llu\n", bound_val, digit, value, bound_check);

      if (value > bound_check) {
        return false;
      }

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

    return true;
  
}

__global__ void check_string_kernel(const char* str, size_t len, bool* result) {
    if (threadIdx.x == 0 && blockIdx.x == 0) {
        result[0] = check_string(str, len);
    }
}

int main() {
    std::string input = "9223372036854775806";
    
    char* d_str;
    bool* d_result;
    hipMalloc(&d_str, input.size());
    hipMalloc(&d_result, sizeof(bool));
    
    hipMemcpy(d_str, input.c_str(), input.size(), hipMemcpyHostToDevice);
    
    check_string_kernel<<<1, 256>>>(d_str, input.size(), d_result);
    hipDeviceSynchronize();
    
    bool h_result;
    hipMemcpy(&h_result, d_result, sizeof(bool), hipMemcpyDeviceToHost);
    
    std::cout << (h_result ? "check passed" : "Got 'false' when checking if \"9223372036854775806\" is a uint64_t. Expect 'true'.") << std::endl;
    
    hipFree(d_str);
    hipFree(d_result);
    
    return 0;
}

[llvmbot]

@llvm/issue-subscribers-backend-amdgpu

Author: Helium (0oyyo0)

①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" />

④Current analysis: In the amdgpu-codegenprepare pass, udiv is converted to fmul.
⑤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;
}

[frederik-h]

①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

It would be helpful to reduce the example as far as possible. There seems to be a lot of code than cannot be related to your problem (e.g. handling '-' although your input does not contain it; and why does your operator() not use its argument?). I suggest to strip it away. This will help others to figure out if this is a compiler issue.

[arsenm]

Was this the issue fixed by 8d2e611?

[0oyyo0]

①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

It would be helpful to reduce the example as far as possible. There seems to be a lot of code than cannot be related to your problem (e.g. handling '-' although your input does not contain it; and why does your operator() not use its argument?). I suggest to strip it away. This will help others to figure out if this is a compiler issue.

Redundant code has been deleted

[0oyyo0]

Was this the issue fixed by 8d2e611?

It doesn't look like the same issue. I tried to modify the code, but the problem still exists