diff --git a/flang/include/flang/Runtime/api-attrs.h b/flang/include/flang/Runtime/api-attrs.h
index 9c8a67ffc34a8..fc3eb42e1b73f 100644
--- a/flang/include/flang/Runtime/api-attrs.h
+++ b/flang/include/flang/Runtime/api-attrs.h
@@ -109,6 +109,18 @@
 #endif
 #endif /* !defined(RT_CONST_VAR_ATTRS) */
 
+/*
+ * RT_VAR_ATTRS is marking non-const/constexpr module scope variables
+ * referenced by Flang runtime.
+ */
+#ifndef RT_VAR_ATTRS
+#if (defined(__CUDACC__) || defined(__CUDA__)) && defined(__CUDA_ARCH__)
+#define RT_VAR_ATTRS __device__
+#else
+#define RT_VAR_ATTRS
+#endif
+#endif /* !defined(RT_VAR_ATTRS) */
+
 /*
  * RT_DEVICE_COMPILATION is defined for any device compilation.
  * Note that it can only be used reliably with compilers that perform
diff --git a/flang/runtime/environment.cpp b/flang/runtime/environment.cpp
index 29196ae8f3105..fe6701d72c9ff 100644
--- a/flang/runtime/environment.cpp
+++ b/flang/runtime/environment.cpp
@@ -23,7 +23,9 @@ extern char **environ;
 
 namespace Fortran::runtime {
 
-ExecutionEnvironment executionEnvironment;
+RT_OFFLOAD_VAR_GROUP_BEGIN
+RT_VAR_ATTRS ExecutionEnvironment executionEnvironment;
+RT_OFFLOAD_VAR_GROUP_END
 
 static void SetEnvironmentDefaults(const EnvironmentDefaultList *envDefaults) {
   if (!envDefaults) {
diff --git a/flang/runtime/environment.h b/flang/runtime/environment.h
index 6da2c7bb3cf78..49c7dbd2940f8 100644
--- a/flang/runtime/environment.h
+++ b/flang/runtime/environment.h
@@ -10,6 +10,7 @@
 #define FORTRAN_RUNTIME_ENVIRONMENT_H_
 
 #include "flang/Decimal/decimal.h"
+#include "flang/Runtime/api-attrs.h"
 #include <optional>
 
 struct EnvironmentDefaultList;
@@ -32,7 +33,11 @@ enum class Convert { Unknown, Native, LittleEndian, BigEndian, Swap };
 std::optional<Convert> GetConvertFromString(const char *, std::size_t);
 
 struct ExecutionEnvironment {
-  constexpr ExecutionEnvironment(){};
+#if !defined(_OPENMP)
+  // FIXME: https://github.com/llvm/llvm-project/issues/84942
+  constexpr
+#endif
+      ExecutionEnvironment(){};
   void Configure(int argc, const char *argv[], const char *envp[],
       const EnvironmentDefaultList *envDefaults);
   const char *GetEnv(
@@ -51,7 +56,10 @@ struct ExecutionEnvironment {
   bool checkPointerDeallocation{true}; // FORT_CHECK_POINTER_DEALLOCATION
 };
 
-extern ExecutionEnvironment executionEnvironment;
+RT_OFFLOAD_VAR_GROUP_BEGIN
+extern RT_VAR_ATTRS ExecutionEnvironment executionEnvironment;
+RT_OFFLOAD_VAR_GROUP_END
+
 } // namespace Fortran::runtime
 
 #endif // FORTRAN_RUNTIME_ENVIRONMENT_H_