... now that it has been enhanced to do so.
This is a follow-up to PR97348 commit 383400a607
"[nvptx] Set -misa=sm_35 by default".
gcc/
PR target/97348
* config/nvptx/nvptx.h (ASM_SPEC): Don't set.
* config/nvptx/nvptx.opt (misa): Adjust comment.
In the docs we have for m64:
...
Ignored, but preserved for backward compatibility. Only 64-bit ABI is
supported.
...
But with --target-help, we have instead:
...
$ gcc --target-help
...
-m64 Generate code for a 64-bit ABI.
...
which could be interpreted as meaning that generating code for a 32-bit ABI is
still possible.
Fix this by instead emitting the same text as in the docs:
...
-m64 Ignored, but preserved for backward compatibility. Only 64-bit
ABI is supported.
...
Tested on nvptx.
gcc/ChangeLog:
2022-03-29 Tom de Vries <tdevries@suse.de>
* config/nvptx/nvptx.opt (m64): Update help text to reflect that it
is ignored.
Say we have an sm_50 board, and we want to run a benchmark using the highest
possible march setting.
Currently there's march=sm_30, march=sm_35, march=sm_53, but no march=sm_50.
So, we'd need to pick march=sm_35.
Likewise, for a test script that handles multiple boards, we'd need a mapping
from native board sm_xx to march, which might have to be updated with newer
gcc releases.
Add an option march-map, such that we can just specify march-map=sm_50, and
let the compiler map this to the appropriate march.
The option is implemented as a list of aliases, such that we have a somewhat
lengthy (17 lines in total):
...
$ gcc --help=target
...
-march-map=sm_30 Same as -misa=sm_30.
-march-map=sm_32 Same as -misa=sm_30.
...
-march-map=sm_87 Same as -misa=sm_80.
-march-map=sm_90 Same as -misa=sm_80.
...
This implementation was chosen in the hope that it'll be easier if
we end up with some misa multilib.
It would be nice to have the mapping list generated from an updated
nvptx-sm.def, but for now it's spelled out in nvptx.opt.
Tested on nvptx.
gcc/ChangeLog:
2022-03-29 Tom de Vries <tdevries@suse.de>
PR target/104714
* config/nvptx/nvptx.opt (march-map=*): Add aliases.
gcc/testsuite/ChangeLog:
2022-03-29 Tom de Vries <tdevries@suse.de>
PR target/104714
* gcc.target/nvptx/march-map.c: New test.
The target option misa has the following description:
...
$ gcc --target-help 2>&1 | grep misa
-misa= Specify the PTX ISA target architecture to use.
...
The name misa is somewhat poorly chosen. It suggests that for a use
-misa=sm_30, sm_30 is the name of a specific Instruction Set Architecture.
Instead, sm_30 is the name of a specific target architecture in the generic
PTX Instruction Set Architecture.
Futhermore, there's mptx, which also has ISA in the description:
...
-mptx= Specify the PTX ISA version to use.
...
Add the more intuitive alias march for misa:
...
$ gcc --target-help 2>&1 | grep march
-march= Alias: Same as -misa=.
...
Tested on nvptx.
gcc/ChangeLog:
2022-03-29 Tom de Vries <tdevries@suse.de>
* config/nvptx/nvptx.opt (march): Add alias of misa.
gcc/testsuite/ChangeLog:
2022-03-29 Tom de Vries <tdevries@suse.de>
* gcc.target/nvptx/main.c: New test.
* gcc.target/nvptx/march.c: New test.
Currently we have:
...
$ gcc --target-help 2>&1 | egrep "misa|mptx"
-misa= Specify the version of the ptx ISA to use.
-mptx= Specify the version of the ptx version to use.
Known PTX ISA versions (for use with the -misa= option):
Known PTX versions (for use with the -mptx= option):
...
As reported in PR104818, the "version of the ptx version" doesn't make much
sense.
Furthermore, the description of misa (and 'Known ISA versions') is misleading
because it does not specify the version of the PTX ISA, but rather the PTX ISA
target architecture.
Fix this by printing instead:
...
$ gcc --target-help 2>&1 | egrep "misa|mptx"
-misa= Specify the PTX ISA target architecture to use.
-mptx= Specify the PTX ISA version to use.
Known PTX ISA target architectures (for use with the -misa= option):
Known PTX ISA versions (for use with the -mptx= option):
...
Tested on nvptx.
gcc/ChangeLog:
2022-03-28 Tom de Vries <tdevries@suse.de>
PR target/104818
* config/nvptx/gen-opt.sh (ptx_isa): Improve help text.
* config/nvptx/nvptx-gen.opt: Regenerate.
* config/nvptx/nvptx.opt (misa, mptx, ptx_version): Improve help text.
* config/nvptx/t-nvptx (s-nvptx-gen-opt): Add missing dependency on
gen-opt.sh.
Add new option -mexperimental.
This allows, rather than developing a new feature to completion in a
development branch, to develop a new feature on trunk, without disturbing
trunk.
The equivalent of the feature branch merge then becomes making the
functionality available for -mno-experimental.
If more features at the same time will be developed, we can do something like
-mexperimental=feature1,feature2 but for now that's not necessary.
For now, has no effect.
gcc/ChangeLog:
2022-03-19 Tom de Vries <tdevries@suse.de>
* config/nvptx/nvptx.opt (mexperimental): New option.
Starting with ptx isa version 6.3, a ptx directive .alias is available.
Use this directive to support symbol aliases, as far as possible.
The alias support is off by default. It can be turned on using a switch
-malias.
Furthermore, for pre-sm_75, it's not effective unless the ptx version is
bumped to 6.3 or higher using -mptx (given that the default for pre-sm_75 is
6.0).
The alias support has the following limitations.
Only function aliases are supported.
Weak aliases are not supported. That is, if I disable the check in
nvptx_asm_output_def_from_decls that disallows this, a weak alias is emitted
and parsed by the driver. But the test gcc.dg/globalalias.c starts failing,
with the behaviour matching the comment about "weird behavior of AIX's .set
pseudo-op": a weak alias may resolve to different functions in different
files.
Aliases to weak symbols are not supported (see gcc.dg/localalias.c). This is
currently not prohibited by the compiler, but with the driver link we run
into: "error: Function test with .weak scope cannot be aliased".
Aliases to aliases are not supported (see libgomp.c-c++-common/pr96390.c).
This is currently not prohibited by the compiler, but with the driver link we
run into: "Internal error: alias to unknown symbol" .
Unreferenced aliases are not emitted (these can occur f.i. when inlining a
call to an alias). This avoids driver link error "Internal error: reference
to deleted section".
When enabling malias by default, libgomp detects alias support and
consequently libgomp.a will contains a few uses of .alias. This however
results in aforementioned "Internal error: reference to deleted section" in
many test-cases. Either there's some error with how .alias is used, or
there's a driver bug. While this issue is not resolved, we keep malias
off-by-default.
At some point we may add support in the nvptx-tools linker for symbol
aliases, and define f.i. malias=ptx and malias=ld to choose between the two in
the compiler.
An example of where this support is useful, is the OvO (OpenMP vs Offload)
testsuite. The testsuite passes already at -O2. But at -O0, there are errors
in some c++ test-cases due to missing symbol alias support. By compiling with
-malias, the whole testsuite passes also at -O0.
This patch causes a regression:
...
-PASS: gcc.dg/pr60797.c (test for errors, line 4)
+FAIL: gcc.dg/pr60797.c (test for errors, line 4)
...
The test-case is skipped for effective target alias, and both without and with
this patch the nvptx target is considered to not support it, so the test-case is
executed. The test-case expects an error message along the lines of "alias
definitions not supported in this configuration", but instead we run into:
...
gcc.dg/pr60797.c:4:12: error: foo aliased to undefined symbol
...
This is probably due to the fact that the nvptx backend now defines macros
ASM_OUTPUT_DEF and ASM_OUTPUT_DEF_FROM_DECLS, so from the point of view of the
common part of the compiler, aliases are supported.
gcc/ChangeLog:
2022-03-18 Tom de Vries <tdevries@suse.de>
PR target/104957
* config/nvptx/nvptx-protos.h (nvptx_asm_output_def_from_decls): Declare.
* config/nvptx/nvptx.cc (write_fn_proto_1): Don't add function marker
for alias.
(SET_ASM_OP, NVPTX_ASM_OUTPUT_DEF): New macro def.
(nvptx_asm_output_def_from_decls): New function.
* config/nvptx/nvptx.h (ASM_OUTPUT_DEF): New macro def, define to
gcc_unreachable ().
(ASM_OUTPUT_DEF_FROM_DECLS): New macro def, define to
nvptx_asm_output_def_from_decls.
* config/nvptx/nvptx.opt (malias): New opt.
gcc/testsuite/ChangeLog:
2022-03-18 Tom de Vries <tdevries@suse.de>
PR target/104957
* gcc.target/nvptx/alias-1.c: New test.
* gcc.target/nvptx/alias-2.c: New test.
* gcc.target/nvptx/alias-3.c: New test.
* gcc.target/nvptx/alias-4.c: New test.
* gcc.target/nvptx/nvptx.exp
(check_effective_target_runtime_ptx_isa_version_6_3): New proc.
With commit 07667c911b ("[nvptx] Build libraries with misa=sm_30") the
intention was that the sm_xx for all libraries was switched back to sm_30
using MULTILIB_EXTRA_OPTS, without changing the default sm_35.
Testing on an sm_30 board revealed that still some libs were build with sm_35,
so fix this by switching back to default sm_30.
Tested on nvptx.
gcc/ChangeLog:
2022-03-07 Tom de Vries <tdevries@suse.de>
PR target/104758
* config/nvptx/nvptx.opt (misa): Set default to sm_30.
* config/nvptx/t-nvptx (MULTILIB_EXTRA_OPTS): Remove misa=sm_30.
Use nvptx-sm.def to generate new files nvptx-gen.h and nvptx-gen.opt, and:
- include nvptx-gen.h in nvptx.h, and
- add nvptx-gen.opt to extra_options (before nvptx.opt, in case that matters).
Tested on nvptx.
gcc/ChangeLog:
2022-02-25 Tom de Vries <tdevries@suse.de>
* config.gcc (nvptx*-*-*): Add nvptx/nvptx-gen.opt to extra_options.
* config/nvptx/gen-copyright.sh: New file.
* config/nvptx/gen-h.sh: New file.
* config/nvptx/gen-opt.sh: New file.
* config/nvptx/nvptx.h (TARGET_SM35, TARGET_SM53, TARGET_SM70)
(TARGET_SM75, TARGET_SM80): Move ...
* config/nvptx/nvptx-gen.h: ... here. New file, generate.
* config/nvptx/nvptx.opt (Enum ptx_isa): Move ...
* config/nvptx/nvptx-gen.opt: ... here. New file, generate.
* config/nvptx/t-nvptx ($(srcdir)/config/nvptx/nvptx-gen.h)
($(srcdir)/config/nvptx/nvptx-gen.opt): New make target.
Add an -mptx=_ value, that indicates the default ptx version.
It can be used to undo an explicit -mptx setting, so this:
...
$ gcc test.c -mptx=3.1 -mptx=_
...
has the same effect as:
...
$ gcc test.c
...
Tested on nvptx.
gcc/ChangeLog:
2022-02-28 Tom de Vries <tdevries@suse.de>
* config/nvptx/nvptx-opts.h (enum ptx_version): Add
PTX_VERSION_default.
* config/nvptx/nvptx.cc (handle_ptx_version_option): Handle
PTX_VERSION_default.
* config/nvptx/nvptx.opt: Add EnumValue "_" / PTX_VERSION_default.
Add -misa=sm_70, and use it to specify the misa value in test-case
gcc.target/nvptx/atomic-store-2.c.
Tested on nvptx.
gcc/ChangeLog:
* config/nvptx/nvptx-c.cc (nvptx_cpu_cpp_builtins): Handle SM70.
* config/nvptx/nvptx.cc (first_ptx_version_supporting_sm):
Likewise.
* config/nvptx/nvptx.opt (misa): Add sm_70 alias PTX_ISA_SM70.
gcc/testsuite/ChangeLog:
2022-02-22 Tom de Vries <tdevries@suse.de>
* gcc.target/nvptx/atomic-store-2.c: Use -misa=sm_70.
* gcc.target/nvptx/uniform-simt-3.c: Same.
Co-Authored-By: Tom de Vries <tdevries@suse.de>
Currently supported internally are 3.1, 6.0, 6.3 and 7.0.
However, -mptx= supports 3.1, 6.3, 7.0 – but not the internal default 6.0.
Add -mptx=6.0 for consistency.
Tested on nvptx.
gcc/ChangeLog:
* config/nvptx/nvptx.opt (mptx): Add 6.0 alias PTX_VERSION_6_0.
* doc/invoke.texi (-mptx): Update for new values and defaults.
Co-Authored-By: Tom de Vries <tdevries@suse.de>
Add functionality that indicates which insns are added by -minit-regs, such
that for instance we have for pr53465.s:
...
// #APP
// 9 "gcc/testsuite/gcc.c-torture/execute/pr53465.c" 1
// Start: Added by -minit-regs=3:
// #NO_APP
mov.u32 %r26, 0;
// #APP
// 9 "gcc/testsuite/gcc.c-torture/execute/pr53465.c" 1
// End: Added by -minit-regs=3:
// #NO_APP
...
Can be switched off using -mno-ptx-comment.
Tested on nvptx.
gcc/ChangeLog:
2022-02-21 Tom de Vries <tdevries@suse.de>
* config/nvptx/nvptx.cc (gen_comment): New function.
(workaround_uninit_method_1, workaround_uninit_method_2)
(workaround_uninit_method_3): : Use gen_comment.
* config/nvptx/nvptx.opt (mptx-comment): New option.
With nvptx target, driver version 510.47.03 and board GT 1030 I, we run into:
...
FAIL: gcc.c-torture/execute/pr53465.c -O1 execution test
FAIL: gcc.c-torture/execute/pr53465.c -O2 execution test
FAIL: gcc.c-torture/execute/pr53465.c -O3 -g execution test
...
while the test-cases pass with nvptx-none-run -O0.
The problem is that the generated ptx contains a read from an uninitialized
ptx register, and the driver JIT doesn't handle this well.
For -O2 and -O3, we can get rid of the FAIL using --param
logical-op-non-short-circuit=0. But not for -O1.
At -O1, the test-case minimizes to:
...
void __attribute__((noinline, noclone))
foo (int y) {
int c;
for (int i = 0; i < y; i++)
{
int d = i + 1;
if (i && d <= c)
__builtin_abort ();
c = d;
}
}
int main () {
foo (2); return 0;
}
...
Note that the test-case does not contain an uninitialized use. In the first
iteration, i is 0 and consequently c is not read. In the second iteration, c
is read, but by that time it's already initialized by 'c = d' from the first
iteration.
AFAICT the problem is introduced as follows: the conditional use of c in the
loop body is translated into an unconditional use of c in the loop header:
...
# c_1 = PHI <c_4(D)(2), c_9(6)>
...
which forwprop1 propagates the 'c_9 = d_7' assignment into:
...
# c_1 = PHI <c_4(D)(2), d_7(6)>
...
which ends up being translated by expand into an unconditional:
...
(insn 13 12 0 (set (reg/v:SI 22 [ c ])
(reg/v:SI 23 [ d ])) -1
(nil))
...
at the start of the loop body, creating an uninitialized read of d on the
path from loop entry.
By disabling coalesce_ssa_name, we get the more usual copies on the incoming
edges. The copy on the loop entry path still does an uninitialized read, but
that one's now initialized by init-regs. The test-case passes, also when
disabling init-regs, so it's possible that the JIT driver doesn't object to
this type of uninitialized read.
Now that we characterized the problem to some degree, we need to fix this,
because either:
- we're violating an undocumented ptx invariant, and this is a compiler bug,
or
- this is is a driver JIT bug and we need to work around it.
There are essentially two strategies to address this:
- stop the compiler from creating uninitialized reads
- patch up uninitialized reads using additional initialization
The former will probably involve:
- making some optimizations more conservative in the presence of
uninitialized reads, and
- disabling some other optimizations (where making them more conservative is
not possible, or cannot easily be achieved).
This will probably will have a cost penalty for code that does not suffer from
the original problem.
The latter has the problem that it may paper over uninitialized reads
in the source code, or indeed over ones that were incorrectly introduced
by the compiler. But it has the advantage that it allows for the problem to
be addressed at a single location.
There's an existing pass, init-regs, which implements a form of the latter,
but it doesn't work for this example because it only inserts additional
initialization for uses that have not a single reaching definition.
Fix this by adding initialization of uninitialized ptx regs in reorg.
Control the new functionality using -minit-regs=<0|1|2|3>, meaning:
- 0: disabled.
- 1: add initialization of all regs at the entry bb
- 2: add initialization of uninitialized regs at the entry bb
- 3: add initialization of uninitialized regs close to the use
and defaulting to 3.
Tested on nvptx.
gcc/ChangeLog:
2022-02-17 Tom de Vries <tdevries@suse.de>
PR target/104440
* config/nvptx/nvptx.cc (workaround_uninit_method_1)
(workaround_uninit_method_2, workaround_uninit_method_3)
(workaround_uninit): New function.
(nvptx_reorg): Use workaround_uninit.
* config/nvptx/nvptx.opt (minit-regs): New option.
While testing with driver version 390.147 I ran into the problem that it
doesn't support ptx isa version 6.3 (the new default), only 6.1.
Furthermore, using the -mptx option is a bit user-unfriendly.
Say we want to compile for sm_80. We can use -misa=sm_80 to specify that, but
then run into errors because the default ptx version is 6.3, which doesn't
support sm_80 yet.
Address both these issues by:
- picking a default -mptx based on the active -misa, and
- ensuring that the default -mptx is at least 6.0 (instead
of 6.3).
Also add an error in case of incompatible options like
"-misa=sm_80 -mptx=6.3":
...
cc1: error: PTX version (-mptx) needs to be at least 7.0 to support \
selected -misa (sm_80)
...
Tested on x86_64-linux with nvptx accelerator.
gcc/ChangeLog:
2022-02-08 Tom de Vries <tdevries@suse.de>
PR target/104283
* config/nvptx/nvptx-opts.h (enum ptx_version): Add PTX_VERSION_3_0
and PTX_VERSION_4_2.
* config/nvptx/nvptx.cc (first_ptx_version_supporting_sm)
(default_ptx_version_option, ptx_version_to_string)
(sm_version_to_string, handle_ptx_version_option): New function.
(nvptx_option_override): Call handle_ptx_version_option.
(nvptx_file_start): Use ptx_version_to_string and sm_version_to_string.
* config/nvptx/nvptx.md (define_insn "nvptx_shuffle<mode>")
(define_insn "nvptx_vote_ballot"): Use TARGET_PTX_6_0.
* config/nvptx/nvptx.opt (mptx): Remove 'Init'.
With the following example, minimized from parallel-dims.c:
...
int
main (void)
{
int vectors_max = -1;
#pragma acc parallel num_gangs (1) num_workers (1) copy (vectors_max)
{
for (int i = 0; i < 2; i++)
for (int j = 0; j < 2; j++)
#pragma acc loop vector reduction (max: vectors_max)
for (int k = 0; k < 32; k++)
vectors_max = k;
}
if (vectors_max != 31)
__builtin_abort ();
return 0;
}
...
I run into (T400, driver version 470.94):
...
FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims.c \
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 \
execution test
...
The FAIL does not happen with GOMP_NVPTX_JIT=-O0.
The problem seems to be that the shfl insns for the vector reduction are not
executed uniformly by the warp. Enforcing this by using shfl.sync fixes the
problem.
Fix this by setting the ptx isa to 6.3 by default, which allows the use of
shfl.sync.
Tested on x86_64 with nvptx accelerator.
gcc/ChangeLog:
2022-01-27 Tom de Vries <tdevries@suse.de>
* config/nvptx/nvptx.opt (mptx): Set to PTX_VERSION_6_3 by default.
Add new target macros TARGET_SM75 and TARGET_SM80. Add support for
__builtin_tanhf, HFmode exp2/tanh and also for HFmode min/max, controlled by
TARGET_SM75 and TARGET_SM80 respectively.
The following has been tested on nvptx-none, hosted on x86_64-pc-linux-gnu
with a "make" and "make -k check" with no new failures.
gcc/ChangeLog:
* config/nvptx/nvptx-opts.h (ptx_isa): PTX_ISA_SM75 and PTX_ISA_SM80
ISA levels.
* config/nvptx/nvptx.opt: Add sm_75 and sm_80 to -misa.
* config/nvptx/nvptx.h (TARGET_SM75, TARGET_SM80):
New helper macros to conditionalize functionality on target ISA.
* config/nvptx/nvptx-c.c (nvptx_cpu_cpp_builtins): Add __PTX_SM__
support for the new ISA levels.
* config/nvptx/nvptx.c (nvptx_file_start): Add support for TARGET_SM75
and TARGET_SM80.
* config/nvptx/nvptx.md (define_c_enum "unspec"): New UNSPEC_TANH.
(define_mode_iterator HSFM): New iterator for HFmode and SFmode.
(exp2hf2): New define_insn controlled by TARGET_SM75.
(tanh<mode>2): New define_insn controlled by TARGET_SM75.
(sminhf3, smaxhf3): New define_isnns controlled by TARGET_SM80.
gcc/testsuite/ChangeLog:
* gcc.target/nvptx/float16-2.c: New test case.
* gcc.target/nvptx/tanh-1.c: New test case.
Add support for ptx isa version 7.0, required for the addition of -misa=sm_75
and -misa=sm_80.
Tested by setting the default ptx isa version to 7.0, and doing a build and
libgomp test run.
gcc/ChangeLog:
* config/nvptx/nvptx-opts.h (enum ptx_version): Add PTX_VERSION_7_0.
* config/nvptx/nvptx.c (nvptx_file_start): Handle TARGET_PTX_7_0.
* config/nvptx/nvptx.h (TARGET_PTX_7_0): New macro.
* config/nvptx/nvptx.opt (ptx_version): Add 7.0.
The recent flurry of activity around HFmode on gcc-patches intrigued me
to investigate adding HFmode support to the nvptx backend. NVidia GPUs
with an SM ISA above 5.3 support IEEE 16-bit floating point instructions.
Hence, this patch adds support for -misa=sm_53, and implements some
backend patterns/insns sufficient for a proof-of-concept prototype.
The following has been tested on nvptx-none, hosted on x86_64-pc-linux-gnu
with a "make" and "make -k check" with no new failures.
gcc/ChangeLog:
* config/nvptx/nvptx-opts.h (ptx_isa): Add PTX_ISA_SM53 ISA level
to enumeration.
* config/nvptx/nvptx.opt: Add sm_53 to -misa.
* config/nvptx/nvptx-modes.def: Add support for HFmode.
* config/nvptx/nvptx.h (TARGET_SM53):
New helper macro to conditionalize functionality on target ISA.
* config/nvptx/nvptx-c.c (nvptx_cpu_cpp_builtins): Add __PTX_SM__
support for the new ISA levels.
* config/nvptx/nvptx.c (nvtx_ptx_type_from_mode): Support new HFmode
with the ".f16" suffix/qualifier.
(nvptx_file_start): Add support for TARGET_SM53.
(nvptx_omp_device_kind_arch_isa): Add support for TARGET_SM53
and tweak TARGET_SM35.
(nvptx_scalar_mode_supported_p): Target hook with conditional
HFmode support on TARGET_SM53 and higher.
(nvptx_libgcc_floating_mode_supported_p): Likewise.
(TARGET_SCALAR_MODE_SUPPORTED_P): Use nvptx_scalar_mode_supported_p.
(TARGET_LIBGCC_FLOATING_MODE_SUPPORTED_P): Likewise, use new hook.
* config/nvptx/nvptx.md (*movhf_insn): New define_insn.
(movhf): New define_expand for HFmode moves.
(addhf3, subhf3, mulhf, extendhf<mode>2, trunc<mode>hf2): New
instructions conditional on TARGET_SM53 (i.e. -misa=sm_53).
gcc/testsuite/ChangeLog:
* gcc.target/nvptx/float16-1.c: New test case.
Add nvptx option -mptx that sets the ptx ISA version. This is currently
hardcoded to 3.1.
Tested libgomp on x86_64-linux with nvptx accelerator, both with default set to
3.1 and 6.3.
gcc/ChangeLog:
2021-05-12 Tom de Vries <tdevries@suse.de>
PR target/96005
* config/nvptx/nvptx-opts.h (enum ptx_version): New enum.
* config/nvptx/nvptx.c (nvptx_file_start): Print .version according
to ptx_version_option.
* config/nvptx/nvptx.h (TARGET_PTX_6_3): Define.
* config/nvptx/nvptx.md (define_insn "nvptx_shuffle<mode>")
(define_insn "nvptx_vote_ballot"): Use sync variant for
TARGET_PTX_6_3.
* config/nvptx/nvptx.opt (ptx_version): Add enum.
(mptx): Add option.
* doc/invoke.texi (Nvidia PTX Options): Add mptx item.
The nvptx port has an -m32 option, but it's not clear whether this
was ever build/tested/used.
Don't expose to user anymore. Tested on nvptx.
gcc/ChangeLog:
2020-10-15 Tom de Vries <tdevries@suse.de>
PR target/97436
* config/nvptx/nvptx.opt (m32): Comment out.
* doc/invoke.texi (NVPTX options): Remove -m32.
gcc/testsuite/ChangeLog:
2020-10-15 Tom de Vries <tdevries@suse.de>
PR target/97436
* gcc.target/nvptx/atomic_fetch-3.c: Remove.
Currently, in order to use the switch -msoft-stack-reserve-local with the
default arg 128, you have to specify '-msoft-stack-reserve-local128'.
Fix the switch format such that you specify '-msoft-stack-reserve-local=128'
instead.
Tested on nvptx.
gcc/ChangeLog:
2020-10-12 Tom de Vries <tdevries@suse.de>
* config/nvptx/nvptx.opt (-msoft-stack-reserve-local): Rename to ...
(-msoft-stack-reserve-local=): ... this.
The nvptx-as assembler verifies the ptx code using ptxas, if there's any
in the PATH.
The default in the nvptx port for -misa=sm_xx is sm_30, but the ptxas of the
latest cuda release (11.1) no longer supports sm_30.
Consequently we cannot build gcc against that release (although we should
still be able to build without any cuda release).
Fix this by setting -misa=sm_35 by default.
Tested check-gcc on nvptx.
Tested libgomp on x86_64-linux with nvpx accelerator.
Both build again cuda 9.1.
gcc/ChangeLog:
2020-10-09 Tom de Vries <tdevries@suse.de>
PR target/97348
* config/nvptx/nvptx.h (ASM_SPEC): Also pass -m to nvptx-as if
default is used.
* config/nvptx/nvptx.opt (misa): Init with PTX_ISA_SM35.
Thomas Schwinge
Tom de Vries
Tobias Burnus
Jakub Jelinek
Roger Sayle
Martin Liska
Jakub Jelinek
Cesar Philippidis
Alexander Monakov
Nathan Sidwell
Martin Sebor
Bernd Schmidt