From 89a1ec7fae77fdfec7b07f35360ab3fbf37e0be5 Mon Sep 17 00:00:00 2001 From: Ryan Padrone Date: Tue, 10 Mar 2026 17:49:15 -0700 Subject: [PATCH] Add validation test for single if clause on data-related constructs --- Tests/acc_single_if_clause_data.F90 | 258 ++++++++++++++++++++++++++++ Tests/acc_single_if_clause_data.c | 218 +++++++++++++++++++++++ Tests/acc_single_if_clause_data.cpp | 218 +++++++++++++++++++++++ 3 files changed, 694 insertions(+) create mode 100644 Tests/acc_single_if_clause_data.F90 create mode 100644 Tests/acc_single_if_clause_data.c create mode 100644 Tests/acc_single_if_clause_data.cpp diff --git a/Tests/acc_single_if_clause_data.F90 b/Tests/acc_single_if_clause_data.F90 new file mode 100644 index 0000000..cb9e481 --- /dev/null +++ b/Tests/acc_single_if_clause_data.F90 @@ -0,0 +1,258 @@ +! acc_single_if_clause_data.F90 +! +! Feature under test (OpenACC 3.4, Sections 2.6.5, 2.6.6, 2.8. March 2026): +! The specification clarifies that the data, enter data, exit data, +! and host_data constructs may contain at most one if clause. +! These tests verify correct behavior when a single valid if clause +! is used on these constructs in Fortran. +! +! Tests: +! T1 – data construct with single if clause: +! Executes a data region with copyin/copyout and a single +! if(dev) clause controlling execution. +! +! T2 – enter data / exit data with single if clause: +! Uses enter data and exit data directives with if(dev) +! to manage device memory and verify correct results. +! +! T3 – host_data construct with single if clause: +! Uses host_data use_device(...) with if(dev) to ensure +! device pointer access functions correctly. +! +! T4 – combined constructs with single if clauses: +! Uses both data and host_data constructs within the +! same region, each containing a single if clause. + +#ifndef T1 +!T1:syntax,data,if-clause,construct-independent,V:3.4- + LOGICAL FUNCTION test1() + USE OPENACC + IMPLICIT NONE + INCLUDE "acc_testsuite.Fh" + INTEGER :: x + INTEGER :: errors = 0 + INTEGER :: dev + REAL(8), DIMENSION(LOOPCOUNT) :: a, b, c + + dev = 1 + SEEDDIM(1) = 1 +# ifdef SEED + SEEDDIM(1) = SEED +# endif + CALL RANDOM_SEED(PUT=SEEDDIM) + CALL RANDOM_NUMBER(a) + CALL RANDOM_NUMBER(b) + c = 0.0 + + !$acc data copyin(a(1:LOOPCOUNT), b(1:LOOPCOUNT)) copyout(c(1:LOOPCOUNT)) if(dev .ne. 0) + !$acc parallel loop present(a(1:LOOPCOUNT), b(1:LOOPCOUNT), c(1:LOOPCOUNT)) + DO x = 1, LOOPCOUNT + c(x) = a(x) + b(x) + END DO + !$acc end parallel loop + !$acc end data + + DO x = 1, LOOPCOUNT + IF (abs(c(x) - (a(x) + b(x))) .gt. PRECISION) THEN + errors = errors + 1 + END IF + END DO + + IF (errors .eq. 0) THEN + test1 = .FALSE. + ELSE + test1 = .TRUE. + END IF + END FUNCTION +#endif + +#ifndef T2 +!T2:syntax,enter-data,exit-data,if-clause,construct-independent,V:3.4- + LOGICAL FUNCTION test2() + USE OPENACC + IMPLICIT NONE + INCLUDE "acc_testsuite.Fh" + INTEGER :: x + INTEGER :: errors = 0 + INTEGER :: dev + REAL(8), DIMENSION(LOOPCOUNT) :: a, b, c + + dev = 1 + SEEDDIM(1) = 1 +# ifdef SEED + SEEDDIM(1) = SEED +# endif + CALL RANDOM_SEED(PUT=SEEDDIM) + CALL RANDOM_NUMBER(a) + CALL RANDOM_NUMBER(b) + c = 0.0 + + !$acc enter data copyin(a(1:LOOPCOUNT), b(1:LOOPCOUNT)) create(c(1:LOOPCOUNT)) if(dev .ne. 0) + + !$acc parallel loop present(a(1:LOOPCOUNT), b(1:LOOPCOUNT), c(1:LOOPCOUNT)) + DO x = 1, LOOPCOUNT + c(x) = a(x) + b(x) + END DO + !$acc end parallel loop + + !$acc exit data copyout(c(1:LOOPCOUNT)) delete(a(1:LOOPCOUNT), b(1:LOOPCOUNT)) if(dev .ne. 0) + + DO x = 1, LOOPCOUNT + IF (abs(c(x) - (a(x) + b(x))) .gt. PRECISION) THEN + errors = errors + 1 + END IF + END DO + + IF (errors .eq. 0) THEN + test2 = .FALSE. + ELSE + test2 = .TRUE. + END IF + END FUNCTION +#endif + +#ifndef T3 +!T3:syntax,host-data,if-clause,construct-independent,V:3.4- + LOGICAL FUNCTION test3() + USE OPENACC + IMPLICIT NONE + INCLUDE "acc_testsuite.Fh" + INTEGER :: x + INTEGER :: errors = 0 + INTEGER :: dev + REAL(8), DIMENSION(LOOPCOUNT), TARGET :: a + REAL(8), POINTER :: seen_ptr(:) + + dev = 1 + NULLIFY(seen_ptr) + + SEEDDIM(1) = 1 +# ifdef SEED + SEEDDIM(1) = SEED +# endif + CALL RANDOM_SEED(PUT=SEEDDIM) + CALL RANDOM_NUMBER(a) + + !$acc data copyin(a(1:LOOPCOUNT)) + !$acc host_data use_device(a) if(dev .ne. 0) + seen_ptr => a + !$acc end host_data + !$acc end data + + IF (.not. ASSOCIATED(seen_ptr)) THEN + errors = errors + 1 + END IF + + IF (errors .eq. 0) THEN + test3 = .FALSE. + ELSE + test3 = .TRUE. + END IF + END FUNCTION +#endif + +#ifndef T4 +!T4:syntax,data,host-data,if-clause,construct-independent,V:3.4- + LOGICAL FUNCTION test4() + USE OPENACC + IMPLICIT NONE + INCLUDE "acc_testsuite.Fh" + INTEGER :: x + INTEGER :: errors = 0 + INTEGER :: host + REAL(8), DIMENSION(LOOPCOUNT) :: a, b + + host = 0 + SEEDDIM(1) = 1 +# ifdef SEED + SEEDDIM(1) = SEED +# endif + CALL RANDOM_SEED(PUT=SEEDDIM) + CALL RANDOM_NUMBER(a) + b = 0.0 + + !$acc data copyin(a(1:LOOPCOUNT)) copyout(b(1:LOOPCOUNT)) if(.TRUE.) + !$acc parallel loop present(a(1:LOOPCOUNT), b(1:LOOPCOUNT)) + DO x = 1, LOOPCOUNT + b(x) = a(x) * 2.0 + END DO + !$acc end parallel loop + + !$acc host_data use_device(b) if(host .ne. 0) + !$acc end host_data + !$acc end data + + DO x = 1, LOOPCOUNT + IF (abs(b(x) - (a(x) * 2.0)) .gt. PRECISION) THEN + errors = errors + 1 + END IF + END DO + + IF (errors .eq. 0) THEN + test4 = .FALSE. + ELSE + test4 = .TRUE. + END IF + END FUNCTION +#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 +#ifndef T4 + LOGICAL :: test4 +#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. + 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. + 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. + END IF +#endif +#ifndef T4 + DO testrun = 1, NUM_TEST_CALLS + failed = failed .or. test4() + END DO + IF (failed) THEN + failcode = failcode + 2 ** 3 + failed = .FALSE. + END IF +#endif + + CALL EXIT(failcode) + END PROGRAM diff --git a/Tests/acc_single_if_clause_data.c b/Tests/acc_single_if_clause_data.c new file mode 100644 index 0000000..6d94626 --- /dev/null +++ b/Tests/acc_single_if_clause_data.c @@ -0,0 +1,218 @@ +// acc_single_if_clause_data.c +// +// Feature under test (OpenACC 3.4, Sections 2.6.5, 2.6.6, and 2.8. March 2026): +// The specification clarifies that the data, enter data, exit data, +// and host_data constructs may contain at most one if clause. +// These tests verify correct behavior when a single valid if clause +// is used on these constructs. +// +// Tests: +// T1 – data construct with single if clause: +// Uses a data region with copyin/copyout and a single if(dev) +// clause controlling device execution. +// +// T2 – enter data / exit data with single if clause: +// Uses enter data and exit data directives with if(dev) to +// manage device memory and verifies correct data movement. +// +// T3 – host_data construct with single if clause: +// Uses host_data use_device(...) with if(dev) to ensure device +// pointer access works correctly when an if clause is present. +// +// T4 – combined constructs with single if clauses: +// Uses both data and host_data constructs within the same region, +// each containing a single if clause, verifying that multiple +// constructs can independently use valid if clauses. + +#include "acc_testsuite.h" +#include + +#ifndef T1 +int test1(){ + 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)); + int dev = 1; + + for (int x = 0; x < n; ++x){ + a[x] = rand() / (real_t)(RAND_MAX / 10); + b[x] = rand() / (real_t)(RAND_MAX / 10); + c[x] = 0.0; + } + + #pragma acc data copyin(a[0:n], b[0:n]) copyout(c[0:n]) if(dev) + { + #pragma acc parallel loop present(a[0:n], b[0:n], c[0:n]) + for (int x = 0; x < n; ++x){ + c[x] = a[x] + b[x]; + } + } + + for (int x = 0; x < n; ++x){ + if (fabs(c[x] - (a[x] + b[x])) > PRECISION){ + err += 1; + } + } + + free(a); + free(b); + free(c); + + return err; +} +#endif + +#ifndef T2 +int test2(){ + 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)); + int dev = 1; + + for (int x = 0; x < n; ++x){ + a[x] = rand() / (real_t)(RAND_MAX / 10); + b[x] = rand() / (real_t)(RAND_MAX / 10); + c[x] = 0.0; + } + + #pragma acc enter data copyin(a[0:n], b[0:n]) create(c[0:n]) if(dev) + + #pragma acc parallel loop present(a[0:n], b[0:n], c[0:n]) + for (int x = 0; x < n; ++x){ + c[x] = a[x] + b[x]; + } + + #pragma acc exit data copyout(c[0:n]) delete(a[0:n], b[0:n]) if(dev) + + for (int x = 0; x < n; ++x){ + if (fabs(c[x] - (a[x] + b[x])) > PRECISION){ + err += 1; + } + } + + free(a); + free(b); + free(c); + + return err; +} +#endif + +#ifndef T3 +int test3(){ + int err = 0; + srand(SEED); + + real_t *a = (real_t *)malloc(n * sizeof(real_t)); + real_t *seen_ptr = NULL; + int dev = 1; + + for (int x = 0; x < n; ++x){ + a[x] = rand() / (real_t)(RAND_MAX / 10); + } + + #pragma acc data copyin(a[0:n]) + { + #pragma acc host_data use_device(a) if(dev) + { + seen_ptr = a; + } + } + + if (seen_ptr == NULL){ + err += 1; + } + + free(a); + + return err; +} +#endif + +#ifndef T4 +int test4(){ + 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)); + int host = 0; + + for (int x = 0; x < n; ++x){ + a[x] = rand() / (real_t)(RAND_MAX / 10); + b[x] = 0.0; + } + + #pragma acc data copyin(a[0:n]) copyout(b[0:n]) if(1) + { + #pragma acc parallel loop present(a[0:n], b[0:n]) + for (int x = 0; x < n; ++x){ + b[x] = a[x] * 2; + } + + #pragma acc host_data use_device(b) if(host) + { + ; + } + } + + for (int x = 0; x < n; ++x){ + if (fabs(b[x] - (a[x] * 2)) > PRECISION){ + err += 1; + } + } + + free(a); + free(b); + + return err; +} +#endif + +int main(){ + int failcode = 0; + int failed; +#ifndef T1 + failed = 0; + for (int x = 0; x < NUM_TEST_CALLS; ++x){ + failed = failed + test1(); + } + if (failed != 0){ + failcode = failcode + (1 << 0); + } +#endif +#ifndef T2 + failed = 0; + for (int x = 0; x < NUM_TEST_CALLS; ++x){ + failed = failed + test2(); + } + if (failed != 0){ + failcode = failcode + (1 << 1); + } +#endif +#ifndef T3 + failed = 0; + for (int x = 0; x < NUM_TEST_CALLS; ++x){ + failed = failed + test3(); + } + if (failed != 0){ + failcode = failcode + (1 << 2); + } +#endif +#ifndef T4 + failed = 0; + for (int x = 0; x < NUM_TEST_CALLS; ++x){ + failed = failed + test4(); + } + if (failed != 0){ + failcode = failcode + (1 << 3); + } +#endif + return failcode; +} diff --git a/Tests/acc_single_if_clause_data.cpp b/Tests/acc_single_if_clause_data.cpp new file mode 100644 index 0000000..a118f64 --- /dev/null +++ b/Tests/acc_single_if_clause_data.cpp @@ -0,0 +1,218 @@ +// acc_single_if_clause_data.cpp +// +// Feature under test (OpenACC 3.4, Sections 2.6.5, 2.6.6, 2.8. March 2026): +// The specification clarifies that the data, enter data, exit data, +// and host_data constructs may contain at most one if clause. +// These tests verify correct behavior when a single valid if clause +// is used on these constructs in a C++ environment. +// +// Tests: +// T1 – data construct with single if clause: +// Uses a data region with copyin/copyout and a single if(dev) +// clause controlling device execution. +// +// T2 – enter data / exit data with single if clause: +// Uses enter data and exit data directives with if(dev) to +// manage device memory and verify correct computation. +// +// T3 – host_data construct with single if clause: +// Uses host_data use_device(...) with if(dev) to confirm +// that device pointer access functions correctly. +// +// T4 – combined constructs with single if clauses: +// Uses both data and host_data constructs within the same +// region to verify that each construct may independently +// contain a single if clause. + +#include "acc_testsuite.h" +#include + +#ifndef T1 +int test1(){ + int err = 0; + srand(SEED); + + real_t *a = new real_t[n]; + real_t *b = new real_t[n]; + real_t *c = new real_t[n]; + int dev = 1; + + for (int x = 0; x < n; ++x){ + a[x] = rand() / (real_t)(RAND_MAX / 10); + b[x] = rand() / (real_t)(RAND_MAX / 10); + c[x] = 0.0; + } + + #pragma acc data copyin(a[0:n], b[0:n]) copyout(c[0:n]) if(dev) + { + #pragma acc parallel loop present(a[0:n], b[0:n], c[0:n]) + for (int x = 0; x < n; ++x){ + c[x] = a[x] + b[x]; + } + } + + for (int x = 0; x < n; ++x){ + if (fabs(c[x] - (a[x] + b[x])) > PRECISION){ + err += 1; + } + } + + delete[] a; + delete[] b; + delete[] c; + + return err; +} +#endif + +#ifndef T2 +int test2(){ + int err = 0; + srand(SEED); + + real_t *a = new real_t[n]; + real_t *b = new real_t[n]; + real_t *c = new real_t[n]; + int dev = 1; + + for (int x = 0; x < n; ++x){ + a[x] = rand() / (real_t)(RAND_MAX / 10); + b[x] = rand() / (real_t)(RAND_MAX / 10); + c[x] = 0.0; + } + + #pragma acc enter data copyin(a[0:n], b[0:n]) create(c[0:n]) if(dev) + + #pragma acc parallel loop present(a[0:n], b[0:n], c[0:n]) + for (int x = 0; x < n; ++x){ + c[x] = a[x] + b[x]; + } + + #pragma acc exit data copyout(c[0:n]) delete(a[0:n], b[0:n]) if(dev) + + for (int x = 0; x < n; ++x){ + if (fabs(c[x] - (a[x] + b[x])) > PRECISION){ + err += 1; + } + } + + delete[] a; + delete[] b; + delete[] c; + + return err; +} +#endif + +#ifndef T3 +int test3(){ + int err = 0; + srand(SEED); + + real_t *a = new real_t[n]; + real_t *seen_ptr = NULL; + int dev = 1; + + for (int x = 0; x < n; ++x){ + a[x] = rand() / (real_t)(RAND_MAX / 10); + } + + #pragma acc data copyin(a[0:n]) + { + #pragma acc host_data use_device(a) if(dev) + { + seen_ptr = a; + } + } + + if (seen_ptr == NULL){ + err += 1; + } + + delete[] a; + + return err; +} +#endif + +#ifndef T4 +int test4(){ + int err = 0; + srand(SEED); + + real_t *a = new real_t[n]; + real_t *b = new real_t[n]; + int host = 0; + + for (int x = 0; x < n; ++x){ + a[x] = rand() / (real_t)(RAND_MAX / 10); + b[x] = 0.0; + } + + #pragma acc data copyin(a[0:n]) copyout(b[0:n]) if(1) + { + #pragma acc parallel loop present(a[0:n], b[0:n]) + for (int x = 0; x < n; ++x){ + b[x] = a[x] * 2; + } + + #pragma acc host_data use_device(b) if(host) + { + ; + } + } + + for (int x = 0; x < n; ++x){ + if (fabs(b[x] - (a[x] * 2)) > PRECISION){ + err += 1; + } + } + + delete[] a; + delete[] b; + + return err; +} +#endif + +int main(){ + int failcode = 0; + int failed; +#ifndef T1 + failed = 0; + for (int x = 0; x < NUM_TEST_CALLS; ++x){ + failed = failed + test1(); + } + if (failed != 0){ + failcode = failcode + (1 << 0); + } +#endif +#ifndef T2 + failed = 0; + for (int x = 0; x < NUM_TEST_CALLS; ++x){ + failed = failed + test2(); + } + if (failed != 0){ + failcode = failcode + (1 << 1); + } +#endif +#ifndef T3 + failed = 0; + for (int x = 0; x < NUM_TEST_CALLS; ++x){ + failed = failed + test3(); + } + if (failed != 0){ + failcode = failcode + (1 << 2); + } +#endif +#ifndef T4 + failed = 0; + for (int x = 0; x < NUM_TEST_CALLS; ++x){ + failed = failed + test4(); + } + if (failed != 0){ + failcode = failcode + (1 << 3); + } +#endif + return failcode; +}