https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104364
--- Comment #3 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Thomas Schwinge from comment #0)
> And, is it correct here to use the non-'atom' replacement, though? '%frame'
> comes from:
>
> .visible .func GOMP_taskwait
> {
> .reg .u64 %stack;
> .reg .u64 %frame;
> .reg .u64 %sspslot;
> .reg .u64 %sspprev;
> {
> .reg .u32 %fstmp0;
> .reg .u64 %fstmp1;
> .reg .u64 %fstmp2;
> mov.u32 %fstmp0,%tid.y;
> mul.wide.u32 %fstmp1,%fstmp0,8;
> mov.u64 %fstmp2,__nvptx_stacks;
> add.u64 %sspslot,%fstmp2,%fstmp1;
> ld.shared.u64 %sspprev,[%sspslot];
> sub.u64 %frame,%sspprev,32;
> sub.u64 %stack,%frame,16;
> st.shared.u64 [%sspslot],%stack;
> }
> [...]
>
> ... which -- at least from my superficial look -- seems to be some kind of
> 'shared' location, thus mandating 'atom' access? (But haven't looked in
> detail, so I may certainly be wrong, of course.)
Yeah, that's a mistake, when using -msoft-stack, we can't assume that
stack-related means local.
