* 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 --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 };