From patchwork Tue Oct 16 12:55:17 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 984743 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-487639-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="LCpBAzz7"; 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 42ZFfH5j2Rz9s9G for ; Tue, 16 Oct 2018 23:55:35 +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:message-id:date:mime-version :content-type; q=dns; s=default; b=GzOF11oLprRLrrRz0eKdfEyIsvYpM uAvAlcuQnWEWW9//Qcqgr2Z0hSZ7gfWPNCCLEZUCd3H4bb5TviBQhvNjxBn+bzv1 LGB0hjsOdQ/E+dH6Y6cxmU+OnXk8k3ejiiSa3ZD9BrSQyTO4sJw9hP1lPyj2wBK2 6ThIL6l6orskuI= 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:message-id:date:mime-version :content-type; s=default; bh=3qhhvJu0G2l935r8IkDaI7iz1Dg=; b=LCp BAzz7XHxviMTjLQ0YjBiiduHaaNG1S21Ltr4bCC7R1Q19018iYn2zZFOTffKCAit f7vO3MdAGMCSKXYsSDhQQuS5TQUXxQSh17zyMzrTsftBKaQo0dnx8FerScJx1EIU 6+auSpmTKlvfjAQbOzJ8mkD7cRglvgC+NJeQJOyo= Received: (qmail 59275 invoked by alias); 16 Oct 2018 12:55:28 -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 59176 invoked by uid 89); 16 Oct 2018 12:55:26 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-24.9 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, RCVD_IN_DNSWL_NONE, SPF_PASS, UPPERCASE_50_75 autolearn=ham version=3.3.2 spammy=H*r:0700 X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Tue, 16 Oct 2018 12:55:24 +0000 Received: from svr-orw-mbx-02.mgc.mentorg.com ([147.34.90.202]) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1gCOsd-00075A-1C from ChungLin_Tang@mentor.com ; Tue, 16 Oct 2018 05:55:23 -0700 Received: from [0.0.0.0] (147.34.91.1) by svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Tue, 16 Oct 2018 05:55:20 -0700 Reply-To: From: Chung-Lin Tang Subject: [PATCH, OpenACC, 1/8] Multi-dimensional dynamic array support for OpenACC data clauses, gomp-constants.h additions To: , Jakub Jelinek , Thomas Schwinge Message-ID: Date: Tue, 16 Oct 2018 20:55:17 +0800 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.13; rv:52.0) Gecko/20100101 Thunderbird/52.9.1 MIME-Version: 1.0 This part defines GOMP_MAP_DYNAMIC_ARRAY_* symbols in include/gomp-constants.h. Basically use the next bit to define GOMP_MAP_FLAG_SPECIAL_3 to achieve this purpose. Thanks, Chung-Lin Tang include/ * gomp-constants.h (GOMP_MAP_FLAG_SPECIAL_3): Define. (enum gomp_map_kind): Add GOMP_MAP_DYNAMIC_ARRAY, GOMP_MAP_DYNAMIC_ARRAY_TO, GOMP_MAP_DYNAMIC_ARRAY_FROM, GOMP_MAP_DYNAMIC_ARRAY_TOFROM, GOMP_MAP_DYNAMIC_ARRAY_FORCE_TO, GOMP_MAP_DYNAMIC_ARRAY_FORCE_FROM, GOMP_MAP_DYNAMIC_ARRAY_FORCE_TOFROM, GOMP_MAP_DYNAMIC_ARRAY_ALLOC, GOMP_MAP_DYNAMIC_ARRAY_FORCE_ALLOC, GOMP_MAP_DYNAMIC_ARRAY_FORCE_PRESENT. (GOMP_MAP_DYNAMIC_ARRAY_P): Define. diff --git a/include/gomp-constants.h b/include/gomp-constants.h index ccfb657..f25169c 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -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). */ @@ -128,6 +129,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 dynamic arrays. */ + GOMP_MAP_DYNAMIC_ARRAY = (GOMP_MAP_FLAG_SPECIAL_3), + GOMP_MAP_DYNAMIC_ARRAY_TO = (GOMP_MAP_DYNAMIC_ARRAY + | GOMP_MAP_TO), + GOMP_MAP_DYNAMIC_ARRAY_FROM = (GOMP_MAP_DYNAMIC_ARRAY + | GOMP_MAP_FROM), + GOMP_MAP_DYNAMIC_ARRAY_TOFROM = (GOMP_MAP_DYNAMIC_ARRAY + | GOMP_MAP_TOFROM), + GOMP_MAP_DYNAMIC_ARRAY_FORCE_TO = (GOMP_MAP_DYNAMIC_ARRAY_TO + | GOMP_MAP_FLAG_FORCE), + GOMP_MAP_DYNAMIC_ARRAY_FORCE_FROM = (GOMP_MAP_DYNAMIC_ARRAY_FROM + | GOMP_MAP_FLAG_FORCE), + GOMP_MAP_DYNAMIC_ARRAY_FORCE_TOFROM = (GOMP_MAP_DYNAMIC_ARRAY_TOFROM + | GOMP_MAP_FLAG_FORCE), + GOMP_MAP_DYNAMIC_ARRAY_ALLOC = (GOMP_MAP_DYNAMIC_ARRAY + | GOMP_MAP_ALLOC), + GOMP_MAP_DYNAMIC_ARRAY_FORCE_ALLOC = (GOMP_MAP_DYNAMIC_ARRAY + | GOMP_MAP_FORCE_ALLOC), + GOMP_MAP_DYNAMIC_ARRAY_FORCE_PRESENT = (GOMP_MAP_DYNAMIC_ARRAY + | GOMP_MAP_FORCE_PRESENT), /* Internal to GCC, not used in libgomp. */ /* Do not map, but pointer assign a pointer instead. */ @@ -156,6 +177,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_DYNAMIC_ARRAY_P(X) \ + ((X) & GOMP_MAP_DYNAMIC_ARRAY) /* Asynchronous behavior. Keep in sync with libgomp/{openacc.h,openacc.f90,openacc_lib.h}:acc_async_t. */ From patchwork Tue Oct 16 12:55:32 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 984744 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-487640-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="b/9QZxOF"; 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 42ZFfb5TFjz9s9G for ; Tue, 16 Oct 2018 23:55:51 +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:message-id:date:mime-version :content-type; q=dns; s=default; b=omQrzDtzHXS6L4t5Zv+lrv1lkRZbR qxwKG4sbvuT0A1x2fxqCcvsS2XqqaMKAlasXyC5KhaTJfFhOl2fgWzlwLzG+HsYh gn2XJZ9xHsIqfWtR4CdHwU5VEFUr1zCDiss9JSGMef3c72tBFhqb2z2x4vPd9QPt yzgNs79wDpmJvE= 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:message-id:date:mime-version :content-type; s=default; bh=aB397W6SndOTl0Ao3OYWafBcSEI=; b=b/9 QZxOFkkmBZEI2CmR0tiBhasX8yfXumtmAkERzzQnFGcZlek22uYdY84ZB0gvgVtY 8rmR/Yp530y2QKqNykIJ9389ZbQLH+1zmDCE0AYMkWNSAYYwkROVsKIEs7whoCCU vAJgz00I/HLzH1JG0FIj/Dc7umpa5EQA2+3r5g6Y= Received: (qmail 61059 invoked by alias); 16 Oct 2018 12:55:43 -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 61048 invoked by uid 89); 16 Oct 2018 12:55:43 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-25.0 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, RCVD_IN_DNSWL_NONE, SPF_PASS autolearn=ham version=3.3.2 spammy=clause, ort, H*r:0700 X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Tue, 16 Oct 2018 12:55:40 +0000 Received: from svr-orw-mbx-02.mgc.mentorg.com ([147.34.90.202]) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1gCOss-00079W-Sk from ChungLin_Tang@mentor.com ; Tue, 16 Oct 2018 05:55:38 -0700 Received: from [0.0.0.0] (147.34.91.1) by svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Tue, 16 Oct 2018 05:55:35 -0700 Reply-To: From: Chung-Lin Tang Subject: [PATCH, OpenACC, 2/8] Multi-dimensional dynamic array support for OpenACC data clauses, C/C++ front-end parts To: , Jakub Jelinek , Thomas Schwinge Message-ID: Date: Tue, 16 Oct 2018 20:55:32 +0800 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.13; rv:52.0) Gecko/20100101 Thunderbird/52.9.1 MIME-Version: 1.0 These are the parts for the C/C++ front-ends. We now allow certain non-contiguous cases under OpenACC, and pass the defined base/length pairs for each array dimension as a TREE_LIST passed to the middle-end through OMP_CLAUSE_SIZE. Thanks, Chung-Lin 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. diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 0f639be..c273435 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -12409,7 +12409,7 @@ c_finish_omp_cancellation_point (location_t loc, tree clauses) 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) @@ -12494,7 +12494,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, } ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types, - maybe_zero_len, first_non_one, ort); + maybe_zero_len, first_non_one, + non_contiguous, ort); if (ret == error_mark_node || ret == NULL_TREE) return ret; @@ -12654,6 +12655,21 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, } } } + + /* For OpenACC, if the low_bound/length suggest this is a subarray, + and is referenced through by a pointer, then mark this as + non-contiguous. */ + if (ort == C_ORT_ACC + && types.length () > 0 + && (TREE_CODE (low_bound) != INTEGER_CST + || integer_nonzerop (low_bound) + || (length && (TREE_CODE (length) != INTEGER_CST + || !tree_int_cst_equal (size, length))))) + { + tree x = types.last (); + if (TREE_CODE (x) == POINTER_TYPE) + non_contiguous = true; + } } else if (length == NULL_TREE) { @@ -12695,13 +12711,16 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, /* If there is a pointer type anywhere but in the very first array-section-subscript, the array section can't be contiguous. */ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND - && TREE_CODE (TREE_CHAIN (t)) == TREE_LIST) + && TREE_CODE (TREE_CHAIN (t)) == TREE_LIST + && ort != C_ORT_ACC) { error_at (OMP_CLAUSE_LOCATION (c), "array section is not contiguous in %qs clause", omp_clause_code_name[OMP_CLAUSE_CODE (c)]); return error_mark_node; } + else if (TREE_CODE (TREE_CHAIN (t)) == TREE_LIST) + non_contiguous = true; } else { @@ -12729,10 +12748,11 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) { bool maybe_zero_len = false; unsigned int first_non_one = 0; + bool non_contiguous = false; auto_vec types; tree first = handle_omp_array_sections_1 (c, OMP_CLAUSE_DECL (c), types, maybe_zero_len, first_non_one, - ort); + non_contiguous, ort); if (first == error_mark_node) return true; if (first == NULL_TREE) @@ -12765,6 +12785,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) unsigned int num = types.length (), i; tree t, side_effects = NULL_TREE, size = NULL_TREE; tree condition = NULL_TREE; + tree da_dims = NULL_TREE; if (int_size_in_bytes (TREE_TYPE (first)) <= 0) maybe_zero_len = true; @@ -12788,6 +12809,13 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) length = fold_convert (sizetype, length); if (low_bound == NULL_TREE) low_bound = integer_zero_node; + + if (non_contiguous) + { + da_dims = tree_cons (low_bound, length, da_dims); + continue; + } + if (!maybe_zero_len && i > first_non_one) { if (integer_nonzerop (low_bound)) @@ -12880,6 +12908,14 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) 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_DYNAMIC_ARRAY); + OMP_CLAUSE_DECL (c) = t; + OMP_CLAUSE_SIZE (c) = da_dims; + return false; + } if (side_effects) size = build2 (COMPOUND_EXPR, sizetype, side_effects, size); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION) diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 85c7cfa..af7a1a6 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -4521,7 +4521,7 @@ omp_privatize_field (tree t, bool shared) 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) @@ -4604,7 +4604,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, && TREE_CODE (TREE_CHAIN (t)) == FIELD_DECL) TREE_CHAIN (t) = omp_privatize_field (TREE_CHAIN (t), false); ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types, - maybe_zero_len, first_non_one, ort); + maybe_zero_len, first_non_one, + non_contiguous, ort); if (ret == error_mark_node || ret == NULL_TREE) return ret; @@ -4776,6 +4777,21 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, } } } + + /* For OpenACC, if the low_bound/length suggest this is a subarray, + and is referenced through by a pointer, then mark this as + non-contiguous. */ + if (ort == C_ORT_ACC + && types.length () > 0 + && (TREE_CODE (low_bound) != INTEGER_CST + || integer_nonzerop (low_bound) + || (length && (TREE_CODE (length) != INTEGER_CST + || !tree_int_cst_equal (size, length))))) + { + tree x = types.last (); + if (TREE_CODE (x) == POINTER_TYPE) + non_contiguous = true; + } } else if (length == NULL_TREE) { @@ -4817,13 +4833,16 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, /* If there is a pointer type anywhere but in the very first array-section-subscript, the array section can't be contiguous. */ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND - && TREE_CODE (TREE_CHAIN (t)) == TREE_LIST) + && TREE_CODE (TREE_CHAIN (t)) == TREE_LIST + && ort != C_ORT_ACC) { error_at (OMP_CLAUSE_LOCATION (c), "array section is not contiguous in %qs clause", omp_clause_code_name[OMP_CLAUSE_CODE (c)]); return error_mark_node; } + else if (TREE_CODE (TREE_CHAIN (t)) == TREE_LIST) + non_contiguous = true; } else { @@ -4851,10 +4870,11 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) { bool maybe_zero_len = false; unsigned int first_non_one = 0; + bool non_contiguous = false; auto_vec types; tree first = handle_omp_array_sections_1 (c, OMP_CLAUSE_DECL (c), types, maybe_zero_len, first_non_one, - ort); + non_contiguous, ort); if (first == error_mark_node) return true; if (first == NULL_TREE) @@ -4888,6 +4908,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) unsigned int num = types.length (), i; tree t, side_effects = NULL_TREE, size = NULL_TREE; tree condition = NULL_TREE; + tree da_dims = NULL_TREE; if (int_size_in_bytes (TREE_TYPE (first)) <= 0) maybe_zero_len = true; @@ -4913,6 +4934,13 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) length = fold_convert (sizetype, length); if (low_bound == NULL_TREE) low_bound = integer_zero_node; + + if (non_contiguous) + { + da_dims = tree_cons (low_bound, length, da_dims); + continue; + } + if (!maybe_zero_len && i > first_non_one) { if (integer_nonzerop (low_bound)) @@ -5000,6 +5028,14 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) } if (!processing_template_decl) { + if (non_contiguous) + { + int kind = OMP_CLAUSE_MAP_KIND (c); + OMP_CLAUSE_SET_MAP_KIND (c, kind | GOMP_MAP_DYNAMIC_ARRAY); + OMP_CLAUSE_DECL (c) = t; + OMP_CLAUSE_SIZE (c) = da_dims; + return false; + } if (side_effects) size = build2 (COMPOUND_EXPR, sizetype, side_effects, size); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION) From patchwork Tue Oct 16 12:55:57 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 984745 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-487641-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="aEW6Wk+4"; 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 42ZFg13BfPz9s5c for ; Tue, 16 Oct 2018 23:56:13 +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:message-id:date:mime-version :content-type; q=dns; s=default; b=AjdCPujhgcg1OcMYiwkR7oynfsMIa ZurLBicgH8BTITylBk8/nrh72hZuwJlQt8pSNl32n15wsSIAMgy3x2Z2Cw+EO9JS 0bYQ8RzF+SHTh/w5yWDMJ2VxLTesOFXO/VwMh0DdXJdLA5bjFJfYHsDLchZ7QVbR XbMm/BlFx0kK/Y= 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:message-id:date:mime-version :content-type; s=default; bh=xU42vDeWxZvsdLFWj1KDoUE0owo=; b=aEW 6Wk+4nQnZNTijCZ9FsRzKQ0Miyz7c4yQ+pr8wt8SznV6eyI76iKdS2J+gSoGfE/4 XO0Y9BKg//lzft0rV4FPU8ggTDZm0NCW9aFT5QW24Za5ms496Z8hHQlniQklYFwL OyPweOdrurJgM2vcHsdeygaTfV8oLoosx4qr0lAo= Received: (qmail 68166 invoked by alias); 16 Oct 2018 12:56:06 -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 68155 invoked by uid 89); 16 Oct 2018 12:56:06 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-25.0 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, RCVD_IN_DNSWL_NONE, SPF_PASS autolearn=ham version=3.3.2 spammy=H*r:0700 X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Tue, 16 Oct 2018 12:56:05 +0000 Received: from svr-orw-mbx-02.mgc.mentorg.com ([147.34.90.202]) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1gCOtH-0007Cc-LL from ChungLin_Tang@mentor.com ; Tue, 16 Oct 2018 05:56:03 -0700 Received: from [0.0.0.0] (147.34.91.1) by svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Tue, 16 Oct 2018 05:56:00 -0700 Reply-To: From: Chung-Lin Tang Subject: [PATCH, OpenACC, 3/8] Multi-dimensional dynamic array support for OpenACC data clauses, gimplify patch To: , Jakub Jelinek , Thomas Schwinge Message-ID: <5dafa510-8741-978b-2183-09cfdf45691c@mentor.com> Date: Tue, 16 Oct 2018 20:55:57 +0800 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.13; rv:52.0) Gecko/20100101 Thunderbird/52.9.1 MIME-Version: 1.0 This gimplify.c patch adds to the omp clause scanning to handle dynamic array cases, mainly to properly handle dimension biases of GOMP_MAP_DYNAMIC_ARRAYs by making sure the bias field is seen in the omp-ctx. Thanks, Chung-Lin gcc/ * gimplify.c (gimplify_scan_omp_clauses): For dynamic array map kinds, make sure bias in each dimension are put into firstprivate variables. diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 781d430..09ef876 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -7901,8 +7901,28 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, 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 (OMP_CLAUSE_SIZE (c) + && TREE_CODE (OMP_CLAUSE_SIZE (c)) == TREE_LIST + && GOMP_MAP_DYNAMIC_ARRAY_P (OMP_CLAUSE_MAP_KIND (c))) + { + tree dims = OMP_CLAUSE_SIZE (c); + for (tree t = dims; t; t = TREE_CHAIN (t)) + { + /* If a dimension bias isn't a constant, we have to ensure + that the value gets transferred to the offload target. */ + tree low_bound = TREE_PURPOSE (t); + if (TREE_CODE (low_bound) != INTEGER_CST) + { + low_bound = get_initialized_tmp_var (low_bound, pre_p, + NULL, false); + omp_add_variable (ctx, low_bound, + GOVD_FIRSTPRIVATE | GOVD_SEEN); + TREE_PURPOSE (t) = low_bound; + } + } + } + else if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, + NULL, is_gimple_val, fb_rvalue) == GS_ERROR) { remove = true; break; From patchwork Tue Oct 16 12:56:14 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 984746 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-487642-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="XhF9hizc"; 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 42ZFgN2b1lz9s5c for ; Tue, 16 Oct 2018 23:56:32 +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:message-id:date:mime-version :content-type; q=dns; s=default; b=A3qvophE4MfVTEGEwgHY5A+cRd+OC oBqfM7TXODr1GGaGsgyCH8WyNaSZ2Vx5u+wSeTonlXo7e7csa2GGpqRcyZdV+cdX T0h6g4NRK5D7TYwrBVY0rZgs80ck/ojmT9PlatleJUPGPv1+6oI1GhVqRceZhvrj jbLiUkvbLyLsPo= 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:message-id:date:mime-version :content-type; s=default; bh=bhaBk1NOFIYZZmFnv14HZz6/E6Q=; b=XhF 9hizchhxLOp/BEQr0DOsQzfAoW8mPL/D7zU1vn7r9pQMU4nTIZP+u04uiSedL7gJ 7CpqS+dF6d0cYhfe6rBEElCCfg/BUQc+mKfVqH49zPb5daTVOSmpEiYcIVzwWNvq a6hutvNHz6axI6qag2X11aIBw6eaLtRe28YDdKxQ= Received: (qmail 75159 invoked by alias); 16 Oct 2018 12:56:25 -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 75147 invoked by uid 89); 16 Oct 2018 12:56:24 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-25.1 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, RCVD_IN_DNSWL_NONE, SPF_PASS autolearn=ham version=3.3.2 spammy=limitation, 1259, dimensions X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Tue, 16 Oct 2018 12:56:22 +0000 Received: from svr-orw-mbx-02.mgc.mentorg.com ([147.34.90.202]) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1gCOtZ-0007H6-0S from ChungLin_Tang@mentor.com ; Tue, 16 Oct 2018 05:56:21 -0700 Received: from [0.0.0.0] (147.34.91.1) by svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Tue, 16 Oct 2018 05:56:17 -0700 Reply-To: From: Chung-Lin Tang Subject: [PATCH, OpenACC, 4/8] Multi-dimensional dynamic array support for OpenACC data clauses, omp-low: dynamic array descriptor creation To: , Jakub Jelinek , Thomas Schwinge Message-ID: <4c8edea9-4685-2c19-d742-9f1a63e89f91@mentor.com> Date: Tue, 16 Oct 2018 20:56:14 +0800 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.13; rv:52.0) Gecko/20100101 Thunderbird/52.9.1 MIME-Version: 1.0 The next two patches are the bulk of the compiler patch in the middle-ends. The first patch here, implements the creation of dynamic array descriptors to pass to the runtime, a different way than completely using map-clauses. Because we support arbitrary number of dimensions, adding more map kind cases may convolute a lot of the compiler/runtime logic handling the long map sequences. This implementation uses a descriptor struct created on stack, and passes the pointer to descriptor through to the libgomp runtime, using the exact same receiver field for the dynamic array. The libgomp runtime then does its stuff to set things up, and properly adjusts the device-side receiver field pointer to the on-device created dynamic array structures. I.e. the same receiver field serves as descriptor address field on the compiler side, and as the actual data address once we get to device code (a pretty important point needed to clarify). Thanks, Chung-Lin gcc/ * omp-low.c (struct omp_context): Add 'hash_map *dynamic_arrays' field, also added include of "tree-hash-traits.h". (append_field_to_record_type): New function. (create_dynamic_array_descr_type): Likewise. (create_dynamic_array_descr_init_code): Likewise. (new_omp_context): Add initialize of dynamic_arrays field. (delete_omp_context): Add delete of dynamic_arrays field. (scan_sharing_clauses): For dynamic array map kinds, check for supported dimension structure, and install dynamic array variable into current omp_context. (lower_omp_target): Add handling for dynamic array map kinds. diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 6a1cb05..4c44800 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -60,6 +60,7 @@ along with GCC; see the file COPYING3. If not see #include "hsa-common.h" #include "stringpool.h" #include "attribs.h" +#include "tree-hash-traits.h" /* Lowering of OMP parallel and workshare constructs proceeds in two phases. The first phase scans the function looking for OMP statements @@ -124,6 +125,9 @@ struct omp_context /* True if this construct can be cancelled. */ bool cancellable; + + /* Hash map of dynamic arrays in this context. */ + hash_map *dynamic_arrays; }; static splay_tree all_contexts; @@ -843,6 +847,136 @@ omp_copy_decl (tree var, copy_body_data *cb) return error_mark_node; } +/* Helper function for create_dynamic_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 dynamic array descriptor. Returns created type, and + returns the number of dimensions in *DIM_NUM. */ + +static tree +create_dynamic_array_descr_type (tree decl, tree dims, int *dim_num) +{ + int n = 0; + tree da_descr_type, name, x; + gcc_assert (TREE_CODE (dims) == TREE_LIST); + + da_descr_type = lang_hooks.types.make_type (RECORD_TYPE); + name = create_tmp_var_name (".omp_dynamic_array_descr_type"); + name = build_decl (UNKNOWN_LOCATION, TYPE_DECL, name, da_descr_type); + DECL_ARTIFICIAL (name) = 1; + DECL_NAMELESS (name) = 1; + TYPE_NAME (da_descr_type) = name; + TYPE_ARTIFICIAL (da_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 (da_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 (da_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 (da_descr_type, get_identifier (fldname), + sizetype); + /* One for the length. */ + ASM_FORMAT_PRIVATE_NAME (fldname, "$dim_length", n); + append_field_to_record_type (da_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 (da_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 (da_descr_type, get_identifier (fldname), + sizetype); + } + + layout_type (da_descr_type); + *dim_num = n; + return da_descr_type; +} + +/* Generate code sequence for initializing dynamic array descriptor. */ + +static void +create_dynamic_array_descr_init_code (tree da_descr, tree da_var, + tree dimensions, int da_dim_num, + gimple_seq *ilist) +{ + tree fld, fldref; + tree da_descr_type = TREE_TYPE (da_descr); + tree dim_type = TREE_TYPE (da_var); + + fld = TYPE_FIELDS (da_descr_type); + fldref = omp_build_component_ref (da_descr, fld); + gimplify_assign (fldref, (TREE_CODE (dim_type) == ARRAY_TYPE + ? build_fold_addr_expr (da_var) : da_var), ilist); + + if (TREE_CODE (dim_type) == REFERENCE_TYPE) + dim_type = TREE_TYPE (dim_type); + + fld = TREE_CHAIN (fld); + fldref = omp_build_component_ref (da_descr, fld); + gimplify_assign (fldref, build_int_cst (sizetype, da_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 (da_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 (da_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 (da_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 (da_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 * @@ -877,6 +1011,8 @@ new_omp_context (gimple *stmt, omp_context *outer_ctx) ctx->cb.decl_map = new hash_map; + ctx->dynamic_arrays = new hash_map; + return ctx; } @@ -951,6 +1087,8 @@ delete_omp_context (splay_tree_value value) if (is_task_ctx (ctx)) finalize_task_copyfn (as_a (ctx->stmt)); + delete ctx->dynamic_arrays; + XDELETE (ctx); } @@ -1256,6 +1394,42 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, install_var_local (decl, ctx); break; } + + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && GOMP_MAP_DYNAMIC_ARRAY_P (OMP_CLAUSE_MAP_KIND (c))) + { + tree da_decl = OMP_CLAUSE_DECL (c); + tree da_dimensions = OMP_CLAUSE_SIZE (c); + tree da_type = TREE_TYPE (da_decl); + bool by_ref = (TREE_CODE (da_type) == ARRAY_TYPE + ? true : false); + + /* Checking code to ensure we only have arrays at top dimension. + This limitation might be lifted in the future. */ + if (TREE_CODE (da_type) == REFERENCE_TYPE) + da_type = TREE_TYPE (da_type); + tree t = da_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" + " dynamic array"); + break; + } + prev_t = t; + t = TREE_TYPE (t); + } + + install_var_field (da_decl, by_ref, 3, ctx); + tree new_var = install_var_local (da_decl, ctx); + + bool existed = ctx->dynamic_arrays->put (new_var, da_dimensions); + gcc_assert (!existed); + break; + } + if (DECL_P (decl)) { if (DECL_SIZE (decl) @@ -7687,6 +7861,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GOMP_MAP_FORCE_PRESENT: case GOMP_MAP_FORCE_DEVICEPTR: case GOMP_MAP_DEVICE_RESIDENT: + case GOMP_MAP_DYNAMIC_ARRAY_TO: + case GOMP_MAP_DYNAMIC_ARRAY_FROM: + case GOMP_MAP_DYNAMIC_ARRAY_TOFROM: + case GOMP_MAP_DYNAMIC_ARRAY_FORCE_TO: + case GOMP_MAP_DYNAMIC_ARRAY_FORCE_FROM: + case GOMP_MAP_DYNAMIC_ARRAY_FORCE_TOFROM: + case GOMP_MAP_DYNAMIC_ARRAY_ALLOC: + case GOMP_MAP_DYNAMIC_ARRAY_FORCE_ALLOC: + case GOMP_MAP_DYNAMIC_ARRAY_FORCE_PRESENT: case GOMP_MAP_LINK: gcc_assert (is_gimple_omp_oacc (stmt)); break; @@ -7749,7 +7932,14 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) 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_DYNAMIC_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 @@ -7993,6 +8183,25 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) avar = build_fold_addr_expr (avar); gimplify_assign (x, avar, &ilist); } + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_DYNAMIC_ARRAY)) + { + int da_dim_num; + tree dimensions = OMP_CLAUSE_SIZE (c); + + tree da_descr_type = + create_dynamic_array_descr_type (OMP_CLAUSE_DECL (c), + dimensions, &da_dim_num); + tree da_descr = + create_tmp_var_raw (da_descr_type, ".$omp_da_descr"); + gimple_add_tmp_var (da_descr); + + create_dynamic_array_descr_init_code + (da_descr, ovar, dimensions, da_dim_num, &ilist); + + gimplify_assign (x, build_fold_addr_expr (da_descr), + &ilist); + } else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE) { gcc_assert (is_gimple_omp_oacc (ctx->stmt)); @@ -8053,6 +8262,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) s = TREE_TYPE (s); s = TYPE_SIZE_UNIT (s); } + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_DYNAMIC_ARRAY)) + s = NULL_TREE; else s = OMP_CLAUSE_SIZE (c); if (s == NULL_TREE) From patchwork Tue Oct 16 12:56:29 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 984747 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-487643-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="oj0G5T3l"; 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 42ZFgg36vwz9s5c for ; Tue, 16 Oct 2018 23:56: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:message-id:date:mime-version :content-type; q=dns; s=default; b=e5Zb5wHTigIkkdez/4FqtrxeeLoqt N09tKZxvrt2rGS1nS3tOViFhVF5IzZ3kXhg7C2aJDAbQWFNc2jA+0EJoD8nxEm8d XazLsR//ezDIwHHF79fEs5jVfEgk4sX6SI1CXJY5Tvmkvtkagw9bzbxHtvdOgBlA P6YmsysihsYLqY= 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:message-id:date:mime-version :content-type; s=default; bh=mu+YZiRfHw9BOqSbAgsIdH/T8O4=; b=oj0 G5T3l3RmJPm9GWMOKEft/dG58w9KDq91SlehQOaQSh2lJLgdkBrucQ/e87vRZS20 QARiVRnqNy3fVyBMSOk76BJWrSyRNl0HNjN5Y4KKbGNS2xtRsmZ+a7l86YmjMOv7 iO9I2ScM17rrovhpbMvcr2eFPNtZS87JIOow8mz0= Received: (qmail 76690 invoked by alias); 16 Oct 2018 12:56:39 -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 76676 invoked by uid 89); 16 Oct 2018 12:56:39 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-25.1 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, RCVD_IN_DNSWL_NONE, SPF_PASS autolearn=ham version=3.3.2 spammy= X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Tue, 16 Oct 2018 12:56:37 +0000 Received: from svr-orw-mbx-02.mgc.mentorg.com ([147.34.90.202]) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1gCOto-0007IF-03 from ChungLin_Tang@mentor.com ; Tue, 16 Oct 2018 05:56:36 -0700 Received: from [0.0.0.0] (147.34.91.1) by svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Tue, 16 Oct 2018 05:56:32 -0700 Reply-To: From: Chung-Lin Tang Subject: [PATCH, OpenACC, 5/8] Multi-dimensional dynamic array support for OpenACC data clauses, omp-low: bias scanning/adjustment during omp-lowering To: , Jakub Jelinek , Thomas Schwinge Message-ID: Date: Tue, 16 Oct 2018 20:56:29 +0800 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.13; rv:52.0) Gecko/20100101 Thunderbird/52.9.1 MIME-Version: 1.0 This part is also in omp-low.c. We scan and adjust the code during omp-lowering, to add the biases for each dimension when a dynamic array access is detected, which is required for generally supporting copying sections of each dimension. The code is a bit sophisticated, and I wonder if this is better implemented in gimplify.c (though probably a non-trivial task as well). Nevertheless, it is currently working. Thanks, Chung-Lin gcc/ * omp-low.c (dynamic_array_lookup): New function. (dynamic_array_reference_start): Likewise. (scan_for_op): Likewise. (scan_for_reference): Likewise. (da_create_bias): Likewise. (da_dimension_peel): Likewise. (lower_omp_1): Add case to look for start of dynamic array reference, and handle bias adjustments for the code sequence. diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 6a1cb05..4c44800 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -8734,6 +8946,201 @@ lower_omp_grid_body (gimple_stmt_iterator *gsi_p, omp_context *ctx) gimple_build_omp_return (false)); } +/* Helper to lookup dynamic array through nested omp contexts. Returns + TREE_LIST of dimensions, and the CTX where it was found in *CTX_P. */ + +static tree +dynamic_array_lookup (tree t, omp_context **ctx_p) +{ + omp_context *c = *ctx_p; + while (c) + { + tree *dims = c->dynamic_arrays->get (t); + if (dims) + { + *ctx_p = c; + return *dims; + } + c = c->outer; + } + return NULL_TREE; +} + +/* Tests if this gimple STMT is the start of a dynamic array access sequence. + Returns true if found, and also returns the gimple operand ptr and + dimensions tree list through *OUT_REF and *OUT_DIMS respectively. */ + +static bool +dynamic_array_reference_start (gimple *stmt, omp_context **ctx_p, + tree **out_ref, tree *out_dims) +{ + if (gimple_code (stmt) == GIMPLE_ASSIGN) + for (unsigned i = 1; i < gimple_num_ops (stmt); i++) + { + tree *op = gimple_op_ptr (stmt, i), dims; + if (TREE_CODE (*op) == ARRAY_REF) + op = &TREE_OPERAND (*op, 0); + if (TREE_CODE (*op) == MEM_REF) + op = &TREE_OPERAND (*op, 0); + if ((dims = dynamic_array_lookup (*op, ctx_p)) != NULL_TREE) + { + *out_ref = op; + *out_dims = dims; + return true; + } + } + return false; +} + +static tree +scan_for_op (tree *tp, int *walk_subtrees, void *data) +{ + struct walk_stmt_info *wi = (struct walk_stmt_info *) data; + tree t = *tp; + tree op = (tree) wi->info; + *walk_subtrees = 1; + if (operand_equal_p (t, op, 0)) + { + wi->info = tp; + return t; + } + return NULL_TREE; +} + +static tree * +scan_for_reference (gimple *stmt, tree op) +{ + struct walk_stmt_info wi; + memset (&wi, 0, sizeof (wi)); + wi.info = op; + if (walk_gimple_op (stmt, scan_for_op, &wi)) + return (tree *) wi.info; + return NULL; +} + +static tree +da_create_bias (tree orig_bias, tree unit_type) +{ + return build2 (MULT_EXPR, sizetype, fold_convert (sizetype, orig_bias), + TYPE_SIZE_UNIT (unit_type)); +} + +/* Main worker for adjusting dynamic array accesses, handles the adjustment + of many cases of statement forms, and called multiple times to 'peel' away + each dimension. */ + +static gimple_stmt_iterator +da_dimension_peel (omp_context *da_ctx, + gimple_stmt_iterator da_gsi, tree orig_da, + tree *da_op_p, tree *da_type_p, tree *da_dims_p) +{ + gimple *stmt = gsi_stmt (da_gsi); + tree lhs = gimple_assign_lhs (stmt); + tree rhs = gimple_assign_rhs1 (stmt); + + if (gimple_num_ops (stmt) == 2 + && TREE_CODE (rhs) == MEM_REF + && operand_equal_p (*da_op_p, TREE_OPERAND (rhs, 0), 0) + && !operand_equal_p (orig_da, TREE_OPERAND (rhs, 0), 0) + && (TREE_OPERAND (rhs, 1) == NULL_TREE + || integer_zerop (TREE_OPERAND (rhs, 1)))) + { + gcc_assert (TREE_CODE (TREE_TYPE (*da_type_p)) == POINTER_TYPE); + *da_type_p = TREE_TYPE (*da_type_p); + } + else + { + gimple *g; + gimple_seq ilist = NULL; + tree bias, t; + tree op = *da_op_p; + tree orig_type = *da_type_p; + tree orig_bias = TREE_PURPOSE (*da_dims_p); + bool by_ref = false; + + if (TREE_CODE (orig_bias) != INTEGER_CST) + orig_bias = lookup_decl (orig_bias, da_ctx); + + if (gimple_num_ops (stmt) == 2) + { + if (TREE_CODE (rhs) == ADDR_EXPR) + { + rhs = TREE_OPERAND (rhs, 0); + *da_dims_p = NULL_TREE; + } + + if (TREE_CODE (rhs) == ARRAY_REF + && TREE_CODE (TREE_OPERAND (rhs, 0)) == MEM_REF + && operand_equal_p (TREE_OPERAND (TREE_OPERAND (rhs, 0), 0), + *da_op_p, 0)) + { + bias = da_create_bias (orig_bias, + TREE_TYPE (TREE_TYPE (orig_type))); + *da_type_p = TREE_TYPE (TREE_TYPE (orig_type)); + } + else if (TREE_CODE (rhs) == ARRAY_REF + && TREE_CODE (TREE_OPERAND (rhs, 0)) == VAR_DECL + && operand_equal_p (TREE_OPERAND (rhs, 0), *da_op_p, 0)) + { + tree ptr_type = build_pointer_type (orig_type); + op = create_tmp_var (ptr_type); + gimplify_assign (op, build_fold_addr_expr (TREE_OPERAND (rhs, 0)), + &ilist); + bias = da_create_bias (orig_bias, TREE_TYPE (orig_type)); + *da_type_p = TREE_TYPE (orig_type); + orig_type = ptr_type; + by_ref = true; + } + else if (TREE_CODE (rhs) == MEM_REF + && operand_equal_p (*da_op_p, TREE_OPERAND (rhs, 0), 0) + && TREE_OPERAND (rhs, 1) != NULL_TREE) + { + bias = da_create_bias (orig_bias, TREE_TYPE (orig_type)); + *da_type_p = TREE_TYPE (orig_type); + } + else if (TREE_CODE (lhs) == MEM_REF + && operand_equal_p (*da_op_p, TREE_OPERAND (lhs, 0), 0)) + { + if (*da_dims_p != NULL_TREE) + { + gcc_assert (TREE_CHAIN (*da_dims_p) == NULL_TREE); + bias = da_create_bias (orig_bias, TREE_TYPE (orig_type)); + *da_type_p = TREE_TYPE (orig_type); + } + else + /* This should be the end of the dynamic array access + sequence. */ + return da_gsi; + } + else + gcc_unreachable (); + } + else if (gimple_num_ops (stmt) == 3 + && gimple_assign_rhs_code (stmt) == POINTER_PLUS_EXPR + && operand_equal_p (*da_op_p, rhs, 0)) + { + bias = da_create_bias (orig_bias, TREE_TYPE (orig_type)); + } + else + gcc_unreachable (); + + bias = fold_build1 (NEGATE_EXPR, sizetype, bias); + bias = fold_build2 (POINTER_PLUS_EXPR, orig_type, op, bias); + + t = create_tmp_var (by_ref ? build_pointer_type (orig_type) : orig_type); + + g = gimplify_assign (t, bias, &ilist); + gsi_insert_seq_before (&da_gsi, ilist, GSI_NEW_STMT); + *da_op_p = gimple_assign_lhs (g); + + if (by_ref) + *da_op_p = build2 (MEM_REF, TREE_TYPE (orig_type), *da_op_p, + build_int_cst (orig_type, 0)); + *da_dims_p = TREE_CHAIN (*da_dims_p); + } + + return da_gsi; +} /* Callback for lower_omp_1. Return non-NULL if *tp needs to be regimplified. If DATA is non-NULL, lower_omp_1 is outside @@ -9009,6 +9416,51 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx) } /* FALLTHRU */ default: + + /* If we detect the start of a dynamic array reference sequence, scan + and do the needed adjustments. */ + tree da_dims, *da_op_p; + omp_context *da_ctx = ctx; + if (da_ctx && dynamic_array_reference_start (stmt, &da_ctx, + &da_op_p, &da_dims)) + { + bool started = false; + tree orig_da = *da_op_p; + tree da_type = TREE_TYPE (orig_da); + tree next_da_op; + + gimple_stmt_iterator da_gsi = *gsi_p, new_gsi; + while (da_op_p) + { + if (!is_gimple_assign (gsi_stmt (da_gsi)) + || ((gimple_assign_single_p (gsi_stmt (da_gsi)) + || gimple_assign_cast_p (gsi_stmt (da_gsi))) + && *da_op_p == gimple_assign_rhs1 (gsi_stmt (da_gsi)))) + break; + + new_gsi = da_dimension_peel (da_ctx, da_gsi, orig_da, + da_op_p, &da_type, &da_dims); + if (!started) + { + /* Point 'stmt' to the start of the newly added + sequence. */ + started = true; + *gsi_p = new_gsi; + stmt = gsi_stmt (*gsi_p); + } + if (!da_dims) + break; + + next_da_op = gimple_assign_lhs (gsi_stmt (da_gsi)); + + do { + gsi_next (&da_gsi); + da_op_p = scan_for_reference (gsi_stmt (da_gsi), next_da_op); + } + while (!da_op_p); + } + } + if ((ctx || task_shared_vars) && walk_gimple_op (stmt, lower_omp_regimplify_p, ctx ? NULL : &wi)) From patchwork Tue Oct 16 12:56:45 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 984748 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-487644-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="w8a6LQqC"; 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 42ZFgy37Jgz9s9G for ; Tue, 16 Oct 2018 23:57:00 +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:message-id:date:mime-version :content-type; q=dns; s=default; b=wU3O7PBAOZ0rvRQgHjVInHZiEigiP K/BLIzPCRvVGwYam8EewuA3zSfsSAUfZ1efjiz1B3FylfpbT47MuRxlaSh1q2qUk 8Z6x67hB1SAq2tLoVSI70h87eYWrtHAu8MT6ReFzCK/UuwX76L/eVGmmdU2LdP6C x9y2VO/ZOvWIHQ= 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:message-id:date:mime-version :content-type; s=default; bh=oeb7Nh6nD/qUYNRhuT2ROgIZsRg=; b=w8a 6LQqCG4Lll0rckWArCqkr7pR8bycSae/1oRktroeKVibijaGSvP3mXZf+JzWp9Ex XxQqQdV++LGNLx5xqU8vhhIayhqcpzk3bfeadzlR8AjYbtPkuK5s8op4h3li+SsY rrHQWpiDt7M+8GCPmOKuKLKNa4oQMDtQ3zd/f8iE= Received: (qmail 78228 invoked by alias); 16 Oct 2018 12:56:54 -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 78211 invoked by uid 89); 16 Oct 2018 12:56:53 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-25.1 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, RCVD_IN_DNSWL_NONE, SPF_PASS autolearn=ham version=3.3.2 spammy=H*r:0700 X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Tue, 16 Oct 2018 12:56:52 +0000 Received: from svr-orw-mbx-02.mgc.mentorg.com ([147.34.90.202]) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1gCOu3-0007JY-3l from ChungLin_Tang@mentor.com ; Tue, 16 Oct 2018 05:56:51 -0700 Received: from [0.0.0.0] (147.34.91.1) by svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Tue, 16 Oct 2018 05:56:48 -0700 Reply-To: From: Chung-Lin Tang Subject: [PATCH, OpenACC, 6/8] Multi-dimensional dynamic array support for OpenACC data clauses, tree pretty-printing additions To: , Jakub Jelinek , Thomas Schwinge Message-ID: <426c4819-b42c-090a-c0df-1dfc3f0924a7@mentor.com> Date: Tue, 16 Oct 2018 20:56:45 +0800 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.13; rv:52.0) Gecko/20100101 Thunderbird/52.9.1 MIME-Version: 1.0 This tree-pretty-print.c patch allows proper dumping of the dynamic arrays case of OMP_CLAUSE_MAP. Thanks, Chung-Lin gcc/ * tree-pretty-print.c (dump_omp_clauses): Add cases for printing GOMP_MAP_DYNAMIC_ARRAY map kinds. diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index 1c7982c..803f76b 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -745,6 +745,33 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_LINK: pp_string (pp, "link"); break; + case GOMP_MAP_DYNAMIC_ARRAY_TO: + pp_string (pp, "to,dynamic_array"); + break; + case GOMP_MAP_DYNAMIC_ARRAY_FROM: + pp_string (pp, "from,dynamic_array"); + break; + case GOMP_MAP_DYNAMIC_ARRAY_TOFROM: + pp_string (pp, "tofrom,dynamic_array"); + break; + case GOMP_MAP_DYNAMIC_ARRAY_FORCE_TO: + pp_string (pp, "force_to,dynamic_array"); + break; + case GOMP_MAP_DYNAMIC_ARRAY_FORCE_FROM: + pp_string (pp, "force_from,dynamic_array"); + break; + case GOMP_MAP_DYNAMIC_ARRAY_FORCE_TOFROM: + pp_string (pp, "force_tofrom,dynamic_array"); + break; + case GOMP_MAP_DYNAMIC_ARRAY_ALLOC: + pp_string (pp, "alloc,dynamic_array"); + break; + case GOMP_MAP_DYNAMIC_ARRAY_FORCE_ALLOC: + pp_string (pp, "force_alloc,dynamic_array"); + break; + case GOMP_MAP_DYNAMIC_ARRAY_FORCE_PRESENT: + pp_string (pp, "force_present,dynamic_array"); + break; default: gcc_unreachable (); } @@ -766,6 +793,10 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_TO_PSET: pp_string (pp, " [pointer set, len: "); break; + case GOMP_MAP_DYNAMIC_ARRAY: + gcc_assert (TREE_CODE (OMP_CLAUSE_SIZE (clause)) == TREE_LIST); + pp_string (pp, " [dimensions: "); + break; default: pp_string (pp, " [len: "); break; From patchwork Tue Oct 16 12:57:00 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 984749 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-487645-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="KRLEw91k"; 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 42ZFhG6JHgz9s9G for ; Tue, 16 Oct 2018 23:57:18 +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:message-id:date:mime-version :content-type; q=dns; s=default; b=o5VBMFSi2mrD6AHw/Pc388x5ZXeT5 ZMCtr3Q7SQRvpXp+vgcw7uOw8ou8djdD7gre1X7VT6GKlN1t9ChJPwZxNkajt8Wh u6NSBA570qErawgo3/48uJuznLOqByi6jErY1Je2RIaDu/uFqW89BIMHh8fw6cC9 7zrX7Z4GLDpNis= 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:message-id:date:mime-version :content-type; s=default; bh=Qx66ls5+xZ+20GZ3Ezma+kTSUHY=; b=KRL Ew91kpTftk/q2S3sAuYJ1YIBefCeojVTuYb46xyOVuzFqVP2LyjWKbCy9VDPXzJd uVaBYIFV46wxaybYVa3brfOelODrpaU5bTkFwDPFdYb2pO+NjAuG5be0Xobrs7Lv 0Qz2QwTBP5D0pBq0zjXBD/0HcmLpm41v2a5kceLw= Received: (qmail 79880 invoked by alias); 16 Oct 2018 12:57:11 -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 79852 invoked by uid 89); 16 Oct 2018 12:57:10 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-22.6 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, RCVD_IN_DNSWL_NONE, SPF_PASS, UNWANTED_LANGUAGE_BODY autolearn=ham version=3.3.2 spammy=tgt, Processing, row, filling X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Tue, 16 Oct 2018 12:57:08 +0000 Received: from svr-orw-mbx-02.mgc.mentorg.com ([147.34.90.202]) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1gCOuJ-0007Lv-3J from ChungLin_Tang@mentor.com ; Tue, 16 Oct 2018 05:57:07 -0700 Received: from [0.0.0.0] (147.34.91.1) by svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Tue, 16 Oct 2018 05:57:04 -0700 Reply-To: From: Chung-Lin Tang Subject: [PATCH, OpenACC, 7/8] Multi-dimensional dynamic array support for OpenACC data clauses, libgomp support To: , Jakub Jelinek , Thomas Schwinge Message-ID: Date: Tue, 16 Oct 2018 20:57:00 +0800 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.13; rv:52.0) Gecko/20100101 Thunderbird/52.9.1 MIME-Version: 1.0 This part is the libgomp runtime handling for OpenACC dynamic arrays. We handle such arrays by creating a "pointer block" that emulates the N-1 dimensions, and then treating each data row of the final Nth dimension as an individual object mapped in the TGT. All the rows are processed as appended after all the other map kind objects. Thanks, Chung-Lin libgomp/ * target.c (struct da_dim): New struct declaration. (struct da_descr_type): Likewise. (struct da_info): Likewise. (gomp_dynamic_array_count_rows): New function. (gomp_dynamic_array_compute_info): Likewise. (gomp_dynamic_array_fill_rows_1): Likewise. (gomp_dynamic_array_fill_rows): Likewise. (gomp_dynamic_array_create_ptrblock): Likewise. (gomp_map_vars): Add code to handle dynamic array map kinds. diff --git a/libgomp/target.c b/libgomp/target.c index 4c9fae0..071dc70 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -490,6 +490,140 @@ gomp_map_val (struct target_mem_desc *tgt, void **hostaddrs, size_t i) return tgt->tgt_start + tgt->list[i].offset; } +/* Dynamic array related data structures, interfaces with the compiler. */ + +struct da_dim { + size_t base; + size_t length; + size_t elem_size; + size_t is_array; +}; + +struct da_descr_type { + void *ptr; + size_t ndims; + struct da_dim dims[]; +}; + +/* Internal dynamic array info struct, used only here inside the runtime. */ + +struct da_info +{ + struct da_descr_type *descr; + size_t map_index; + size_t ptrblock_size; + size_t data_row_num; + size_t data_row_size; +}; + +static size_t +gomp_dynamic_array_count_rows (struct da_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 +gomp_dynamic_array_compute_info (struct da_info *da) +{ + size_t d, n = 1; + struct da_descr_type *descr = da->descr; + + da->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); + da->ptrblock_size += dim_ptrblock_size; + n *= dim_count; + } + da->data_row_num = n; + da->data_row_size = descr->dims[d].length; +} + +static void +gomp_dynamic_array_fill_rows_1 (struct da_descr_type *descr, void *da, + 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 = da + 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); + gomp_dynamic_array_fill_rows_1 (descr, ptr, d + 1, row_ptr, count); + } + } + else + { + **row_ptr = da + descr->dims[d].base; + *row_ptr += 1; + *count += 1; + } +} + +static size_t +gomp_dynamic_array_fill_rows (struct da_descr_type *descr, void *rows[]) +{ + size_t count = 0; + void **p = rows; + gomp_dynamic_array_fill_rows_1 (descr, descr->ptr, 0, &p, &count); + return count; +} + +static void * +gomp_dynamic_array_create_ptrblock (struct da_info *da, + void *tgt_addr, void *tgt_data_rows[]) +{ + struct da_descr_type *descr = da->descr; + void *ptrblock = gomp_malloc (da->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 *); + + 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_addr + (ptr - ptrblock); + 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]; + } + void *addr = &curr_dim_ptrblock[b * curr_dim_num + i]; + assert (ptrblock <= addr && addr < ptrblock + da->ptrblock_size); + } + + n *= curr_dim_num; + curr_dim_ptrblock = next_dim_ptrblock; + } + assert (n == da->data_row_num); + return ptrblock; +} + 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, @@ -501,9 +635,29 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, const int typemask = short_mapkind ? 0xff : 0x7; 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; + struct target_mem_desc *tgt; + + size_t da_data_row_num = 0, row_start = 0; + size_t da_info_num = 0, da_index; + struct da_info *da_info = NULL; + struct target_var_desc *row_desc; + uintptr_t target_row_addr; + void **host_data_rows = NULL, **target_data_rows = NULL; + void *row; + + for (i = 0; i < mapnum; i++) + { + int kind = get_kind (short_mapkind, kinds, i); + if (GOMP_MAP_DYNAMIC_ARRAY_P (kind & typemask)) + { + da_data_row_num += gomp_dynamic_array_count_rows (hostaddrs[i]); + da_info_num += 1; + } + } + + tgt = gomp_malloc (sizeof (*tgt) + + sizeof (tgt->list[0]) * (mapnum + da_data_row_num)); + tgt->list_count = mapnum + da_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; @@ -515,6 +669,14 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, return tgt; } + if (da_info_num) + da_info = gomp_alloca (sizeof (struct da_info) * da_info_num); + if (da_data_row_num) + { + host_data_rows = gomp_malloc (sizeof (void *) * da_data_row_num); + target_data_rows = gomp_malloc (sizeof (void *) * da_data_row_num); + } + tgt_align = sizeof (void *); tgt_size = 0; cbuf.chunks = NULL; @@ -546,7 +708,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, return NULL; } - for (i = 0; i < mapnum; i++) + for (i = 0, da_index = 0; i < mapnum; i++) { int kind = get_kind (short_mapkind, kinds, i); if (hostaddrs[i] == NULL @@ -619,6 +781,20 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, has_firstprivate = true; continue; } + else if (GOMP_MAP_DYNAMIC_ARRAY_P (kind & typemask)) + { + /* Ignore dynamic arrays for now, we process them together + later. */ + tgt->list[i].key = NULL; + tgt->list[i].offset = 0; + not_found_cnt++; + + struct da_info *da = &da_info[da_index++]; + da->descr = (struct da_descr_type *) hostaddrs[i]; + da->map_index = i; + 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]; @@ -687,6 +863,55 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, } } + /* For dynamic arrays. Each data row is one target item, separated from + the normal map clause items, hence we order them after mapnum. */ + for (i = 0, da_index = 0, row_start = 0; i < mapnum; i++) + { + int kind = get_kind (short_mapkind, kinds, i); + if (!GOMP_MAP_DYNAMIC_ARRAY_P (kind & typemask)) + continue; + + struct da_info *da = &da_info[da_index++]; + struct da_descr_type *descr = da->descr; + size_t nr; + + gomp_dynamic_array_compute_info (da); + + /* We have allocated space in host/target_data_rows to place all the + row data block pointers, now we can start filling them in. */ + nr = gomp_dynamic_array_fill_rows (descr, &host_data_rows[row_start]); + assert (nr == da->data_row_num); + + size_t align = (size_t) 1 << (kind >> rshift); + if (tgt_align < align) + tgt_align = align; + tgt_size = (tgt_size + align - 1) & ~(align - 1); + tgt_size += da->ptrblock_size; + + for (size_t j = 0; j < da->data_row_num; j++) + { + row = host_data_rows[row_start + j]; + row_desc = &tgt->list[mapnum + row_start + j]; + + cur_node.host_start = (uintptr_t) row; + cur_node.host_end = cur_node.host_start + da->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, n, &cur_node, row_desc, + kind & typemask, /* TODO: cbuf? */ NULL); + } + else + { + tgt_size = (tgt_size + align - 1) & ~(align - 1); + tgt_size += da->data_row_size; + not_found_cnt++; + } + } + row_start += da->data_row_num; + } + if (devaddrs) { if (mapnum != 1) @@ -830,6 +1055,15 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, default: break; } + + if (GOMP_MAP_DYNAMIC_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)) @@ -976,6 +1210,108 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, array++; } } + + /* Processing of dynamic array rows. */ + for (i = 0, da_index = 0, row_start = 0; i < mapnum; i++) + { + int kind = get_kind (short_mapkind, kinds, i); + if (!GOMP_MAP_DYNAMIC_ARRAY_P (kind & typemask)) + continue; + + struct da_info *da = &da_info[da_index++]; + assert (da->descr == hostaddrs[i]); + + /* The map for the dynamic array itself is never copied from during + unmapping, its the data rows that count. Set copy from flags are + set to false here. */ + tgt->list[i].copy_from = false; + tgt->list[i].always_copy_from = false; + + size_t align = (size_t) 1 << (kind >> rshift); + tgt_size = (tgt_size + align - 1) & ~(align - 1); + + /* For the map of the dynamic array itself, adjust so that the passed + device address points to the beginning of the ptrblock. */ + tgt->list[i].key->tgt_offset = tgt_size; + + void *target_ptrblock = (void*) tgt->tgt_start + tgt_size; + tgt_size += da->ptrblock_size; + + /* Add splay key for each data row in current DA. */ + for (size_t j = 0; j < da->data_row_num; j++) + { + row = host_data_rows[row_start + j]; + row_desc = &tgt->list[mapnum + row_start + j]; + + cur_node.host_start = (uintptr_t) row; + cur_node.host_end = cur_node.host_start + da->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, n, &cur_node, row_desc, + kind & typemask, cbufp); + target_row_addr = n->tgt->tgt_start + n->tgt_offset; + } + else + { + tgt->refcount++; + + splay_tree_key k = &array->key; + k->host_start = (uintptr_t) row; + k->host_end = k->host_start + da->data_row_size; + + k->tgt = tgt; + k->refcount = 1; + k->link_key = NULL; + tgt_size = (tgt_size + align - 1) & ~(align - 1); + target_row_addr = tgt->tgt_start + tgt_size; + k->tgt_offset = tgt_size; + tgt_size += da->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 = da->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, + (void *) tgt->tgt_start + k->tgt_offset, + (void *) k->host_start, + da->data_row_size, cbufp); + array++; + } + target_data_rows[row_start + j] = (void *) target_row_addr; + } + + /* 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 (da->ptrblock_size) + { + void *ptrblock = gomp_dynamic_array_create_ptrblock + (da, target_ptrblock, target_data_rows + row_start); + gomp_copy_host2dev (devicep, target_ptrblock, ptrblock, + da->ptrblock_size, cbufp); + free (ptrblock); + } + + row_start += da->data_row_num; + } + assert (row_start == da_data_row_num && da_index == da_info_num); + } + + if (da_data_row_num) + { + free (host_data_rows); + free (target_data_rows); } if (pragma_kind == GOMP_MAP_VARS_TARGET) From patchwork Tue Oct 16 12:57:16 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 984750 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-487646-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="B8EWkUhs"; 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 42ZFhZ3YZWz9s8r for ; Tue, 16 Oct 2018 23:57:34 +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:message-id:date:mime-version :content-type; q=dns; s=default; b=bvV4kjUlX1AFrWMZnSvKLeVH+zKLb Kn0Y4Nz3Q57v+LUuU58T1hadhax9UkpG/Zjp6yV7r3xRiDp1OQtFNYXww9EM1s95 F50VSOsuwBzwn4HemjEqQAqC7oGNqEUUfJobgvd5rOjJCwk6KvwRfI4RV+22yBqD pmMQoG7VolypQ4= 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:message-id:date:mime-version :content-type; s=default; bh=rvMuQCE8lLWP+OQwhdmKifZ9yNw=; b=B8E WkUhsYpzJfuNoiLYxiYpTFEthFkYjNm6NsX0vmY/+k4oam2B5j+Qbcn98VpFEcLb kQ74k0asRH1bBNEmGbzZ1ovIQB6rnQ4jxjXSZEg33Ge0v5RdjA/j43DqN7zvXzJZ N/RDg7Q4t+YZCSF9p3sz6MKlhBAE1SPNAkh+usPE= Received: (qmail 81627 invoked by alias); 16 Oct 2018 12:57:26 -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 81614 invoked by uid 89); 16 Oct 2018 12:57:25 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-25.2 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, RCVD_IN_DNSWL_NONE, SPF_PASS autolearn=ham version=3.3.2 spammy=exercises, gang, Array X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Tue, 16 Oct 2018 12:57:23 +0000 Received: from svr-orw-mbx-02.mgc.mentorg.com ([147.34.90.202]) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1gCOuY-0007NJ-51 from ChungLin_Tang@mentor.com for gcc-patches@gcc.gnu.org; Tue, 16 Oct 2018 05:57:22 -0700 Received: from [0.0.0.0] (147.34.91.1) by svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Tue, 16 Oct 2018 05:57:19 -0700 Reply-To: From: Chung-Lin Tang Subject: [PATCH, OpenACC, 8/8] Multi-dimensional dynamic array support for OpenACC data clauses, libgomp testsuite additions To: , Thomas Schwinge Message-ID: <40c6cbaf-8599-b958-ad37-9c61769c09b8@mentor.com> Date: Tue, 16 Oct 2018 20:57:16 +0800 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.13; rv:52.0) Gecko/20100101 Thunderbird/52.9.1 MIME-Version: 1.0 These are the added cases for testing the OpenACC dynamic (sub)arrays functionality. Thanks, Chung-Lin libgomp/ * testsuite/libgomp.oacc-c-c++-common/da-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/da-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/da-3.c: New test. * testsuite/libgomp.oacc-c-c++-common/da-4.c: New test. * testsuite/libgomp.oacc-c-c++-common/da-utils.h: New test. diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/da-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/da-1.c new file mode 100644 index 0000000..c1c205d --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/da-1.c @@ -0,0 +1,103 @@ +/* { dg-do run { target { ! openacc_host_selected } } } */ + +#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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/da-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/da-2.c new file mode 100644 index 0000000..6ee7855 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/da-2.c @@ -0,0 +1,37 @@ +/* { dg-do run { target { ! openacc_host_selected } } } */ + +#include +#include "da-utils.h" + +int +main (void) +{ + int n = 10; + int ***a = (int ***) create_da (sizeof (int), n, 3); + int ***b = (int ***) create_da (sizeof (int), n, 3); + int ***c = (int ***) create_da (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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/da-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/da-3.c new file mode 100644 index 0000000..877c6df --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/da-3.c @@ -0,0 +1,45 @@ +/* { dg-do run { target { ! openacc_host_selected } } } */ + +#include +#include "da-utils.h" + +int main (void) +{ + int n = 20, x = 5, y = 12; + int *****a = (int *****) create_da (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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/da-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/da-4.c new file mode 100644 index 0000000..2059c5f --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/da-4.c @@ -0,0 +1,36 @@ +/* { dg-do run { target { ! openacc_host_selected } } } */ + +#include +#include "da-utils.h" + +int main (void) +{ + int n = 128; + double ***a = (double ***) create_da (sizeof (double), n, 3); + double ***b = (double ***) create_da (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 dynamic 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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/da-utils.h b/libgomp/testsuite/libgomp.oacc-c-c++-common/da-utils.h new file mode 100644 index 0000000..2f87795 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/da-utils.h @@ -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_da (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; +}