diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index 0204ad3..df915b9 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -114,6 +114,9 @@ static unsigned worker_red_align;
 #define worker_red_name "__worker_red"
 static GTY(()) rtx worker_red_sym;
 
+/* True if any function references __nvptx_stacks.  */
+static bool need_softstack_decl;
+
 /* Allocate a new, cleared machine_function structure.  */
 
 static struct machine_function *
@@ -689,15 +692,46 @@ nvptx_declare_function_name (FILE *file, const char *name, const_tree decl)
 
   /* Declare a local variable for the frame.  */
   sz = get_frame_size ();
-  if (sz > 0 || cfun->machine->has_call_with_sc)
+  if (sz == 0 && cfun->machine->has_call_with_sc)
+    sz = 1;
+  if (sz > 0)
     {
       int alignment = crtl->stack_alignment_needed / BITS_PER_UNIT;
 
-      fprintf (file, "\t.reg.u%d %%frame;\n"
-	       "\t.local.align %d .b8 %%farray[" HOST_WIDE_INT_PRINT_DEC"];\n",
-	       BITS_PER_WORD, alignment, sz == 0 ? 1 : sz);
-      fprintf (file, "\tcvta.local.u%d %%frame, %%farray;\n",
-	       BITS_PER_WORD);
+      fprintf (file, "\t.reg.u%d %%frame;\n", BITS_PER_WORD);
+      if (TARGET_SOFT_STACK)
+	{
+	  /* Maintain 64-bit stack alignment.  */
+	  int keep_align = BIGGEST_ALIGNMENT / BITS_PER_UNIT;
+	  sz = (sz + keep_align - 1) & ~(keep_align - 1);
+	  int bits = BITS_PER_WORD;
+	  fprintf (file, "\t.reg.u32 %%fstmp0;\n");
+	  fprintf (file, "\t.reg.u%d %%fstmp1;\n", bits);
+	  fprintf (file, "\t.reg.u%d %%fstmp2;\n", bits);
+	  fprintf (file, "\tmov.u32 %%fstmp0, %%tid.y;\n");
+	  fprintf (file, "\tmul%s.u32 %%fstmp1, %%fstmp0, %d;\n",
+	           bits == 64 ? ".wide" : "", bits);
+	  fprintf (file, "\tmov.u%d %%fstmp2, __nvptx_stacks;\n", bits);
+	  /* fstmp2 = &__nvptx_stacks[tid.y];  */
+	  fprintf (file, "\tadd.u%d %%fstmp2, %%fstmp2, %%fstmp1;\n", bits);
+	  fprintf (file, "\tld.shared.u%d %%fstmp1, [%%fstmp2];\n", bits);
+	  fprintf (file, "\tsub.u%d %%frame, %%fstmp1, "
+	           HOST_WIDE_INT_PRINT_DEC ";\n", bits, sz);
+	  if (alignment > keep_align)
+	    fprintf (file, "\tand.b%d %%frame, %%frame, %d;\n",
+		     bits, -alignment);
+	  if (!crtl->is_leaf)
+	    fprintf (file, "\tst.shared.u%d [%%fstmp2], %%frame;\n", bits);
+	  need_softstack_decl = true;
+	}
+      else
+	{
+	  fprintf (file, "\t.local.align %d "
+		   ".b8 %%farray[" HOST_WIDE_INT_PRINT_DEC"];\n",
+		   alignment, sz);
+	  fprintf (file, "\tcvta.local.u%d %%frame, %%farray;\n",
+		   BITS_PER_WORD);
+	}
     }
 
   if (cfun->machine->has_call_with_varargs)
@@ -734,6 +768,13 @@ nvptx_output_return (void)
 {
   machine_mode mode = (machine_mode)cfun->machine->ret_reg_mode;
 
+  if (TARGET_SOFT_STACK
+      && !crtl->is_leaf
+      && (get_frame_size () > 0 || cfun->machine->has_call_with_sc))
+    {
+      int bits = BITS_PER_WORD;
+      fprintf (asm_out_file, "\tst.shared.u%d [%%fstmp2], %%fstmp1;\n", bits);
+    }
   if (mode != VOIDmode)
     {
       mode = arg_promotion (mode);
@@ -3278,6 +3319,11 @@ nvptx_file_end (void)
 	       worker_red_align,
 	       worker_red_name, worker_red_size);
     }
+
+  if (need_softstack_decl)
+    {
+      fprintf (asm_out_file, ".extern .shared .u64 __nvptx_stacks[];\n;");
+    }
 }
 
 /* Expander for the shuffle builtins.  */
diff --git a/gcc/config/nvptx/nvptx.opt b/gcc/config/nvptx/nvptx.opt
index 8017046..7ab09b9 100644
--- a/gcc/config/nvptx/nvptx.opt
+++ b/gcc/config/nvptx/nvptx.opt
@@ -28,3 +28,7 @@ Generate code for a 64-bit ABI.
 mmainkernel
 Target Report RejectNegative
 Link in code for a __main kernel.
+
+msoft-stack
+Target Report Mask(SOFT_STACK)
+Use custom stacks instead of local memory for automatic storage.
diff --git a/gcc/doc/invoke.texi b/gcc/doc/invoke.texi
index 587e30e..6e45fb6 100644
--- a/gcc/doc/invoke.texi
+++ b/gcc/doc/invoke.texi
@@ -18935,6 +18935,13 @@ Generate code for 32-bit or 64-bit ABI.
 Link in code for a __main kernel.  This is for stand-alone instead of
 offloading execution.
 
+@item -msoft-stack
+@opindex msoft-stack
+Do not use @code{.local} memory for automatic storage.  Instead, use pointer
+in shared memory array @code{char *__nvptx_stacks[]} at position @code{tid.y}
+as the stack pointer.  This is for placing automatic variables into storage
+that can be accessed from other threads, or modified with atomic instructions.
+
 @end table
 
 @node PDP-11 Options
diff --git a/libgcc/config/nvptx/crt0.s b/libgcc/config/nvptx/crt0.s
index 38327ed..7a42e87 100644
--- a/libgcc/config/nvptx/crt0.s
+++ b/libgcc/config/nvptx/crt0.s
@@ -22,6 +22,9 @@
         exit;
 }
 
+.visible .shared .u64 __nvptx_stacks[1];
+.global .u64 %__softstack[16384];
+
 .extern .func (.param.u32 retval) main (.param.u32 argc, .param.u64 argv);
 
 .visible .entry __main (.param .u64 __retval, .param.u32 __argc, .param.u64 __argv)
@@ -34,6 +37,12 @@
         ld.param.u64    %rd0, [__retval];
         st.global.u64   [%__exitval], %rd0;
 
+        .reg .u64 %stackptr;
+        mov.u64	%stackptr, %__softstack;
+        cvta.global.u64	%stackptr, %stackptr;
+        add.u64	%stackptr, %stackptr, 131072;
+        st.shared.u64	[__nvptx_stacks], %stackptr;
+
 	ld.param.u32	%r1, [__argc];
 	ld.param.u64	%rd1, [__argv];
 	st.param.u32	[%argc], %r1;
