<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/61172>61172</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
Bug in a offload multi thread team distribute reduction test on AMD and Nvidia GPUs systems
</td>
</tr>
<tr>
<th>Labels</th>
<td>
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
fel-cab
</td>
</tr>
</table>
<pre>
The test fails on AMD and Nvidia GPUs systems, when executed with more than 1 thread (usually greater than 8).
I compile it with `clang -g -fopenmp --offload-arch=native -lomptarget`
The test is based on (https://github.com/SOLLVE/sollve_vv/blob/master/tests/4.5/target_teams_distribute/test_target_teams_distribute_reduction_add.c)
but creating many threads.
```
#include <omp.h>
#include <stdio.h>
#include <stdlib.h>
#define N 1024
int test_add() {
int a[N];
int b[N];
int total = 0;
int host_total = 0;
int errors = 0;
int num_teams[N];
for (int x = 0; x < N; ++x) {
a[x] = 1;
b[x] = x;
num_teams[x] = -1;
}
#pragma omp target teams distribute reduction(+:total) defaultmap(tofrom:scalar) num_teams(16)
for (int x = 0; x < N; ++x) {
num_teams[x] = omp_get_num_teams();
total += a[x] + b[x];
}
for (int x = 0; x < N; ++x) {
host_total += a[x] + b[x];
}
if(num_teams[0] == 1)
printf("Test operated with one team\n");
if(num_teams[0] <= 0)
printf("Test reported invalid number of teams.\n");
if(host_total != total) {
printf("Error in the reduction, host val = %d, target val = %d\n",host_total,total);
errors++;
}
return errors;
}
int main(int argc, char** argv)
{
if(argc != 2) {
printf("please enter the number of threads\n");
return 1;
}
int threads = atoi(argv[1]);
int total_errors = 0;
#pragma omp parallel num_threads(threads)
{
#pragma omp task shared(total_errors)
{
printf("task %d\n",omp_get_thread_num());
int errors = test_add();
#pragma omp atomic
total_errors += errors;
}
}
if(total_errors > 0)
printf("test FAIL\n");
else
printf("test PASSED\n");
return total_errors;
}
```
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJykVl1vqzgQ_TWTl1EiYvJBHvKQj2Z1pW73Sr27r5GBAbxrMLJNbvvvV8aEQJpWurpS1YDHnplz5sxgbozIK6ItLPewPE54YwultxnJacLjSazS9-2PgtCSsZhxIQ2qCnd_HpFXKb5cRCo4_vH9b4Pm3VgqDbAD_iyoQnqjpLGU4k9hCyyVJrQFr3COttDEUwQWNabhUr5jrolb0n5DBGwzQwiOEOy-YaLKWkhCYb0jWAWJ5FWO0xynmaqpKmucTlWWScXTKddJAeGx4lZcCKdSlbXlOicLq8C79P97TMJgzA2lDhawqLC2NhDugJ2AnXJhiyaeJaoEdnr96_n5nydgJ6OkvND5cgF2iqWKgZ1KbixpYCfn0wA7LWZL99aGPlvipTmnwlgt4sZSt-_8ifmsKW0SK1R15mk6S4BtfNJxYzFxVIkqx5JX7x2VZjaE5pD6P__KQlElskkJITyosp4VED49shmbCvWVVYp4aL5uSikTFeELzgO28Kuisi29DgAwV1GE9d7bEJ2Vw3L_AssjhKPl-PGyVZZLhPCIwZ2lUI7IT82ktdLmoalqSk_9fcjrpkxpJwm396330D4e8MU9AtsD27_docMW2xssj-2h-SAstvh609vINMyn3zIdHIf18Y75WvO85KjKGr2WsHWANy1hr6W2DHsIdy1ZLueUMt5IW_IaWGRVplUJ4c4kXHLt7LeEWDRf9Sr8DWIeQlRlfXZtMAzngg256Qrs8j8O6GX7ntAHNP1GokNV_WpQkQGLhkiDDqlXw5VGCDaItRaVzVrA7IcbR6omzfupqSpqKwrLQwWMjVn5LM7BA_06jqZaaRdHVBcuRepKE5NGlXkJzT4POeJm7oL1iuo5_BD0yfUhigptMdLkoaUaL137AlumbrFT82j5mtDhlgCwwzX2oHs32LW9r-4g_0GVNNlGV9eN_eFxiznhlFxUnYa4zhOXXVK4BtkB27mly43qAfyWKnfgShK7JwiCzZChWhI3hFT5DyENK-Kn_IOSeDcdlvlo-YZk005Q76Olk1slfHYXWO7nTshjn_3IPT-cn879aPjUXHMpSfoG77JlUf-0GaQ1yvx-hpn_0BRcU9pOpFsCIxcfvYyZbN2MJXOdMT4jN2q6IfOByw79APf4O_Zx-xgDt6oUyd2WMZd-oNwrrwN2fFzB7I4QhPBp3OT3JLguP-2-PT9WDUlDXx79vnt9fTo-PtzpbVSgTzqov4xM0m2YbsINn9B2vlqvo_WKRYtJsV3E0SrK0mTNs81qFWfreRivlutonqRss4zCidiygIVBGIQsWCwWixlLs1WYLJZJQKuUBwtYBFRyIWdSXsqZ0vlEGNPQdjWfr9lE8pikud5v9dZtmsZNbmARSGGsuR2zwkra7pvcDSqO3Z0Sy0Zacb21uun48Pvqr5NfX44njZbbL66YLpPuZ1pr9S8lFtipReMulS2g_wMAAP__RFhehg">