Min/max footgun in __clang_hip_math.h

Clang Frontend
1

There is a discrepancy between the __host__ and __device__ implementation of min and max in __clang_hip_math.h, which can result in potential footgun situations.

While the __device__ overload set contains function templates that would match non-int integral types before the int overload, the __host__ overload set does not provide these templates. This can result in erroneous narrowing conversions to int in host code, when called with a larger integral type. E.g. consider the following example:

#include <cstdint>
#include <iostream>

#include <hip/hip_runtime.h>

int main()
{
    std::size_t i = std::size_t{1} << 35;
    std::size_t j = std::size_t{15};
    auto result = min(i, j);
    std::cout << result << std::endl; // prints 0 instead of 15
}

The fact that the device implementation doesn’t suffer from this problem just makes related bugs to remain undetected even more so.

  • Does someone has information why this API was designed this way? Is there some requirement unknown to me that can only be fulfilled like this?
1 Like