Fix correctness of test unveiled by new device rtl behavior: number of threads participating in a parallel region can be <= thread_limit clause value
#define TEAMS 2
int main(){
- int gpu_results[THREADS];
- int correct_results[THREADS] = {2,2};
- #pragma omp target teams thread_limit(THREADS) num_teams(TEAMS) map(from:gpu_results)
+ int gpu_results[TEAMS];
+ int correct_results[TEAMS];
+ int actual_num_threads = -1;
+
+ // the runtime is allowed to use <=THREADS in the parallel regions and
+ // it actually chooses 1 (new runtime) or 2 (old runtime)
+ #pragma omp target teams thread_limit(THREADS) num_teams(TEAMS) \
+ map(from:gpu_results, actual_num_threads)
{
int dist[THREADS];
// Uncomment line below to trigger generic kernel before fix was in place
@@ -25,15 +30,20 @@ int main(){
#pragma omp barrier
if(thread == 0) {
+ if (team == 0)
+ actual_num_threads = omp_get_num_threads();
for(int i = 1; i < THREADS; i++)
dist[0] += dist[i];
gpu_results[team] = dist[0];
}
}
}
- int status = memcmp(correct_results, gpu_results, THREADS * sizeof(int));
- if (status != 0){
+ for(int i = 0; i < TEAMS; i++)
+ correct_results[i] = actual_num_threads;
+ int status = memcmp(correct_results, gpu_results, TEAMS * sizeof(int));
+
+ if (status != 0 || actual_num_threads > 2) {
printf("FAIL\n");
return 1;
}
