https://gcc.gnu.org/bugzilla/show_bug.cgi?id=83590
--- Comment #1 from Jakub Jelinek <jakub at gcc dot gnu.org> --- With that change, we end up with weird array, TYPE_SIZE_UNIT/TYPE_SIZE on its type is constant, but DECL_SIZE_UNIT/DECL_SIZE on the decl is non-constant (pedantically it is a VLA, but we know the sizes at compile time). So, gimplify_vla_decl is invoked on it. Perhaps just omp-low.c (is_variable_sized) should be changed: static inline bool is_variable_sized (const_tree expr) { + if (DECL_P (expr)) + return TREE_CODE (DECL_SIZE_UNIT (expr)) != INTEGER_CST; return !TREE_CONSTANT (TYPE_SIZE_UNIT (TREE_TYPE (expr))); } I don't have nvptx offloading setup yet on my WS after reinstalling it, can you check if following works with -fopenmp -O2 or not? const int n = 100; int main () { int s = 0; int array[n]; for (int i = 0; i < n; i++) array[i] = i + 1; #pragma omp target parallel for map(tofrom:s) reduction (+:s) for (int i = 0; i < n; i++) s += array[i]; #pragma omp target parallel for map(tofrom:s) reduction (+:s) map(to:array) for (int i = 0; i < n; i++) s += array[i]; if (s != n * (n + 1)) __builtin_abort (); return 0; }