<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/62532>62532</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
D146051 breaks some reductions in AMD GPU offload if -O0
</td>
</tr>
<tr>
<th>Labels</th>
<td>
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
jdenny-ornl
</td>
</tr>
</table>
<pre>
I expect the following reproducer to print i=18:
```
#include <stdio.h>
int main() {
int i = 10;
#pragma omp target teams num_teams(1) reduction(+:i)
{
#pragma omp distribute parallel for num_threads(1) reduction(+: i)
for (int x = 0; x < 4; ++x)
i += 1;
#pragma omp distribute parallel for num_threads(1) reduction(+: i)
for (int x = 0; x < 4; ++x)
i += 1;
}
printf("i=%d\n", i);
return i != 18;
}
```
Working at 8b8466fd31e5 (a main commit from today) while offloading to a gfx906, I saw the wrong result:
```
$ clang -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa test.c
$ ./a.out
i=10
```
The test also reproduced the bug without num_threads(1) and num_teams(1), which I added just to see if parallelism mattered.
I found some workarounds that caused the test to print i=18 as expected:
- Add -O1 or higher to the clang command line.
- Offload to nvptx64 or x86_64 instead.
- Revert fb5683449e97 (D146051), a commit from March that I identified via git bisect. That commit just relocates DeadArgumentEliminationPass
I configured as follows, and check-all succeeded with or without the revert:
```
$ cmake \
-DCMAKE_BUILD_TYPE=Release \
-DLLVM_ENABLE_PROJECTS='clang;lld' \
-DLLVM_ENABLE_RUNTIMES=openmp \
-DLLVM_TARGETS_TO_BUILD='host;AMDGPU' \
-DCMAKE_C_COMPILER=gcc \
-DCMAKE_CXX_COMPILER=g++ \
../llvm
```
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzMVU9z4rgT_TTi0gVl5D_YBw6AyRS_X5ikMmR39kTJVtvWRLYoSQ7Mt9-SDIRkUrN73FQwkuh-3a_1us2MEXWHOCfxksT5iPW2UXr-g2PX_Rwr3clRofjP-QbwdMDSgm0QKiWlOoquBo0HrXhfogar4KBFZ0GQMJ-mJFyQICfB5ZkE5_9hS0PRlbLnCCRcGcuFmjQkXN-6OKyWiY7QlNAMyGw5nAP4KEDCHKYBCa_HhIYHzeqWgWoPYJmu0YJF1hro-nbvV4SmU4emkfelFWpAX5JwIQjNrkizJQBcdh-RuTBWi6K3CAemmZQooVJ6CNJoZPw3YeA2Dng_QlPH6OQZOUJ-uYLILZ0XXZ7enIQ_ctxvqP-3UvR_n-VJZvnFxGul8jGpUwyhMSfxqiOUEroaUnjjp9H2uvOYU4-ZXn-8Yn6UmH_-qfSLEyqzkBZplCQVD6cYO0LMqwtK1bbCQqVVC1Zx9tMV5dgIiaCqSirGnb9VwKCuTlmQuPQ2YNjR98JRK98Hppf2nzQfQSlZV8O4Ugfs2sN1MR7EakiYs5bXZTdmLXefxjCwaOykfMOYEHrHJqq35zZx7Rb8pgS7Bj0GMGnUW8dyn37R13AUtlG9_UwcrOO_9I7jf2xE2cAGGOfI4UdvrKuQQQRRXRUnTAstsxY18sltRhuoVN9xMKpFOCr9wrTbG7ANs1Cy3pyz82l_GCzAzHkWIf9Q8DEsOIfxwxSUhkbUzTCWHNJQeHfXjpIUHU4uPg_DLTvL7vVgT0nk3E9psk8iEJ2xyPjV-AlfUVuoijhJwyjKMJs5LeXTKAniS3XYO1FtmS6bgdoGBMfOikogh1fBoBYWCmGwtBOAnWc_ePqSapSqZBYN5Mj4Qtd9i51dS9GKjrmWfWTGvC9sqbpK1L1G7uo0DGrjc-o4lA2WL2MmJZi-LBHd3bnLd3wvInDF0p7kv1Bzy14QSLw6t-k4X20X_1_vl8-b-3y_--txTcL8CSUyc2sG4_z-_o_tfv11sbxf7x-fHv63Xu2--SEw8xdFwqWUnNDZO_Bbp6fnr7vNdu2czr30K_5u8fRlvfu23z0MGQ0BGmUsCZeLbf7l8fl9iCuB1X71sH3c3K-fSJjXZfmpzffv76yGKXhjOXGNKuVrCx8KOOLzkGdhxkY4nyZpOI0imk1HzTybpWGcZCzOaJoGRcRplCSsCnkY8AorNhJzGtAwiIOQBkEUB5OYpwFmWZQUUVJMiymJAmyZkBMXd6J0PRLG9DhPaBzSkWQFSnN51eu5MxoXfW1IFEhhrHlzs8JKnJ9lDYVG9mKGfr2-LwyIDhbbHL48Pl9GpWv_8UMw6rWcN9YejBMRvSP0rha26YtJqdpzVc5f44NWP7C0hN75VA2hdz7bvwMAAP__G4SH-w">