From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1534) id E8C363858412; Wed, 20 Sep 2023 14:04:27 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org E8C363858412 DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1695218667; bh=IomhOA8NtDeLAI3RjbrOVks6AJuym/0zwLFTnHpdJJE=; h=From:To:Subject:Date:From; b=TcSpXskLFuUsxc/RJyC9wmo1PCj9SxInm4mXSIvri7ZsHwG4joE9e6J5S0XsXVyxO I3xMjgZwd4/DKDSBAjv6CHLZaqov6b1Fzl5n3gcdvkhovJ5gvjD+gwZqYT8VF/upDV T0FqQxA8AeTbPKrE4MuzCqr7iWWbU3SmkP45IhqA= MIME-Version: 1.0 Content-Transfer-Encoding: 7bit Content-Type: text/plain; charset="utf-8" From: Tobias Burnus To: gcc-cvs@gcc.gnu.org Subject: [gcc r14-4176] OpenMP: Add ME support for 'omp allocate' stack variables X-Act-Checkin: gcc X-Git-Author: Tobias Burnus X-Git-Refname: refs/heads/master X-Git-Oldrev: b9cb735fc1bb4ca2339ab900e2d07667d7c0f6b4 X-Git-Newrev: 1a554a2c9f33fdb3c170f1c37274037ece050114 Message-Id: <20230920140427.E8C363858412@sourceware.org> Date: Wed, 20 Sep 2023 14:04:27 +0000 (GMT) List-Id: https://gcc.gnu.org/g:1a554a2c9f33fdb3c170f1c37274037ece050114 commit r14-4176-g1a554a2c9f33fdb3c170f1c37274037ece050114 Author: Tobias Burnus Date: Wed Sep 20 16:03:19 2023 +0200 OpenMP: Add ME support for 'omp allocate' stack variables Call GOMP_alloc/free for 'omp allocate' allocated variables. This is for C only as C++ and Fortran show a sorry already in the FE. Note that this only applies to stack variables as the C FE shows a sorry for static variables. gcc/ChangeLog: * gimplify.cc (gimplify_bind_expr): Call GOMP_alloc/free for 'omp allocate' variables; move stack cleanup after other cleanup. (omp_notice_variable): Process original decl when decl of the value-expression for a 'omp allocate' variable is passed. * omp-low.cc (scan_omp_1_op): Handle 'omp allocate' variables libgomp/ChangeLog: * libgomp.texi (OpenMP 5.1 Impl.): Mark 'omp allocate' as implemented for C only. * testsuite/libgomp.c/allocate-4.c: New test. * testsuite/libgomp.c/allocate-5.c: New test. * testsuite/libgomp.c/allocate-6.c: New test. gcc/testsuite/ChangeLog: * c-c++-common/gomp/allocate-11.c: Remove C-only dg-message for 'sorry, unimplemented'. * c-c++-common/gomp/allocate-12.c: Likewise. * c-c++-common/gomp/allocate-15.c: Likewise. * c-c++-common/gomp/allocate-9.c: Likewise. * c-c++-common/gomp/allocate-10.c: New test. * c-c++-common/gomp/allocate-17.c: New test. Diff: --- gcc/gimplify.cc | 108 +++++++-- gcc/omp-low.cc | 28 ++- gcc/testsuite/c-c++-common/gomp/allocate-10.c | 49 ++++ gcc/testsuite/c-c++-common/gomp/allocate-11.c | 3 - gcc/testsuite/c-c++-common/gomp/allocate-12.c | 3 - gcc/testsuite/c-c++-common/gomp/allocate-15.c | 4 +- gcc/testsuite/c-c++-common/gomp/allocate-17.c | 37 +++ gcc/testsuite/c-c++-common/gomp/allocate-9.c | 2 - libgomp/libgomp.texi | 5 +- libgomp/testsuite/libgomp.c/allocate-4.c | 84 +++++++ libgomp/testsuite/libgomp.c/allocate-5.c | 126 ++++++++++ libgomp/testsuite/libgomp.c/allocate-6.c | 319 ++++++++++++++++++++++++++ 12 files changed, 733 insertions(+), 35 deletions(-) diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index a0e8cc2199d..9f4722f7458 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -36,6 +36,7 @@ along with GCC; see the file COPYING3. If not see #include "cgraph.h" #include "tree-pretty-print.h" #include "diagnostic-core.h" +#include "diagnostic.h" /* For errorcount. */ #include "alias.h" #include "fold-const.h" #include "calls.h" @@ -1372,6 +1373,7 @@ gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p) && (attr = lookup_attribute ("omp allocate", DECL_ATTRIBUTES (t))) != NULL_TREE) { + gcc_assert (!DECL_HAS_VALUE_EXPR_P (t)); tree alloc = TREE_PURPOSE (TREE_VALUE (attr)); tree align = TREE_VALUE (TREE_VALUE (attr)); /* Allocate directives that appear in a target region must specify @@ -1396,12 +1398,56 @@ gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p) error_at (DECL_SOURCE_LOCATION (t), "% directive for %qD inside a target " "region must specify an % clause", t); - else if (align != NULL_TREE - || alloc == NULL_TREE - || !integer_onep (alloc)) - sorry_at (DECL_SOURCE_LOCATION (t), - "OpenMP % directive, used for %qD, not " - "yet supported", t); + /* Skip for omp_default_mem_alloc (= 1), + unless align is present. */ + else if (!errorcount + && (align != NULL_TREE + || alloc == NULL_TREE + || !integer_onep (alloc))) + { + tree tmp = build_pointer_type (TREE_TYPE (t)); + tree v = create_tmp_var (tmp, get_name (t)); + DECL_IGNORED_P (v) = 0; + tmp = remove_attribute ("omp allocate", DECL_ATTRIBUTES (t)); + DECL_ATTRIBUTES (v) + = tree_cons (get_identifier ("omp allocate var"), + build_tree_list (NULL_TREE, t), tmp); + tmp = build_fold_indirect_ref (v); + TREE_THIS_NOTRAP (tmp) = 1; + SET_DECL_VALUE_EXPR (t, tmp); + DECL_HAS_VALUE_EXPR_P (t) = 1; + tree sz = TYPE_SIZE_UNIT (TREE_TYPE (t)); + if (alloc == NULL_TREE) + alloc = build_zero_cst (ptr_type_node); + if (align == NULL_TREE) + align = build_int_cst (size_type_node, DECL_ALIGN_UNIT (t)); + else + align = build_int_cst (size_type_node, + MAX (tree_to_uhwi (align), + DECL_ALIGN_UNIT (t))); + tmp = builtin_decl_explicit (BUILT_IN_GOMP_ALLOC); + tmp = build_call_expr_loc (DECL_SOURCE_LOCATION (t), tmp, + 3, align, sz, alloc); + tmp = fold_build2_loc (DECL_SOURCE_LOCATION (t), MODIFY_EXPR, + TREE_TYPE (v), v, + fold_convert (TREE_TYPE (v), tmp)); + gcc_assert (BIND_EXPR_BODY (bind_expr) != NULL_TREE + && (TREE_CODE (BIND_EXPR_BODY (bind_expr)) + == STATEMENT_LIST)); + tree_stmt_iterator e = tsi_start (BIND_EXPR_BODY (bind_expr)); + while (!tsi_end_p (e)) + { + if ((TREE_CODE (*e) == DECL_EXPR + && TREE_OPERAND (*e, 0) == t) + || (TREE_CODE (*e) == CLEANUP_POINT_EXPR + && TREE_CODE (TREE_OPERAND (*e, 0)) == DECL_EXPR + && TREE_OPERAND (TREE_OPERAND (*e, 0), 0) == t)) + break; + ++e; + } + gcc_assert (!tsi_end_p (e)); + tsi_link_before (&e, tmp, TSI_SAME_STMT); + } } /* Mark variable as local. */ @@ -1486,22 +1532,6 @@ gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p) cleanup = NULL; stack_save = NULL; - /* If the code both contains VLAs and calls alloca, then we cannot reclaim - the stack space allocated to the VLAs. */ - if (gimplify_ctxp->save_stack && !gimplify_ctxp->keep_stack) - { - gcall *stack_restore; - - /* Save stack on entry and restore it on exit. Add a try_finally - block to achieve this. */ - build_stack_save_restore (&stack_save, &stack_restore); - - gimple_set_location (stack_save, start_locus); - gimple_set_location (stack_restore, end_locus); - - gimplify_seq_add_stmt (&cleanup, stack_restore); - } - /* Add clobbers for all variables that go out of scope. */ for (t = BIND_EXPR_VARS (bind_expr); t ; t = DECL_CHAIN (t)) { @@ -1509,6 +1539,17 @@ gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p) && !is_global_var (t) && DECL_CONTEXT (t) == current_function_decl) { + if (flag_openmp + && DECL_HAS_VALUE_EXPR_P (t) + && TREE_USED (t) + && lookup_attribute ("omp allocate", DECL_ATTRIBUTES (t))) + { + tree tmp = builtin_decl_explicit (BUILT_IN_GOMP_FREE); + tmp = build_call_expr_loc (end_locus, tmp, 2, + TREE_OPERAND (DECL_VALUE_EXPR (t), 0), + build_zero_cst (ptr_type_node)); + gimplify_and_add (tmp, &cleanup); + } if (!DECL_HARD_REGISTER (t) && !TREE_THIS_VOLATILE (t) && !DECL_HAS_VALUE_EXPR_P (t) @@ -1565,6 +1606,22 @@ gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p) gimplify_ctxp->live_switch_vars->remove (t); } + /* If the code both contains VLAs and calls alloca, then we cannot reclaim + the stack space allocated to the VLAs. */ + if (gimplify_ctxp->save_stack && !gimplify_ctxp->keep_stack) + { + gcall *stack_restore; + + /* Save stack on entry and restore it on exit. Add a try_finally + block to achieve this. */ + build_stack_save_restore (&stack_save, &stack_restore); + + gimple_set_location (stack_save, start_locus); + gimple_set_location (stack_restore, end_locus); + + gimplify_seq_add_stmt (&cleanup, stack_restore); + } + if (ret_clauses) { gomp_target *stmt; @@ -7894,6 +7951,13 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) if (error_operand_p (decl)) return false; + if (DECL_ARTIFICIAL (decl)) + { + tree attr = lookup_attribute ("omp allocate var", DECL_ATTRIBUTES (decl)); + if (attr) + decl = TREE_VALUE (TREE_VALUE (attr)); + } + if (ctx->region_type == ORT_NONE) return lang_hooks.decls.omp_disregard_value_expr (decl, false); diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 5d7c32dac39..b0c3ef7a9cc 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -3951,6 +3951,7 @@ scan_omp_1_op (tree *tp, int *walk_subtrees, void *data) struct walk_stmt_info *wi = (struct walk_stmt_info *) data; omp_context *ctx = (omp_context *) wi->info; tree t = *tp; + tree tmp; switch (TREE_CODE (t)) { @@ -3960,12 +3961,37 @@ scan_omp_1_op (tree *tp, int *walk_subtrees, void *data) case RESULT_DECL: if (ctx) { + tmp = NULL_TREE; + if (TREE_CODE (t) == VAR_DECL + && (tmp = lookup_attribute ("omp allocate var", + DECL_ATTRIBUTES (t))) != NULL_TREE) + t = TREE_VALUE (TREE_VALUE (tmp)); tree repl = remap_decl (t, &ctx->cb); gcc_checking_assert (TREE_CODE (repl) != ERROR_MARK); - *tp = repl; + if (tmp != NULL_TREE && t != repl) + *tp = build_fold_addr_expr (repl); + else if (tmp == NULL_TREE) + *tp = repl; } break; + case INDIRECT_REF: + case MEM_REF: + if (ctx + && TREE_CODE (TREE_OPERAND (t, 0)) == VAR_DECL + && ((tmp = lookup_attribute ("omp allocate var", + DECL_ATTRIBUTES (TREE_OPERAND (t, 0)))) + != NULL_TREE)) + { + tmp = TREE_VALUE (TREE_VALUE (tmp)); + tree repl = remap_decl (tmp, &ctx->cb); + gcc_checking_assert (TREE_CODE (repl) != ERROR_MARK); + if (tmp != repl) + *tp = repl; + break; + } + gcc_fallthrough (); + default: if (ctx && TYPE_P (t)) *tp = remap_type (t, &ctx->cb); diff --git a/gcc/testsuite/c-c++-common/gomp/allocate-10.c b/gcc/testsuite/c-c++-common/gomp/allocate-10.c new file mode 100644 index 00000000000..7e8f579871c --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/allocate-10.c @@ -0,0 +1,49 @@ +/* TODO: enable for C++ once implemented. */ +/* { dg-do compile { target c } } */ +/* { dg-additional-options "-Wall -fdump-tree-gimple" } */ + +typedef enum omp_allocator_handle_t +#if __cplusplus >= 201103L +: __UINTPTR_TYPE__ +#endif +{ + omp_default_mem_alloc = 1, + __omp_allocator_handle_t_max__ = __UINTPTR_MAX__ +} omp_allocator_handle_t; + +void +f() +{ + int n; + int A[n]; /* { dg-warning "'n' is used uninitialized" } */ + /* { dg-warning "unused variable 'A'" "" { target *-*-* } .-1 } */ +} + +void +h1() +{ + omp_allocator_handle_t my_handle; + int B1[3]; /* { dg-warning "'my_handle' is used uninitialized" } */ + /* { dg-warning "variable 'B1' set but not used" "" { target *-*-* } .-1 } */ + #pragma omp allocate(B1) allocator(my_handle) + B1[0] = 5; + /* { dg-final { scan-tree-dump-times "__builtin_GOMP_alloc" 1 "gimple" } } */ + /* { dg-final { scan-tree-dump-times "B1.\[0-9\]+ = __builtin_GOMP_alloc \\(4, 12, my_handle\\);" 1 "gimple" } } */ + /* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(B1.\[0-9\]+, 0B\\);" 1 "gimple" } } */ +} + +void +h2() +{ + omp_allocator_handle_t my_handle; + int B2[3]; /* { dg-warning "unused variable 'B2'" } */ + #pragma omp allocate(B2) allocator(my_handle) /* No warning as 'B2' is unused */ +} + +void +h3() +{ + omp_allocator_handle_t my_handle; + int B3[3] = {1,2,3}; /* { dg-warning "unused variable 'B3'" } */ + #pragma omp allocate(B3) allocator(my_handle) /* No warning as 'B3' is unused */ +} diff --git a/gcc/testsuite/c-c++-common/gomp/allocate-11.c b/gcc/testsuite/c-c++-common/gomp/allocate-11.c index f9ad50abb7f..dceb97f8c5f 100644 --- a/gcc/testsuite/c-c++-common/gomp/allocate-11.c +++ b/gcc/testsuite/c-c++-common/gomp/allocate-11.c @@ -10,7 +10,6 @@ f (int i) switch (i) /* { dg-note "switch starts here" } */ { int j; /* { dg-note "'j' declared here" } */ - /* { dg-message "sorry, unimplemented: OpenMP 'allocate' directive" "" { target *-*-* } .-1 } */ #pragma omp allocate(j) case 42: /* { dg-error "switch jumps over OpenMP 'allocate' allocation" } */ bar (); @@ -30,9 +29,7 @@ h (int i2) return 5; int k2; /* { dg-note "'k2' declared here" } */ - /* { dg-message "sorry, unimplemented: OpenMP 'allocate' directive" "" { target *-*-* } .-1 } */ int j2 = 4; /* { dg-note "'j2' declared here" } */ - /* { dg-message "sorry, unimplemented: OpenMP 'allocate' directive" "" { target *-*-* } .-1 } */ #pragma omp allocate(k2, j2) label: /* { dg-note "label 'label' defined here" } */ k2 = 4; diff --git a/gcc/testsuite/c-c++-common/gomp/allocate-12.c b/gcc/testsuite/c-c++-common/gomp/allocate-12.c index 3c7c3bb3a2b..1b77db9bd6f 100644 --- a/gcc/testsuite/c-c++-common/gomp/allocate-12.c +++ b/gcc/testsuite/c-c++-common/gomp/allocate-12.c @@ -17,7 +17,6 @@ f () omp_allocator_handle_t my_allocator; int n = 5; /* { dg-note "to be allocated variable declared here" } */ my_allocator = omp_default_mem_alloc; /* { dg-note "modified here" } */ - /* { dg-message "sorry, unimplemented: OpenMP 'allocate' directive" "" { target *-*-* } .-2 } */ #pragma omp allocate(n) allocator(my_allocator) /* { dg-error "variable 'my_allocator' used in the 'allocator' clause must not be modified between declaration of 'n' and its 'allocate' directive" } */ n = 7; return n; @@ -28,7 +27,6 @@ int g () { int n = 5; /* { dg-note "to be allocated variable declared here" } */ - /* { dg-message "sorry, unimplemented: OpenMP 'allocate' directive" "" { target *-*-* } .-1 } */ omp_allocator_handle_t my_allocator = omp_low_lat_mem_alloc; /* { dg-note "declared here" } */ #pragma omp allocate(n) allocator(my_allocator) /* { dg-error "variable 'my_allocator' used in the 'allocator' clause must be declared before 'n'" } */ n = 7; @@ -42,7 +40,6 @@ h () see gomp/allocate-10.c. */ omp_allocator_handle_t my_allocator; int n = 5; - /* { dg-message "sorry, unimplemented: OpenMP 'allocate' directive" "" { target *-*-* } .-1 } */ #pragma omp allocate(n) allocator(my_allocator) n = 7; return n; diff --git a/gcc/testsuite/c-c++-common/gomp/allocate-15.c b/gcc/testsuite/c-c++-common/gomp/allocate-15.c index d9600f96c46..15105b9102e 100644 --- a/gcc/testsuite/c-c++-common/gomp/allocate-15.c +++ b/gcc/testsuite/c-c++-common/gomp/allocate-15.c @@ -8,7 +8,7 @@ void f () { - int var; /* { dg-message "sorry, unimplemented: OpenMP 'allocate' directive, used for 'var', not yet supported" } */ + int var; #pragma omp allocate(var) var = 5; } @@ -21,7 +21,7 @@ h () #pragma omp parallel #pragma omp serial { - int var2[5]; /* { dg-message "sorry, unimplemented: OpenMP 'allocate' directive, used for 'var2', not yet supported" } */ + int var2[5]; #pragma omp allocate(var2) var2[0] = 7; } diff --git a/gcc/testsuite/c-c++-common/gomp/allocate-17.c b/gcc/testsuite/c-c++-common/gomp/allocate-17.c new file mode 100644 index 00000000000..f75af0c2d93 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/allocate-17.c @@ -0,0 +1,37 @@ +/* This file has a syntax error but should not ICE. + Namely, a '}' is missing in one(). */ + +typedef enum omp_allocator_handle_t +#if __cplusplus >= 201103L +: __UINTPTR_TYPE__ +#endif +{ + omp_default_mem_alloc = 1, + omp_low_lat_mem_alloc = 5, + __omp_allocator_handle_t_max__ = __UINTPTR_MAX__ +} omp_allocator_handle_t; + +#include + +void +one () +{ /* { dg-note "to match this '\{'" "" { target c++ } } */ + int result = 0, n = 3; + #pragma omp target map(tofrom: result) firstprivate(n) + { + int var = 5; //, var2[n]; + #pragma omp allocate(var) align(128) allocator(omp_low_lat_mem_alloc) /* { dg-message "sorry, unimplemented: '#pragma omp allocate' not yet supported" "" { target c++ } } */ + var = 7; +} + +void +two () +{ /* { dg-error "a function-definition is not allowed here before '\{' token" "" { target c++ } } */ + int scalar = 44; + #pragma omp allocate(scalar) + + #pragma omp parallel firstprivate(scalar) + scalar = 33; +} +/* { dg-error "expected declaration or statement at end of input" "" { target c } .-1 } */ +/* { dg-error "expected '\}' at end of input" "" { target c++ } .-2 } */ diff --git a/gcc/testsuite/c-c++-common/gomp/allocate-9.c b/gcc/testsuite/c-c++-common/gomp/allocate-9.c index 8e010419a5f..3c11080dd16 100644 --- a/gcc/testsuite/c-c++-common/gomp/allocate-9.c +++ b/gcc/testsuite/c-c++-common/gomp/allocate-9.c @@ -86,8 +86,6 @@ int g() /* { dg-note "declared here" "" { target c } .-8 } */ /* { dg-message "sorry, unimplemented: '#pragma omp allocate' not yet supported" "" { target c++ } .-2 } */ return c2+a2+b2; - /* { dg-message "sorry, unimplemented: OpenMP 'allocate' directive" "" { target c } .-5 } */ - /* { dg-message "sorry, unimplemented: OpenMP 'allocate' directive" "" { target c } .-12 } */ } } diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi index f5cb5b643a2..f21557c3c52 100644 --- a/libgomp/libgomp.texi +++ b/libgomp/libgomp.texi @@ -225,7 +225,7 @@ The OpenMP 4.5 specification is fully supported. @item Predefined memory spaces, memory allocators, allocator traits @tab Y @tab See also @ref{Memory allocation} @item Memory management routines @tab Y @tab -@item @code{allocate} directive @tab N @tab +@item @code{allocate} directive @tab P @tab Only C, only stack variables @item @code{allocate} clause @tab P @tab Initial support @item @code{use_device_addr} clause on @code{target data} @tab Y @tab @item @code{ancestor} modifier on @code{device} clause @tab Y @tab @@ -296,7 +296,8 @@ The OpenMP 4.5 specification is fully supported. @item Loop transformation constructs @tab N @tab @item @code{strict} modifier in the @code{grainsize} and @code{num_tasks} clauses of the @code{taskloop} construct @tab Y @tab -@item @code{align} clause in @code{allocate} directive @tab N @tab +@item @code{align} clause in @code{allocate} directive @tab P + @tab Only C (and only stack variables) @item @code{align} modifier in @code{allocate} clause @tab Y @tab @item @code{thread_limit} clause to @code{target} construct @tab Y @tab @item @code{has_device_addr} clause to @code{target} construct @tab Y @tab diff --git a/libgomp/testsuite/libgomp.c/allocate-4.c b/libgomp/testsuite/libgomp.c/allocate-4.c new file mode 100644 index 00000000000..e81cc4093aa --- /dev/null +++ b/libgomp/testsuite/libgomp.c/allocate-4.c @@ -0,0 +1,84 @@ +/* TODO: move to ../libgomp.c-c++-common once C++ is implemented. */ +/* NOTE: { target c } is unsupported with with the C compiler. */ + +/* { dg-do run } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +#include +#include + +/* { dg-final { scan-tree-dump-times "__builtin_GOMP_alloc \\(" 5 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(" 5 "gimple" } } */ + + +int one () +{ + int sum = 0; + #pragma omp allocate(sum) + /* { dg-final { scan-tree-dump-times "sum\\.\[0-9\]+ = __builtin_GOMP_alloc \\(4, 4, 0B\\);" 1 "gimple" } } */ + /* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(sum\\.\[0-9\]+, 0B\\);" 1 "gimple" } } */ + + /* NOTE: Initializer cannot be omp_init_allocator - as 'A' is + in the same scope and the auto-omp_free comes later than + any omp_destroy_allocator. */ + omp_allocator_handle_t my_allocator = omp_low_lat_mem_alloc; + int n = 25; + int A[n]; + #pragma omp allocate(A) align(128) allocator(my_allocator) + /* { dg-final { scan-tree-dump-times "A\\.\[0-9\]+ = __builtin_GOMP_alloc \\(128, _\[0-9\]+, my_allocator\\);" 1 "gimple" } } */ + /* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(A\\.\[0-9\]+, 0B\\);" 1 "gimple" } } */ + + if (((intptr_t)A) % 128 != 0) + __builtin_abort (); + for (int i = 0; i < n; ++i) + A[i] = i; + + omp_alloctrait_t traits[1] = { { omp_atk_alignment, 64 } }; + my_allocator = omp_init_allocator(omp_low_lat_mem_space,1,traits); + { + int B[n] = { }; + int C[5] = {1,2,3,4,5}; + #pragma omp allocate(B,C) allocator(my_allocator) + /* { dg-final { scan-tree-dump-times "B\\.\[0-9\]+ = __builtin_GOMP_alloc \\(\[0-9\]+, _\[0-9\]+, my_allocator\\);" 1 "gimple" } } */ + /* { dg-final { scan-tree-dump-times "C\\.\[0-9\]+ = __builtin_GOMP_alloc \\(\[0-9\]+, 20, my_allocator\\);" 1 "gimple" } } */ + /* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(B\\.\[0-9\]+, 0B\\);" 1 "gimple" } } */ + /* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(C\\.\[0-9\]+, 0B\\);" 1 "gimple" } } */ + + int D[5] = {11,22,33,44,55}; + #pragma omp allocate(D) align(256) + /* { dg-final { scan-tree-dump-times "D\\.\[0-9\]+ = __builtin_GOMP_alloc \\(256, 20, 0B\\);" 1 "gimple" } } */ + /* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(D\\.\[0-9\]+, 0B\\);" 1 "gimple" } } */ + + if (((intptr_t) B) % 64 != 0) + __builtin_abort (); + if (((intptr_t) C) % 64 != 0) + __builtin_abort (); + if (((intptr_t) D) % 64 != 0) + __builtin_abort (); + + for (int i = 0; i < 5; ++i) + { + if (C[i] != i+1) + __builtin_abort (); + if (D[i] != i+1 + 10*(i+1)) + __builtin_abort (); + } + + for (int i = 0; i < n; ++i) + { + if (B[i] != 0) + __builtin_abort (); + sum += A[i]+B[i]+C[i%5]+D[i%5]; + } + } + omp_destroy_allocator (my_allocator); + return sum; +} + +int +main () +{ + if (one () != 1200) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/allocate-5.c b/libgomp/testsuite/libgomp.c/allocate-5.c new file mode 100644 index 00000000000..beaf16440e1 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/allocate-5.c @@ -0,0 +1,126 @@ +/* TODO: move to ../libgomp.c-c++-common once C++ is implemented. */ +/* NOTE: { target c } is unsupported with with the C compiler. */ + +/* { dg-do run } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +#include +#include + +/* { dg-final { scan-tree-dump-not "__builtin_stack_save" "gimple" } } */ +/* { dg-final { scan-tree-dump-not "__builtin_alloca" "gimple" } } */ +/* { dg-final { scan-tree-dump-not "__builtin_stack_restore" "gimple" } } */ + +/* { dg-final { scan-tree-dump-times "__builtin_GOMP_alloc \\(" 5 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(" 5 "gimple" } } */ + +void +one () +{ + int result = 0, n = 3; + #pragma omp target map(tofrom: result) firstprivate(n) + { + int var = 5, var2[n]; + #pragma omp allocate(var,var2) align(128) allocator(omp_low_lat_mem_alloc) +/* { dg-final { scan-tree-dump-times "var\\.\[0-9\]+ = __builtin_GOMP_alloc \\(128, 4, 5\\);" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "var2\\.\[0-9\]+ = __builtin_GOMP_alloc \\(128, D\\.\[0-9\]+, 5\\);" 1 "gimple" } } */ + +/* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(var\\.\[0-9\]+, 0B\\);" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(var2\\.\[0-9\]+, 0B\\);" 1 "gimple" } } */ + + if ((intptr_t) &var % 128 != 0) + __builtin_abort (); + if ((intptr_t) var2 % 128 != 0) + __builtin_abort (); + if (var != 5) + __builtin_abort (); + + #pragma omp parallel for + for (int i = 0; i < n; ++i) + var2[i] = (i+33); + + #pragma omp loop reduction(+:result) + for (int i = 0; i < n; ++i) + result += var + var2[i]; + } + if (result != (3*5 + 33 + 34 + 35)) + __builtin_abort (); +} + +void +two () +{ + struct st { + int a, b; + }; + int scalar = 44, array[5] = {1,2,3,4,5}; + struct st s = {.a=11, .b=56}; + #pragma omp allocate(scalar, array, s) +/* { dg-final { scan-tree-dump-times "scalar\\.\[0-9\]+ = __builtin_GOMP_alloc \\(4, 4, 0B\\);" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "array\\.\[0-9\]+ = __builtin_GOMP_alloc \\(4, 20, 0B\\);" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "s\\.\[0-9\]+ = __builtin_GOMP_alloc \\(4, 8, 0B\\);" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(scalar\\.\[0-9\]+, 0B\\);" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(array\\.\[0-9\]+, 0B\\);" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(s\\.\[0-9\]+, 0B\\);" 1 "gimple" } } */ + + #pragma omp parallel firstprivate(scalar) firstprivate(array) firstprivate(s) + { + if (scalar != 44) + __builtin_abort (); + scalar = 33; + for (int i = 0; i < 5; ++i) + if (array[i] != i+1) + __builtin_abort (); + for (int i = 0; i < 5; ++i) + array[i] = 10*(i+1); + if (s.a != 11 || s.b != 56) + __builtin_abort (); + s.a = 74; + s.b = 674; + } + if (scalar != 44) + __builtin_abort (); + for (int i = 0; i < 5; ++i) + if (array[i] != i+1) + __builtin_abort (); + if (s.a != 11 || s.b != 56) + __builtin_abort (); + + #pragma omp target defaultmap(firstprivate : scalar) defaultmap(none : aggregate) defaultmap(none : pointer) + { + if (scalar != 44) + __builtin_abort (); + scalar = 33; + } + if (scalar != 44) + __builtin_abort (); + + #pragma omp target defaultmap(none : scalar) defaultmap(firstprivate : aggregate) defaultmap(none : pointer) + { + for (int i = 0; i < 5; ++i) + if (array[i] != i+1) + __builtin_abort (); + for (int i = 0; i < 5; ++i) + array[i] = 10*(i+1); + } + for (int i = 0; i < 5; ++i) + if (array[i] != i+1) + __builtin_abort (); + #pragma omp target defaultmap(none : scalar) defaultmap(firstprivate : aggregate) defaultmap(none : pointer) + { + if (s.a != 11 || s.b != 56) + __builtin_abort (); + s.a = 74; + s.b = 674; + } + if (s.a != 11 || s.b != 56) + __builtin_abort (); +} + +int +main () +{ + one (); + two (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/allocate-6.c b/libgomp/testsuite/libgomp.c/allocate-6.c new file mode 100644 index 00000000000..6d7278ce571 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/allocate-6.c @@ -0,0 +1,319 @@ +/* TODO: move to ../libgomp.c-c++-common once C++ is implemented. */ +/* NOTE: { target c } is unsupported with with the C compiler. */ + +/* { dg-do run } */ +/* { dg-additional-options "-fdump-tree-omplower" } */ + +/* For the 4 vars in omp_parallel, 4 in omp_target and 1 of 2 in no_alloc2_func. */ +/* { dg-final { scan-tree-dump-times "__builtin_GOMP_alloc \\(" 9 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(" 9 "omplower" } } */ + +#include + +void +check_int (int *x, int y) +{ + if (*x != y) + __builtin_abort (); +} + +void +check_ptr (int **x, int *y) +{ + if (*x != y) + __builtin_abort (); +} + + +int +no_alloc_func () +{ + /* There is no __builtin_GOMP_alloc / __builtin_GOMP_free as + allocator == omp_default_mem_alloc (known at compile time. */ + int no_alloc; + #pragma omp allocate(no_alloc) allocator(omp_default_mem_alloc) + no_alloc = 7; + return no_alloc; +} + +int +no_alloc2_func() +{ + /* There is no __builtin_GOMP_alloc / __builtin_GOMP_free as + no_alloc2 is TREE_UNUSED. But there is for is_alloc2. */ + int no_alloc2, is_alloc2; + #pragma omp allocate(no_alloc2, is_alloc2) + is_alloc2 = 7; + return is_alloc2; +} + + +void +omp_parallel () +{ + int n = 6; + int iii = 5, jjj[5], kkk[n]; + int *ptr = (int *) 0x1234; + #pragma omp allocate(iii, jjj, kkk, ptr) + + for (int i = 0; i < 5; i++) + jjj[i] = 3*i; + for (int i = 0; i < 6; i++) + kkk[i] = 7*i; + + #pragma omp parallel default(none) firstprivate(iii, jjj, kkk, ptr) if(0) + { + if (iii != 5) + __builtin_abort(); + iii = 7; + check_int (&iii, 7); + for (int i = 0; i < 5; i++) + if (jjj[i] != 3*i) + __builtin_abort (); + for (int i = 0; i < 6; i++) + if (kkk[i] != 7*i) + __builtin_abort (); + for (int i = 0; i < 5; i++) + jjj[i] = 4*i; + for (int i = 0; i < 6; i++) + kkk[i] = 8*i; + for (int i = 0; i < 5; i++) + check_int (&jjj[i], 4*i); + for (int i = 0; i < 6; i++) + check_int (&kkk[i], 8*i); + if (ptr != (int *) 0x1234) + __builtin_abort (); + ptr = (int *) 0xabcd; + if (ptr != (int *) 0xabcd) + __builtin_abort (); + check_ptr (&ptr, (int *) 0xabcd); + } + if (iii != 5) + __builtin_abort (); + check_int (&iii, 5); + for (int i = 0; i < 5; i++) + { + if (jjj[i] != 3*i) + __builtin_abort (); + check_int (&jjj[i], 3*i); + } + for (int i = 0; i < 6; i++) + { + if (kkk[i] != 7*i) + __builtin_abort (); + check_int (&kkk[i], 7*i); + } + if (ptr != (int *) 0x1234) + __builtin_abort (); + check_ptr (&ptr, (int *) 0x1234); + + #pragma omp parallel default(firstprivate) if(0) + { + if (iii != 5) + __builtin_abort(); + iii = 7; + check_int (&iii, 7); + for (int i = 0; i < 5; i++) + if (jjj[i] != 3*i) + __builtin_abort (); + for (int i = 0; i < 6; i++) + if (kkk[i] != 7*i) + __builtin_abort (); + for (int i = 0; i < 5; i++) + jjj[i] = 4*i; + for (int i = 0; i < 6; i++) + kkk[i] = 8*i; + for (int i = 0; i < 5; i++) + check_int (&jjj[i], 4*i); + for (int i = 0; i < 6; i++) + check_int (&kkk[i], 8*i); + if (ptr != (int *) 0x1234) + __builtin_abort (); + ptr = (int *) 0xabcd; + if (ptr != (int *) 0xabcd) + __builtin_abort (); + check_ptr (&ptr, (int *) 0xabcd); + } + if (iii != 5) + __builtin_abort (); + check_int (&iii, 5); + for (int i = 0; i < 5; i++) + { + if (jjj[i] != 3*i) + __builtin_abort (); + check_int (&jjj[i], 3*i); + } + for (int i = 0; i < 6; i++) + { + if (kkk[i] != 7*i) + __builtin_abort (); + check_int (&kkk[i], 7*i); + } + if (ptr != (int *) 0x1234) + __builtin_abort (); + check_ptr (&ptr, (int *) 0x1234); +} + + + +void +omp_target () +{ + int n = 6; + int iii = 5, jjj[5], kkk[n]; + int *ptr = (int *) 0x1234; + #pragma omp allocate(iii, jjj, kkk, ptr) + + for (int i = 0; i < 5; i++) + jjj[i] = 3*i; + for (int i = 0; i < 6; i++) + kkk[i] = 7*i; + + #pragma omp target defaultmap(none) firstprivate(iii, jjj, kkk, ptr) + { + if (iii != 5) + __builtin_abort(); + iii = 7; + check_int (&iii, 7); + for (int i = 0; i < 5; i++) + if (jjj[i] != 3*i) + __builtin_abort (); + for (int i = 0; i < 6; i++) + if (kkk[i] != 7*i) + __builtin_abort (); + for (int i = 0; i < 5; i++) + jjj[i] = 4*i; + for (int i = 0; i < 6; i++) + kkk[i] = 8*i; + for (int i = 0; i < 5; i++) + check_int (&jjj[i], 4*i); + for (int i = 0; i < 6; i++) + check_int (&kkk[i], 8*i); + if (ptr != (int *) 0x1234) + __builtin_abort (); + ptr = (int *) 0xabcd; + if (ptr != (int *) 0xabcd) + __builtin_abort (); + check_ptr (&ptr, (int *) 0xabcd); + } + if (iii != 5) + __builtin_abort (); + check_int (&iii, 5); + for (int i = 0; i < 5; i++) + { + if (jjj[i] != 3*i) + __builtin_abort (); + check_int (&jjj[i], 3*i); + } + for (int i = 0; i < 6; i++) + { + if (kkk[i] != 7*i) + __builtin_abort (); + check_int (&kkk[i], 7*i); + } + if (ptr != (int *) 0x1234) + __builtin_abort (); + check_ptr (&ptr, (int *) 0x1234); + + #pragma omp target defaultmap(firstprivate) + { + if (iii != 5) + __builtin_abort(); + iii = 7; + check_int (&iii, 7); + for (int i = 0; i < 5; i++) + if (jjj[i] != 3*i) + __builtin_abort (); + for (int i = 0; i < 6; i++) + if (kkk[i] != 7*i) + __builtin_abort (); + for (int i = 0; i < 5; i++) + jjj[i] = 4*i; + for (int i = 0; i < 6; i++) + kkk[i] = 8*i; + for (int i = 0; i < 5; i++) + check_int (&jjj[i], 4*i); + for (int i = 0; i < 6; i++) + check_int (&kkk[i], 8*i); + if (ptr != (int *) 0x1234) + __builtin_abort (); + ptr = (int *) 0xabcd; + if (ptr != (int *) 0xabcd) + __builtin_abort (); + check_ptr (&ptr, (int *) 0xabcd); + } + if (iii != 5) + __builtin_abort (); + check_int (&iii, 5); + for (int i = 0; i < 5; i++) + { + if (jjj[i] != 3*i) + __builtin_abort (); + check_int (&jjj[i], 3*i); + } + for (int i = 0; i < 6; i++) + { + if (kkk[i] != 7*i) + __builtin_abort (); + check_int (&kkk[i], 7*i); + } + if (ptr != (int *) 0x1234) + __builtin_abort (); + check_ptr (&ptr, (int *) 0x1234); + + #pragma omp target defaultmap(tofrom) + { + if (iii != 5) + __builtin_abort(); + iii = 7; + check_int (&iii, 7); + for (int i = 0; i < 5; i++) + if (jjj[i] != 3*i) + __builtin_abort (); + for (int i = 0; i < 6; i++) + if (kkk[i] != 7*i) + __builtin_abort (); + for (int i = 0; i < 5; i++) + jjj[i] = 4*i; + for (int i = 0; i < 6; i++) + kkk[i] = 8*i; + for (int i = 0; i < 5; i++) + check_int (&jjj[i], 4*i); + for (int i = 0; i < 6; i++) + check_int (&kkk[i], 8*i); + if (ptr != (int *) 0x1234) + __builtin_abort (); + ptr = (int *) 0xabcd; + if (ptr != (int *) 0xabcd) + __builtin_abort (); + check_ptr (&ptr, (int *) 0xabcd); + } + + if (iii != 7) + __builtin_abort (); + check_int (&iii, 7); + for (int i = 0; i < 5; i++) + { + if (jjj[i] != 4*i) + __builtin_abort (); + check_int (&jjj[i], 4*i); + } + for (int i = 0; i < 6; i++) + { + if (kkk[i] != 8*i) + __builtin_abort (); + check_int (&kkk[i], 8*i); + } + if (ptr != (int *) 0xabcd) + __builtin_abort (); + check_ptr (&ptr, (int *) 0xabcd); +} + + +int +main () +{ + omp_parallel (); + omp_target (); + return 0; +}