https://gcc.gnu.org/g:0cefec2e1f9348b80d7e2c54de517523a0968da2
commit 0cefec2e1f9348b80d7e2c54de517523a0968da2 Author: Andrew Stubbs <[email protected]> Date: Thu Oct 29 18:17:57 2026 +0000 amdgcn: Align "omp simd array" variables, and cap vector alignment The vectorizer accesses "omp simd array" variables (such as the per-lane partial results of SIMD reductions) as whole vectors, typically loading and storing them on every loop iteration, but their type is an array of scalars, so they only get the alignment of the element type. Experiments with BabelStream's "dot" kernel, whose accumulator is such an array, show that wide vector accesses to the stack are significantly faster when the address is aligned to 256 bytes than when it is aligned to only 64 or 128 bytes (the exact cause is unknown), but that larger alignments give no further benefit. This patch implements LOCAL_DECL_ALIGNMENT to over-align these arrays to the largest power of two that divides their size, up to 256 bytes, or not at all when that would be less than 128 bytes. The stack itself remains 64-bit aligned, so the extra alignment is achieved at run time, and costs stack space only in functions that have such arrays. On an MI210 (gfx90a), the double-precision dot bandwidth improves from about 565 GB/s to 630 GB/s, and the single-precision from about 438 GB/s to 460 GB/s. Some of the single-precision improvement comes from the run-time aligned address being calculated once, outside the loop, rather than on every iteration. The other kernels are unaffected. The patch also implements TARGET_VECTOR_ALIGNMENT to cap the alignment of vector types at 256 bytes; the default is their natural alignment, which can be up to 1024 bytes, and wastes stack space for no benefit. Assisted-By: Claude Opus 5.5 Claude did some of the benchmark analysis and generated the testcases, as well as some prototype patches my final patches are derived from. Backport from mainline: gcc/ChangeLog: * config/gcn/gcn-protos.h (gcn_local_decl_alignment): New prototype. * config/gcn/gcn.cc (GCN_VECTOR_STACK_ALIGNMENT): New define. (gcn_vector_alignment): New function. (gcn_local_decl_alignment): New function. (TARGET_VECTOR_ALIGNMENT): Define. * config/gcn/gcn.h (LOCAL_DECL_ALIGNMENT): Define. gcc/testsuite/ChangeLog: * gcc.target/gcn/simd-array-align-1.c: New test. * gcc.target/gcn/vector-align-1.c: New test. Diff: --- gcc/config/gcn/gcn-protos.h | 1 + gcc/config/gcn/gcn.cc | 54 ++++++++++++++++++++++- gcc/config/gcn/gcn.h | 1 + gcc/testsuite/gcc.target/gcn/simd-array-align-1.c | 28 ++++++++++++ gcc/testsuite/gcc.target/gcn/vector-align-1.c | 17 +++++++ 5 files changed, 100 insertions(+), 1 deletion(-) diff --git a/gcc/config/gcn/gcn-protos.h b/gcc/config/gcn/gcn-protos.h index 958fd935b911..e03dc337648e 100644 --- a/gcc/config/gcn/gcn-protos.h +++ b/gcc/config/gcn/gcn-protos.h @@ -66,6 +66,7 @@ extern void gcn_hsa_declare_function_name (FILE *file, const char *name, extern HOST_WIDE_INT gcn_initial_elimination_offset (int, int); extern REAL_VALUE_TYPE gcn_dconst1over2pi (void); extern bool gcn_inline_constant64_p (rtx, bool); +extern unsigned int gcn_local_decl_alignment (tree); extern bool gcn_inline_constant_p (rtx); extern int gcn_inline_fp_constant_p (rtx, bool); extern rtx gcn_oacc_dim_pos (int dim); diff --git a/gcc/config/gcn/gcn.cc b/gcc/config/gcn/gcn.cc index 92d7b939cfdd..9c305398c782 100644 --- a/gcc/config/gcn/gcn.cc +++ b/gcc/config/gcn/gcn.cc @@ -5666,13 +5666,63 @@ gcn_related_vector_mode (machine_mode vector_mode, return VnMODE (n, element_mode); } +/* Known GCN devices have 64-byte cache lines, but experimentation shows that + 256-byte alignments are beneficial for some other part of the memory + system. */ + +#define GCN_VECTOR_STACK_ALIGNMENT (256 * BITS_PER_UNIT) + +/* Implement TARGET_VECTOR_ALIGNMENT. + + Vector types default to their natural alignment, which is up to 1024 + bytes for V64TImode. Cap it at GCN_VECTOR_STACK_ALIGNMENT. */ + +static HOST_WIDE_INT +gcn_vector_alignment (const_tree type) +{ + return MIN (default_vector_alignment (type), GCN_VECTOR_STACK_ALIGNMENT); +} + +/* Implement LOCAL_DECL_ALIGNMENT. + + For most cases, the default alignment is probably fine, but "omp simd array" + variables are performance critical so this hook makes them a special case. + In particular, these are used for the per-lane accumulators for SIMD + reductions. Aligning these arrays as if they were vectors means that + they're optimized for when they're are accessed as vectors. + + FIXME can/should this be applied to all arrays? */ + +unsigned int +gcn_local_decl_alignment (tree decl) +{ + unsigned int align = LOCAL_ALIGNMENT (TREE_TYPE (decl), DECL_ALIGN (decl)); + + if (VAR_P (decl) + && lookup_attribute ("omp simd array", DECL_ATTRIBUTES (decl)) + && DECL_SIZE (decl) + && tree_fits_uhwi_p (DECL_SIZE (decl))) + { + unsigned HOST_WIDE_INT size = tree_to_uhwi (DECL_SIZE (decl)); + unsigned HOST_WIDE_INT want = MIN (least_bit_hwi (size), + GCN_VECTOR_STACK_ALIGNMENT); + if (want >= 128 * BITS_PER_UNIT && want > align) + align = want; + } + + return align; +} + /* Implement TARGET_VECTORIZE_PREFERRED_VECTOR_ALIGNMENT. Returns the preferred alignment in bits for accesses to vectors of type type in vectorized code. This might be less than or greater than the ABI-defined value returned by TARGET_VECTOR_ALIGNMENT. It can be equal to the alignment of a single element, in which case the vectorizer will not try to optimize - for alignment. */ + for alignment. + + Note: Changing this to GCN_VECTOR_STACK_ALIGNMENT did not give a measurable + improvement running BabelStream. FIXME Might other benchmarks benefit? */ static poly_uint64 gcn_preferred_vector_alignment (const_tree type) @@ -8753,6 +8803,8 @@ gcn_output_ensure_scalar (rtx *operands, machine_mode mode) #undef TARGET_VECTORIZE_VECTOR_ALIGNMENT_REACHABLE #define TARGET_VECTORIZE_VECTOR_ALIGNMENT_REACHABLE \ gcn_vector_alignment_reachable +#undef TARGET_VECTOR_ALIGNMENT +#define TARGET_VECTOR_ALIGNMENT gcn_vector_alignment #undef TARGET_VECTOR_MODE_SUPPORTED_P #define TARGET_VECTOR_MODE_SUPPORTED_P gcn_vector_mode_supported_p diff --git a/gcc/config/gcn/gcn.h b/gcc/config/gcn/gcn.h index 581e77ed04c7..a758280a8403 100644 --- a/gcc/config/gcn/gcn.h +++ b/gcc/config/gcn/gcn.h @@ -122,6 +122,7 @@ extern const struct gcn_device_def { #define DATA_ALIGNMENT(TYPE,ALIGN) ((ALIGN) > 128 ? (ALIGN) : 128) #define LOCAL_ALIGNMENT(TYPE,ALIGN) ((ALIGN) > 64 ? (ALIGN) : 64) +#define LOCAL_DECL_ALIGNMENT(DECL) gcn_local_decl_alignment (DECL) #define STACK_SLOT_ALIGNMENT(TYPE,MODE,ALIGN) ((ALIGN) > 64 ? (ALIGN) : 64) #define STRICT_ALIGNMENT 1 diff --git a/gcc/testsuite/gcc.target/gcn/simd-array-align-1.c b/gcc/testsuite/gcc.target/gcn/simd-array-align-1.c new file mode 100644 index 000000000000..994323484970 --- /dev/null +++ b/gcc/testsuite/gcc.target/gcn/simd-array-align-1.c @@ -0,0 +1,28 @@ +/* { dg-do compile } */ +/* { dg-options "-O2 -fopenmp-simd -ftree-vectorize" } */ + +/* Check that the "omp simd array" variables holding the per-lane partial + sums are aligned at run time, to 256 bytes (or 128 bytes for a 32-lane + float array). */ + +double +f (double *a, double *b, int n) +{ + double sum = 0; +#pragma omp simd reduction(+:sum) + for (int i = 0; i < n; i++) + sum += a[i] * b[i]; + return sum; +} + +float +g (float *a, float *b, int n) +{ + float sum = 0; +#pragma omp simd reduction(+:sum) + for (int i = 0; i < n; i++) + sum += a[i] * b[i]; + return sum; +} + +/* { dg-final { scan-assembler-times {_and_b32\t[^\n]*, -(128|256), } 2 } } */ diff --git a/gcc/testsuite/gcc.target/gcn/vector-align-1.c b/gcc/testsuite/gcc.target/gcn/vector-align-1.c new file mode 100644 index 000000000000..d766bbec3793 --- /dev/null +++ b/gcc/testsuite/gcc.target/gcn/vector-align-1.c @@ -0,0 +1,17 @@ +/* { dg-do compile } */ +/* { dg-options "-O2" } */ + +/* Check that vector type alignment is capped at 256 bytes. */ + +typedef int v32si __attribute__ ((vector_size (128))); +typedef double v32df __attribute__ ((vector_size (256))); +typedef double v64df __attribute__ ((vector_size (512))); +typedef double v128df __attribute__ ((vector_size (1024))); + +v32si a; +v32df b; +v64df c; +v128df d; + +/* { dg-final { scan-assembler-times "\\.align\t128\n" 1 } } */ +/* { dg-final { scan-assembler-times "\\.align\t256\n" 3 } } */
