[Issue]: __hadd(int x, int y) and __rhadd(int x, int y) is wrong

[kwu130]

Problem Description

The calculation result of __hadd(int x, int y) and __rhadd(int x, int y) is inconsistent with NV.

For example, to calculate __hadd(123, -789)，the result on Rocm is 1, but the result on NV is -666.

The following are the descriptions of the __hadd function by Rocm and NV respectively.
Rocm

NV

The functions described are the same, but the calculation results are very different

Operating System

Centos7

CPU

AMD Ryzen 9 5900HX with Radeon Graphics

GPU

AMD Instinct MI100

ROCm Version

ROCm 5.7.0

ROCm Component

clr

Steps to Reproduce

kernel function as follows

__global__
void
kernel() {
    int x = 123;
    int y = -789;
    auto res = __hadd(x, y);
    printf("res %d\n", res);
}

(Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support

No response

Additional Information

Used cuda 11.8 to compare

[cjatin]

Hmm strange, thanks for reporting.
Will look into this.

[kwu130]

Hmm strange, thanks for reporting. Will look into this.
Please take a look at the following functions, looking forward to your reply. @cjatin
__hadd
__rhadd
__uhadd
__urhadd

[cjatin]

I found a few bugs. Raised a patch internally to fix it.