diff mbox

[gomp4,14/14] libgomp: use more generic implementations on nvptx

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

Commit Message

Alexander Monakov Oct. 20, 2015, 6:34 p.m. UTC 
This patch removes 0-size libgomp stubs where generic implementations can be
compiled for the NVPTX target.

It also removes non-stub critical.c, which contains assembly implementations
for GOMP_atomic_{start,end}, but does not contain implementations for
GOMP_critical_*.  My understanding is that OpenACC offloading uses
GOMP_atomic_* routines (by virtue of OpenMP lowering using them).  Linking in
GOMP_critical_* and dependencies would be pointless for OpenACC.

If OpenACC indeed uses GOMP_atomic_*, then it makes sense to split them out
into a separate file (atomic.c?).

After this patch, a few 0-size stubs remain in libgomp/config/nvptx.  They
fall roughly into these categories:

  - Files which must remain a 0-size stub, like env.c, because they implement
    something that doesn't make sense on an accelerator-only target.

  - Files that are 0-size because all functionality is implemented in the
    corresponding header for now, like mutex.c.

  - Fiels that are 0-size, but probably should be changed to use generic
    implementations soon, like sections.c

  - Files that are 0-size and will need to have a custom implementation for
    nvptx, like time.c
---
 libgomp/config/nvptx/alloc.c    |  0
 libgomp/config/nvptx/barrier.c  |  0
 libgomp/config/nvptx/critical.c | 57 -----------------------------------------
 libgomp/config/nvptx/error.c    |  0
 libgomp/config/nvptx/iter.c     |  0
 libgomp/config/nvptx/iter_ull.c |  0
 libgomp/config/nvptx/loop.c     |  0
 libgomp/config/nvptx/loop_ull.c |  0
 libgomp/config/nvptx/ordered.c  |  0
 libgomp/config/nvptx/parallel.c |  0
 libgomp/config/nvptx/single.c   |  0
 libgomp/config/nvptx/task.c     |  0
 libgomp/config/nvptx/work.c     |  0
 13 files changed, 57 deletions(-)
 delete mode 100644 libgomp/config/nvptx/alloc.c
 delete mode 100644 libgomp/config/nvptx/barrier.c
 delete mode 100644 libgomp/config/nvptx/critical.c
 delete mode 100644 libgomp/config/nvptx/error.c
 delete mode 100644 libgomp/config/nvptx/iter.c
 delete mode 100644 libgomp/config/nvptx/iter_ull.c
 delete mode 100644 libgomp/config/nvptx/loop.c
 delete mode 100644 libgomp/config/nvptx/loop_ull.c
 delete mode 100644 libgomp/config/nvptx/ordered.c
 delete mode 100644 libgomp/config/nvptx/parallel.c
 delete mode 100644 libgomp/config/nvptx/single.c
 delete mode 100644 libgomp/config/nvptx/task.c
 delete mode 100644 libgomp/config/nvptx/work.c

diff --git a/libgomp/config/nvptx/error.c b/libgomp/config/nvptx/error.c
deleted file mode 100644
index e69de29..0000000
diff --git a/libgomp/config/nvptx/iter.c b/libgomp/config/nvptx/iter.c
deleted file mode 100644
index e69de29..0000000
diff --git a/libgomp/config/nvptx/iter_ull.c b/libgomp/config/nvptx/iter_ull.c
deleted file mode 100644
index e69de29..0000000
diff --git a/libgomp/config/nvptx/loop.c b/libgomp/config/nvptx/loop.c
deleted file mode 100644
index e69de29..0000000
diff --git a/libgomp/config/nvptx/loop_ull.c b/libgomp/config/nvptx/loop_ull.c
deleted file mode 100644
index e69de29..0000000
diff --git a/libgomp/config/nvptx/ordered.c b/libgomp/config/nvptx/ordered.c
deleted file mode 100644
index e69de29..0000000
diff --git a/libgomp/config/nvptx/parallel.c b/libgomp/config/nvptx/parallel.c
deleted file mode 100644
index e69de29..0000000
diff --git a/libgomp/config/nvptx/single.c b/libgomp/config/nvptx/single.c
deleted file mode 100644
index e69de29..0000000
diff --git a/libgomp/config/nvptx/task.c b/libgomp/config/nvptx/task.c
deleted file mode 100644
index e69de29..0000000
diff --git a/libgomp/config/nvptx/work.c b/libgomp/config/nvptx/work.c
deleted file mode 100644
index e69de29..0000000

