https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104769
Bug ID: 104769
Summary: [nvptx] mptx/misa multilibs
Product: gcc
Version: 12.0
Status: UNCONFIRMED
Severity: enhancement
Priority: P3
Component: target
Assignee: unassigned at gcc dot gnu.org
Reporter: vries at gcc dot gnu.org
Target Milestone: ---
The current situation is:
- default: -misa=sm_35 -mptx=6.0
- libraries: -misa=sm_30 -mptx=3.1
There's an open question on whether we need or want multilibs for different
values of misa and/or mptx.
First, let's look at drivers. The libraries use the lowest (supported by gcc)
versions, which are supported by the latest production branch driver 510.x. So
from that point of view, there's no need.
Then let's look at misa. Potentially, having a multilib for each misa value is
the most optimal solution. But is it faster in practice? I'd expect the
performance critical code to reside in the applications, which can be using
their own optimal -misa settings.
Now, let's look at mptx. Newer ptx isa versions unlock new features, which
could enable code that is currently not compiling. F.i., alloca, support for
.alias. Again here, the question is whether we need this only in the
applications or also in the libraries. We'll assess this each time we start
using a new feature.
Finally, let's look at mgomp/libgomp.a.
This has code like:
...
@ %r73 membar.sys;
@ %r73 atom.add.u64 %r25,[%r35+72],%r32;
@ %r73 membar.sys;
{
.reg .b32 act;
vote.ballot.b32 act,1;
.reg .pred uni;
setp.eq.b32 uni,act,0xffffffff;
@ ! uni trap;
@ ! uni exit;
...
Arguably, when compiling with -mptx=6.0, which is the default, we want to use
bar.warp.sync instead.
So, perhaps a default mptx=3.1 and a multilib mptx=6.0?
