On NVPTX, we don't need most of target.c functionality, except for GOMP_teams.
Provide it as a copy of the generic implementation for now (it most likely
will need to change down the line: on NVPTX we do need to spawn several
thread blocks for #pragma omp teams).

Alternatively, it might make sense to split GOMP_teams out of target.c into
its own file (teams.c?), leaving target.c a 0-size stub in config/nvptx.

        * config/nvptx/target.c: New.
---
 libgomp/config/nvptx/target.c | 39 +++++++++++++++++++++++++++++++++++++++
 1 file changed, 39 insertions(+)

diff --git a/libgomp/config/nvptx/target.c b/libgomp/config/nvptx/target.c
index e69de29..ad36013 100644
--- a/libgomp/config/nvptx/target.c
+++ b/libgomp/config/nvptx/target.c
@@ -0,0 +1,39 @@
+/* Copyright (C) 2013-2015 Free Software Foundation, Inc.
+   Contributed by Jakub Jelinek <ja...@redhat.com>.
+
+   This file is part of the GNU Offloading and Multi Processing Library
+   (libgomp).
+
+   Libgomp is free software; you can redistribute it and/or modify it
+   under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3, or (at your option)
+   any later version.
+
+   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
+   more details.
+
+   Under Section 7 of GPL version 3, you are granted additional
+   permissions described in the GCC Runtime Library Exception, version
+   3.1, as published by the Free Software Foundation.
+
+   You should have received a copy of the GNU General Public License and
+   a copy of the GCC Runtime Library Exception along with this program;
+   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+   <http://www.gnu.org/licenses/>.  */
+
+#include "libgomp.h"
+#include <limits.h>
+
+void
+GOMP_teams (unsigned int num_teams, unsigned int thread_limit)
+{
+  if (thread_limit)
+    {
+      struct gomp_task_icv *icv = gomp_icv (true);
+      icv->thread_limit_var
+       = thread_limit > INT_MAX ? UINT_MAX : thread_limit;
+    }
+  (void) num_teams;
+}

Reply via email to