From 95feccdbf1243fb409bc35688a588dbb12d4a098 Mon Sep 17 00:00:00 2001 From: Ryan Padrone Date: Sat, 14 Feb 2026 14:45:26 -0800 Subject: [PATCH 1/3] =?UTF-8?q?Add=20V&V=20tests=20for=20evaluation=20of?= =?UTF-8?q?=20expressions=20in=20directive=20arguments=20(OpenACC=203.4=20?= =?UTF-8?q?=C2=A72.1)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- Tests/acc_expr_eval_directive_args.F90 | 280 +++++++++++++++++++++++++ Tests/acc_expr_eval_directive_args.c | 236 +++++++++++++++++++++ Tests/acc_expr_eval_directive_args.cpp | 234 +++++++++++++++++++++ 3 files changed, 750 insertions(+) create mode 100644 Tests/acc_expr_eval_directive_args.F90 create mode 100644 Tests/acc_expr_eval_directive_args.c create mode 100644 Tests/acc_expr_eval_directive_args.cpp 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..227417e --- /dev/null +++ b/Tests/acc_expr_eval_directive_args.c @@ -0,0 +1,236 @@ +// 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 + +/* Pure helpers: deterministic, no side effects */ +static int foo_pure(int x){ + return (x % 64) + 1; +} + +static int bar_pure(int x){ + return (x % 32) + 1; +} + +/* Side-effecting helper used only to exercise "may be elided"; we never assert call count */ +static volatile int size_calls = 0; + +static int size_maybe_elided(int nval){ + size_calls = size_calls + 1; + return nval; +} + +#ifndef T1 +//T1:syntax,expressions,runtime,construct-independent,V:3.4- +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 +//T2:syntax,expressions,runtime,construct-independent,V:3.4- +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; + } + + /* Ensure 'a' is NOT present on device: we do not enter/create any data for it. */ + size_calls = 0; + #pragma acc update device(a[0:size_maybe_elided(n)]) if_present + + /* User-visible behavior: no crash; host values unchanged. */ + for (int i = 0; i < n; ++i){ + if (fabs(a[i] - (real_t)i) > PRECISION){ + err = err + 1; + } + } + + /* DO NOT assert anything about size_calls (may be 0, 1, or more). */ + + free(a); + return err; +} +#endif + +#ifndef T3 +//T3:syntax,expressions,runtime,construct-independent,V:3.4- +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; + } + + /* Side-effecting expressions in directive arguments. + Per Section 2.1, evaluation order/number is unspecified. + We do NOT use i afterward (do not rely on side effects). */ + 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; + } + } + + /* DO NOT assert anything about i. */ + + 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..e8b2ba9 --- /dev/null +++ b/Tests/acc_expr_eval_directive_args.cpp @@ -0,0 +1,234 @@ +// 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 +//T1:syntax,expressions,runtime,construct-independent,V:3.4- +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 +//T2:syntax,expressions,runtime,construct-independent,V:3.4- +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; + } + + // Ensure 'a' is NOT present on device: do NOT enter/create any data. + size_calls = 0; + #pragma acc update device(a[0:size_maybe_elided(n)]) if_present + + // User-visible behavior: no crash; host values unchanged. + for (int i = 0; i < n; ++i){ + if (std::fabs(a[i] - (real_t)i) > PRECISION){ + err = err + 1; + } + } + + // DO NOT assert anything about size_calls (may be 0, 1, or more). + + std::free(a); + return err; +} +#endif + +#ifndef T3 +//T3:syntax,expressions,runtime,construct-independent,V:3.4- +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; + } + + // Side-effecting expressions in directive arguments. + // Per Section 2.1, evaluation order/number is unspecified. + // We do NOT use i afterward (do not rely on side effects). + 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; + } + } + + // DO NOT assert anything about i. + + 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; +} From 7662256b6e368f0878aa4a4c70cc333dc8f1a2c0 Mon Sep 17 00:00:00 2001 From: Ryanpadrone <159075564+Ryanpadrone@users.noreply.github.com> Date: Wed, 18 Feb 2026 19:57:55 -0500 Subject: [PATCH 2/3] Update acc_expr_eval_directive_args.c --- Tests/acc_expr_eval_directive_args.c | 14 +------------- 1 file changed, 1 insertion(+), 13 deletions(-) diff --git a/Tests/acc_expr_eval_directive_args.c b/Tests/acc_expr_eval_directive_args.c index 227417e..594a6af 100644 --- a/Tests/acc_expr_eval_directive_args.c +++ b/Tests/acc_expr_eval_directive_args.c @@ -1,4 +1,5 @@ // 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 @@ -16,7 +17,6 @@ #include #include -/* Pure helpers: deterministic, no side effects */ static int foo_pure(int x){ return (x % 64) + 1; } @@ -25,7 +25,6 @@ static int bar_pure(int x){ return (x % 32) + 1; } -/* Side-effecting helper used only to exercise "may be elided"; we never assert call count */ static volatile int size_calls = 0; static int size_maybe_elided(int nval){ @@ -34,7 +33,6 @@ static int size_maybe_elided(int nval){ } #ifndef T1 -//T1:syntax,expressions,runtime,construct-independent,V:3.4- int test1(){ int err = 0; int cond = 0; @@ -100,7 +98,6 @@ int test1(){ #endif #ifndef T2 -//T2:syntax,expressions,runtime,construct-independent,V:3.4- int test2(){ int err = 0; @@ -113,26 +110,21 @@ int test2(){ a[i] = (real_t)i; } - /* Ensure 'a' is NOT present on device: we do not enter/create any data for it. */ size_calls = 0; #pragma acc update device(a[0:size_maybe_elided(n)]) if_present - /* User-visible behavior: no crash; host values unchanged. */ for (int i = 0; i < n; ++i){ if (fabs(a[i] - (real_t)i) > PRECISION){ err = err + 1; } } - /* DO NOT assert anything about size_calls (may be 0, 1, or more). */ - free(a); return err; } #endif #ifndef T3 -//T3:syntax,expressions,runtime,construct-independent,V:3.4- int test3(){ int err = 0; @@ -161,9 +153,6 @@ int test3(){ c[k] = 0; } - /* Side-effecting expressions in directive arguments. - Per Section 2.1, evaluation order/number is unspecified. - We do NOT use i afterward (do not rely on side effects). */ int i = 0; #pragma acc data copyin(a[0:n], b[0:n]) copy(c[0:n]) @@ -188,7 +177,6 @@ int test3(){ } } - /* DO NOT assert anything about i. */ free(a); free(b); From 11e59813e462a3698acd986fee2d050dc622acf5 Mon Sep 17 00:00:00 2001 From: Ryanpadrone <159075564+Ryanpadrone@users.noreply.github.com> Date: Wed, 18 Feb 2026 19:59:33 -0500 Subject: [PATCH 3/3] Update acc_expr_eval_directive_args.cpp --- Tests/acc_expr_eval_directive_args.cpp | 11 +---------- 1 file changed, 1 insertion(+), 10 deletions(-) diff --git a/Tests/acc_expr_eval_directive_args.cpp b/Tests/acc_expr_eval_directive_args.cpp index e8b2ba9..f9eb0ee 100644 --- a/Tests/acc_expr_eval_directive_args.cpp +++ b/Tests/acc_expr_eval_directive_args.cpp @@ -1,4 +1,5 @@ // 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 @@ -32,7 +33,6 @@ static int size_maybe_elided(int nval){ } #ifndef T1 -//T1:syntax,expressions,runtime,construct-independent,V:3.4- int test1(){ int err = 0; int cond = 0; @@ -98,7 +98,6 @@ int test1(){ #endif #ifndef T2 -//T2:syntax,expressions,runtime,construct-independent,V:3.4- int test2(){ int err = 0; @@ -111,18 +110,15 @@ int test2(){ a[i] = (real_t)i; } - // Ensure 'a' is NOT present on device: do NOT enter/create any data. size_calls = 0; #pragma acc update device(a[0:size_maybe_elided(n)]) if_present - // User-visible behavior: no crash; host values unchanged. for (int i = 0; i < n; ++i){ if (std::fabs(a[i] - (real_t)i) > PRECISION){ err = err + 1; } } - // DO NOT assert anything about size_calls (may be 0, 1, or more). std::free(a); return err; @@ -130,7 +126,6 @@ int test2(){ #endif #ifndef T3 -//T3:syntax,expressions,runtime,construct-independent,V:3.4- int test3(){ int err = 0; @@ -159,9 +154,6 @@ int test3(){ c[k] = 0; } - // Side-effecting expressions in directive arguments. - // Per Section 2.1, evaluation order/number is unspecified. - // We do NOT use i afterward (do not rely on side effects). int i = 0; #pragma acc data copyin(a[0:n], b[0:n]) copy(c[0:n]) @@ -186,7 +178,6 @@ int test3(){ } } - // DO NOT assert anything about i. std::free(a); std::free(b);