Comments

Jakub Jelinek Oct. 21, 2015, 10:12 a.m. UTC | #1 
On Tue, Oct 20, 2015 at 09:34:36PM +0300, Alexander Monakov wrote:
> This patch removes 0-size libgomp stubs where generic implementations can be
> compiled for the NVPTX target.
> 
> It also removes non-stub critical.c, which contains assembly implementations
> for GOMP_atomic_{start,end}, but does not contain implementations for
> GOMP_critical_*.  My understanding is that OpenACC offloading uses
> GOMP_atomic_* routines (by virtue of OpenMP lowering using them).  Linking in
> GOMP_critical_* and dependencies would be pointless for OpenACC.
> 
> If OpenACC indeed uses GOMP_atomic_*, then it makes sense to split them out
> into a separate file (atomic.c?).

Splitting atomic stuff into atomic.c and keeping critical in critical.c is
fine.

Note, when you actually start supporting multiple teams, likely
both GOMP_critical_* and GOMP_atomic_* will need to change for NVPTX
- the locks should be per-CTA variables in that case, rather than global
vars.  In OpenMP 4.5, 8, 16, 32 and 64-bit atomics (on aligned vars) are
required to be per-device, but that should be fine for PTX, which (I'd hope)
should expand all those inline using atomic instructions.  GOMP_atomic_*
is then used either for other types, where it should be private to the
contention group (CTA for PTX), and for reductions if HW atomics aren't
used (weird type or many different reductions) but for parallel/for/sections
reductions different CTAs shouldn't really fight for the same lock, each
should have its own.  Critical is similar thing.
I bet for PTX we'll need for named critical to adjust the compiler side.

And for teams reduction, PTX will need a different approach.

	Jakub
diff mbox

Patch

diff --git a/libgomp/config/nvptx/alloc.c b/libgomp/config/nvptx/alloc.c
deleted file mode 100644
index e69de29..0000000
diff --git a/libgomp/config/nvptx/barrier.c b/libgomp/config/nvptx/barrier.c
deleted file mode 100644
index e69de29..0000000
diff --git a/libgomp/config/nvptx/critical.c b/libgomp/config/nvptx/critical.c
deleted file mode 100644
index 1f55aad..0000000
--- a/libgomp/config/nvptx/critical.c
+++ /dev/null
@@ -1,57 +0,0 @@ 
-/* GOMP atomic routines
-
-   Copyright (C) 2014-2015 Free Software Foundation, Inc.
-
-   Contributed by Mentor Embedded.
-
-   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/>.  */
-
-__asm__ ("// BEGIN VAR DEF: libgomp_ptx_lock\n"
-	 ".global .align 4 .u32 libgomp_ptx_lock;\n"
-	 "\n"
-	 "// BEGIN GLOBAL FUNCTION DECL: GOMP_atomic_start\n"
-	 ".visible .func GOMP_atomic_start;\n"
-	 "// BEGIN GLOBAL FUNCTION DEF: GOMP_atomic_start\n"
-	 ".visible .func GOMP_atomic_start\n"
-	 "{\n"
-	 "	.reg .pred 	%p<2>;\n"
-	 "	.reg .s32 	%r<2>;\n"
-	 "	.reg .s64 	%rd<2>;\n"
-	 "BB5_1:\n"
-	 "	mov.u64 	%rd1, libgomp_ptx_lock;\n"
-	 "	atom.global.cas.b32 	%r1, [%rd1], 0, 1;\n"
-	 "	setp.ne.s32	%p1, %r1, 0;\n"
-	 "	@%p1 bra 	BB5_1;\n"
-	 "	ret;\n"
-	 "	}\n"
-	 "// BEGIN GLOBAL FUNCTION DECL: GOMP_atomic_end\n"
-	 ".visible .func GOMP_atomic_end;\n"
-	 "// BEGIN GLOBAL FUNCTION DEF: GOMP_atomic_end\n"
-	 ".visible .func GOMP_atomic_end\n"
-	 "{\n"
-	 "	.reg .s32 	%r<2>;\n"
-	 "	.reg .s64 	%rd<2>;\n"
-	 "	mov.u64 	%rd1, libgomp_ptx_lock;\n"
-	 "	atom.global.exch.b32 	%r1, [%rd1], 0;\n"
-	 "	ret;\n"
-	 "	}");