On Thu, Jan 20, 2022 at 09:26:50AM +0100, Thomas Schwinge wrote:
> That's what we need to look into, in particular: if we decompose (GIMPLE
> sequence) an OpenACC 'kernels' region into parts, how to move or
> otherwise handle any 'GIMPLE_DEBUG's.
I admit I haven't looked at the pass except now for the toplevel comment.
It says that OpenACC constructs in the region are perhaps adjusted but
their body is unchanged, so that suggests that debug stmts inside of those
bodies should be kept as is.
Next it says that sequential code in between those loops/whatever are
put into some sequential construct, so I guess if you decide so because
of some non-debug stmts, you can just move the debug stmts into that
construct as well, including those debug stmts before the first such
non-debug stmt and debug stmts after the last such non-debug stmts.
It is not a perfect solution, because normally debug stmts before
loops would affect also what is in the loop unless overridden, but
what the pass does seems terribly destructive for debug experience anyway.
There is then another case, only debug stmts e.g. in between or before
the loops or after them and nothing else. Perhaps throwing them away at
this point is the best thing to do (but, all of this only after the pass
decides that it will change something).
Another thing is, this is apparently a very early pass, so most real
debug stmts don't exist, they are typically created later.
I'd expect you mostly see gimple_debug_begin_stmt_p stmts.
Those can be removed more easily, it doesn't mean var has this value
for the following code until stated otherwise, but it just said here was
the start of some source code statement. So, if you drop them, all that
will work worse is break some_line.
So citing from e.g. PR100400:
void foo ()
{
# DEBUG BEGIN_STMT // Outside of region, don't touch this
#pragma omp target oacc_kernels map(force_tofrom:p [len: 8])
{
int c.0;
# DEBUG BEGIN_STMT // Drop this
try
{
# DEBUG BEGIN_STMT // If p = &c; is moved somewhere, move the
surrounding DEBUG BEGIN_STMTs with it
# DEBUG BEGIN_STMT
p = &c;
# DEBUG BEGIN_STMT // Up to here
#pragma acc loop independent private(c.0) private(c)
for (c.0 = 0; c.0 < 1; c.0 = c.0 + 1)
{
c = c.0;
# DEBUG BEGIN_STMT // Keep this in the body
}
}
finally
{
c = {CLOBBER};
}
}
}
If you don't have time for it right now, after deciding you are
going to transform it just gsi_remove gimple_debug_begin_stmt_p stmts
you don't know how to handle.
> With these things now hopfully clarified, is the attached
> "Catch 'GIMPLE_DEBUG' misbehavior in OpenACC 'kernels' decomposition
> [PR100400, PR103836, PR104061]" OK to push? It's of course not the final
> fix, but it at least makes obvious any current silent miscompilation, and
> incremental improvement over the current status.
No, users really don't want to see sorry messages just because they turned
-g on their code. They might be ok with their kernels not being easily
debuggable, but they surely will not be ok with not being able to debug
the host code in the same TU.
Jakub
