diff --git a/Tests/acc_expr_eval_directive_args.F90 b/Tests/acc_expr_eval_directive_args.F90 new file mode 100644 index 0000000..f686fda --- /dev/null +++ b/Tests/acc_expr_eval_directive_args.F90 @@ -0,0 +1,280 @@ +! acc_expr_eval_directive_args.F90 +! +! Feature under test (OpenACC 3.4, Section 2.1, Feb 2026): +! - Clarified user-visible behavior of evaluation of expressions in directive arguments. +! A program must not depend on the order/number of evaluations of expressions in +! clause/construct/directive arguments, nor on any side effects of those evaluations. +! +! Notes: +! - T1: uses complex but side-effect-free expressions in directive arguments; checks correctness. +! - T2: uses update ... if_present where data is not present; must be a no-op and not crash. +! The section length expression may be elided; we do NOT assert call count. +! - T3: uses side-effecting function calls in directive arguments; we do NOT rely on side effects; +! checks correctness only. + +MODULE acc_expr_eval_helpers + IMPLICIT NONE + INTEGER, SAVE :: size_calls = 0 + INTEGER, SAVE :: sidefx_calls = 0 +CONTAINS + + INTEGER FUNCTION foo_pure(x) + IMPLICIT NONE + INTEGER, INTENT(IN) :: x + foo_pure = MOD(x, 64) + 1 + END FUNCTION foo_pure + + INTEGER FUNCTION bar_pure(x) + IMPLICIT NONE + INTEGER, INTENT(IN) :: x + bar_pure = MOD(x, 32) + 1 + END FUNCTION bar_pure + + INTEGER FUNCTION size_maybe_elided(nval) + IMPLICIT NONE + INTEGER, INTENT(IN) :: nval + size_calls = size_calls + 1 + size_maybe_elided = nval + END FUNCTION size_maybe_elided + + INTEGER FUNCTION sidefx_inc(x) + IMPLICIT NONE + INTEGER, INTENT(IN) :: x + sidefx_calls = sidefx_calls + 1 + sidefx_inc = x + END FUNCTION sidefx_inc + +END MODULE acc_expr_eval_helpers + + +#ifndef T1 +!T1:syntax,expressions,runtime,construct-independent,V:3.4- +LOGICAL FUNCTION test1() + USE OPENACC + USE acc_expr_eval_helpers + IMPLICIT NONE + INCLUDE "acc_testsuite.Fh" + + INTEGER :: i + INTEGER :: cond + INTEGER :: errors + REAL(8), DIMENSION(LOOPCOUNT) :: a, b, c + + errors = 0 + + SEEDDIM(1) = 1 +# ifdef SEED + SEEDDIM(1) = SEED +# endif + CALL RANDOM_SEED(PUT=SEEDDIM) + + CALL RANDOM_NUMBER(a) + CALL RANDOM_NUMBER(b) + + DO i = 1, LOOPCOUNT + c(i) = 0.0D0 + END DO + + cond = 0 + IF (LOOPCOUNT .GT. 0) THEN + cond = 1 + ELSE + cond = 0 + END IF + + !$acc data copyin(a(1:LOOPCOUNT), b(1:LOOPCOUNT)) copy(c(1:LOOPCOUNT)) + !$acc parallel & + !$acc& if((cond .EQ. 1) .AND. ((LOOPCOUNT/2) .GT. 0)) & + !$acc& num_gangs(foo_pure(LOOPCOUNT + 7)) & + !$acc& async(MOD(LOOPCOUNT, 3) + 1) + + !$acc loop + DO i = 1, LOOPCOUNT + c(i) = a(i) + b(i) + END DO + + !$acc end parallel + !$acc wait + !$acc end data + + DO i = 1, LOOPCOUNT + IF (ABS(c(i) - (a(i) + b(i))) .GT. PRECISION) THEN + errors = errors + 1 + ELSE + errors = errors + 0 + END IF + END DO + + IF (errors .EQ. 0) THEN + test1 = .FALSE. + ELSE + test1 = .TRUE. + END IF +END FUNCTION test1 +#endif + + +#ifndef T2 +!T2:syntax,expressions,runtime,construct-independent,V:3.4- +LOGICAL FUNCTION test2() + USE OPENACC + USE acc_expr_eval_helpers + IMPLICIT NONE + INCLUDE "acc_testsuite.Fh" + + INTEGER :: i + INTEGER :: errors + REAL(8), DIMENSION(LOOPCOUNT) :: a + + errors = 0 + + DO i = 1, LOOPCOUNT + a(i) = DBLE(i) + END DO + + ! Ensure 'a' is NOT present on device: do NOT enter/create any data. + size_calls = 0 + !$acc update device(a(1:size_maybe_elided(LOOPCOUNT))) if_present + + ! User-visible behavior: no crash; host values unchanged. + DO i = 1, LOOPCOUNT + IF (ABS(a(i) - DBLE(i)) .GT. PRECISION) THEN + errors = errors + 1 + ELSE + errors = errors + 0 + END IF + END DO + + ! DO NOT assert anything about size_calls (may be 0, 1, or more). + + IF (errors .EQ. 0) THEN + test2 = .FALSE. + ELSE + test2 = .TRUE. + END IF +END FUNCTION test2 +#endif + + +#ifndef T3 +!T3:syntax,expressions,runtime,construct-independent,V:3.4- +LOGICAL FUNCTION test3() + USE OPENACC + USE acc_expr_eval_helpers + IMPLICIT NONE + INCLUDE "acc_testsuite.Fh" + + INTEGER :: i + INTEGER :: k + INTEGER :: tmp + INTEGER :: errors + REAL(8), DIMENSION(LOOPCOUNT) :: a, b, c + + errors = 0 + + SEEDDIM(1) = 1 +# ifdef SEED + SEEDDIM(1) = SEED +# endif + CALL RANDOM_SEED(PUT=SEEDDIM) + + CALL RANDOM_NUMBER(a) + CALL RANDOM_NUMBER(b) + + DO i = 1, LOOPCOUNT + c(i) = 0.0D0 + END DO + + sidefx_calls = 0 + tmp = 7 + + !$acc data copyin(a(1:LOOPCOUNT), b(1:LOOPCOUNT)) copy(c(1:LOOPCOUNT)) + !$acc parallel & + !$acc& num_gangs(foo_pure(sidefx_inc(tmp))) & + !$acc& async(MOD(sidefx_inc(tmp + 2), 3) + 1) + + !$acc loop + DO k = 1, LOOPCOUNT + c(k) = a(k) + b(k) + END DO + + !$acc end parallel + !$acc wait + !$acc end data + + DO i = 1, LOOPCOUNT + IF (ABS(c(i) - (a(i) + b(i))) .GT. PRECISION) THEN + errors = errors + 1 + ELSE + errors = errors + 0 + END IF + END DO + + ! DO NOT assert anything about sidefx_calls. + + IF (errors .EQ. 0) THEN + test3 = .FALSE. + ELSE + test3 = .TRUE. + END IF +END FUNCTION test3 +#endif + + +PROGRAM main + IMPLICIT NONE + INTEGER :: failcode, testrun + LOGICAL :: failed + INCLUDE "acc_testsuite.Fh" + +#ifndef T1 + LOGICAL :: test1 +#endif +#ifndef T2 + LOGICAL :: test2 +#endif +#ifndef T3 + LOGICAL :: test3 +#endif + + failcode = 0 + failed = .FALSE. + +#ifndef T1 + DO testrun = 1, NUM_TEST_CALLS + failed = failed .OR. test1() + END DO + IF (failed) THEN + failcode = failcode + 2 ** 0 + failed = .FALSE. + ELSE + failed = .FALSE. + END IF +#endif + +#ifndef T2 + DO testrun = 1, NUM_TEST_CALLS + failed = failed .OR. test2() + END DO + IF (failed) THEN + failcode = failcode + 2 ** 1 + failed = .FALSE. + ELSE + failed = .FALSE. + END IF +#endif + +#ifndef T3 + DO testrun = 1, NUM_TEST_CALLS + failed = failed .OR. test3() + END DO + IF (failed) THEN + failcode = failcode + 2 ** 2 + failed = .FALSE. + ELSE + failed = .FALSE. + END IF +#endif + + CALL EXIT(failcode) +END PROGRAM main diff --git a/Tests/acc_expr_eval_directive_args.c b/Tests/acc_expr_eval_directive_args.c new file mode 100644 index 0000000..594a6af --- /dev/null +++ b/Tests/acc_expr_eval_directive_args.c @@ -0,0 +1,224 @@ +// acc_expr_eval_directive_args.c +// +// Feature under test (OpenACC 3.4, Section 2.1, Feb 2026): +// - Clarified user-visible behavior of evaluation of expressions in directive arguments. +// A program must not depend on the order/number of evaluations of expressions in +// clause/construct/directive arguments, nor on any side effects of those evaluations. +// +// Notes: +// - T1: uses complex but side-effect-free expressions in directive arguments; checks correctness. +// - T2: uses update ... if_present where data is not present; must be a no-op and not crash. +// The section length expression may be elided; we do NOT assert call count. +// - T3: uses side-effecting expressions in directive arguments; we do NOT rely on the side effects; +// checks correctness only. + +#include "acc_testsuite.h" +#include +#include +#include + +static int foo_pure(int x){ + return (x % 64) + 1; +} + +static int bar_pure(int x){ + return (x % 32) + 1; +} + +static volatile int size_calls = 0; + +static int size_maybe_elided(int nval){ + size_calls = size_calls + 1; + return nval; +} + +#ifndef T1 +int test1(){ + int err = 0; + int cond = 0; + + srand(SEED); + + real_t *a = (real_t *)malloc(n * sizeof(real_t)); + real_t *b = (real_t *)malloc(n * sizeof(real_t)); + real_t *c = (real_t *)malloc(n * sizeof(real_t)); + + if (a == NULL || b == NULL || c == NULL){ + if (a != NULL){ + free(a); + } + if (b != NULL){ + free(b); + } + if (c != NULL){ + free(c); + } + return 1; + } + + for (int i = 0; i < n; ++i){ + a[i] = rand() / (real_t)(RAND_MAX / 10); + b[i] = rand() / (real_t)(RAND_MAX / 10); + c[i] = 0; + } + + if (n > 0){ + cond = 1; + } + else{ + cond = 0; + } + + #pragma acc data copyin(a[0:n], b[0:n]) copy(c[0:n]) + { + #pragma acc parallel \ + if((cond == 1) && ((n / 2) > 0)) \ + num_gangs(foo_pure(n + 7)) \ + vector_length(((n % 128) + 1)) + { + #pragma acc loop + for (int i = 0; i < n; ++i){ + c[i] = a[i] + b[i]; + } + } + } + + for (int i = 0; i < n; ++i){ + if (fabs(c[i] - (a[i] + b[i])) > PRECISION){ + err = err + 1; + } + } + + free(a); + free(b); + free(c); + + return err; +} +#endif + +#ifndef T2 +int test2(){ + int err = 0; + + real_t *a = (real_t *)malloc(n * sizeof(real_t)); + if (a == NULL){ + return 1; + } + + for (int i = 0; i < n; ++i){ + a[i] = (real_t)i; + } + + size_calls = 0; + #pragma acc update device(a[0:size_maybe_elided(n)]) if_present + + for (int i = 0; i < n; ++i){ + if (fabs(a[i] - (real_t)i) > PRECISION){ + err = err + 1; + } + } + + free(a); + return err; +} +#endif + +#ifndef T3 +int test3(){ + int err = 0; + + srand(SEED); + + real_t *a = (real_t *)malloc(n * sizeof(real_t)); + real_t *b = (real_t *)malloc(n * sizeof(real_t)); + real_t *c = (real_t *)malloc(n * sizeof(real_t)); + + if (a == NULL || b == NULL || c == NULL){ + if (a != NULL){ + free(a); + } + if (b != NULL){ + free(b); + } + if (c != NULL){ + free(c); + } + return 1; + } + + for (int k = 0; k < n; ++k){ + a[k] = rand() / (real_t)(RAND_MAX / 10); + b[k] = rand() / (real_t)(RAND_MAX / 10); + c[k] = 0; + } + + int i = 0; + + #pragma acc data copyin(a[0:n], b[0:n]) copy(c[0:n]) + { + #pragma acc parallel \ + num_gangs(foo_pure(++i)) \ + num_workers(bar_pure(i++)) \ + async(((i + 1) % 3) + 1) + { + #pragma acc loop + for (int k = 0; k < n; ++k){ + c[k] = a[k] + b[k]; + } + } + + #pragma acc wait + } + + for (int k = 0; k < n; ++k){ + if (fabs(c[k] - (a[k] + b[k])) > PRECISION){ + err = err + 1; + } + } + + + free(a); + free(b); + free(c); + + return err; +} +#endif + +int main(){ + int failcode = 0; + int failed = 0; + +#ifndef T1 + failed = 0; + for (int t = 0; t < NUM_TEST_CALLS; ++t){ + failed = failed + test1(); + } + if (failed != 0){ + failcode = failcode + (1 << 0); + } +#endif + +#ifndef T2 + failed = 0; + for (int t = 0; t < NUM_TEST_CALLS; ++t){ + failed = failed + test2(); + } + if (failed != 0){ + failcode = failcode + (1 << 1); + } +#endif + +#ifndef T3 + failed = 0; + for (int t = 0; t < NUM_TEST_CALLS; ++t){ + failed = failed + test3(); + } + if (failed != 0){ + failcode = failcode + (1 << 2); + } +#endif + + return failcode; +} diff --git a/Tests/acc_expr_eval_directive_args.cpp b/Tests/acc_expr_eval_directive_args.cpp new file mode 100644 index 0000000..f9eb0ee --- /dev/null +++ b/Tests/acc_expr_eval_directive_args.cpp @@ -0,0 +1,225 @@ +// acc_expr_eval_directive_args.cpp +// +// Feature under test (OpenACC 3.4, Section 2.1, Feb 2026): +// - Clarified user-visible behavior of evaluation of expressions in directive arguments. +// A program must not depend on the order/number of evaluations of expressions in +// clause/construct/directive arguments, nor on any side effects of those evaluations. +// +// Notes: +// - T1: uses complex but side-effect-free expressions in directive arguments; checks correctness. +// - T2: uses update ... if_present where data is not present; must be a no-op and not crash. +// The section length expression may be elided; we do NOT assert call count. +// - T3: uses side-effecting expressions in directive arguments; we do NOT rely on side effects; +// checks correctness only. + +#include "acc_testsuite.h" +#include +#include +#include + +static int foo_pure(int x){ + return (x % 64) + 1; +} + +static int bar_pure(int x){ + return (x % 32) + 1; +} + +static volatile int size_calls = 0; + +static int size_maybe_elided(int nval){ + size_calls = size_calls + 1; + return nval; +} + +#ifndef T1 +int test1(){ + int err = 0; + int cond = 0; + + std::srand(SEED); + + real_t *a = (real_t *)std::malloc(n * sizeof(real_t)); + real_t *b = (real_t *)std::malloc(n * sizeof(real_t)); + real_t *c = (real_t *)std::malloc(n * sizeof(real_t)); + + if (a == NULL || b == NULL || c == NULL){ + if (a != NULL){ + std::free(a); + } + if (b != NULL){ + std::free(b); + } + if (c != NULL){ + std::free(c); + } + return 1; + } + + for (int i = 0; i < n; ++i){ + a[i] = std::rand() / (real_t)(RAND_MAX / 10); + b[i] = std::rand() / (real_t)(RAND_MAX / 10); + c[i] = 0; + } + + if (n > 0){ + cond = 1; + } + else{ + cond = 0; + } + + #pragma acc data copyin(a[0:n], b[0:n]) copy(c[0:n]) + { + #pragma acc parallel \ + if((cond == 1) && ((n / 2) > 0)) \ + num_gangs(foo_pure(n + 7)) \ + vector_length(((n % 128) + 1)) + { + #pragma acc loop + for (int i = 0; i < n; ++i){ + c[i] = a[i] + b[i]; + } + } + } + + for (int i = 0; i < n; ++i){ + if (std::fabs(c[i] - (a[i] + b[i])) > PRECISION){ + err = err + 1; + } + } + + std::free(a); + std::free(b); + std::free(c); + + return err; +} +#endif + +#ifndef T2 +int test2(){ + int err = 0; + + real_t *a = (real_t *)std::malloc(n * sizeof(real_t)); + if (a == NULL){ + return 1; + } + + for (int i = 0; i < n; ++i){ + a[i] = (real_t)i; + } + + size_calls = 0; + #pragma acc update device(a[0:size_maybe_elided(n)]) if_present + + for (int i = 0; i < n; ++i){ + if (std::fabs(a[i] - (real_t)i) > PRECISION){ + err = err + 1; + } + } + + + std::free(a); + return err; +} +#endif + +#ifndef T3 +int test3(){ + int err = 0; + + std::srand(SEED); + + real_t *a = (real_t *)std::malloc(n * sizeof(real_t)); + real_t *b = (real_t *)std::malloc(n * sizeof(real_t)); + real_t *c = (real_t *)std::malloc(n * sizeof(real_t)); + + if (a == NULL || b == NULL || c == NULL){ + if (a != NULL){ + std::free(a); + } + if (b != NULL){ + std::free(b); + } + if (c != NULL){ + std::free(c); + } + return 1; + } + + for (int k = 0; k < n; ++k){ + a[k] = std::rand() / (real_t)(RAND_MAX / 10); + b[k] = std::rand() / (real_t)(RAND_MAX / 10); + c[k] = 0; + } + + int i = 0; + + #pragma acc data copyin(a[0:n], b[0:n]) copy(c[0:n]) + { + #pragma acc parallel \ + num_gangs(foo_pure(++i)) \ + num_workers(bar_pure(i++)) \ + async(((i + 1) % 3) + 1) + { + #pragma acc loop + for (int k = 0; k < n; ++k){ + c[k] = a[k] + b[k]; + } + } + + #pragma acc wait + } + + for (int k = 0; k < n; ++k){ + if (std::fabs(c[k] - (a[k] + b[k])) > PRECISION){ + err = err + 1; + } + } + + + std::free(a); + std::free(b); + std::free(c); + + return err; +} +#endif + +int main(){ + int failcode = 0; + int failed = 0; + +#ifndef T1 + failed = 0; + for (int t = 0; t < NUM_TEST_CALLS; ++t){ + failed = failed + test1(); + } + if (failed != 0){ + failcode = failcode + (1 << 0); + } +#endif + +#ifndef T2 + failed = 0; + for (int t = 0; t < NUM_TEST_CALLS; ++t){ + failed = failed + test2(); + } + if (failed != 0){ + failcode = failcode + (1 << 1); + } +#endif + +#ifndef T3 + failed = 0; + for (int t = 0; t < NUM_TEST_CALLS; ++t){ + failed = failed + test3(); + } + if (failed != 0){ + failcode = failcode + (1 << 2); + } +#endif + + return failcode; +}