diff --git a/Makefile b/Makefile index a215f088c..0aa4f3690 100644 --- a/Makefile +++ b/Makefile @@ -322,13 +322,12 @@ $(CURDIR)/tests/4.5/application_kernels/qmcpack_target_static_lib.c.o: $(CURDIR) $(@:.run=.o) $(VERBOSE) $(if $(LOG),$(RECORD)$(notdir $(@:.run=.log)) \ && echo "PASS" > $(LOGTEMPFILE) || echo "FAIL" > $(LOGTEMPFILE)), \ $(if $(findstring _env_,$@), \ - $(call CHECK_OUTPUT, $(call loadModules,$(C_COMPILER_MODULE)) $(BSRUN)$(RUN_TEST) --env \ $(shell echo "$@" | sed -e 's@.*/@@' -e 's@test_\(.*\)_env_.*@\1@' | tr 'a-z' 'A-Z') \ $(shell echo "$@" | sed -e 's@.*/@@' -e 's@.*_env_\([^.]*\).*@\1@') \ $(@:.run=.o) $(VERBOSE) $(if $(LOG),$(RECORD)$(notdir $(@:.run=.log))) ), \ $(call CHECK_OUTPUT, $(call loadModules,$(C_COMPILER_MODULE)) $(BSRUN)$(RUN_TEST) $(@:.run=.o) $(VERBOSE) $(if $(LOG),$(RECORD)$(notdir $(@:.run=.log))) ) \ - ) + )) -$(call log_section_footer,"RUN",$(SYSTEM),$$(cat $(LOGTEMPFILE)),$(LOG_NOTE),$(notdir $(@:.run=.log))) -@$(if $(LOG), rm $(LOGTEMPFILE)) diff --git a/ompvv/ompvv.h b/ompvv/ompvv.h index 70f78df3b..cc8c37ab7 100644 --- a/ompvv/ompvv.h +++ b/ompvv/ompvv.h @@ -80,12 +80,12 @@ _Pragma("omp target map (from: _ompvv_isOffloadingOn)") \ #define OMPVV_REPORT(err) { \ OMPVV_INFOMSG("The value of " #err " is %d.", err); \ if (_ompvv_isOffloadingOn == -1) \ - if (err == -667) \ + if (err == OMPVV_SKIPPED_EXIT_CODE) \ printf("[OMPVV_RESULT: %s] Test %s.\n", __FILENAME__, "skipped"); \ else \ printf("[OMPVV_RESULT: %s] Test %s.\n", __FILENAME__, ((err) == 0)? "passed":"failed"); \ else \ - if (err == -667) \ + if (err == OMPVV_SKIPPED_EXIT_CODE) \ printf("[OMPVV_RESULT: %s] Test %s on the %s.\n", __FILENAME__, "skipped", (_ompvv_isOffloadingOn)? "device" : "host"); \ else \ printf("[OMPVV_RESULT: %s] Test %s on the %s.\n", __FILENAME__, ((err) == 0)? "passed":"failed", (_ompvv_isOffloadingOn)? "device" : "host"); \ @@ -124,6 +124,12 @@ _Pragma("omp target map (from: _ompvv_isOffloadingOn) map(to: _ompvv_isSharedEnv var2set = (_ompvv_isOffloadingOn && _ompvv_isSharedEnv == 1);\ } +// Macro for skipping a test +#define OMPVV_CHECK_TO_SKIP(condition) { \ + int SKIPPED_EXIT_CODE = OMPVV_SKIPPED_EXIT_CODE; \ + if (condition) {OMPVV_REPORT_AND_RETURN(SKIPPED_EXIT_CODE)}; \ + } + // Macros to provide thread and team nums if they are not specified #ifndef OMPVV_NUM_THREADS_DEVICE #define OMPVV_NUM_THREADS_DEVICE 8 @@ -141,10 +147,6 @@ _Pragma("omp target map (from: _ompvv_isOffloadingOn) map(to: _ompvv_isSharedEnv #define OMPVV_NUM_THREADS_HOST 8 #endif -#ifndef OMPVV_NUM_THREADS_HOST - #define OMPVV_NUM_THREADS_HOST 8 -#endif - #ifndef OMPVV_SKIPPED_EXIT_CODE #define OMPVV_SKIPPED_EXIT_CODE -667 #endif diff --git a/tests/6.0/declare_target/test_declare_target_directive_local.c b/tests/6.0/declare_target/test_declare_target_directive_local.c new file mode 100644 index 000000000..210077322 --- /dev/null +++ b/tests/6.0/declare_target/test_declare_target_directive_local.c @@ -0,0 +1,83 @@ +//----------------test_declare_target_directive_local.c--------------------------------// +// OpenMP API Version 6.0 November 2024 +// Pg. 902, line 4 +// *********** +// DIRECTIVE: declare_target +// CLAUSE: local +// The declare_target directive is used with the local clause to make data +// persist on multiple devices, each with their own copy of the referenced +// variables in the clause. If the values retrieved from each of the available +// devices add up to the expected amount, the test will pass. +//-------------------------------------------------------------------------------------// + +#include "ompvv.h" +#include + +#define N 100 + +int sum; +int x[N]; + +#pragma omp declare_target local(sum, x) +#pragma omp begin declare_target +void init_x(int dev_id) { + for (int j = 0; j < N; ++j) { + x[j] = j + dev_id; + } +} + +void foo(void) { + #pragma omp for reduction(+ : sum) + for (int i = 0; i < N; i++) { + sum += x[i]; + } +} +#pragma omp end declare target + +int test_declare_target_local() { + int errors = 0; + int TotGpus = omp_get_num_devices(); + int errors_arr[TotGpus]; + + for (int i = 0; i < TotGpus; i++) { + errors_arr[i] = 0; + { + init_x(i); + sum = 0; + } + } + for (int i = 0; i < TotGpus; i++) { + #pragma omp target parallel device(i) + { + foo(); + } + } + #pragma omp taskwait + + for (int i = 0; i < TotGpus; i++) { + #pragma omp target map(tofrom : errors_arr[i]) device(i) + { + if ((N) * (N - 1) / 2 + (N)*i != sum) { + ++errors_arr[i]; + } + } + } + for (int i = 0; i < TotGpus; i++) { + errors += errors_arr[i]; + } + OMPVV_TEST_AND_SET(errors, errors != 0); + + return errors; +} + +int main() { + int errors = 0, test_exit_code = 0; + int TotGpus = omp_get_num_devices(); + OMPVV_WARNING_IF(TotGpus < 1, "Test requires non-host devices, but none were " + "found. \n This test will be skipped.\n"); + OMPVV_CHECK_TO_SKIP(TotGpus < 1) + OMPVV_TEST_OFFLOADING; + + OMPVV_TEST_AND_SET(errors, test_declare_target_local() != 0); + OMPVV_REPORT_AND_RETURN(errors); +}