[nvptx] Update bar.sync for ptx isa 6.0

[nvptx] Update bar.sync for ptx isa 6.0

In ptx isa 6.0, a new barrier instruction was added, and bar.sync was
redefined as barrier.sync.aligned.

The aligned modifier indicates that all threads in a CTA will execute the same
barrier instruction.

The seems fine for a form "bar.sync 0".

But a "bar.sync %rx,64" (as used for vector length > 32) may execute a
diffferent barrier depending on the value of %rx, so we can't assume it's
aligned.

Fix this by using "barrier.sync %rx,64" instead.

Tested on x86_64 with nvptx accelerator.

gcc/ChangeLog:

2022-01-27  Tom de Vries  <tdevries@suse.de>

	* config/nvptx/nvptx-opts.h (enum ptx_version): Add PTX_VERSION_6_0.
	* config/nvptx/nvptx.h (TARGET_PTX_6_0): New macro.
	* config/nvptx/nvptx.md (define_insn "nvptx_barsync"): Use barrier
	insn for TARGET_PTX_6_0.