this adds a check that the device number of teams and number of threads are less than or equal to the numbers specified in num_teams and thread_limit clauses
@@ -21,21 +21,40 @@ void checkHost(int gpu_error, int* errors, long long a){
void reduction(int num_teams, int num_threads, int* errors){
long long result = 0;
int gpu_error = 0;
- #pragma omp target teams num_teams(num_teams) thread_limit(num_threads) map(tofrom: result)
+ int device_teams = 0;
+ int device_threads = 0;
+ #pragma omp target teams num_teams(num_teams) thread_limit(num_threads) map(tofrom: result) map(from:device_teams,device_threads)
{
long long a, i;
a = 0;
#pragma omp parallel for reduction(+:a)
for (i = 0; i < N; i++) {
a += i;
+ device_threads = omp_get_num_threads();
}
result = a;
if (a != SUM && omp_get_team_num() <= 50){ //limit teams that print
printf ("GPU - Incorrect result = %lld, expected = %lld!\n", a, SUM);
gpu_error = 1;
}
+ device_teams = omp_get_num_teams();
} //end of target
+ // Spec says you cannot have more teams than num_teams clause
+ if ( device_teams > num_teams ) {
+ (*errors)++;
+ gpu_error++;
+ printf("ERROR: num_teams requested:%d actual teams on device:%d\n",
+ num_teams, device_teams);
+ }
+ // Spec says you cannot have more threads than thread_limit clause
+ if ( device_threads > num_threads ) {
+ (*errors)++;
+ gpu_error++;
+ printf("ERROR: num_threads limit:%d Actual threads on device:%d\n",
+ num_threads, device_threads);
+ }
+
checkHost(gpu_error, errors, result);
}
