diff mbox

[gomp4,13/14] libgomp: provide minimal GOMP_teams

Message ID 1445366076-16082-14-git-send-email-amonakov@ispras.ru
State New
Headers show

Commit Message

Alexander Monakov Oct. 20, 2015, 6:34 p.m. UTC 
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(+)

Comments

Jakub Jelinek Oct. 21, 2015, 10:03 a.m. UTC | #1 
On Tue, Oct 20, 2015 at 09:34:35PM +0300, Alexander Monakov wrote:
> 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.

This is fine.  I bet it will need changes on the compiler side too,
GOMP_teams has been written with rough idea of the dynamic parallelism,
we'll need to do something with private (per-CTA) vars, shared (global) vars
etc., dunno if with tweaks just in the late omp lowering pass after it
is lowered/expanded as OpenMP already, or if we'll need Martin's HSA
approach.

	Jakub
diff mbox

Patch

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 <jakub@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;
+}