From patchwork Tue Nov 26 14:49:21 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 1201095 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-514614-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=mentor.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="oDIlEhAY"; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 47Mmyg1gLyz9sP3 for ; Wed, 27 Nov 2019 01:49:47 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :reply-to:from:subject:to:cc:message-id:date:mime-version :content-type; q=dns; s=default; b=cW4UvaVMmuYQBpgmSB7rvCBlNhpWr N2QgK5SH6sK+yoljEIUMa2+V+pgJzT4xIiWjyRY564uh9sepcF8GMpygJOqaGn7q av76rHkYbna5nIIF9jYleb79Ah4uAFESVxGHO08OcrUWmDfO8wwdYgDQneK70xFD h6TfQvIiQ9IsOk= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :reply-to:from:subject:to:cc:message-id:date:mime-version :content-type; s=default; bh=tkWqGV5w90ZxbCfX7dx2GuUoXyQ=; b=oDI lEhAYZnrmsuQl3jpBCEev2CUmCPIzGoHNNQ2mdydN/giHI/b1X1ha1GEFc5MIE02 fXagU/ih+mHQPbLHH2rPhCYq/HhdSeUGqAoEZQWrN3h8Uhs1rzB8CQ1fgibfgvdu RuXWRXeck9x00VKD3rexRtI7QidjOF3sEmZX6g6Y= Received: (qmail 36423 invoked by alias); 26 Nov 2019 14:49:38 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 36415 invoked by uid 89); 26 Nov 2019 14:49:37 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-12.6 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, KAM_ASCII_DIVIDERS, SPF_PASS autolearn=ham version=3.3.1 spammy= X-HELO: esa3.mentor.iphmx.com Received: from esa3.mentor.iphmx.com (HELO esa3.mentor.iphmx.com) (68.232.137.180) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Tue, 26 Nov 2019 14:49:30 +0000 IronPort-SDR: C2cHtJeLNah1VAm7bxeqfmYFaF0769CSiwP+e5acnDt+NpYVytDEDEcxEgh0ofCASwTXBoODAj qHda7zZOtedCqf/ImCHoRWHs569Idz8yYQoEeShkWdm5xnU7132a0jNPoirCQJeEfYG6PsJ4px 7OJeK6GcyVOc6MzgSSgzPxHz24M9Lxlmjx6m31ADgQXYYpuSjO1Tu02s2VSvIu0MNgQsGQXtoS 9mVIzYH1qt+LWGLo+0ASpoy3TrEvtgOih75VH5tbox9ZaO3aDZLWMeSWSeaOkXpGodB7HRHaI5 alk= Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 26 Nov 2019 06:49:28 -0800 IronPort-SDR: ayERaf+bEHVpF3YktRCbIqTQliieZq8N8PXWgaeElDl+8FkzNTW5zeTxhZT2XZaMncOFfbEshM Xin0Rj8z7t1ZYhW2+np7j34T6ebNTaRSsNSVwh7bNisCXsPQl4RiRqXNoJRaL4kNw488p1kG+v pHM6h2j1l342F62du5d85yrFSicv0VmoFVwpTudIno0VpoBGPAkj+8OdB9Q7fPtvHkSseDK6HZ in7SqGJoesvWwVJc8GcqunVR9ALYBBQvyE6jwk8N7xl6qE7CsjJ0KqF5yd9NrBE0kXEsYfnpmd dPg= Reply-To: From: Chung-Lin Tang Subject: [PATCH, OpenACC, v3] Non-contiguous array support for OpenACC data clauses To: gcc-patches , Thomas Schwinge CC: Jakub Jelinek Message-ID: <6b17767f-9b7d-50a7-c52f-e80bf1be991e@mentor.com> Date: Tue, 26 Nov 2019 22:49:21 +0800 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.13; rv:60.0) Gecko/20100101 Thunderbird/60.9.1 MIME-Version: 1.0 Hi Thomas, this is a reorg of the last non-contiguous arrays patch. You'll notice that: (1) A large part of the code has been pulled into oacc-parallel.c, with most of the data structure declarations in oacc-int.h. (2) target.c only contains relatively little code from gomp_map_vars_internal that processes what GOACC_parallel_keyed/data_start gives it. (3) Instead of directly passed in the map pointer, the array descriptor pointers are now passed to GOACC_parallel_keyed/data_start using varargs. (I believe the adding of '...' to GOACC_data_start does not break any compatiblity) (4) Along the way, I've added a 'gomp_map_vars_openacc' for specializing our uses, which should shave off quite some code through inlining. The GOMP_MAP_NONCONTIG_ARRAY_P maps are still placed at the beginning of the recieved map sequence in this patch. It should still be relatively easy to use a GOACC_FLAG_* to do so if deemed better before committing. Thanks, Chung-Lin PR other/76739 gcc/c/ * c-typeck.c (handle_omp_array_sections_1): Add 'bool &non_contiguous' parameter, adjust recursive call site, add cases for allowing pointer based multi-dimensional arrays for OpenACC. (handle_omp_array_sections): Adjust handle_omp_array_sections_1 call, handle non-contiguous case to create dynamic array map. gcc/cp/ * semantics.c (handle_omp_array_sections_1): Add 'bool &non_contiguous' parameter, adjust recursive call site, add cases for allowing pointer based multi-dimensional arrays for OpenACC. (handle_omp_array_sections): Adjust handle_omp_array_sections_1 call, handle non-contiguous case to create dynamic array map. gcc/fortran/ * f95-lang.c (DEF_FUNCTION_TYPE_VAR_5): New symbol. * types.def (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR): New type. gcc/ * builtin-types.def (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR): New type. * omp-builtins.def (BUILT_IN_GOACC_DATA_START): Adjust function type to new BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR. * gimplify.c (gimplify_scan_omp_clauses): Skip gimplification of OMP_CLAUSE_SIZE of non-contiguous array maps (which is a TREE_LIST). * omp-expand.c (expand_omp_target): Add non-contiguous array descriptor pointers to variadic arguments. * omp-low.c (append_field_to_record_type): New function. (create_noncontig_array_descr_type): Likewise. (create_noncontig_array_descr_init_code): Likewise. (scan_sharing_clauses): For non-contiguous array map kinds, check for supported dimension structure, and install non-contiguous array variable into current omp_context. (reorder_noncontig_array_clauses): New function. (scan_omp_target): Call reorder_noncontig_array_clauses to place non-contiguous array map clauses at beginning of clause sequence. (lower_omp_target): Add handling for non-contiguous array map kinds, add all created non-contiguous array descriptors to gimple_omp_target_data_arg. gcc/testsuite/ * c-c++-common/goacc/noncontig_array-1.c: New test. libgomp/ * libgomp_g.h (GOACC_data_start): Add variadic '...' to declaration. * libgomp.h (gomp_map_vars_openacc): New function declaration. * oacc-int.h (struct goacc_ncarray_dim): New struct declaration. (struct goacc_ncarray_descr_type): Likewise. (struct goacc_ncarray): Likewise. (struct goacc_ncarray_info): Likewise. (goacc_noncontig_array_create_ptrblock): New function declaration. * oacc-parallel.c (goacc_noncontig_array_count_rows): New function. (goacc_noncontig_array_compute_sizes): Likewise. (goacc_noncontig_array_fill_rows_1): Likewise. (goacc_noncontig_array_fill_rows): Likewise. (goacc_process_noncontiguous_arrays): Likewise. (goacc_noncontig_array_create_ptrblock): Likewise. (GOACC_parallel_keyed): Use goacc_process_noncontiguous_arrays to handle non-contiguous array descriptors at end of varargs, adjust to use gomp_map_vars_openacc. (GOACC_data_start): Likewise. Adjust function type to accept varargs. * target.c (gomp_map_vars_internal): Add struct goacc_ncarray_info * nca_info parameter, add handling code for non-contiguous arrays. (gomp_map_vars_openacc): Add new function for specialization of gomp_map_vars_internal for OpenACC structured region usage. * testsuite/libgomp.oacc-c-c++-common/noncontig_array-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/noncontig_array-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/noncontig_array-3.c: New test. * testsuite/libgomp.oacc-c-c++-common/noncontig_array-4.c: New test. * testsuite/libgomp.oacc-c-c++-common/noncontig_array-utils.h: Support header for new tests. Index: include/gomp-constants.h =================================================================== --- include/gomp-constants.h (revision 278656) +++ include/gomp-constants.h (working copy) @@ -40,6 +40,7 @@ #define GOMP_MAP_FLAG_SPECIAL_0 (1 << 2) #define GOMP_MAP_FLAG_SPECIAL_1 (1 << 3) #define GOMP_MAP_FLAG_SPECIAL_2 (1 << 4) +#define GOMP_MAP_FLAG_SPECIAL_3 (1 << 5) #define GOMP_MAP_FLAG_SPECIAL (GOMP_MAP_FLAG_SPECIAL_1 \ | GOMP_MAP_FLAG_SPECIAL_0) /* Flag to force a specific behavior (or else, trigger a run-time error). */ @@ -127,6 +128,26 @@ enum gomp_map_kind /* Decrement usage count and deallocate if zero. */ GOMP_MAP_RELEASE = (GOMP_MAP_FLAG_SPECIAL_2 | GOMP_MAP_DELETE), + /* Mapping kinds for non-contiguous arrays. */ + GOMP_MAP_NONCONTIG_ARRAY = (GOMP_MAP_FLAG_SPECIAL_3), + GOMP_MAP_NONCONTIG_ARRAY_TO = (GOMP_MAP_NONCONTIG_ARRAY + | GOMP_MAP_TO), + GOMP_MAP_NONCONTIG_ARRAY_FROM = (GOMP_MAP_NONCONTIG_ARRAY + | GOMP_MAP_FROM), + GOMP_MAP_NONCONTIG_ARRAY_TOFROM = (GOMP_MAP_NONCONTIG_ARRAY + | GOMP_MAP_TOFROM), + GOMP_MAP_NONCONTIG_ARRAY_FORCE_TO = (GOMP_MAP_NONCONTIG_ARRAY_TO + | GOMP_MAP_FLAG_FORCE), + GOMP_MAP_NONCONTIG_ARRAY_FORCE_FROM = (GOMP_MAP_NONCONTIG_ARRAY_FROM + | GOMP_MAP_FLAG_FORCE), + GOMP_MAP_NONCONTIG_ARRAY_FORCE_TOFROM = (GOMP_MAP_NONCONTIG_ARRAY_TOFROM + | GOMP_MAP_FLAG_FORCE), + GOMP_MAP_NONCONTIG_ARRAY_ALLOC = (GOMP_MAP_NONCONTIG_ARRAY + | GOMP_MAP_ALLOC), + GOMP_MAP_NONCONTIG_ARRAY_FORCE_ALLOC = (GOMP_MAP_NONCONTIG_ARRAY + | GOMP_MAP_FORCE_ALLOC), + GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT = (GOMP_MAP_NONCONTIG_ARRAY + | GOMP_MAP_FORCE_PRESENT), /* Internal to GCC, not used in libgomp. */ /* Do not map, but pointer assign a pointer instead. */ @@ -155,6 +176,8 @@ enum gomp_map_kind #define GOMP_MAP_ALWAYS_P(X) \ (GOMP_MAP_ALWAYS_TO_P (X) || ((X) == GOMP_MAP_ALWAYS_FROM)) +#define GOMP_MAP_NONCONTIG_ARRAY_P(X) \ + ((X) & GOMP_MAP_NONCONTIG_ARRAY) /* Asynchronous behavior. Keep in sync with libgomp/{openacc.h,openacc.f90,openacc_lib.h}:acc_async_t. */ Index: gcc/builtin-types.def =================================================================== --- gcc/builtin-types.def (revision 278656) +++ gcc/builtin-types.def (working copy) @@ -822,6 +822,9 @@ DEF_FUNCTION_TYPE_VAR_5 (BT_FN_INT_STRING_SIZE_INT DEF_FUNCTION_TYPE_VAR_5 (BT_FN_INT_INT_INT_INT_INT_INT_VAR, BT_INT, BT_INT, BT_INT, BT_INT, BT_INT, BT_INT) +DEF_FUNCTION_TYPE_VAR_5 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR, + BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR, BT_PTR) + DEF_FUNCTION_TYPE_VAR_6 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR, BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, BT_PTR, BT_PTR, BT_PTR) Index: gcc/c/c-typeck.c =================================================================== --- gcc/c/c-typeck.c (revision 278656) +++ gcc/c/c-typeck.c (working copy) @@ -12876,12 +12876,14 @@ c_finish_omp_cancellation_point (location_t loc, t <= FIRST_NON_ONE we diagnose non-contiguous arrays if low bound isn't 0 or length isn't the array domain max + 1, for > FIRST_NON_ONE we can if MAYBE_ZERO_LEN is false. MAYBE_ZERO_LEN will be true in the above - case though, as some lengths could be zero. */ + case though, as some lengths could be zero. + NON_CONTIGUOUS will be true if this is an OpenACC non-contiguous array + section. */ static tree handle_omp_array_sections_1 (tree c, tree t, vec &types, bool &maybe_zero_len, unsigned int &first_non_one, - enum c_omp_region_type ort) + bool &non_contiguous, enum c_omp_region_type ort) { tree ret, low_bound, length, type; if (TREE_CODE (t) != TREE_LIST) @@ -12966,7 +12968,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec types; tree *tp = &OMP_CLAUSE_DECL (c); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND @@ -13218,7 +13229,7 @@ handle_omp_array_sections (tree c, enum c_omp_regi tp = &TREE_VALUE (*tp); tree first = handle_omp_array_sections_1 (c, *tp, types, maybe_zero_len, first_non_one, - ort); + non_contiguous, ort); if (first == error_mark_node) return true; if (first == NULL_TREE) @@ -13251,6 +13262,7 @@ handle_omp_array_sections (tree c, enum c_omp_regi unsigned int num = types.length (), i; tree t, side_effects = NULL_TREE, size = NULL_TREE; tree condition = NULL_TREE; + tree ncarray_dims = NULL_TREE; if (int_size_in_bytes (TREE_TYPE (first)) <= 0) maybe_zero_len = true; @@ -13274,6 +13286,13 @@ handle_omp_array_sections (tree c, enum c_omp_regi length = fold_convert (sizetype, length); if (low_bound == NULL_TREE) low_bound = integer_zero_node; + + if (non_contiguous) + { + ncarray_dims = tree_cons (low_bound, length, ncarray_dims); + continue; + } + if (!maybe_zero_len && i > first_non_one) { if (integer_nonzerop (low_bound)) @@ -13370,6 +13389,14 @@ handle_omp_array_sections (tree c, enum c_omp_regi size = size_binop (MULT_EXPR, size, l); } } + if (non_contiguous) + { + int kind = OMP_CLAUSE_MAP_KIND (c); + OMP_CLAUSE_SET_MAP_KIND (c, kind | GOMP_MAP_NONCONTIG_ARRAY); + OMP_CLAUSE_DECL (c) = t; + OMP_CLAUSE_SIZE (c) = ncarray_dims; + return false; + } if (side_effects) size = build2 (COMPOUND_EXPR, sizetype, side_effects, size); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION Index: gcc/cp/semantics.c =================================================================== --- gcc/cp/semantics.c (revision 278656) +++ gcc/cp/semantics.c (working copy) @@ -4735,12 +4735,14 @@ omp_privatize_field (tree t, bool shared) <= FIRST_NON_ONE we diagnose non-contiguous arrays if low bound isn't 0 or length isn't the array domain max + 1, for > FIRST_NON_ONE we can if MAYBE_ZERO_LEN is false. MAYBE_ZERO_LEN will be true in the above - case though, as some lengths could be zero. */ + case though, as some lengths could be zero. + NON_CONTIGUOUS will be true if this is an OpenACC non-contiguous array + section. */ static tree handle_omp_array_sections_1 (tree c, tree t, vec &types, bool &maybe_zero_len, unsigned int &first_non_one, - enum c_omp_region_type ort) + bool &non_contiguous, enum c_omp_region_type ort) { tree ret, low_bound, length, type; if (TREE_CODE (t) != TREE_LIST) @@ -4825,7 +4827,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec types; tree *tp = &OMP_CLAUSE_DECL (c); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND @@ -5100,7 +5111,7 @@ handle_omp_array_sections (tree c, enum c_omp_regi tp = &TREE_VALUE (*tp); tree first = handle_omp_array_sections_1 (c, *tp, types, maybe_zero_len, first_non_one, - ort); + non_contiguous, ort); if (first == error_mark_node) return true; if (first == NULL_TREE) @@ -5134,6 +5145,7 @@ handle_omp_array_sections (tree c, enum c_omp_regi unsigned int num = types.length (), i; tree t, side_effects = NULL_TREE, size = NULL_TREE; tree condition = NULL_TREE; + tree ncarray_dims = NULL_TREE; if (int_size_in_bytes (TREE_TYPE (first)) <= 0) maybe_zero_len = true; @@ -5159,6 +5171,13 @@ handle_omp_array_sections (tree c, enum c_omp_regi length = fold_convert (sizetype, length); if (low_bound == NULL_TREE) low_bound = integer_zero_node; + + if (non_contiguous) + { + ncarray_dims = tree_cons (low_bound, length, ncarray_dims); + continue; + } + if (!maybe_zero_len && i > first_non_one) { if (integer_nonzerop (low_bound)) @@ -5250,6 +5269,14 @@ handle_omp_array_sections (tree c, enum c_omp_regi } if (!processing_template_decl) { + if (non_contiguous) + { + int kind = OMP_CLAUSE_MAP_KIND (c); + OMP_CLAUSE_SET_MAP_KIND (c, kind | GOMP_MAP_NONCONTIG_ARRAY); + OMP_CLAUSE_DECL (c) = t; + OMP_CLAUSE_SIZE (c) = ncarray_dims; + return false; + } if (side_effects) size = build2 (COMPOUND_EXPR, sizetype, side_effects, size); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION Index: gcc/fortran/f95-lang.c =================================================================== --- gcc/fortran/f95-lang.c (revision 278656) +++ gcc/fortran/f95-lang.c (working copy) @@ -632,6 +632,8 @@ gfc_init_builtin_functions (void) #define DEF_FUNCTION_TYPE_VAR_0(NAME, RETURN) NAME, #define DEF_FUNCTION_TYPE_VAR_1(NAME, RETURN, ARG1) NAME, #define DEF_FUNCTION_TYPE_VAR_2(NAME, RETURN, ARG1, ARG2) NAME, +#define DEF_FUNCTION_TYPE_VAR_5(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5) \ + NAME, #define DEF_FUNCTION_TYPE_VAR_6(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \ ARG6) NAME, #define DEF_FUNCTION_TYPE_VAR_7(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \ Index: gcc/fortran/types.def =================================================================== --- gcc/fortran/types.def (revision 278656) +++ gcc/fortran/types.def (working copy) @@ -269,6 +269,9 @@ DEF_FUNCTION_TYPE_VAR_7 (BT_FN_VOID_INT_SIZE_PTR_P BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR, BT_PTR, BT_INT, BT_INT) +DEF_FUNCTION_TYPE_VAR_5 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR, + BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR, BT_PTR) + DEF_FUNCTION_TYPE_VAR_6 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR, BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, BT_PTR, BT_PTR, BT_PTR) Index: gcc/gimplify.c =================================================================== --- gcc/gimplify.c (revision 278656) +++ gcc/gimplify.c (working copy) @@ -8641,9 +8641,17 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_se if (OMP_CLAUSE_SIZE (c) == NULL_TREE) OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl) : TYPE_SIZE_UNIT (TREE_TYPE (decl)); - if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, - NULL, is_gimple_val, fb_rvalue) == GS_ERROR) + if (GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c))) { + gcc_assert (OMP_CLAUSE_SIZE (c) + && TREE_CODE (OMP_CLAUSE_SIZE (c)) == TREE_LIST); + /* For non-contiguous array maps, OMP_CLAUSE_SIZE is a TREE_LIST + of the individual array dimensions, which gimplify_expr doesn't + handle, so skip the call to gimplify_expr here. */ + } + else if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, + NULL, is_gimple_val, fb_rvalue) == GS_ERROR) + { remove = true; break; } Index: gcc/omp-builtins.def =================================================================== --- gcc/omp-builtins.def (revision 278656) +++ gcc/omp-builtins.def (working copy) @@ -32,7 +32,7 @@ along with GCC; see the file COPYING3. If not see DEF_GOACC_BUILTIN (BUILT_IN_ACC_GET_DEVICE_TYPE, "acc_get_device_type", BT_FN_INT, ATTR_NOTHROW_LIST) DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_START, "GOACC_data_start", - BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST) + BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR, ATTR_NOTHROW_LIST) DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_END, "GOACC_data_end", BT_FN_VOID, ATTR_NOTHROW_LIST) DEF_GOACC_BUILTIN (BUILT_IN_GOACC_ENTER_EXIT_DATA, "GOACC_enter_exit_data", Index: gcc/omp-expand.c =================================================================== --- gcc/omp-expand.c (revision 278656) +++ gcc/omp-expand.c (working copy) @@ -8458,6 +8458,21 @@ expand_omp_target (struct omp_region *region) if (tagging) /* Push terminal marker - zero. */ args.safe_push (oacc_launch_pack (0, NULL_TREE, 0)); + + /* We assume index >= 3 in gimple_omp_target_data_arg are non-contiguous + array descriptor pointer arguments. */ + if (TREE_VEC_LENGTH (t) > 3 + && (start_ix == BUILT_IN_GOACC_DATA_START + || start_ix == BUILT_IN_GOACC_PARALLEL)) + { + gcc_assert ((c = omp_find_clause (clauses, OMP_CLAUSE_MAP)) + && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c))); + //if (tagging) + //args.safe_push (oacc_launch_pack (GOMP_LAUNCH_NONCONTIG_ARRAYS, + // NULL_TREE, 0)); + for (int i = 3; i < TREE_VEC_LENGTH (t); i++) + args.safe_push (TREE_VEC_ELT (t, i)); + } g = gimple_build_call_vec (builtin_decl_explicit (start_ix), args); gimple_set_location (g, gimple_location (entry_stmt)); Index: gcc/omp-low.c =================================================================== --- gcc/omp-low.c (revision 278656) +++ gcc/omp-low.c (working copy) @@ -903,6 +903,141 @@ omp_copy_decl (tree var, copy_body_data *cb) return error_mark_node; } +/* Helper function for create_noncontig_array_descr_type(), to append a new field + to a record type. */ + +static void +append_field_to_record_type (tree record_type, tree fld_ident, tree fld_type) +{ + tree *p, fld = build_decl (UNKNOWN_LOCATION, FIELD_DECL, fld_ident, fld_type); + DECL_CONTEXT (fld) = record_type; + + for (p = &TYPE_FIELDS (record_type); *p; p = &DECL_CHAIN (*p)) + ; + *p = fld; +} + +/* Create type for non-contiguous array descriptor. Returns created type, and + returns the number of dimensions in *DIM_NUM. */ + +static tree +create_noncontig_array_descr_type (tree decl, tree dims, int *dim_num) +{ + int n = 0; + tree array_descr_type, name, x; + gcc_assert (TREE_CODE (dims) == TREE_LIST); + + array_descr_type = lang_hooks.types.make_type (RECORD_TYPE); + name = create_tmp_var_name (".omp_noncontig_array_descr_type"); + name = build_decl (UNKNOWN_LOCATION, TYPE_DECL, name, array_descr_type); + DECL_ARTIFICIAL (name) = 1; + DECL_NAMELESS (name) = 1; + TYPE_NAME (array_descr_type) = name; + TYPE_ARTIFICIAL (array_descr_type) = 1; + + /* Main starting pointer/array. */ + /* + tree main_var_type = TREE_TYPE (decl); + if (TREE_CODE (main_var_type) == REFERENCE_TYPE) + main_var_type = TREE_TYPE (main_var_type); + append_field_to_record_type (array_descr_type, DECL_NAME (decl), + (TREE_CODE (TREE_TYPE (decl)) == POINTER_TYPE + ? main_var_type + : build_pointer_type (main_var_type))); + */ + + /* Number of dimensions. */ + append_field_to_record_type (array_descr_type, get_identifier ("__dim_num"), + sizetype); + + for (x = dims; x; x = TREE_CHAIN (x), n++) + { + char *fldname; + /* One for the start index. */ + ASM_FORMAT_PRIVATE_NAME (fldname, "__dim_base", n); + append_field_to_record_type (array_descr_type, get_identifier (fldname), + sizetype); + /* One for the length. */ + ASM_FORMAT_PRIVATE_NAME (fldname, "__dim_length", n); + append_field_to_record_type (array_descr_type, get_identifier (fldname), + sizetype); + /* One for the element size. */ + ASM_FORMAT_PRIVATE_NAME (fldname, "__dim_elem_size", n); + append_field_to_record_type (array_descr_type, get_identifier (fldname), + sizetype); + /* One for is_array flag. */ + ASM_FORMAT_PRIVATE_NAME (fldname, "__dim_is_array", n); + append_field_to_record_type (array_descr_type, get_identifier (fldname), + sizetype); + } + + layout_type (array_descr_type); + *dim_num = n; + return array_descr_type; +} + +/* Generate code sequence for initializing non-contiguous array descriptor. */ + +static void +create_noncontig_array_descr_init_code (tree array_descr, tree array_var, + tree dimensions, int dim_num, + gimple_seq *ilist) +{ + tree fld, fldref; + tree array_descr_type = TREE_TYPE (array_descr); + tree dim_type = TREE_TYPE (array_var); + + fld = TYPE_FIELDS (array_descr_type); + /* + fldref = omp_build_component_ref (array_descr, fld); + gimplify_assign (fldref, (TREE_CODE (dim_type) == ARRAY_TYPE + ? build_fold_addr_expr (array_var) : array_var), + ilist); + */ + if (TREE_CODE (dim_type) == REFERENCE_TYPE) + dim_type = TREE_TYPE (dim_type); + + //fld = TREE_CHAIN (fld); + fldref = omp_build_component_ref (array_descr, fld); + gimplify_assign (fldref, build_int_cst (sizetype, dim_num), ilist); + + while (dimensions) + { + tree dim_base = fold_convert (sizetype, TREE_PURPOSE (dimensions)); + tree dim_length = fold_convert (sizetype, TREE_VALUE (dimensions)); + tree dim_elem_size = TYPE_SIZE_UNIT (TREE_TYPE (dim_type)); + tree dim_is_array = (TREE_CODE (dim_type) == ARRAY_TYPE + ? integer_one_node : integer_zero_node); + /* Set base. */ + fld = TREE_CHAIN (fld); + fldref = omp_build_component_ref (array_descr, fld); + dim_base = fold_build2 (MULT_EXPR, sizetype, dim_base, dim_elem_size); + gimplify_assign (fldref, dim_base, ilist); + + /* Set length. */ + fld = TREE_CHAIN (fld); + fldref = omp_build_component_ref (array_descr, fld); + dim_length = fold_build2 (MULT_EXPR, sizetype, dim_length, dim_elem_size); + gimplify_assign (fldref, dim_length, ilist); + + /* Set elem_size. */ + fld = TREE_CHAIN (fld); + fldref = omp_build_component_ref (array_descr, fld); + dim_elem_size = fold_convert (sizetype, dim_elem_size); + gimplify_assign (fldref, dim_elem_size, ilist); + + /* Set is_array flag. */ + fld = TREE_CHAIN (fld); + fldref = omp_build_component_ref (array_descr, fld); + dim_is_array = fold_convert (sizetype, dim_is_array); + gimplify_assign (fldref, dim_is_array, ilist); + + dimensions = TREE_CHAIN (dimensions); + dim_type = TREE_TYPE (dim_type); + } + gcc_assert (TREE_CHAIN (fld) == NULL_TREE); +} + /* Create a new context, with OUTER_CTX being the surrounding context. */ static omp_context * @@ -1385,6 +1520,38 @@ scan_sharing_clauses (tree clauses, omp_context *c install_var_local (decl, ctx); break; } + + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c))) + { + tree array_decl = OMP_CLAUSE_DECL (c); + tree array_type = TREE_TYPE (array_decl); + bool by_ref = (TREE_CODE (array_type) == ARRAY_TYPE + ? true : false); + + /* Checking code to ensure we only have arrays at top dimension. + This limitation might be lifted in the future. See PR76639. */ + if (TREE_CODE (array_type) == REFERENCE_TYPE) + array_type = TREE_TYPE (array_type); + tree t = array_type, prev_t = NULL_TREE; + while (t) + { + if (TREE_CODE (t) == ARRAY_TYPE && prev_t) + { + error_at (gimple_location (ctx->stmt), "array types are" + " only allowed at outermost dimension of" + " non-contiguous array"); + break; + } + prev_t = t; + t = TREE_TYPE (t); + } + + install_var_field (array_decl, by_ref, 3, ctx); + install_var_local (array_decl, ctx); + break; + } + if (DECL_P (decl)) { if (DECL_SIZE (decl) @@ -2697,6 +2864,50 @@ scan_omp_single (gomp_single *stmt, omp_context *o layout_type (ctx->record_type); } +/* Reorder clauses so that non-contiguous array map clauses are placed at the very + front of the chain. */ + +static void +reorder_noncontig_array_clauses (tree *clauses_ptr) +{ + tree c, clauses = *clauses_ptr; + tree prev_clause = NULL_TREE, next_clause; + tree array_clauses = NULL_TREE, array_clauses_tail = NULL_TREE; + + for (c = clauses; c; c = next_clause) + { + next_clause = OMP_CLAUSE_CHAIN (c); + + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c))) + { + /* Unchain c from clauses. */ + if (c == clauses) + clauses = next_clause; + + /* Link on to array_clauses. */ + if (array_clauses_tail) + OMP_CLAUSE_CHAIN (array_clauses_tail) = c; + else + array_clauses = c; + array_clauses_tail = c; + + if (prev_clause) + OMP_CLAUSE_CHAIN (prev_clause) = next_clause; + continue; + } + + prev_clause = c; + } + + /* Place non-contiguous array clauses at the start of the clause list. */ + if (array_clauses) + { + OMP_CLAUSE_CHAIN (array_clauses_tail) = clauses; + *clauses_ptr = array_clauses; + } +} + /* Scan a GIMPLE_OMP_TARGET. */ static void @@ -2705,7 +2916,6 @@ scan_omp_target (gomp_target *stmt, omp_context *o omp_context *ctx; tree name; bool offloaded = is_gimple_omp_offloaded (stmt); - tree clauses = gimple_omp_target_clauses (stmt); ctx = new_omp_context (stmt, outer_ctx); ctx->field_map = splay_tree_new (splay_tree_compare_pointers, 0, 0); @@ -2724,6 +2934,14 @@ scan_omp_target (gomp_target *stmt, omp_context *o gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn); } + /* If is OpenACC construct, put non-contiguous array clauses (if any) + in front of clause chain. The runtime can then test the first to see + if the additional map processing for them is required. */ + if (is_gimple_omp_oacc (stmt)) + reorder_noncontig_array_clauses (gimple_omp_target_clauses_ptr (stmt)); + + tree clauses = gimple_omp_target_clauses (stmt); + scan_sharing_clauses (clauses, ctx); scan_omp (gimple_omp_body_ptr (stmt), ctx); @@ -11438,6 +11656,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp case GOMP_MAP_FORCE_PRESENT: case GOMP_MAP_FORCE_DEVICEPTR: case GOMP_MAP_DEVICE_RESIDENT: + case GOMP_MAP_NONCONTIG_ARRAY_TO: + case GOMP_MAP_NONCONTIG_ARRAY_FROM: + case GOMP_MAP_NONCONTIG_ARRAY_TOFROM: + case GOMP_MAP_NONCONTIG_ARRAY_FORCE_TO: + case GOMP_MAP_NONCONTIG_ARRAY_FORCE_FROM: + case GOMP_MAP_NONCONTIG_ARRAY_FORCE_TOFROM: + case GOMP_MAP_NONCONTIG_ARRAY_ALLOC: + case GOMP_MAP_NONCONTIG_ARRAY_FORCE_ALLOC: + case GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT: case GOMP_MAP_LINK: gcc_assert (is_gimple_omp_oacc (stmt)); break; @@ -11500,7 +11727,14 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp if (offloaded && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_IN_REDUCTION (c))) { - x = build_receiver_ref (var, true, ctx); + tree var_type = TREE_TYPE (var); + bool rcv_by_ref = + (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c)) + && TREE_CODE (var_type) != ARRAY_TYPE + ? false : true); + + x = build_receiver_ref (var, rcv_by_ref, ctx); tree new_var = lookup_decl (var, ctx); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP @@ -11674,6 +11908,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp vec_alloc (vkind, map_cnt); unsigned int map_idx = 0; + vec nca_descrs = vNULL; + for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) switch (OMP_CLAUSE_CODE (c)) { @@ -11750,6 +11986,28 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp avar = build_fold_addr_expr (avar); gimplify_assign (x, avar, &ilist); } + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c))) + { + int dim_num; + tree dimensions = OMP_CLAUSE_SIZE (c); + + tree array_descr_type = + create_noncontig_array_descr_type (OMP_CLAUSE_DECL (c), + dimensions, &dim_num); + tree array_descr = + create_tmp_var_raw (array_descr_type, + ".omp_noncontig_array_descr"); + gimple_add_tmp_var (array_descr); + + create_noncontig_array_descr_init_code + (array_descr, ovar, dimensions, dim_num, &ilist); + nca_descrs.safe_push (build_fold_addr_expr (array_descr)); + + gimplify_assign (x, (TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE + ? build_fold_addr_expr (ovar) : ovar), + &ilist); + } else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE) { gcc_assert (is_gimple_omp_oacc (ctx->stmt)); @@ -11821,6 +12079,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp s = TREE_TYPE (s); s = TYPE_SIZE_UNIT (s); } + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c))) + s = NULL_TREE; else s = OMP_CLAUSE_SIZE (c); if (s == NULL_TREE) @@ -12072,6 +12333,19 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp gcc_assert (map_idx == map_cnt); + unsigned nca_num = nca_descrs.length (); + if (nca_num > 0) + { + tree nca, t = gimple_omp_target_data_arg (stmt); + int i, oldlen = TREE_VEC_LENGTH (t); + tree nt = make_tree_vec (oldlen + nca_num); + for (i = 0; i < oldlen; i++) + TREE_VEC_ELT (nt, i) = TREE_VEC_ELT (t, i); + for (i = 0; nca_descrs.iterate (i, &nca); i++) + TREE_VEC_ELT (nt, oldlen + i) = nca; + gimple_omp_target_set_data_arg (stmt, nt); + } + DECL_INITIAL (TREE_VEC_ELT (t, 1)) = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 1)), vsize); DECL_INITIAL (TREE_VEC_ELT (t, 2)) Index: gcc/testsuite/c-c++-common/goacc/noncontig_array-1.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/noncontig_array-1.c (nonexistent) +++ gcc/testsuite/c-c++-common/goacc/noncontig_array-1.c (working copy) @@ -0,0 +1,25 @@ +/* { dg-do compile } */ + +void foo (void) +{ + int array_of_array[10][10]; + int **ptr_to_ptr; + int *array_of_ptr[10]; + int (*ptr_to_array)[10]; + + #pragma acc parallel copy (array_of_array[2:4][0:10]) + array_of_array[5][5] = 1; + + #pragma acc parallel copy (ptr_to_ptr[2:4][1:7]) + ptr_to_ptr[5][5] = 1; + + #pragma acc parallel copy (array_of_ptr[2:4][1:7]) + array_of_ptr[5][5] = 1; + + #pragma acc parallel copy (ptr_to_array[2:4][1:7]) /* { dg-error "array section is not contiguous in 'map' clause" } */ + ptr_to_array[5][5] = 1; +} +/* { dg-final { scan-tree-dump-times {#pragma omp target oacc_parallel map\(tofrom:array_of_array} 1 gimple } } */ +/* { dg-final { scan-tree-dump-times {#pragma omp target oacc_parallel map\(tofrom,noncontig_array:ptr_to_ptr \[dimensions: 2 4, 1 7\]} 1 gimple } } */ +/* { dg-final { scan-tree-dump-times {#pragma omp target oacc_parallel map\(tofrom,noncontig_array:array_of_ptr \[dimensions: 2 4, 1 7\]} 1 gimple } } */ +/* { dg-final { scan-tree-dump-times {#pragma omp target oacc_parallel map\(tofrom,noncontig_array:ptr_to_array \[dimensions: 2 4, 1 7\]} 1 gimple { xfail *-*-* } } } */ Index: gcc/tree-pretty-print.c =================================================================== --- gcc/tree-pretty-print.c (revision 278656) +++ gcc/tree-pretty-print.c (working copy) @@ -849,6 +849,33 @@ dump_omp_clause (pretty_printer *pp, tree clause, case GOMP_MAP_LINK: pp_string (pp, "link"); break; + case GOMP_MAP_NONCONTIG_ARRAY_TO: + pp_string (pp, "to,noncontig_array"); + break; + case GOMP_MAP_NONCONTIG_ARRAY_FROM: + pp_string (pp, "from,noncontig_array"); + break; + case GOMP_MAP_NONCONTIG_ARRAY_TOFROM: + pp_string (pp, "tofrom,noncontig_array"); + break; + case GOMP_MAP_NONCONTIG_ARRAY_FORCE_TO: + pp_string (pp, "force_to,noncontig_array"); + break; + case GOMP_MAP_NONCONTIG_ARRAY_FORCE_FROM: + pp_string (pp, "force_from,noncontig_array"); + break; + case GOMP_MAP_NONCONTIG_ARRAY_FORCE_TOFROM: + pp_string (pp, "force_tofrom,noncontig_array"); + break; + case GOMP_MAP_NONCONTIG_ARRAY_ALLOC: + pp_string (pp, "alloc,noncontig_array"); + break; + case GOMP_MAP_NONCONTIG_ARRAY_FORCE_ALLOC: + pp_string (pp, "force_alloc,noncontig_array"); + break; + case GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT: + pp_string (pp, "force_present,noncontig_array"); + break; default: gcc_unreachable (); } @@ -859,8 +886,15 @@ dump_omp_clause (pretty_printer *pp, tree clause, if (OMP_CLAUSE_SIZE (clause)) { switch (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP - ? OMP_CLAUSE_MAP_KIND (clause) : GOMP_MAP_TO) + ? (GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (clause)) + ? GOMP_MAP_NONCONTIG_ARRAY + : OMP_CLAUSE_MAP_KIND (clause)) + : GOMP_MAP_TO) { + case GOMP_MAP_NONCONTIG_ARRAY: + gcc_assert (TREE_CODE (OMP_CLAUSE_SIZE (clause)) == TREE_LIST); + pp_string (pp, " [dimensions: "); + break; case GOMP_MAP_POINTER: case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: Index: libgomp/libgomp.h =================================================================== --- libgomp/libgomp.h (revision 278656) +++ libgomp/libgomp.h (working copy) @@ -1167,6 +1167,10 @@ extern struct target_mem_desc *gomp_map_vars_async size_t, void **, void **, size_t *, void *, bool, enum gomp_map_vars_kind); +extern struct target_mem_desc *gomp_map_vars_openacc (struct gomp_device_descr *, + struct goacc_asyncqueue *, + size_t, void **, size_t *, + unsigned short *, void *); extern void gomp_unmap_tgt (struct target_mem_desc *); extern void gomp_unmap_vars (struct target_mem_desc *, bool); extern void gomp_unmap_vars_async (struct target_mem_desc *, bool, Index: libgomp/libgomp_g.h =================================================================== --- libgomp/libgomp_g.h (revision 278656) +++ libgomp/libgomp_g.h (working copy) @@ -364,7 +364,7 @@ extern void GOACC_parallel_keyed (int, void (*) (v extern void GOACC_parallel (int, void (*) (void *), size_t, void **, size_t *, unsigned short *, int, int, int, int, int, ...); extern void GOACC_data_start (int, size_t, void **, size_t *, - unsigned short *); + unsigned short *, ...); extern void GOACC_data_end (void); extern void GOACC_enter_exit_data (int, size_t, void **, size_t *, unsigned short *, int, int, ...); Index: libgomp/oacc-int.h =================================================================== --- libgomp/oacc-int.h (revision 278656) +++ libgomp/oacc-int.h (working copy) @@ -164,6 +164,57 @@ bool _goacc_profiling_setup_p (struct goacc_thread void goacc_profiling_dispatch (acc_prof_info *, acc_event_info *, acc_api_info *); +/* Definitions for data structures describing OpenACC non-contiguous arrays + (Note: interfaces with compiler) + + The compiler generates a descriptor for each such array, places the + descriptor on stack, and passes the address of the descriptor to the libgomp + runtime as a normal map argument. The runtime then processes the array + data structure setup, and replaces the argument with the new actual + array address for the child function. + + Care must be taken such that the struct field and layout assumptions + of struct goacc_ncarray_dim, goacc_ncarray_descr_type inside the compiler + be consistant with the below declarations. */ + +struct goacc_ncarray_dim { + size_t base; + size_t length; + size_t elem_size; + size_t is_array; +}; + +struct goacc_ncarray_descr_type +{ + size_t ndims; + struct goacc_ncarray_dim dims[]; +}; + +/* Internal non-contiguous array info struct, used only here inside the runtime. */ + +struct goacc_ncarray +{ + struct goacc_ncarray_descr_type *descr; + void *ptr; + size_t map_index; + size_t ptrblock_size; + void **data_rows; + void **tgt_data_rows; + size_t data_row_num; + size_t data_row_size; +}; + +struct goacc_ncarray_info +{ + size_t num_data_rows, num_ncarray; + void **data_rows; + void **tgt_data_rows; + struct goacc_ncarray ncarray[]; +}; + +extern void *goacc_noncontig_array_create_ptrblock (struct goacc_ncarray *, void *); + + #ifdef HAVE_ATTRIBUTE_VISIBILITY # pragma GCC visibility pop #endif Index: libgomp/oacc-parallel.c =================================================================== --- libgomp/oacc-parallel.c (revision 278656) +++ libgomp/oacc-parallel.c (working copy) @@ -36,8 +36,8 @@ #include #include #include +#include - /* In the ABI, the GOACC_FLAGs are encoded as an inverted bitmask, so that we continue to support the following two legacy values. */ _Static_assert (GOACC_FLAGS_UNMARSHAL (GOMP_DEVICE_ICV) == 0, @@ -113,7 +113,174 @@ handle_ftn_pointers (size_t mapnum, void **hostadd static void goacc_wait (int async, int num_waits, va_list *ap); +static size_t +goacc_noncontig_array_count_rows (struct goacc_ncarray_descr_type *descr) +{ + size_t nrows = 1; + for (size_t d = 0; d < descr->ndims - 1; d++) + nrows *= descr->dims[d].length / sizeof (void *); + return nrows; +} +static void +goacc_noncontig_array_compute_sizes (struct goacc_ncarray *nca) +{ + size_t d, n = 1; + struct goacc_ncarray_descr_type *descr = nca->descr; + + nca->ptrblock_size = 0; + for (d = 0; d < descr->ndims - 1; d++) + { + size_t dim_count = descr->dims[d].length / descr->dims[d].elem_size; + size_t dim_ptrblock_size = (descr->dims[d + 1].is_array + ? 0 : descr->dims[d].length * n); + nca->ptrblock_size += dim_ptrblock_size; + n *= dim_count; + } + nca->data_row_num = n; + nca->data_row_size = descr->dims[d].length; + +} + +static void +goacc_noncontig_array_fill_rows_1 (struct goacc_ncarray_descr_type *descr, void *nca, + size_t d, void ***row_ptr, size_t *count) +{ + if (d < descr->ndims - 1) + { + size_t elsize = descr->dims[d].elem_size; + size_t n = descr->dims[d].length / elsize; + void *p = nca + descr->dims[d].base; + for (size_t i = 0; i < n; i++) + { + void *ptr = p + i * elsize; + /* Deref if next dimension is not array. */ + if (!descr->dims[d + 1].is_array) + ptr = *((void **) ptr); + goacc_noncontig_array_fill_rows_1 (descr, ptr, d + 1, row_ptr, count); + } + } + else + { + **row_ptr = nca + descr->dims[d].base; + *row_ptr += 1; + *count += 1; + } +} + +static size_t +goacc_noncontig_array_fill_rows (struct goacc_ncarray *nca) +{ + size_t count = 0; + void **p = nca->data_rows; + goacc_noncontig_array_fill_rows_1 (nca->descr, nca->ptr, 0, &p, &count); + return count; +} + +static struct goacc_ncarray_info * +goacc_process_noncontiguous_arrays (size_t mapnum, void **hostaddrs, + unsigned short *kinds, va_list* ap) +{ + size_t i, nr, num_data_rows = 0, num_ncarray = 0, curr_row_start = 0; + struct goacc_ncarray_descr_type *descr; + + /* We need to go over *ap twice, so preserve *ap state here. */ + va_list itr; + va_copy (itr, *ap); + for (i = 0; i < mapnum; i++) + if (GOMP_MAP_NONCONTIG_ARRAY_P (kinds[i] & 0xff)) + { + descr = va_arg (itr, struct goacc_ncarray_descr_type *); + num_data_rows += goacc_noncontig_array_count_rows (descr); + num_ncarray += 1; + } + else + break; + + /* Allocate the entire info struct, array entries, and row pointer + arrays in one large block. */ + struct goacc_ncarray_info *nca_info + = gomp_malloc (sizeof (struct goacc_ncarray_info) + + sizeof (struct goacc_ncarray) * num_ncarray + + sizeof (void *) * num_data_rows * 2); + nca_info->num_data_rows = num_data_rows; + nca_info->num_ncarray = num_ncarray; + nca_info->data_rows = (void **) (nca_info->ncarray + num_ncarray); + nca_info->tgt_data_rows = nca_info->data_rows + num_data_rows; + + struct goacc_ncarray *curr_ncarray = nca_info->ncarray; + for (i = 0; i < mapnum; i++) + if (GOMP_MAP_NONCONTIG_ARRAY_P (kinds[i] & 0xff)) + { + descr = va_arg (*ap, struct goacc_ncarray_descr_type *); + curr_ncarray->descr = descr; + curr_ncarray->ptr = hostaddrs[i]; + curr_ncarray->map_index = i; + + goacc_noncontig_array_compute_sizes (curr_ncarray); + + curr_ncarray->data_rows = nca_info->data_rows + curr_row_start; + curr_ncarray->tgt_data_rows = nca_info->tgt_data_rows + curr_row_start; + + nr = goacc_noncontig_array_fill_rows (curr_ncarray); + assert (nr == curr_ncarray->data_row_num); + curr_row_start += nr; + curr_ncarray += 1; + } + else + break; + + return nca_info; +} + +void * +goacc_noncontig_array_create_ptrblock (struct goacc_ncarray *nca, + void *tgt_ptrblock_addr) +{ + struct goacc_ncarray_descr_type *descr = nca->descr; + void **tgt_data_rows = nca->tgt_data_rows; + void *ptrblock = gomp_malloc (nca->ptrblock_size); + void **curr_dim_ptrblock = (void **) ptrblock; + size_t n = 1; + + for (size_t d = 0; d < descr->ndims - 1; d++) + { + int curr_dim_len = descr->dims[d].length; + int next_dim_len = descr->dims[d + 1].length; + int curr_dim_num = curr_dim_len / sizeof (void *); + size_t next_dim_bias = descr->dims[d + 1].base; + + void *next_dim_ptrblock + = (void *)(curr_dim_ptrblock + n * curr_dim_num); + + for (int b = 0; b < n; b++) + for (int i = 0; i < curr_dim_num; i++) + { + if (d < descr->ndims - 2) + { + void *ptr = (next_dim_ptrblock + + b * curr_dim_num * next_dim_len + + i * next_dim_len); + void *tgt_ptr = (tgt_ptrblock_addr + + (ptr - ptrblock) - next_dim_bias); + curr_dim_ptrblock[b * curr_dim_num + i] = tgt_ptr; + } + else + { + curr_dim_ptrblock[b * curr_dim_num + i] + = tgt_data_rows[b * curr_dim_num + i] - next_dim_bias; + } + void *addr = &curr_dim_ptrblock[b * curr_dim_num + i]; + assert (ptrblock <= addr && addr < ptrblock + nca->ptrblock_size); + } + + n *= curr_dim_num; + curr_dim_ptrblock = next_dim_ptrblock; + } + assert (n == nca->data_row_num); + return ptrblock; +} + /* Launch a possibly offloaded function with FLAGS. FN is the host fn address. MAPNUM, HOSTADDRS, SIZES & KINDS describe the memory blocks to be copied to/from the device. Varadic arguments are @@ -138,6 +305,7 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (voi int async = GOMP_ASYNC_SYNC; unsigned dims[GOMP_DIM_MAX]; unsigned tag; + struct goacc_ncarray_info *nca_info = NULL; #ifdef HAVE_INTTYPES_H gomp_debug (0, "%s: mapnum=%"PRIu64", hostaddrs=%p, size=%p, kinds=%p\n", @@ -270,11 +438,20 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (voi break; } + /*case GOMP_LAUNCH_NONCONTIG_ARRAYS: + nca_info = goacc_process_noncontiguous_arrays (mapnum, hostaddrs, + kinds, &ap); + break;*/ + default: gomp_fatal ("unrecognized offload code '%d'," " libgomp is too old", GOMP_LAUNCH_CODE (tag)); } } + + if (mapnum > 0 && GOMP_MAP_NONCONTIG_ARRAY_P (kinds[0] & 0xff)) + nca_info = goacc_process_noncontiguous_arrays (mapnum, hostaddrs, kinds, &ap); + va_end (ap); if (!(acc_dev->capabilities & GOMP_OFFLOAD_CAP_NATIVE_EXEC)) @@ -311,8 +488,10 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (voi goacc_aq aq = get_goacc_asyncqueue (async); - tgt = gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, kinds, - true, GOMP_MAP_VARS_OPENACC); + tgt = gomp_map_vars_openacc (acc_dev, aq, mapnum, hostaddrs, sizes, kinds, + nca_info); + free (nca_info); + if (profiling_p) { prof_info.event_type = acc_ev_enter_data_end; @@ -390,7 +569,7 @@ GOACC_parallel (int flags_m, void (*fn) (void *), void GOACC_data_start (int flags_m, size_t mapnum, - void **hostaddrs, size_t *sizes, unsigned short *kinds) + void **hostaddrs, size_t *sizes, unsigned short *kinds, ...) { int flags = GOACC_FLAGS_UNMARSHAL (flags_m); @@ -480,8 +659,7 @@ GOACC_data_start (int flags_m, size_t mapnum, { prof_info.device_type = acc_device_host; api_info.device_type = prof_info.device_type; - tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true, - GOMP_MAP_VARS_OPENACC); + tgt = gomp_map_vars_openacc (NULL, NULL, 0, NULL, NULL, NULL, NULL); tgt->prev = thr->mapped_data; thr->mapped_data = tgt; @@ -488,9 +666,19 @@ GOACC_data_start (int flags_m, size_t mapnum, goto out_prof; } + struct goacc_ncarray_info *nca_info = NULL; + if (mapnum > 0 && GOMP_MAP_NONCONTIG_ARRAY_P (kinds[0] & 0xff)) + { + va_list ap; + va_start (ap, kinds); + nca_info = goacc_process_noncontiguous_arrays (mapnum, hostaddrs, kinds, &ap); + va_end (ap); + } + gomp_debug (0, " %s: prepare mappings\n", __FUNCTION__); - tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true, - GOMP_MAP_VARS_OPENACC); + tgt = gomp_map_vars_openacc (acc_dev, NULL, mapnum, hostaddrs, sizes, kinds, + nca_info); + free (nca_info); gomp_debug (0, " %s: mappings prepared\n", __FUNCTION__); tgt->prev = thr->mapped_data; thr->mapped_data = tgt; @@ -701,7 +889,6 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, if (num_waits) { va_list ap; - va_start (ap, num_waits); goacc_wait (async, num_waits, &ap); va_end (ap); Index: libgomp/target.c =================================================================== --- libgomp/target.c (revision 278656) +++ libgomp/target.c (working copy) @@ -524,10 +524,11 @@ static inline __attribute__((always_inline)) struc gomp_map_vars_internal (struct gomp_device_descr *devicep, struct goacc_asyncqueue *aq, size_t mapnum, void **hostaddrs, void **devaddrs, size_t *sizes, - void *kinds, bool short_mapkind, - enum gomp_map_vars_kind pragma_kind) + void *kinds, struct goacc_ncarray_info *nca_info, + bool short_mapkind, enum gomp_map_vars_kind pragma_kind) { size_t i, tgt_align, tgt_size, not_found_cnt = 0; + size_t nca_data_row_num = (nca_info ? nca_info->num_data_rows : 0); bool has_firstprivate = false; const int rshift = short_mapkind ? 8 : 3; const int typemask = short_mapkind ? 0xff : 0x7; @@ -534,8 +535,9 @@ gomp_map_vars_internal (struct gomp_device_descr * struct splay_tree_s *mem_map = &devicep->mem_map; struct splay_tree_key_s cur_node; struct target_mem_desc *tgt - = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum); - tgt->list_count = mapnum; + = gomp_malloc (sizeof (*tgt) + + sizeof (tgt->list[0]) * (mapnum + nca_data_row_num)); + tgt->list_count = mapnum + nca_data_row_num; tgt->refcount = pragma_kind == GOMP_MAP_VARS_ENTER_DATA ? 0 : 1; tgt->device_descr = devicep; struct gomp_coalesce_buf cbuf, *cbufp = NULL; @@ -667,6 +669,27 @@ gomp_map_vars_internal (struct gomp_device_descr * has_firstprivate = true; continue; } + else if (GOMP_MAP_NONCONTIG_ARRAY_P (kind & typemask)) + { + /* Ignore non-contiguous arrays for now, we process them together + later. */ + tgt->list[i].key = NULL; + tgt->list[i].offset = 0; + not_found_cnt++; + + /* The map for the non-contiguous array itself is never copied from + during unmapping, its the data rows that count. Set copy-from + flags to false here. */ + tgt->list[i].copy_from = false; + tgt->list[i].always_copy_from = false; + + size_t align = (size_t) 1 << (kind >> rshift); + if (tgt_align < align) + tgt_align = align; + + continue; + } + cur_node.host_start = (uintptr_t) hostaddrs[i]; if (!GOMP_MAP_POINTER_P (kind & typemask)) cur_node.host_end = cur_node.host_start + sizes[i]; @@ -735,6 +758,44 @@ gomp_map_vars_internal (struct gomp_device_descr * } } + /* For non-contiguous arrays. Each data row is one target item, separated + from the normal map clause items, hence we order them after mapnum. */ + if (nca_info) + { + struct target_var_desc *next_var_desc = &tgt->list[mapnum]; + for (i = 0; i < nca_info->num_ncarray; i++) + { + struct goacc_ncarray *nca = &nca_info->ncarray[i]; + int kind = get_kind (short_mapkind, kinds, nca->map_index); + size_t align = (size_t) 1 << (kind >> rshift); + tgt_size = (tgt_size + align - 1) & ~(align - 1); + tgt_size += nca->ptrblock_size; + + for (size_t j = 0; j < nca->data_row_num; j++) + { + struct target_var_desc *row_desc = next_var_desc++; + void *row = nca->data_rows[j]; + cur_node.host_start = (uintptr_t) row; + cur_node.host_end = cur_node.host_start + nca->data_row_size; + splay_tree_key n = splay_tree_lookup (mem_map, &cur_node); + if (n) + { + assert (n->refcount != REFCOUNT_LINK); + gomp_map_vars_existing (devicep, aq, n, &cur_node, row_desc, + kind & typemask, + /* TODO: cbuf? */ NULL); + } + else + { + tgt_size = (tgt_size + align - 1) & ~(align - 1); + tgt_size += nca->data_row_size; + not_found_cnt++; + } + } + } + assert (next_var_desc == &tgt->list[mapnum + nca_info->num_data_rows]); + } + if (devaddrs) { if (mapnum != 1) @@ -895,6 +956,15 @@ gomp_map_vars_internal (struct gomp_device_descr * default: break; } + + if (GOMP_MAP_NONCONTIG_ARRAY_P (kind & typemask)) + { + tgt->list[i].key = &array->key; + tgt->list[i].key->tgt = tgt; + array++; + continue; + } + splay_tree_key k = &array->key; k->host_start = (uintptr_t) hostaddrs[i]; if (!GOMP_MAP_POINTER_P (kind & typemask)) @@ -1044,6 +1114,98 @@ gomp_map_vars_internal (struct gomp_device_descr * array++; } } + + /* Processing of non-contiguous array rows. */ + if (nca_info) + { + struct target_var_desc *next_var_desc = &tgt->list[mapnum]; + for (i = 0; i < nca_info->num_ncarray; i++) + { + struct goacc_ncarray *nca = &nca_info->ncarray[i]; + int kind = get_kind (short_mapkind, kinds, nca->map_index); + size_t align = (size_t) 1 << (kind >> rshift); + tgt_size = (tgt_size + align - 1) & ~(align - 1); + + assert (nca->ptr == hostaddrs[nca->map_index]); + + /* For the map of the non-contiguous array itself, adjust so that + the passed device address points to the beginning of the + ptrblock. Remember to adjust the first-dimension's bias here. */ + tgt->list[nca->map_index].key->tgt_offset + = tgt_size - nca->descr->dims[0].base; + + void *target_ptrblock = (void*) tgt->tgt_start + tgt_size; + tgt_size += nca->ptrblock_size; + + /* Add splay key for each data row in current non-contiguous + array. */ + for (size_t j = 0; j < nca->data_row_num; j++) + { + struct target_var_desc *row_desc = next_var_desc++; + void *row = nca->data_rows[j]; + cur_node.host_start = (uintptr_t) row; + cur_node.host_end = cur_node.host_start + nca->data_row_size; + splay_tree_key k = splay_tree_lookup (mem_map, &cur_node); + if (k) + { + assert (k->refcount != REFCOUNT_LINK); + gomp_map_vars_existing (devicep, aq, k, &cur_node, row_desc, + kind & typemask, + cbufp); + } + else + { + tgt->refcount++; + tgt_size = (tgt_size + align - 1) & ~(align - 1); + + k = &array->key; + k->host_start = (uintptr_t) row; + k->host_end = k->host_start + nca->data_row_size; + + k->tgt = tgt; + k->refcount = 1; + k->dynamic_refcount = 0; + k->link_key = NULL; + k->tgt_offset = tgt_size; + + tgt_size += nca->data_row_size; + + row_desc->key = k; + row_desc->copy_from + = GOMP_MAP_COPY_FROM_P (kind & typemask); + row_desc->always_copy_from + = GOMP_MAP_COPY_FROM_P (kind & typemask); + row_desc->offset = 0; + row_desc->length = nca->data_row_size; + + array->left = NULL; + array->right = NULL; + splay_tree_insert (mem_map, array); + + if (GOMP_MAP_COPY_TO_P (kind & typemask)) + gomp_copy_host2dev (devicep, aq, + (void *) tgt->tgt_start + k->tgt_offset, + (void *) k->host_start, + nca->data_row_size, cbufp); + array++; + } + nca->tgt_data_rows[j] + = (void *) (k->tgt->tgt_start + k->tgt_offset); + } + + /* Now we have the target memory allocated, and target offsets of all + row blocks assigned and calculated, we can construct the + accelerator side ptrblock and copy it in. */ + if (nca->ptrblock_size) + { + void *ptrblock = goacc_noncontig_array_create_ptrblock + (nca, target_ptrblock); + gomp_copy_host2dev (devicep, aq, target_ptrblock, ptrblock, + nca->ptrblock_size, cbufp); + free (ptrblock); + } + } + } } if (pragma_kind == GOMP_MAP_VARS_TARGET) @@ -1086,12 +1248,25 @@ gomp_map_vars_internal (struct gomp_device_descr * } attribute_hidden struct target_mem_desc * +gomp_map_vars_openacc (struct gomp_device_descr *devicep, + struct goacc_asyncqueue *aq, size_t mapnum, + void **hostaddrs, size_t *sizes, unsigned short *kinds, + void *nca_info) +{ + return gomp_map_vars_internal (devicep, aq, mapnum, hostaddrs, NULL, + sizes, (void *) kinds, + (struct goacc_ncarray_info *) nca_info, + true, GOMP_MAP_VARS_OPENACC); +} + +attribute_hidden struct target_mem_desc * gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, void **devaddrs, size_t *sizes, void *kinds, bool short_mapkind, enum gomp_map_vars_kind pragma_kind) { return gomp_map_vars_internal (devicep, NULL, mapnum, hostaddrs, devaddrs, - sizes, kinds, short_mapkind, pragma_kind); + sizes, kinds, NULL, short_mapkind, + pragma_kind); } attribute_hidden struct target_mem_desc * @@ -1102,7 +1277,7 @@ gomp_map_vars_async (struct gomp_device_descr *dev enum gomp_map_vars_kind pragma_kind) { return gomp_map_vars_internal (devicep, aq, mapnum, hostaddrs, devaddrs, - sizes, kinds, short_mapkind, pragma_kind); + sizes, kinds, NULL, short_mapkind, pragma_kind); } attribute_hidden void Index: libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-1.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-1.c (nonexistent) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-1.c (working copy) @@ -0,0 +1,103 @@ +/* { dg-do run } */ + +#include +#include + +#define n 100 +#define m 100 + +int b[n][m]; + +void +test1 (void) +{ + int i, j, *a[100]; + + /* Array of pointers form test. */ + for (i = 0; i < n; i++) + { + a[i] = (int *)malloc (sizeof (int) * m); + for (j = 0; j < m; j++) + b[i][j] = j - i; + } + + #pragma acc parallel loop copyout(a[0:n][0:m]) copyin(b) + for (i = 0; i < n; i++) + #pragma acc loop + for (j = 0; j < m; j++) + a[i][j] = b[i][j]; + + for (i = 0; i < n; i++) + { + for (j = 0; j < m; j++) + assert (a[i][j] == b[i][j]); + /* Clean up. */ + free (a[i]); + } +} + +void +test2 (void) +{ + int i, j, **a = (int **) malloc (sizeof (int *) * n); + + /* Separately allocated blocks. */ + for (i = 0; i < n; i++) + { + a[i] = (int *)malloc (sizeof (int) * m); + for (j = 0; j < m; j++) + b[i][j] = j - i; + } + + #pragma acc parallel loop copyout(a[0:n][0:m]) copyin(b) + for (i = 0; i < n; i++) + #pragma acc loop + for (j = 0; j < m; j++) + a[i][j] = b[i][j]; + + for (i = 0; i < n; i++) + { + for (j = 0; j < m; j++) + assert (a[i][j] == b[i][j]); + /* Clean up. */ + free (a[i]); + } + free (a); +} + +void +test3 (void) +{ + int i, j, **a = (int **) malloc (sizeof (int *) * n); + a[0] = (int *) malloc (sizeof (int) * n * m); + + /* Rows allocated in one contiguous block. */ + for (i = 0; i < n; i++) + { + a[i] = *a + i * m; + for (j = 0; j < m; j++) + b[i][j] = j - i; + } + + #pragma acc parallel loop copyout(a[0:n][0:m]) copyin(b) + for (i = 0; i < n; i++) + #pragma acc loop + for (j = 0; j < m; j++) + a[i][j] = b[i][j]; + + for (i = 0; i < n; i++) + for (j = 0; j < m; j++) + assert (a[i][j] == b[i][j]); + + free (a[0]); + free (a); +} + +int +main (void) +{ + test1 (); + test2 (); + test3 (); + return 0; +} Index: libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-2.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-2.c (nonexistent) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-2.c (working copy) @@ -0,0 +1,37 @@ +/* { dg-do run } */ + +#include +#include "noncontig_array-utils.h" + +int +main (void) +{ + int n = 10; + int ***a = (int ***) create_ncarray (sizeof (int), n, 3); + int ***b = (int ***) create_ncarray (sizeof (int), n, 3); + int ***c = (int ***) create_ncarray (sizeof (int), n, 3); + + for (int i = 0; i < n; i++) + for (int j = 0; j < n; j++) + for (int k = 0; k < n; k++) + { + a[i][j][k] = i + j * k + k; + b[i][j][k] = j + k * i + i * j; + c[i][j][k] = a[i][j][k]; + } + + #pragma acc parallel copy (a[0:n][0:n][0:n]) copyin (b[0:n][0:n][0:n]) + { + for (int i = 0; i < n; i++) + for (int j = 0; j < n; j++) + for (int k = 0; k < n; k++) + a[i][j][k] += b[k][j][i] + i + j + k; + } + + for (int i = 0; i < n; i++) + for (int j = 0; j < n; j++) + for (int k = 0; k < n; k++) + assert (a[i][j][k] == c[i][j][k] + b[k][j][i] + i + j + k); + + return 0; +} Index: libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-3.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-3.c (nonexistent) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-3.c (working copy) @@ -0,0 +1,45 @@ +/* { dg-do run } */ + +#include +#include "noncontig_array-utils.h" + +int main (void) +{ + int n = 20, x = 5, y = 12; + int *****a = (int *****) create_ncarray (sizeof (int), n, 5); + + int sum1 = 0, sum2 = 0, sum3 = 0; + + for (int i = 0; i < n; i++) + for (int j = 0; j < n; j++) + for (int k = 0; k < n; k++) + for (int l = 0; l < n; l++) + for (int m = 0; m < n; m++) + { + a[i][j][k][l][m] = 1; + sum1++; + } + + #pragma acc parallel copy (a[x:y][x:y][x:y][x:y][x:y]) copy(sum2) + { + for (int i = x; i < x + y; i++) + for (int j = x; j < x + y; j++) + for (int k = x; k < x + y; k++) + for (int l = x; l < x + y; l++) + for (int m = x; m < x + y; m++) + { + a[i][j][k][l][m] = 0; + sum2++; + } + } + + for (int i = 0; i < n; i++) + for (int j = 0; j < n; j++) + for (int k = 0; k < n; k++) + for (int l = 0; l < n; l++) + for (int m = 0; m < n; m++) + sum3 += a[i][j][k][l][m]; + + assert (sum1 == sum2 + sum3); + return 0; +} Index: libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-4.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-4.c (nonexistent) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-4.c (working copy) @@ -0,0 +1,36 @@ +/* { dg-do run } */ + +#include +#include "noncontig_array-utils.h" + +int main (void) +{ + int n = 128; + double ***a = (double ***) create_ncarray (sizeof (double), n, 3); + double ***b = (double ***) create_ncarray (sizeof (double), n, 3); + + for (int i = 0; i < n; i++) + for (int j = 0; j < n; j++) + for (int k = 0; k < n; k++) + a[i][j][k] = i + j + k + i * j * k; + + /* This test exercises async copyout of non-contiguous array rows. */ + #pragma acc parallel copyin(a[0:n][0:n][0:n]) copyout(b[0:n][0:n][0:n]) async(5) + { + #pragma acc loop gang + for (int i = 0; i < n; i++) + #pragma acc loop vector + for (int j = 0; j < n; j++) + for (int k = 0; k < n; k++) + b[i][j][k] = a[i][j][k] * 2.0; + } + + #pragma acc wait (5) + + for (int i = 0; i < n; i++) + for (int j = 0; j < n; j++) + for (int k = 0; k < n; k++) + assert (b[i][j][k] == a[i][j][k] * 2.0); + + return 0; +} Index: libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-utils.h =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-utils.h (nonexistent) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-utils.h (working copy) @@ -0,0 +1,44 @@ +#include +#include +#include +#include + +/* Allocate and create a pointer based NDIMS-dimensional array, + each dimension DIMLEN long, with ELSIZE sized data elements. */ +void * +create_ncarray (size_t elsize, int dimlen, int ndims) +{ + size_t blk_size = 0; + size_t n = 1; + + for (int i = 0; i < ndims - 1; i++) + { + n *= dimlen; + blk_size += sizeof (void *) * n; + } + size_t data_rows_num = n; + size_t data_rows_offset = blk_size; + blk_size += elsize * n * dimlen; + + void *blk = (void *) malloc (blk_size); + memset (blk, 0, blk_size); + void **curr_dim = (void **) blk; + n = 1; + + for (int d = 0; d < ndims - 1; d++) + { + uintptr_t next_dim = (uintptr_t) (curr_dim + n * dimlen); + size_t next_dimlen = dimlen * (d < ndims - 2 ? sizeof (void *) : elsize); + + for (int b = 0; b < n; b++) + for (int i = 0; i < dimlen; i++) + if (d < ndims - 1) + curr_dim[b * dimlen + i] + = (void*) (next_dim + b * dimlen * next_dimlen + i * next_dimlen); + + n *= dimlen; + curr_dim = (void**) next_dim; + } + assert (n == data_rows_num); + return blk; +}