declare_target local

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))
 

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"); \
@@ -102,6 +102,15 @@ _Pragma("omp target map (from: _ompvv_isOffloadingOn)") \
   OMPVV_RETURN(err); \
 }
 
+// Macro to check if it is a shared data environment
+#define OMPVV_TEST_SHARED_ENVIRONMENT_PROBE \
+  int _ompvv_isSharedEnv = 0; \
+  _ompvv_isOffloadingOn = 0; \
+_Pragma("omp target map (from: _ompvv_isOffloadingOn) map(to: _ompvv_isSharedEnv)") \
+  {  _ompvv_isOffloadingOn = !omp_is_initial_device();  \
+     _ompvv_isSharedEnv = 1; \
+  }
+
 // Macro to check if it is a shared data environment
 #define OMPVV_TEST_SHARED_ENVIRONMENT_PROBE \
   int _ompvv_isSharedEnv = 0; \
@@ -124,6 +133,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 +156,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

tests/6.0/declare_target/test_declare_target_directive_local.c

@@ -0,0 +1,91 @@
+//--------------- 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 <omp.h>
+
+#define N 100
+
+int sum;
+int x[N];
+
+#pragma omp declare_target to(sum, x) // 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(int dev_id) {
+  int i;
+  for (int j = 0; j < N; ++j) {
+    x[j] = j + dev_id;
+  }
+  #pragma omp for reduction(+ : sum)
+  for (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];
+  int sum_array[TotGpus];
+
+  for (int i = 0; i < TotGpus; i++) {
+    sum_array[i] = 0;
+    errors_arr[i] = 0;
+    // #pragma omp target device(i) map(tofrom : sum_array)
+    // {
+    //   // init_x(i);
+    //   // sum = 0;
+    //   sum_array[i] = 0;
+    // }
+  }
+
+  for (int i = 0; i < TotGpus; i++) {
+    #pragma omp target parallel device(i) nowait
+    {
+      foo(i);
+    }
+  }
+  #pragma omp taskwait
+
+  for (int i = 0; i < TotGpus; i++) {
+    #pragma omp target map(tofrom : errors_arr) device(i)
+    {
+      if ((N) * (N - 1) / 2 + (N)*i != sum_array[i]) {
+        ++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);
+}