diff options
| author | Tom de Vries <tdevries@suse.de> | 2022-01-26 14:16:42 +0100 |
|---|---|---|
| committer | Tom de Vries <tdevries@suse.de> | 2022-02-01 19:28:48 +0100 |
| commit | 57f971f99209cc950d7e706b7b52f4c9ef1d10b0 (patch) | |
| tree | 18dbffcb88d4a184356cf9e98e66f0a0844fc353 /gcc | |
| parent | 456de10c549379b74d4858f00d4b8817035a73fc (diff) | |
| download | gcc-57f971f99209cc950d7e706b7b52f4c9ef1d10b0.zip gcc-57f971f99209cc950d7e706b7b52f4c9ef1d10b0.tar.gz gcc-57f971f99209cc950d7e706b7b52f4c9ef1d10b0.tar.bz2 | |
[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.
Diffstat (limited to 'gcc')
| -rw-r--r-- | gcc/config/nvptx/nvptx-opts.h | 1 | ||||
| -rw-r--r-- | gcc/config/nvptx/nvptx.h | 1 | ||||
| -rw-r--r-- | gcc/config/nvptx/nvptx.md | 8 |
3 files changed, 8 insertions, 2 deletions
diff --git a/gcc/config/nvptx/nvptx-opts.h b/gcc/config/nvptx/nvptx-opts.h index daae72f..c754a51 100644 --- a/gcc/config/nvptx/nvptx-opts.h +++ b/gcc/config/nvptx/nvptx-opts.h @@ -32,6 +32,7 @@ enum ptx_isa enum ptx_version { PTX_VERSION_3_1, + PTX_VERSION_6_0, PTX_VERSION_6_3, PTX_VERSION_7_0 }; diff --git a/gcc/config/nvptx/nvptx.h b/gcc/config/nvptx/nvptx.h index 9fda2f0..065d7aa 100644 --- a/gcc/config/nvptx/nvptx.h +++ b/gcc/config/nvptx/nvptx.h @@ -91,6 +91,7 @@ #define TARGET_SM75 (ptx_isa_option >= PTX_ISA_SM75) #define TARGET_SM80 (ptx_isa_option >= PTX_ISA_SM80) +#define TARGET_PTX_6_0 (ptx_version_option >= PTX_VERSION_6_0) #define TARGET_PTX_6_3 (ptx_version_option >= PTX_VERSION_6_3) #define TARGET_PTX_7_0 (ptx_version_option >= PTX_VERSION_7_0) diff --git a/gcc/config/nvptx/nvptx.md b/gcc/config/nvptx/nvptx.md index 9cbbd95..b391165 100644 --- a/gcc/config/nvptx/nvptx.md +++ b/gcc/config/nvptx/nvptx.md @@ -1968,9 +1968,13 @@ "" { if (INTVAL (operands[1]) == 0) - return "\\tbar.sync\\t%0;"; + return (TARGET_PTX_6_0 + ? "\\tbarrier.sync.aligned\\t%0;" + : "\\tbar.sync\\t%0;"); else - return "\\tbar.sync\\t%0, %1;"; + return (TARGET_PTX_6_0 + ? "\\tbarrier.sync\\t%0, %1;" + : "\\tbar.sync\\t%0, %1;"); } [(set_attr "predicable" "false")]) |