diff mbox

[gomp-nvptx,2/7] nvptx libgcc: use attribute shared

Message ID 1458323327-9908-3-git-send-email-amonakov@ispras.ru
State New
Headers show

Commit Message

Alexander Monakov March 18, 2016, 5:48 p.m. UTC 
* config/nvptx/crt0.c (__nvptx_stacks): Define in C.  Use it...
	(__nvptx_uni): Ditto.
	(__main): ...here instead of inline asm.
	* config/nvptx/stacks.c (__nvptx_stacks): Define in C.
	(__nvptx_uni): Ditto.
---
 libgcc/ChangeLog.gomp-nvptx  |  8 ++++++++
 libgcc/config/nvptx/crt0.c   | 10 ++++------
 libgcc/config/nvptx/stacks.c |  9 ++-------
 3 files changed, 14 insertions(+), 13 deletions(-)
diff mbox

Patch

diff --git a/libgcc/config/nvptx/crt0.c b/libgcc/config/nvptx/crt0.c
index 5e04b0f..9e9a25e 100644
--- a/libgcc/config/nvptx/crt0.c
+++ b/libgcc/config/nvptx/crt0.c
@@ -41,10 +41,8 @@  abort (void)
   exit (255);
 }
 
-asm ("// BEGIN GLOBAL VAR DECL: __nvptx_stacks");
-asm (".extern .shared .u64 __nvptx_stacks[32];");
-asm ("// BEGIN GLOBAL VAR DECL: __nvptx_uni");
-asm (".extern .shared .u32 __nvptx_uni[32];");
+extern char *__nvptx_stacks[32] __attribute__((shared));
+extern unsigned __nvptx_uni[32] __attribute__((shared));
 
 extern int main (int argc, char *argv[]);
 
@@ -54,8 +52,8 @@  __main (int *__retval, int __argc, char *__argv[])
   __exitval = __retval;
 
   static char gstack[131072] __attribute__((aligned(8)));
-  asm ("st.shared.u64 [__nvptx_stacks], %0;" : : "r" (gstack + sizeof gstack));
-  asm ("st.shared.u32 [__nvptx_uni], %0;" : : "r" (0));
+  __nvptx_stacks[0] = gstack + sizeof gstack;
+  __nvptx_uni[0] = 0;
 
   exit (main (__argc, __argv));
 }
diff --git a/libgcc/config/nvptx/stacks.c b/libgcc/config/nvptx/stacks.c
index a7e640a..4640fc9 100644
--- a/libgcc/config/nvptx/stacks.c
+++ b/libgcc/config/nvptx/stacks.c
@@ -21,10 +21,5 @@ 
    see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
    <http://www.gnu.org/licenses/>.  */
 
-/* __shared__ char *__nvptx_stacks[32];  */
-asm ("// BEGIN GLOBAL VAR DEF: __nvptx_stacks");
-asm (".visible .shared .u64 __nvptx_stacks[32];");
-
-/* __shared__ unsigned __nvptx_uni[32];  */
-asm ("// BEGIN GLOBAL VAR DEF: __nvptx_uni");
-asm (".visible .shared .u32 __nvptx_uni[32];");
+char *__nvptx_stacks[32] __attribute__((shared)) = { 0 };
+unsigned __nvptx_uni[32] __attribute__((shared)) = { 0 };