/* "Function scope" (top-level block scope) 'static' variables ... inside OpenACC compute construct regions as well as OpenACC 'routine'. This is to document/verify aspects of GCC's observed behavior, not necessarily as it's (intended to be?) restricted by the OpenACC specification. See also PR84991, PR84992, PR90779 etc., and "C/C++ 'static' variables" (only visible to members of the GitHub OpenACC organization). */ /* { dg-additional-options "-fopt-info-note-omp" } { dg-additional-options "--param=openacc-privatization=noisy" } { dg-additional-options "-foffload=-fopt-info-note-omp" } { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } for testing/documenting aspects of that functionality. */ /* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting aspects of that functionality. */ #undef NDEBUG #include #include #include #include #define IF_DEBUG if (0) /* Without explicit 'num_gangs'. */ static void t0_c(void) { IF_DEBUG __builtin_printf ("%s\n", __FUNCTION__); const int i_limit = 11; const int var_init = 16; for (int i = 0; i < i_limit; ++i) { int result = 0; int num_gangs_actual = -1; #pragma acc parallel \ reduction(max:num_gangs_actual) \ reduction(max:result) /* { dg-note {variable 'var' declared in block isn't candidate for adjusting OpenACC privatization level: static} "" { target *-*-* } .-3 } */ { num_gangs_actual = 1 + __builtin_goacc_parlevel_id(GOMP_DIM_GANG); static int var = var_init; #pragma acc atomic capture result = ++var; /* Irrespective of the order in which the gang-redundant threads execute, 'var' has now been incremented 'num_gangs_actual' times, and the final value captured as 'result'. */ } /* Without an explicit 'num_gangs' clause GCC assigns 'num_gangs(1)' because it doesn't see any use of gang-level parallelism inside the region. */ assert(num_gangs_actual == 1); assert(result == var_init + num_gangs_actual * (1 + i)); } } /* Call a gang-level routine. */ static const int t0_r_var_init = 61; #pragma acc routine gang /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'atomic'" { xfail *-*-* } .+4 } */ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .+3 } */ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .+2 } */ __attribute__((noinline)) static int t0_r_r(void) { static int var = t0_r_var_init; int tmp; #pragma acc atomic capture tmp = ++var; return tmp; } static void t0_r(void) { IF_DEBUG __builtin_printf ("%s\n", __FUNCTION__); const int i_limit = 11; for (int i = 0; i < i_limit; ++i) { int result = 0; int num_gangs_actual = -1; #pragma acc parallel \ reduction(max:num_gangs_actual) \ reduction(max:result) { num_gangs_actual = 1 + __builtin_goacc_parlevel_id(GOMP_DIM_GANG); result = t0_r_r(); /* Irrespective of the order in which the gang-redundant threads execute, 'var' has now been incremented 'num_gangs_actual' times, and the final value captured as 'result'. */ } /* The number of gangs selected by the implemention ought to but must not be bigger than one. */ IF_DEBUG __builtin_printf ("%d: num_gangs_actual: %d\n", i, num_gangs_actual); assert(num_gangs_actual >= 1); assert(result == t0_r_var_init + num_gangs_actual * (1 + i)); } } /* Explicit 'num_gangs'. */ static void t1_c(void) { IF_DEBUG __builtin_printf ("%s\n", __FUNCTION__); const int i_limit = 22; const int num_gangs_request = 444; const int var_init = 5; for (int i = 0; i < i_limit; ++i) { int result = 0; int num_gangs_actual = -1; /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'atomic'" { xfail *-*-* } .+1 } */ #pragma acc parallel \ num_gangs(num_gangs_request) \ reduction(max:num_gangs_actual) \ reduction(max:result) /* { dg-note {variable 'var' declared in block isn't candidate for adjusting OpenACC privatization level: static} "" { target *-*-* } .-4 } */ { num_gangs_actual = 1 + __builtin_goacc_parlevel_id(GOMP_DIM_GANG); static int var = var_init; #pragma acc atomic capture result = ++var; /* Irrespective of the order in which the gang-redundant threads execute, 'var' has now been incremented 'num_gangs_actual' times, and the final value captured as 'result'. */ } if (acc_get_device_type() == acc_device_host) assert(num_gangs_actual == 1); else assert(num_gangs_actual == num_gangs_request); assert(result == var_init + num_gangs_actual * (1 + i)); } } /* Check the same routine called from two compute constructs. */ static const int t1_r2_var_init = 166; #pragma acc routine gang /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'atomic'" { xfail *-*-* } .+4 } */ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .+3 } */ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .+2 } */ __attribute__((noinline)) static int t1_r2_r(void) { static int var = t1_r2_var_init; int tmp; #pragma acc atomic capture tmp = ++var; return tmp; } static void t1_r2(void) { IF_DEBUG __builtin_printf ("%s\n", __FUNCTION__); const int i_limit = 71; /* The checking assumes the same 'num_gangs' for all compute constructs. */ const int num_gangs_request = 333; int num_gangs_actual = -1; if (acc_get_device_type() == acc_device_host) num_gangs_actual = 1; else { /* We're assuming that the implementation is able to accomodate the 'num_gangs' requested (which really ought to be true for 'num_gangs'). */ num_gangs_actual = num_gangs_request; } for (int i = 0; i < i_limit; ++i) { int result_1 = 0; #pragma acc parallel \ num_gangs(num_gangs_request) \ reduction(max:result_1) { result_1 = t1_r2_r(); /* Irrespective of the order in which the gang-redundant threads execute, 'var' has now been incremented 'num_gangs_actual' times, and the final value captured as 'result_1'. */ } IF_DEBUG __builtin_printf ("%d: result_1: %d\n", i, result_1); assert(result_1 == t1_r2_var_init + num_gangs_actual * (1 + (i * 3 + 0))); int result_2 = 0; #pragma acc parallel \ num_gangs(num_gangs_request) \ reduction(max:result_2) { result_2 = t1_r2_r() + t1_r2_r(); /* Irrespective of the order in which the gang-redundant threads execute, 'var' has now been incremented '2 * num_gangs_actual' times. However, the order of the two 't1_r2_r' function calls is not synchronized (between different gang-redundant threads). We thus cannot verify the actual 'result_2' values in this case. */ } IF_DEBUG __builtin_printf ("%d: result_2: %d\n", i, result_2); if (num_gangs_actual == 1) /* Per the rationale above, only in this case we can check the actual result. */ assert(result_2 == (t1_r2_var_init + num_gangs_actual * (1 + (i * 3 + 1)) + t1_r2_var_init + num_gangs_actual * (1 + (i * 3 + 2)))); /* But we can generally check low and high limits. */ { /* Must be bigger than '2 * result_1'. */ int c = 2 * result_1; IF_DEBUG __builtin_printf (" > %d\n", c); assert(result_2 > c); } { /* ..., but limited by the base value for next 'i'. */ int c = 2 * (t1_r2_var_init + num_gangs_actual * (0 + ((i + 1) * 3 + 0))); IF_DEBUG __builtin_printf (" < %d\n", c); assert(result_2 < c); } } } /* Asynchronous execution. */ static const int t2_var_init_2 = -55; #pragma acc routine gang /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'atomic'" { xfail *-*-* } .+4 } */ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .+3 } */ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .+2 } */ __attribute__((noinline)) static int t2_r(void) { static int var = t2_var_init_2; int tmp; #pragma acc atomic capture tmp = ++var; return tmp; } static void t2(void) { IF_DEBUG __builtin_printf ("%s\n", __FUNCTION__); const int i_limit = 12; const int num_gangs_request_1 = 14; const int var_init_1 = 5; int results_1[i_limit][num_gangs_request_1]; memset (results_1, 0, sizeof results_1); const int num_gangs_request_2 = 5; int results_2[i_limit][num_gangs_request_2]; memset (results_2, 0, sizeof results_2); const int num_gangs_request_3 = 34; const int var_init_3 = 1250; int results_3[i_limit][num_gangs_request_3]; memset (results_3, 0, sizeof results_3); #pragma acc data \ copy(results_1, results_2, results_3) /* { dg-note {variable 'num_gangs_request_1\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target { c && { ! __OPTIMIZE__ } } } .-2 } */ /* { dg-note {variable 'num_gangs_request_2\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target { c && { ! __OPTIMIZE__ } } } .-3 } */ /* { dg-note {variable 'num_gangs_request_3\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target { c && { ! __OPTIMIZE__ } } } .-4 } */ /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 } */ { for (int i = 0; i < i_limit; ++i) { /* The following 'async' clauses effect asynchronous execution, but using the same async-argument for each compute construct implies that the respective compute constructs' execution is synchronized with itself, meaning that all 'i = 0' execution has finished (on the device) before 'i = 1' is started (on the device), etc. */ /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'atomic'" { xfail *-*-* } .+1 } */ #pragma acc parallel \ present(results_1) \ num_gangs(num_gangs_request_1) \ async(1) /* { dg-note {variable 'var' declared in block isn't candidate for adjusting OpenACC privatization level: static} "" { target *-*-* } .-4 } */ /* { dg-note {variable 'tmp' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 } */ { static int var = var_init_1; int tmp; #pragma acc atomic capture tmp = ++var; results_1[i][__builtin_goacc_parlevel_id(GOMP_DIM_GANG)] += tmp; } #pragma acc parallel \ present(results_2) \ num_gangs(num_gangs_request_2) \ async(2) { results_2[i][__builtin_goacc_parlevel_id(GOMP_DIM_GANG)] += t2_r(); } /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'atomic'" { xfail *-*-* } .+1 } */ #pragma acc parallel \ present(results_3) \ num_gangs(num_gangs_request_3) \ async(3) /* { dg-note {variable 'var' declared in block isn't candidate for adjusting OpenACC privatization level: static} "" { target *-*-* } .-4 } */ /* { dg-note {variable 'tmp' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 } */ { static int var = var_init_3; int tmp; #pragma acc atomic capture tmp = ++var; results_3[i][__builtin_goacc_parlevel_id(GOMP_DIM_GANG)] += tmp; } } #pragma acc wait } int num_gangs_actual_1; int num_gangs_actual_2; int num_gangs_actual_3; if (acc_get_device_type() == acc_device_host) { num_gangs_actual_1 = 1; num_gangs_actual_2 = 1; num_gangs_actual_3 = 1; } else { /* We're assuming that the implementation is able to accomodate the 'num_gangs' requested (which really ought to be true for 'num_gangs'). */ num_gangs_actual_1 = num_gangs_request_1; num_gangs_actual_2 = num_gangs_request_2; num_gangs_actual_3 = num_gangs_request_3; } /* For 'i = 0', 'results_*[i][0..num_gangs_actual_*]' are expected to each contain one value of '(1 + var_init_*)..(var_init_* + num_gangs_actual_*)', and so on for increasing 'i'. Their order however is unspecified due to the gang-redundant execution. (Thus checking that their sums match.) */ int result_1 = 0; int result_2 = 0; int result_3 = 0; for (int i = 0; i < i_limit; ++i) { int result_1_ = 0; for (int g = 0; g < num_gangs_actual_1; ++g) { IF_DEBUG __builtin_printf ("results_1[%d][%d]: %d\n", i, g, results_1[i][g]); result_1_ += results_1[i][g]; } IF_DEBUG __builtin_printf ("%d result_1_: %d\n", i, result_1_); assert (result_1_ == (((var_init_1 + num_gangs_actual_1 * (1 + i)) * (1 + var_init_1 + num_gangs_actual_1 * (1 + i)) / 2) - ((var_init_1 + num_gangs_actual_1 * (0 + i)) * (1 + var_init_1 + num_gangs_actual_1 * (0 + i)) / 2))); result_1 += result_1_; int result_2_ = 0; for (int g = 0; g < num_gangs_actual_2; ++g) { IF_DEBUG __builtin_printf ("results_2[%d][%d]: %d\n", i, g, results_2[i][g]); result_2_ += results_2[i][g]; } IF_DEBUG __builtin_printf ("%d result_2_: %d\n", i, result_2_); assert (result_2_ == (((t2_var_init_2 + num_gangs_actual_2 * (1 + i)) * (1 + t2_var_init_2 + num_gangs_actual_2 * (1 + i)) / 2) - ((t2_var_init_2 + num_gangs_actual_2 * (0 + i)) * (1 + t2_var_init_2 + num_gangs_actual_2 * (0 + i)) / 2))); result_2 += result_2_; int result_3_ = 0; for (int g = 0; g < num_gangs_actual_3; ++g) { IF_DEBUG __builtin_printf ("results_3[%d][%d]: %d\n", i, g, results_3[i][g]); result_3_ += results_3[i][g]; } IF_DEBUG __builtin_printf ("%d result_3_: %d\n", i, result_3_); assert (result_3_ == (((var_init_3 + num_gangs_actual_3 * (1 + i)) * (1 + var_init_3 + num_gangs_actual_3 * (1 + i)) / 2) - ((var_init_3 + num_gangs_actual_3 * (0 + i)) * (1 + var_init_3 + num_gangs_actual_3 * (0 + i)) / 2))); result_3 += result_3_; } IF_DEBUG __builtin_printf ("result_1: %d\n", result_1); assert (result_1 == (((var_init_1 + num_gangs_actual_1 * i_limit) * (1 + var_init_1 + num_gangs_actual_1 * i_limit) / 2) - (var_init_1 * (var_init_1 + 1) / 2))); IF_DEBUG __builtin_printf ("result_2: %d\n", result_2); assert (result_2 == (((t2_var_init_2 + num_gangs_actual_2 * i_limit) * (1 + t2_var_init_2 + num_gangs_actual_2 * i_limit) / 2) - (t2_var_init_2 * (t2_var_init_2 + 1) / 2))); IF_DEBUG __builtin_printf ("result_3: %d\n", result_3); assert (result_3 == (((var_init_3 + num_gangs_actual_3 * i_limit) * (1 + var_init_3 + num_gangs_actual_3 * i_limit) / 2) - (var_init_3 * (var_init_3 + 1) / 2))); } #pragma acc routine seq __attribute__((noinline)) static int pr84991_1_r_s(int n) { static const int test[] = {1,2,3,4}; return test[n]; } static void pr84991_1(void) { int n[1]; n[0] = 3; #pragma acc parallel copy(n) { n[0] = pr84991_1_r_s(n[0]); } assert(n[0] == 4); } static void pr84992_1(void) { int n[1]; n[0] = 3; #pragma acc parallel copy(n) /* { dg-note {variable 'test' declared in block isn't candidate for adjusting OpenACC privatization level: static} "" { target *-*-* } .-1 } */ { static const int test[] = {1,2,3,4}; n[0] = test[n[0]]; } assert(n[0] == 4); } int main(void) { t0_c(); t0_r(); t1_c(); t1_r2(); t2(); pr84991_1(); pr84992_1(); return 0; }