From b76cfbbdf939969a4360e5b3e93a747900d9de8d Mon Sep 17 00:00:00 2001 From: rjenaa Date: Thu, 8 Aug 2024 09:56:02 -0700 Subject: [PATCH 01/11] Initial test for assumed array size --- .../target/test_target_assumed_array_size.c | 49 +++++++++++++++++++ 1 file changed, 49 insertions(+) create mode 100644 tests/6.0/target/test_target_assumed_array_size.c diff --git a/tests/6.0/target/test_target_assumed_array_size.c b/tests/6.0/target/test_target_assumed_array_size.c new file mode 100644 index 000000000..08955202d --- /dev/null +++ b/tests/6.0/target/test_target_assumed_array_size.c @@ -0,0 +1,49 @@ +//===--- test_target_assumed_array_size.c ------------------------------------===// +// +// OpenMP API Version 6.0 +// +// +////===----------------------------------------------------------------------===// + + +#include +#include +#include "ompvv.h" + +#define N 1024 + +#pragma omp begin declare target + +void update_array(int* array){ + + #pragma omp target parallel map(tofrom: array[:N]) + for (int i = 0; i < N; i++){ + array[i] = i * 2; + } + +} + +#pragma omp end declare target + +int test_assumed_array_size(){ + int errors = 0; + int *array = (int *)malloc(N * sizeof(int)); + + for (int i = 0; i < N; i++){ + array[i] = i; + } + + update_array(array); + for(int i = 0; i < N; i++){ + OMPVV_TEST_AND_SET(errors, array[i] != i * 2); + } + free(array); + return errors; +} + +int main(){ + OMPVV_TEST_OFFLOADING; + int errors = 0; + OMPVV_TEST_AND_SET_VERBOSE(errors, test_assumed_array_size() != 0); + OMPVV_REPORT_AND_RETURN(errors); +} From 54fd2434cb4b4a8521691602effaac99a35ae01d Mon Sep 17 00:00:00 2001 From: rjenaa Date: Thu, 15 Aug 2024 08:24:08 -0700 Subject: [PATCH 02/11] Modified the array sections --- .../target/test_target_assumed_array_size.c | 35 ++++++++----------- 1 file changed, 15 insertions(+), 20 deletions(-) diff --git a/tests/6.0/target/test_target_assumed_array_size.c b/tests/6.0/target/test_target_assumed_array_size.c index 08955202d..72ca8e0dc 100644 --- a/tests/6.0/target/test_target_assumed_array_size.c +++ b/tests/6.0/target/test_target_assumed_array_size.c @@ -1,7 +1,7 @@ //===--- test_target_assumed_array_size.c ------------------------------------===// // // OpenMP API Version 6.0 -// +// Tests the target directive with the map clause on an assumed-size array. // ////===----------------------------------------------------------------------===// @@ -12,32 +12,27 @@ #define N 1024 -#pragma omp begin declare target - -void update_array(int* array){ - - #pragma omp target parallel map(tofrom: array[:N]) - for (int i = 0; i < N; i++){ - array[i] = i * 2; - } - -} - -#pragma omp end declare target - int test_assumed_array_size(){ int errors = 0; - int *array = (int *)malloc(N * sizeof(int)); + int array[N]; + int *pointer; for (int i = 0; i < N; i++){ - array[i] = i; + array[i] = 0; } - update_array(array); - for(int i = 0; i < N; i++){ - OMPVV_TEST_AND_SET(errors, array[i] != i * 2); + pointer = &array[0]; + #pragma omp target data map(tofrom: array[0:N]) + #pragma omp target parallel for map(pointer[:]) + for (int i = 0; i < N; i++){ + *pointer[i] = i; } - free(array); + + + for (int i = 0; i < N; i++){ + OMPVV_TEST_AND_SET_VERBOSE(errors, array[i] != i); + } + return errors; } From 91ec29d82ffcefb3b8ceae2f5c57e324fecf5301 Mon Sep 17 00:00:00 2001 From: rjenaa Date: Thu, 15 Aug 2024 10:28:42 -0700 Subject: [PATCH 03/11] edited for loop to access array --- tests/6.0/target/test_target_assumed_array_size.c | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/6.0/target/test_target_assumed_array_size.c b/tests/6.0/target/test_target_assumed_array_size.c index 72ca8e0dc..dd950c201 100644 --- a/tests/6.0/target/test_target_assumed_array_size.c +++ b/tests/6.0/target/test_target_assumed_array_size.c @@ -25,7 +25,7 @@ int test_assumed_array_size(){ #pragma omp target data map(tofrom: array[0:N]) #pragma omp target parallel for map(pointer[:]) for (int i = 0; i < N; i++){ - *pointer[i] = i; + pointer[i] = i; } From f00176a98f0b2b6a463afe03c352749d701851ac Mon Sep 17 00:00:00 2001 From: rjenaa Date: Thu, 31 Oct 2024 10:02:53 -0700 Subject: [PATCH 04/11] Initial test for groupprivate --- .../target/test_target_teams_groupprivate.c | 51 +++++++++++++++++++ 1 file changed, 51 insertions(+) create mode 100644 tests/6.0/target/test_target_teams_groupprivate.c diff --git a/tests/6.0/target/test_target_teams_groupprivate.c b/tests/6.0/target/test_target_teams_groupprivate.c new file mode 100644 index 000000000..77426fcad --- /dev/null +++ b/tests/6.0/target/test_target_teams_groupprivate.c @@ -0,0 +1,51 @@ +//===--- test_target_groupprivate.c -----------------------------------------===// +// +// OpenMP API Version 6.0 +// Tests the target directive with the groupprivate directive ensuring proper behavior +// +////===----------------------------------------------------------------------===// + + +#include +#include +#include "ompvv.h" + +#define N 10 +#define NUM_TEAMS 4 + +int group_sum; + +int test_target_groupprivate(){ + int errors = 0; + int host_sum = 0; + int team_sum[NUM_TEAMS]; + for (int i = 0; i < NUM_TEAMS; i++){ + team_sum[i] = 0; + } + for (int i = 0; i < N; i++){ + host_sum += i; + } + + #pragma omp target teams num_teams(NUM_TEAMS) map(tofrom: team_sum[:NUM_TEAMS]) groupprivate(group_sum) + { + group_sum = 0; + for (int i = 0; i < N; i++){ + group_sum += i; + } + + team_sum[omp_get_team_num()] = group_sum; + } + + for (int i = 0; i < NUM_TEAMS; i++){ + OMPVV_TEST_AND_SET_VERBOSE(errors, team_sum[i] != host_sum); + } + return errors; +} + +int main(){ + OMPVV_TEST_OFFLOADING; + int errors = 0; + OMPVV_TEST_AND_SET_VERBOSE(errors, test_target_groupprivate() != 0); + OMPVV_REPORT_AND_RETURN(errors); +} + From b7b433d695825480d05a8d3fb8a5f9f5a88d96e8 Mon Sep 17 00:00:00 2001 From: rjenaa Date: Thu, 31 Oct 2024 10:05:10 -0700 Subject: [PATCH 05/11] removed unwanted file --- .../target/test_target_assumed_array_size.c | 44 ------------------- 1 file changed, 44 deletions(-) delete mode 100644 tests/6.0/target/test_target_assumed_array_size.c diff --git a/tests/6.0/target/test_target_assumed_array_size.c b/tests/6.0/target/test_target_assumed_array_size.c deleted file mode 100644 index dd950c201..000000000 --- a/tests/6.0/target/test_target_assumed_array_size.c +++ /dev/null @@ -1,44 +0,0 @@ -//===--- test_target_assumed_array_size.c ------------------------------------===// -// -// OpenMP API Version 6.0 -// Tests the target directive with the map clause on an assumed-size array. -// -////===----------------------------------------------------------------------===// - - -#include -#include -#include "ompvv.h" - -#define N 1024 - -int test_assumed_array_size(){ - int errors = 0; - int array[N]; - int *pointer; - - for (int i = 0; i < N; i++){ - array[i] = 0; - } - - pointer = &array[0]; - #pragma omp target data map(tofrom: array[0:N]) - #pragma omp target parallel for map(pointer[:]) - for (int i = 0; i < N; i++){ - pointer[i] = i; - } - - - for (int i = 0; i < N; i++){ - OMPVV_TEST_AND_SET_VERBOSE(errors, array[i] != i); - } - - return errors; -} - -int main(){ - OMPVV_TEST_OFFLOADING; - int errors = 0; - OMPVV_TEST_AND_SET_VERBOSE(errors, test_assumed_array_size() != 0); - OMPVV_REPORT_AND_RETURN(errors); -} From 664750706173accd2343e8b1d2f4a153f75a5273 Mon Sep 17 00:00:00 2001 From: rjenaa Date: Thu, 31 Oct 2024 10:52:00 -0700 Subject: [PATCH 06/11] remove loop --- .../target/test_target_teams_groupprivate.c | 23 ++++++++----------- 1 file changed, 9 insertions(+), 14 deletions(-) diff --git a/tests/6.0/target/test_target_teams_groupprivate.c b/tests/6.0/target/test_target_teams_groupprivate.c index 77426fcad..b4c0be853 100644 --- a/tests/6.0/target/test_target_teams_groupprivate.c +++ b/tests/6.0/target/test_target_teams_groupprivate.c @@ -15,30 +15,25 @@ int group_sum; + int test_target_groupprivate(){ int errors = 0; int host_sum = 0; - int team_sum[NUM_TEAMS]; - for (int i = 0; i < NUM_TEAMS; i++){ - team_sum[i] = 0; - } - for (int i = 0; i < N; i++){ - host_sum += i; - } + int team_sum = 0; - #pragma omp target teams num_teams(NUM_TEAMS) map(tofrom: team_sum[:NUM_TEAMS]) groupprivate(group_sum) + #pragma omp groupprivate(group_sum) device_type(nohost) + #pragma omp target teams num_teams(NUM_TEAMS) map(tofrom: team_sum) reduction(+: team_sum) { - group_sum = 0; - for (int i = 0; i < N; i++){ - group_sum += i; - } + group_sum = omp_get_team_num(); - team_sum[omp_get_team_num()] = group_sum; + team_sum += group_sum; + } for (int i = 0; i < NUM_TEAMS; i++){ - OMPVV_TEST_AND_SET_VERBOSE(errors, team_sum[i] != host_sum); + host_sum += i; } + OMPVV_TEST_AND_SET_VERBOSE(errors, team_sum != host_sum); return errors; } From fd9be89e0e9a6f9032b31c004693aeae0677159e Mon Sep 17 00:00:00 2001 From: rjenaa Date: Thu, 31 Oct 2024 10:53:37 -0700 Subject: [PATCH 07/11] remove unnecessary N --- tests/6.0/target/test_target_teams_groupprivate.c | 1 - 1 file changed, 1 deletion(-) diff --git a/tests/6.0/target/test_target_teams_groupprivate.c b/tests/6.0/target/test_target_teams_groupprivate.c index b4c0be853..0eab46050 100644 --- a/tests/6.0/target/test_target_teams_groupprivate.c +++ b/tests/6.0/target/test_target_teams_groupprivate.c @@ -10,7 +10,6 @@ #include #include "ompvv.h" -#define N 10 #define NUM_TEAMS 4 int group_sum; From e42eaeffb64b128bb827d668844863c8b49a31e1 Mon Sep 17 00:00:00 2001 From: rjenaa Date: Thu, 7 Nov 2024 05:06:52 -0800 Subject: [PATCH 08/11] Moved groupprivate pragma --- tests/6.0/target/test_target_teams_groupprivate.c | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/tests/6.0/target/test_target_teams_groupprivate.c b/tests/6.0/target/test_target_teams_groupprivate.c index 0eab46050..06fbc00a9 100644 --- a/tests/6.0/target/test_target_teams_groupprivate.c +++ b/tests/6.0/target/test_target_teams_groupprivate.c @@ -1,7 +1,8 @@ //===--- test_target_groupprivate.c -----------------------------------------===// // // OpenMP API Version 6.0 -// Tests the target directive with the groupprivate directive ensuring proper behavior +// Tests the target directive with the groupprivate directive ensuring proper +// behavior. // ////===----------------------------------------------------------------------===// @@ -13,14 +14,13 @@ #define NUM_TEAMS 4 int group_sum; - +#pragma omp groupprivate(group_sum) int test_target_groupprivate(){ int errors = 0; int host_sum = 0; int team_sum = 0; - #pragma omp groupprivate(group_sum) device_type(nohost) #pragma omp target teams num_teams(NUM_TEAMS) map(tofrom: team_sum) reduction(+: team_sum) { group_sum = omp_get_team_num(); From d26ac12ff5ec6521c74b2afb910cddcfc86b5f96 Mon Sep 17 00:00:00 2001 From: rjenaa Date: Thu, 14 Nov 2024 10:12:11 -0800 Subject: [PATCH 09/11] Added any and host device type --- .../test_teams_groupprivate_devicetype_host.c | 44 ++++++++++++++++++ ...target_teams_groupprivate_devicetype_any.c | 45 +++++++++++++++++++ 2 files changed, 89 insertions(+) create mode 100644 tests/6.0/groupprivate/test_teams_groupprivate_devicetype_host.c create mode 100644 tests/6.0/target/test_target_teams_groupprivate_devicetype_any.c diff --git a/tests/6.0/groupprivate/test_teams_groupprivate_devicetype_host.c b/tests/6.0/groupprivate/test_teams_groupprivate_devicetype_host.c new file mode 100644 index 000000000..2c6509e6f --- /dev/null +++ b/tests/6.0/groupprivate/test_teams_groupprivate_devicetype_host.c @@ -0,0 +1,44 @@ +//===--- test_groupprivate_devicetype_host.c ---------------------------===// +// +// OpenMP API Version 6.0 +// Tests the groupprivate directive ensuring proper +// behavior on the host. +// +////===----------------------------------------------------------------------===// + + +#include +#include +#include "ompvv.h" + +#define NUM_TEAMS 4 + +int group_sum; +#pragma omp groupprivate(group_sum) device_type(host) + +int test_groupprivate_devicetype_host(){ + int errors = 0; + int host_sum = 0; + int team_sum = 0; + + #pragma omp teams num_teams(NUM_TEAMS) reduction(+: team_sum) + { + group_sum = omp_get_team_num(); + + team_sum += group_sum; + + } + + for (int i = 0; i < NUM_TEAMS; i++){ + host_sum += i; + } + OMPVV_TEST_AND_SET_VERBOSE(errors, team_sum != host_sum); + return errors; +} + +int main(){ + int errors = 0; + OMPVV_TEST_AND_SET_VERBOSE(errors, test_groupprivate_devicetype_host() != 0); + OMPVV_REPORT_AND_RETURN(errors); +} + diff --git a/tests/6.0/target/test_target_teams_groupprivate_devicetype_any.c b/tests/6.0/target/test_target_teams_groupprivate_devicetype_any.c new file mode 100644 index 000000000..9ffa07318 --- /dev/null +++ b/tests/6.0/target/test_target_teams_groupprivate_devicetype_any.c @@ -0,0 +1,45 @@ +//===--- test_target_groupprivate_devicetype_any.c ---------------------------===// +// +// OpenMP API Version 6.0 +// Tests the target directive with the groupprivate directive ensuring proper +// behavior inside the target region with device type being any. +// +////===----------------------------------------------------------------------===// + + +#include +#include +#include "ompvv.h" + +#define NUM_TEAMS 4 + +int group_sum; +#pragma omp groupprivate(group_sum) device_type(any) + +int test_target_groupprivate_devicetype_any(){ + int errors = 0; + int host_sum = 0; + int team_sum = 0; + + #pragma omp target teams num_teams(NUM_TEAMS) map(tofrom: team_sum) reduction(+: team_sum) + { + group_sum = omp_get_team_num(); + + team_sum += group_sum; + + } + + for (int i = 0; i < NUM_TEAMS; i++){ + host_sum += i; + } + OMPVV_TEST_AND_SET_VERBOSE(errors, team_sum != host_sum); + return errors; +} + +int main(){ + OMPVV_TEST_OFFLOADING; + int errors = 0; + OMPVV_TEST_AND_SET_VERBOSE(errors, test_target_groupprivate_devicetype_any() != 0); + OMPVV_REPORT_AND_RETURN(errors); +} + From 80e07033975b68fd1f281e63da0ff7bc2b5154ea Mon Sep 17 00:00:00 2001 From: rjenaa Date: Thu, 14 Nov 2024 10:55:56 -0800 Subject: [PATCH 10/11] Added nohost version --- ...get_teams_groupprivate_devicetype_nohost.c | 49 +++++++++++++++++++ 1 file changed, 49 insertions(+) create mode 100644 tests/6.0/target/test_target_teams_groupprivate_devicetype_nohost.c diff --git a/tests/6.0/target/test_target_teams_groupprivate_devicetype_nohost.c b/tests/6.0/target/test_target_teams_groupprivate_devicetype_nohost.c new file mode 100644 index 000000000..476030e3f --- /dev/null +++ b/tests/6.0/target/test_target_teams_groupprivate_devicetype_nohost.c @@ -0,0 +1,49 @@ +//===--- test_target_groupprivate_devicetype_nohost.c ---------------------------===// +// +// OpenMP API Version 6.0 +// Tests the target directive with the groupprivate directive ensuring proper +// behavior inside the target region with device type nohost. +// +////===----------------------------------------------------------------------===// + + +#include +#include +#include "ompvv.h" + +#define NUM_TEAMS 4 + +int group_sum; +//#pragma omp groupprivate(group_sum) device_type(nohost) + +int test_target_groupprivate_devicetype_nohost(){ + int errors = 0; + int host_sum = 0; + int team_sum = 0; + + #pragma omp target teams num_teams(NUM_TEAMS) map(tofrom: team_sum, errors) reduction(+: team_sum) + + { + if (omp_get_team_num() == 0 && omp_is_initial_device()){ + errors++; + } + group_sum = omp_get_team_num(); + + team_sum += group_sum; + + } + + for (int i = 0; i < NUM_TEAMS; i++){ + host_sum += i; + } + OMPVV_TEST_AND_SET_VERBOSE(errors, team_sum != host_sum); + return errors; +} + +int main(){ + OMPVV_TEST_OFFLOADING; + int errors = 0; + OMPVV_TEST_AND_SET_VERBOSE(errors, test_target_groupprivate_devicetype_nohost() != 0); + OMPVV_REPORT_AND_RETURN(errors); +} + From 373f0b2f5f9cc2569dc47c9f5d2f962c9ffb95fa Mon Sep 17 00:00:00 2001 From: rjenaa Date: Thu, 14 Nov 2024 10:57:17 -0800 Subject: [PATCH 11/11] Fixed groupprivate pragma --- .../target/test_target_teams_groupprivate_devicetype_nohost.c | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/6.0/target/test_target_teams_groupprivate_devicetype_nohost.c b/tests/6.0/target/test_target_teams_groupprivate_devicetype_nohost.c index 476030e3f..251c41e5d 100644 --- a/tests/6.0/target/test_target_teams_groupprivate_devicetype_nohost.c +++ b/tests/6.0/target/test_target_teams_groupprivate_devicetype_nohost.c @@ -14,7 +14,7 @@ #define NUM_TEAMS 4 int group_sum; -//#pragma omp groupprivate(group_sum) device_type(nohost) +#pragma omp groupprivate(group_sum) device_type(nohost) int test_target_groupprivate_devicetype_nohost(){ int errors = 0;