https://gcc.gnu.org/g:74688b45014c5197ddbaa4b0bf88918d54917845
commit 74688b45014c5197ddbaa4b0bf88918d54917845 Author: Andrew Stubbs <[email protected]> Date: Tue Oct 13 15:10:07 2026 +0000 amdgcn: Enable SIMT for OpenACC Implement the basic SIMT fork/join operations, and the vector reductions. This is enough to run BabelStream, but the additional features are missing, for now. The unimplemented omp_simt_vote_any and omp_simt_xchg_idx patterns now output "sorry, unimplemented" so that any use fails visibly at assembly time, rather than silently emitting a comment. gcc/ChangeLog: * config/gcn/gcn-tree.cc (struct oacc_attrs): New. (gcn_goacc_reduction_init): Add SIMT vector reduction case. (gcn_goacc_reduction_fini): Likewise, using IFN_GOMP_SIMT_REDUCE. (gcn_goacc_reduction_teardown): Read the worker reduction buffer for SIMT vector reductions where the vector length is not 64. (gcn_goacc_reduction): Don't use default_goacc_reduction for vector reductions. * config/gcn/gcn.cc (gcn_oacc_dim_size): Return 64 for the vector dimension in SIMT mode. * config/gcn/gcn.md (UNSPECV_SIMT_FORK, UNSPECV_SIMT_JOIN): Delete. (oacc_fork): Emit omp_simt_enter for the vector dimension in SIMT mode. (oacc_join): Emit omp_simt_exit for the vector dimension in SIMT mode. (*oacc_fork_insn, *oacc_join_insn): Delete. (omp_simt_exit): Remove predicate and constraint from operand 0. (omp_simt_vote_any, omp_simt_xchg_idx): Output "sorry, unimplemented". libgomp/ChangeLog: * testsuite/libgomp.c-c++-common/simt-reduction-amdgcn-1.c: New test. * testsuite/libgomp.c-c++-common/simt-reduction-amdgcn-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/simt-reduction-amdgcn-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/simt-reduction-amdgcn-2.c: New test. Diff: --- gcc/config/gcn/gcn-tree.cc | 199 ++++++++++++++++----- gcc/config/gcn/gcn.cc | 5 +- gcc/config/gcn/gcn.md | 76 ++++---- .../libgomp.c-c++-common/simt-reduction-amdgcn-1.c | 87 +++++++++ .../libgomp.c-c++-common/simt-reduction-amdgcn-2.c | 107 +++++++++++ .../simt-reduction-amdgcn-1.c | 84 +++++++++ .../simt-reduction-amdgcn-2.c | 100 +++++++++++ 7 files changed, 569 insertions(+), 89 deletions(-) diff --git a/gcc/config/gcn/gcn-tree.cc b/gcc/config/gcn/gcn-tree.cc index d31b0230b3b3..54044b8df081 100644 --- a/gcc/config/gcn/gcn-tree.cc +++ b/gcc/config/gcn/gcn-tree.cc @@ -443,6 +443,50 @@ gcn_reduction_update (location_t loc, gimple_stmt_iterator *gsi, return gcn_lockfull_update (loc, gsi, ptr, var, op); } +/* OpenACC offloading function attributes. */ + +struct oacc_attrs +{ + oacc_attrs () + { + tree attr = oacc_get_fn_attrib (current_function_decl); + tree dims = TREE_VALUE (attr); + unsigned ix; + + this->mask = 0; + + for (ix = 0; ix != GOMP_DIM_MAX; ix++, dims = TREE_CHAIN (dims)) + { + tree t = TREE_VALUE (dims); + int size = (t == NULL_TREE) ? -1 : TREE_INT_CST_LOW (t); + tree allowed = TREE_PURPOSE (dims); + + if (size != 1 && !(allowed && integer_zerop (allowed))) + this->mask |= GOMP_DIM_MASK (ix); + + switch (ix) + { + case GOMP_DIM_GANG: + this->num_gangs = size; + break; + + case GOMP_DIM_WORKER: + this->num_workers = size; + break; + + case GOMP_DIM_VECTOR: + this->vector_length = size; + break; + } + } + } + + unsigned mask; + int num_gangs; + int num_workers; + int vector_length; +}; + /* Return a temporary variable decl to use for an OpenACC worker reduction. */ static tree @@ -837,16 +881,71 @@ gcn_goacc_reduction_init (gcall *call) init = omp_reduction_init_op (gimple_location (call), rcode, TREE_TYPE (var)); - if (level == GOMP_DIM_GANG) + if (level == GOMP_DIM_VECTOR + && flag_simt_mode + && oacc_attrs().vector_length == 64) { - /* If there's no receiver object, propagate the incoming VAR. */ - tree ref_to_res = gimple_call_arg (call, 1); - if (integer_zerop (ref_to_res)) - init = var; + /* Initialize vector-non-zeroes to INIT_VAL (OP). */ + tree tid = make_ssa_name (integer_type_node); + tree dim_vector = gimple_call_arg (call, 3); + gimple *tid_call = gimple_build_call_internal (IFN_GOACC_DIM_POS, 1, + dim_vector); + gimple *cond_stmt = gimple_build_cond (NE_EXPR, tid, integer_zero_node, + NULL_TREE, NULL_TREE); + + gimple_call_set_lhs (tid_call, tid); + gimple_seq_add_stmt (&seq, tid_call); + gimple_seq_add_stmt (&seq, cond_stmt); + + /* Split the block just after the call. */ + edge init_edge = split_block (gsi_bb (gsi), call); + basic_block init_bb = init_edge->dest; + basic_block call_bb = init_edge->src; + + /* Fixup flags from call_bb to init_bb. */ + init_edge->flags ^= EDGE_FALLTHRU | EDGE_TRUE_VALUE; + init_edge->probability = profile_probability::even (); + + /* Set the initialization stmts. */ + gimple_seq init_seq = NULL; + tree init_var = make_ssa_name (TREE_TYPE (var)); + gimplify_assign (init_var, init, &init_seq); + gsi = gsi_start_bb (init_bb); + gsi_insert_seq_before (&gsi, init_seq, GSI_SAME_STMT); + + /* Split block just after the init stmt. */ + gsi_prev (&gsi); + edge inited_edge = split_block (gsi_bb (gsi), gsi_stmt (gsi)); + basic_block dst_bb = inited_edge->dest; + + /* Create false edge from call_bb to dst_bb. */ + edge nop_edge = make_edge (call_bb, dst_bb, EDGE_FALSE_VALUE); + nop_edge->probability = profile_probability::even (); + + /* Create phi node in dst block. */ + gphi *phi = create_phi_node (lhs, dst_bb); + add_phi_arg (phi, init_var, inited_edge, gimple_location (call)); + add_phi_arg (phi, var, nop_edge, gimple_location (call)); + + /* Reset dominator of dst bb. */ + set_immediate_dominator (CDI_DOMINATORS, dst_bb, call_bb); + + /* Reset the gsi. */ + gsi = gsi_for_stmt (call); } + else + { + if (level == GOMP_DIM_GANG) + { + /* If there's no receiver object, propagate the incoming VAR. */ + tree ref_to_res = gimple_call_arg (call, 1); + if (integer_zerop (ref_to_res)) + init = var; + } - if (lhs) - gimplify_assign (lhs, init, &seq); + if (lhs != NULL_TREE) + gimplify_assign (lhs, init, &seq); + } pop_gimplify_context (NULL); @@ -876,43 +975,58 @@ gcn_goacc_reduction_fini (gcall *call) push_gimplify_context (true); - tree accum = NULL_TREE; - - if (level == GOMP_DIM_WORKER) + if (level == GOMP_DIM_VECTOR + && flag_simt_mode + && oacc_attrs().vector_length == 64) { - tree offset = gimple_call_arg (call, 5); - tree decl; - if (array_p) - { - tree array_type = TREE_TYPE (TREE_TYPE (array_addr)); - decl = gcn_goacc_get_worker_array_reduction_buffer - (array_type, array_max_idx, &seq); - } - else - { - tree var_type = TREE_TYPE (var); - decl = gcn_goacc_get_worker_red_decl (var_type, offset); - } - accum = build_fold_addr_expr (decl); + /* Implement SIMT reductions using the OpenMP code-path. */ + tree opcode = build_int_cst (unsigned_type_node, op); + gimple *call = gimple_build_call_internal (IFN_GOMP_SIMT_REDUCE, 2, var, + opcode); + r = make_ssa_name (TREE_TYPE (var)); + gimple_call_set_lhs (call, r); + gimple_seq_add_stmt (&seq, call); } - else if (integer_zerop (ref_to_res)) - r = var; else - accum = ref_to_res; - - if (accum) { - /* UPDATE the accumulator. */ - gsi_insert_seq_before (&gsi, seq, GSI_SAME_STMT); - seq = NULL; - if (array_p) + tree accum = NULL_TREE; + + if (level == GOMP_DIM_WORKER) { - gcn_reduction_update (gimple_location (call), &gsi, accum, - array_addr, op, array_max_idx); - r = var; + tree offset = gimple_call_arg (call, 5); + tree decl; + if (array_p) + { + tree array_type = TREE_TYPE (TREE_TYPE (array_addr)); + decl = gcn_goacc_get_worker_array_reduction_buffer + (array_type, array_max_idx, &seq); + } + else + { + tree var_type = TREE_TYPE (var); + decl = gcn_goacc_get_worker_red_decl (var_type, offset); + } + accum = build_fold_addr_expr (decl); } + else if (integer_zerop (ref_to_res)) + r = var; else - r = gcn_reduction_update (gimple_location (call), &gsi, accum, var, op); + accum = ref_to_res; + + if (accum) + { + /* UPDATE the accumulator. */ + gsi_insert_seq_before (&gsi, seq, GSI_SAME_STMT); + seq = NULL; + if (array_p) + { + gcn_reduction_update (gimple_location (call), &gsi, accum, + array_addr, op, array_max_idx); + r = var; + } + else + r = gcn_reduction_update (gimple_location (call), &gsi, accum, var, op); + } } if (lhs) @@ -943,7 +1057,10 @@ gcn_goacc_reduction_teardown (gcall *call) push_gimplify_context (true); - if (level == GOMP_DIM_WORKER) + if (level == GOMP_DIM_WORKER + || (level == GOMP_DIM_VECTOR + && flag_simt_mode + && oacc_attrs().vector_length != 64)) { tree offset = gimple_call_arg (call, 5); if (array_p) @@ -1037,12 +1154,6 @@ gcn_goacc_reduction (gcall *call) { int level = TREE_INT_CST_LOW (gimple_call_arg (call, 3)); - if (level == GOMP_DIM_VECTOR) - { - default_goacc_reduction (call); - return; - } - unsigned code = (unsigned) TREE_INT_CST_LOW (gimple_call_arg (call, 0)); switch (code) diff --git a/gcc/config/gcn/gcn.cc b/gcc/config/gcn/gcn.cc index 90c16b4e0a8b..92d7b939cfdd 100644 --- a/gcc/config/gcn/gcn.cc +++ b/gcc/config/gcn/gcn.cc @@ -7062,9 +7062,10 @@ gcn_oacc_dim_size (int dim) if (dim < 0 || dim > 2) error ("offload dimension out of range (%d)", dim); - /* Vectors are a special case. */ + /* Vectors are a special case. In SIMD mode think of this as 1 times 64. + In SIMT mode each lane is a separate vector position. */ if (dim == 2) - return const1_rtx; /* Think of this as 1 times 64. */ + return flag_simt_mode ? GEN_INT (64) : const1_rtx; static int offset[] = { /* Offsets into dispatch packet. */ diff --git a/gcc/config/gcn/gcn.md b/gcc/config/gcn/gcn.md index cb7030e0efbf..5b9668016792 100644 --- a/gcc/config/gcn/gcn.md +++ b/gcc/config/gcn/gcn.md @@ -63,7 +63,6 @@ UNSPECV_BARRIER UNSPECV_ATOMIC UNSPECV_ICACHE_INV - UNSPECV_SIMT_FORK UNSPECV_SIMT_JOIN UNSPECV_SIMT_ENTER UNSPECV_SIMT_EXIT UNSPECV_SIMT_EXEC_BARRIER UNSPECV_SIMT_REDUCE]) @@ -2633,59 +2632,50 @@ [(set_attr "type" "sopp")]) (define_expand "oacc_fork" - [(set (match_operand:SI 0 "register_operand") - (unspec_volatile [(match_operand:SI 1 "") - (match_operand:SI 2 "")] - UNSPECV_SIMT_FORK))] + [(match_operand:SI 0 "register_operand") + (match_operand:SI 1 "") + (match_operand:SI 2 "")] "" { - /* We need to have oacc_fork/oacc_join named patterns as a pair, - but the fork isn't actually used. */ - if (!flag_simt_mode) + if (flag_simt_mode) { - gcc_unreachable (); - DONE; + int dim = INTVAL (operands[2]); + + if (dim == GOMP_DIM_VECTOR) + emit_insn (gen_omp_simt_enter (operands[0], operands[1], + operands[2])); + /* Else do nothing. */ } - }) + else + /* We need to have oacc_fork/oacc_join named patterns as a pair, + but the fork isn't actually used in SIMD mode. */ + gcc_unreachable (); -; This pseudo instruction should be replaced by the simt_vect pass -(define_insn "*oacc_fork_insn" - [(set (match_operand:SI 0 "register_operand") - (unspec_volatile [(match_operand:SI 1 "") - (match_operand:SI 2 "")] - UNSPECV_SIMT_FORK))] - "" - "; simt_fork %0 %1 %2 ") + DONE; + }) (define_expand "oacc_join" - [(unspec_volatile [(match_operand:SI 0 "") - (match_operand:SI 1 "") - (match_operand:SI 2 "")] - UNSPECV_SIMT_JOIN)] + [(match_operand:SI 0 "") + (match_operand:SI 1 "") + (match_operand:SI 2 "")] "" { int dim = INTVAL (operands[2]); + if (dim == GOMP_DIM_WORKER) + emit_insn (gen_gcn_wavefront_barrier ()); + else if (flag_simt_mode) { - emit_insn (gen_gcn_wavefront_barrier ()); - DONE; + if (dim == GOMP_DIM_VECTOR) + emit_insn (gen_omp_simt_exit (operands[1])); + /* Else do nothing. */ } - else if (!flag_simt_mode) - { - gcc_unreachable (); - DONE; - } - /* else use default expansion. */ - }) + else + /* Other uses should not occur with SIMD mode. */ + gcc_unreachable (); -; This pseudo instruction should be replaced by the simt_vect pass -(define_insn "*oacc_join_insn" - [(unspec_volatile [(match_operand:SI 0 "") - (match_operand:SI 1 "") - (match_operand:SI 2 "")] - UNSPECV_SIMT_JOIN)] - "" - "; simt_join %0 %1 %2 ") + DONE; + }) ; This pseudo instruction should be replaced by the simt_vect pass (define_insn "omp_simt_enter" @@ -2701,7 +2691,7 @@ ; This pseudo instruction should be replaced by the simt_vect pass (define_insn "omp_simt_exit" - [(unspec_volatile [(match_operand 0 "register_operand" "r")] + [(unspec_volatile [(match_operand 0 "" "")] UNSPECV_SIMT_EXIT) (set (reg:DI EXEC_REG) (const_int 1)) (set (reg:BI EXECZ_REG) (const_int 0))] @@ -2748,7 +2738,7 @@ (unspec:SI [(match_operand:SI 1 "register_operand" "")] UNSPEC_SIMT_VOTE_ANY))] "flag_simt_mode" - "; %0 = simt_vote_any %1") + "sorry, unimplemented %0 = simt_vote_any %1") ;; Implement IFN_GOMP_SIMT_LAST_LANE: ;; set operand 0 to the value of the last active lane @@ -2768,7 +2758,7 @@ (match_operand:SI 2 "nonmemory_operand" "")] UNSPEC_SIMT_XCHG_IDX))] "flag_simt_mode" - "; %0 = simt_xchg_idx %1") + "sorry, unimplemented %0 = simt_xchg_idx %1") ;; Ensure that scalar values are visible to all lanes at the SIMT_ENTER ;; transition point. This could be a vec_duplicate, or just ensuring it's diff --git a/libgomp/testsuite/libgomp.c-c++-common/simt-reduction-amdgcn-1.c b/libgomp/testsuite/libgomp.c-c++-common/simt-reduction-amdgcn-1.c new file mode 100644 index 000000000000..4137187fbc63 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/simt-reduction-amdgcn-1.c @@ -0,0 +1,87 @@ +/* Check that reductions on a "#pragma omp simd" loop nested in a + "#pragma omp target" region are computed correctly when GCC maps the + simd loop onto AMD GCN SIMT lanes (one OpenMP thread == one 64-lane + wavefront, see gcn-simt-vectorize.cc / gcn-simt-cbranch.cc). + + Exercise both: + - a trip count that is an exact multiple of the 64-lane wavefront + width, so every SIMT group has all 64 lanes live, and + - a trip count that is not, so the final SIMT group is executed + with fewer than 64 lanes live (EXEC partially masked). */ + +/* { dg-do run } */ +/* { dg-require-effective-target offload_device_gcn } */ +/* { dg-additional-options "--param=offload-gpu-mode=simt" { target offload_target_amdgcn } } */ + +#include <stdlib.h> + +#define FULL_N (64 * 5) /* Multiple of 64: only full groups. */ +#define PART_N (64 * 3 + 17) /* Last group has only 17 lanes live. */ + +static void +test_sum (int n) +{ + long sum = 0, expected = 0; + int i; + + #pragma omp target map(tofrom: sum) + #pragma omp simd reduction(+:sum) + for (i = 0; i < n; i++) + sum += i; + + for (i = 0; i < n; i++) + expected += i; + + if (sum != expected) + abort (); +} + +static void +test_minmax (int n) +{ + int rmin = n, rmax = -1; + int i; + + #pragma omp target map(tofrom: rmin, rmax) + #pragma omp simd reduction(min:rmin) reduction(max:rmax) + for (i = 0; i < n; i++) + { + if (i < rmin) + rmin = i; + if (i > rmax) + rmax = i; + } + + if (rmin != 0 || rmax != n - 1) + abort (); +} + +static void +test_xor (int n) +{ + unsigned int x = 0, expected = 0; + int i; + + #pragma omp target map(tofrom: x) + #pragma omp simd reduction(^:x) + for (i = 0; i < n; i++) + x ^= i; + + for (i = 0; i < n; i++) + expected ^= i; + + if (x != expected) + abort (); +} + +int +main (void) +{ + test_sum (FULL_N); + test_sum (PART_N); + test_minmax (FULL_N); + test_minmax (PART_N); + test_xor (FULL_N); + test_xor (PART_N); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/simt-reduction-amdgcn-2.c b/libgomp/testsuite/libgomp.c-c++-common/simt-reduction-amdgcn-2.c new file mode 100644 index 000000000000..67acdf20df93 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/simt-reduction-amdgcn-2.c @@ -0,0 +1,107 @@ +/* Companion to simt-reduction-amdgcn-1.c: check reductions on a + "#pragma omp simd" loop nested inside team/thread parallelism + ("#pragma omp target teams distribute parallel for"), which is the + intended way to use SIMT lanes (one team's one thread == one 64-lane + wavefront), as opposed to the bare "target" + "simd" with no enclosing + teams/parallel exercised by simt-reduction-amdgcn-1.c. + + num_teams(1) num_threads(1) forces a single team and a single thread, + i.e. a single wavefront, so this isolates the within-wavefront + reduction tree from any cross-team/cross-thread combination -- the + same thing simt-reduction-amdgcn-1.c checks, just reached through the + nested construct instead of the bare one. + + As of this writing, this nested case does not even compile, for two + independent reasons, triggered regardless of the num_teams/num_threads + values since the limitation is structural rather than a function of + the actual team/thread counts: + + - The GCN backend's lane-index translation for a SIMT loop inside + team/thread parallelism is an unimplemented stub + ("omp_simt_xchg_idx_<mode>" in gcn-valu.md, "sorry, placeholder + ... simt_xchg_idx"), hit when generating the amdgcn offload + image. + + - Separately, generating the host fallback version of the same + region ICEs in expand_GOMP_SIMT_LANE (internal-fn.cc) with + "'--param=offload-gpu-mode=simt-directed|simt' is incompatible + with '-no-msimt'" -- the incompatibility is correctly diagnosed, + but compilation does not stop there and crashes instead. */ + +/* { dg-do link } */ +/* { dg-additional-options "--param=offload-gpu-mode=simt" { target offload_target_amdgcn } } */ +/* { dg-excess-errors "sorry, unimplemented: simt_xchg_idx; ICE in expand_GOMP_SIMT_LANE for the host fallback; lto1, mkoffload and lto-wrapper fatal errors" { target offload_target_amdgcn } } */ + +#include <stdlib.h> + +#define FULL_N (64 * 5) /* Multiple of 64: only full groups. */ +#define PART_N (64 * 3 + 17) /* Last group has only 17 lanes live. */ + +static void +test_sum (int n) +{ + long sum = 0, expected = 0; + int i; + + #pragma omp target teams distribute parallel for simd \ + num_teams(1) num_threads(1) reduction(+:sum) map(tofrom: sum) + for (i = 0; i < n; i++) + sum += i; + + for (i = 0; i < n; i++) + expected += i; + + if (sum != expected) + abort (); +} + +static void +test_minmax (int n) +{ + int rmin = n, rmax = -1; + int i; + + #pragma omp target teams distribute parallel for simd \ + num_teams(1) num_threads(1) \ + reduction(min:rmin) reduction(max:rmax) map(tofrom: rmin, rmax) + for (i = 0; i < n; i++) + { + if (i < rmin) + rmin = i; + if (i > rmax) + rmax = i; + } + + if (rmin != 0 || rmax != n - 1) + abort (); +} + +static void +test_xor (int n) +{ + unsigned int x = 0, expected = 0; + int i; + + #pragma omp target teams distribute parallel for simd \ + num_teams(1) num_threads(1) reduction(^:x) map(tofrom: x) + for (i = 0; i < n; i++) + x ^= i; + + for (i = 0; i < n; i++) + expected ^= i; + + if (x != expected) + abort (); +} + +int +main (void) +{ + test_sum (FULL_N); + test_sum (PART_N); + test_minmax (FULL_N); + test_minmax (PART_N); + test_xor (FULL_N); + test_xor (PART_N); + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/simt-reduction-amdgcn-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/simt-reduction-amdgcn-1.c new file mode 100644 index 000000000000..e8e9fad9f902 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/simt-reduction-amdgcn-1.c @@ -0,0 +1,84 @@ +/* Check that reductions on a "#pragma acc loop vector" nested in a + "#pragma acc parallel" region are computed correctly on AMD GCN, where + the vector dimension is mapped onto the 64 SIMT lanes of a wavefront. + + Exercise both: + - a trip count that is an exact multiple of the 64-lane wavefront + width, so every SIMT group has all 64 lanes live, and + - a trip count that is not, so the final SIMT group is executed + with fewer than 64 lanes live (EXEC partially masked). */ + +/* { dg-do run { target openacc_radeon_accel_selected } } */ + +#include <stdlib.h> + +#define FULL_N (64 * 5) /* Multiple of 64: only full groups. */ +#define PART_N (64 * 3 + 17) /* Last group has only 17 lanes live. */ + +static void +test_sum (int n) +{ + long sum = 0, expected = 0; + int i; + + #pragma acc parallel copy(sum) + #pragma acc loop vector reduction(+:sum) + for (i = 0; i < n; i++) + sum += i; + + for (i = 0; i < n; i++) + expected += i; + + if (sum != expected) + abort (); +} + +static void +test_minmax (int n) +{ + int rmin = n, rmax = -1; + int i; + + #pragma acc parallel copy(rmin, rmax) + #pragma acc loop vector reduction(min:rmin) reduction(max:rmax) + for (i = 0; i < n; i++) + { + if (i < rmin) + rmin = i; + if (i > rmax) + rmax = i; + } + + if (rmin != 0 || rmax != n - 1) + abort (); +} + +static void +test_xor (int n) +{ + unsigned int x = 0, expected = 0; + int i; + + #pragma acc parallel copy(x) + #pragma acc loop vector reduction(^:x) + for (i = 0; i < n; i++) + x ^= i; + + for (i = 0; i < n; i++) + expected ^= i; + + if (x != expected) + abort (); +} + +int +main (void) +{ + test_sum (FULL_N); + test_sum (PART_N); + test_minmax (FULL_N); + test_minmax (PART_N); + test_xor (FULL_N); + test_xor (PART_N); + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/simt-reduction-amdgcn-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/simt-reduction-amdgcn-2.c new file mode 100644 index 000000000000..b59fa9f94011 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/simt-reduction-amdgcn-2.c @@ -0,0 +1,100 @@ +/* Companion to simt-reduction-amdgcn-1.c: check reductions on a + "#pragma acc loop vector" nested inside gang/worker parallelism + ("#pragma acc parallel" + "#pragma acc loop worker"), which is the + OpenACC analogue of simt-reduction-amdgcn-2.c's OpenMP + "teams distribute parallel for simd" case, as opposed to the bare + "#pragma acc parallel" + "#pragma acc loop vector" with no enclosing + worker loop exercised by simt-reduction-amdgcn-1.c. + + num_gangs(1) num_workers(1) forces a single gang and a single worker, + i.e. a single wavefront, so this isolates the within-wavefront + reduction tree from any cross-gang/cross-worker combination -- the + same thing simt-reduction-amdgcn-1.c checks, just reached through the + nested construct instead of the bare one. */ + +/* { dg-do run { target openacc_radeon_accel_selected } } */ + +#include <stdlib.h> + +#define FULL_N (64 * 5) /* Multiple of 64: only full groups. */ +#define PART_N (64 * 3 + 17) /* Last group has only 17 lanes live. */ + +static void +test_sum (int n) +{ + long sum = 0, expected = 0; + int i; + + #pragma acc parallel num_gangs(1) num_workers(1) copy(sum) + #pragma acc loop worker + for (int w = 0; w < 1; w++) + { + #pragma acc loop vector reduction(+:sum) + for (i = 0; i < n; i++) + sum += i; + } + + for (i = 0; i < n; i++) + expected += i; + + if (sum != expected) + abort (); +} + +static void +test_minmax (int n) +{ + int rmin = n, rmax = -1; + int i; + + #pragma acc parallel num_gangs(1) num_workers(1) copy(rmin, rmax) + #pragma acc loop worker + for (int w = 0; w < 1; w++) + { + #pragma acc loop vector reduction(min:rmin) reduction(max:rmax) + for (i = 0; i < n; i++) + { + if (i < rmin) + rmin = i; + if (i > rmax) + rmax = i; + } + } + + if (rmin != 0 || rmax != n - 1) + abort (); +} + +static void +test_xor (int n) +{ + unsigned int x = 0, expected = 0; + int i; + + #pragma acc parallel num_gangs(1) num_workers(1) copy(x) + #pragma acc loop worker + for (int w = 0; w < 1; w++) + { + #pragma acc loop vector reduction(^:x) + for (i = 0; i < n; i++) + x ^= i; + } + + for (i = 0; i < n; i++) + expected ^= i; + + if (x != expected) + abort (); +} + +int +main (void) +{ + test_sum (FULL_N); + test_sum (PART_N); + test_minmax (FULL_N); + test_minmax (PART_N); + test_xor (FULL_N); + test_xor (PART_N); + return 0; +}
