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/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"
- " }");
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
