From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 1111 invoked by alias); 13 Feb 2014 13:15:59 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 1090 invoked by uid 89); 13 Feb 2014 13:15:58 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.0 required=5.0 tests=AWL,BAYES_00,RP_MATCHES_RCVD,SPF_HELO_PASS autolearn=ham version=3.3.2 X-Spam-User: qpsmtpd, 2 recipients X-HELO: mailout2.w1.samsung.com Received: from mailout2.w1.samsung.com (HELO mailout2.w1.samsung.com) (210.118.77.12) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (DES-CBC3-SHA encrypted) ESMTPS; Thu, 13 Feb 2014 13:15:54 +0000 Received: from eucpsbgm2.samsung.com (unknown [203.254.199.245]) by mailout2.w1.samsung.com (Oracle Communications Messaging Server 7u4-24.01(7.0.4.24.0) 64bit (built Nov 17 2011)) with ESMTP id <0N0X00806RI9RFA0@mailout2.w1.samsung.com>; Thu, 13 Feb 2014 13:15:45 +0000 (GMT) Received: from eusync3.samsung.com ( [203.254.199.213]) by eucpsbgm2.samsung.com (EUCPMTA) with SMTP id 75.0A.18565.585CCF25; Thu, 13 Feb 2014 13:15:49 +0000 (GMT) Received: from [106.109.130.115] by eusync3.samsung.com (Oracle Communications Messaging Server 7u4-23.01(7.0.4.23.0) 64bit (built Aug 10 2011)) with ESMTPA id <0N0X007B6RICIV10@eusync3.samsung.com>; Thu, 13 Feb 2014 13:15:49 +0000 (GMT) Message-id: <52FCC583.2050709@samsung.com> Date: Thu, 13 Feb 2014 13:15:00 -0000 From: Ilmir Usmanov User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:24.0) Gecko/20100101 Thunderbird/24.0 MIME-version: 1.0 To: Thomas Schwinge Cc: Evgeny Gavrin , GarbuzovViacheslav , Dmitri Botcharnikov , gcc-patches@gcc.gnu.org, jakub@redhat.com, fortran@gcc.gnu.org Subject: Re: [PATCH 4/6] [GOMP4] OpenACC 1.0+ support in fortran front-end References: <52E158EF.9050009@samsung.com> <877g9pqmt2.fsf@schwinge.name> <52E65B24.9070403@samsung.com> <87iot5pgqb.fsf@schwinge.name> <52EB8437.3060602@samsung.com> <52EB84C3.4010407@samsung.com> <52EB8537.4090708@samsung.com> <52EB8596.2010808@samsung.com> <52EB85F7.5020807@samsung.com> <8761olk2yk.fsf@schwinge.name> In-reply-to: <8761olk2yk.fsf@schwinge.name> Content-type: multipart/mixed; boundary=------------010805090902090604010307 X-IsSubscribed: yes X-SW-Source: 2014-02/txt/msg00836.txt.bz2 This is a multi-part message in MIME format. --------------010805090902090604010307 Content-Type: text/plain; charset=UTF-8; format=flowed Content-Transfer-Encoding: 7bit Content-length: 714 Hi Thomas! Thanks a lot for your review! I agree with all your notes. On 11.02.2014 20:51, Thomas Schwinge wrote: > For ChangeLog files updates (on gomp-4_0-branch, use the respective > ChangeLog.gomp files, by the way), should just you be listed as the > author, or also your colleagues? Thank you for the notice, I added Evgeny and Dmitry as authors for this part (see attached ChangeLog entry). > With these issues addressed, this patch is ready for commit to > gomp-4_0-branch. Use your own judgement; if you feel confident, just > commit it, or otherwise post it again for a final review -- as you > prefer. I fixed patch according to your review and ready to commit it. OK for GOMP4 branch? -- Ilmir. --------------010805090902090604010307 Content-Type: text/x-diff; name="0001-OpenACC-GENERIC-nodes.patch" Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="0001-OpenACC-GENERIC-nodes.patch" Content-length: 29011 >From bf14158b1a28c2c5b29c41071fa62c011d9f4f65 Mon Sep 17 00:00:00 2001 From: Ilmir Usmanov Date: Thu, 13 Feb 2014 15:58:28 +0400 Subject: [PATCH] OpenACC GENERIC nodes --- gcc/doc/generic.texi | 45 ++++++++++++++++++ gcc/gimplify.c | 62 +++++++++++++++++++++++++ gcc/omp-low.c | 96 ++++++++++++++++++++++++++++++++------ gcc/tree-core.h | 61 ++++++++++++++++++++++--- gcc/tree-pretty-print.c | 119 ++++++++++++++++++++++++++++++++++++++++++++++++ gcc/tree.c | 44 +++++++++++++++++- gcc/tree.def | 42 +++++++++++++++++ gcc/tree.h | 61 ++++++++++++++++++++++++- 8 files changed, 507 insertions(+), 23 deletions(-) diff --git a/gcc/doc/generic.texi b/gcc/doc/generic.texi index a56715b..ce14620 100644 --- a/gcc/doc/generic.texi +++ b/gcc/doc/generic.texi @@ -2052,6 +2052,15 @@ edge. Rethrowing the exception is represented using @code{RESX_EXPR}. @node OpenMP @subsection OpenMP @tindex OACC_PARALLEL +@tindex OACC_KERNELS +@tindex OACC_DATA +@tindex OACC_HOST_DATA +@tindex OACC_DECLARE +@tindex OACC_UPDATE +@tindex OACC_ENTER_DATA +@tindex OACC_EXIT_DATA +@tindex OACC_WAIT +@tindex OACC_CACHE @tindex OMP_PARALLEL @tindex OMP_FOR @tindex OMP_SECTIONS @@ -2073,6 +2082,42 @@ clauses used by the OpenMP API @w{@uref{http://www.openmp.org/}}. Represents @code{#pragma acc parallel [clause1 @dots{} clauseN]}. +@item OACC_KERNELS + +Represents @code{#pragma acc kernels [clause1 @dots{} clauseN]}. + +@item OACC_DATA + +Represents @code{#pragma acc data [clause1 @dots{} clauseN]}. + +@item OACC_HOST_DATA + +Represents @code{#pragma acc host_data [clause1 @dots{} clauseN]}. + +@item OACC_DECLARE + +Represents @code{#pragma acc declare [clause1 @dots{} clauseN]}. + +@item OACC_UPDATE + +Represents @code{#pragma acc update [clause1 @dots{} clauseN]}. + +@item OACC_ENTER_DATA + +Represents @code{#pragma acc enter data [clause1 @dots{} clauseN]}. + +@item OACC_EXIT_DATA + +Represents @code{#pragma acc exit data [clause1 @dots{} clauseN]}. + +@item OACC_WAIT + +Represents @code{#pragma acc wait [(num @dots{})]}. + +@item OACC_CACHE + +Represents @code{#pragma acc cache (var @dots{})}. + @item OMP_PARALLEL Represents @code{#pragma omp parallel [clause1 @dots{} clauseN]}. It diff --git a/gcc/gimplify.c b/gcc/gimplify.c index d20f07f..06d7790 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -4333,6 +4333,15 @@ is_gimple_stmt (tree t) case ASM_EXPR: case STATEMENT_LIST: case OACC_PARALLEL: + case OACC_KERNELS: + case OACC_DATA: + case OACC_HOST_DATA: + case OACC_DECLARE: + case OACC_UPDATE: + case OACC_ENTER_DATA: + case OACC_EXIT_DATA: + case OACC_WAIT: + case OACC_CACHE: case OMP_PARALLEL: case OMP_FOR: case OMP_SIMD: @@ -6157,6 +6166,23 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, remove = true; break; + case OMP_CLAUSE_HOST: + case OMP_CLAUSE_OACC_DEVICE: + case OMP_CLAUSE_DEVICE_RESIDENT: + case OMP_CLAUSE_USE_DEVICE: + case OMP_CLAUSE_GANG: + case OMP_CLAUSE_WAIT: + case OMP_NO_CLAUSE_CACHE: + case OMP_CLAUSE_INDEPENDENT: + case OMP_CLAUSE_ASYNC: + case OMP_CLAUSE_WORKER: + case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_NUM_GANGS: + case OMP_CLAUSE_NUM_WORKERS: + case OMP_CLAUSE_VECTOR_LENGTH: + remove = true; + break; + case OMP_CLAUSE_NOWAIT: case OMP_CLAUSE_ORDERED: case OMP_CLAUSE_UNTIED: @@ -6498,6 +6524,20 @@ gimplify_adjust_omp_clauses (tree *list_p) case OMP_CLAUSE_DEPEND: break; + case OMP_CLAUSE_HOST: + case OMP_CLAUSE_OACC_DEVICE: + case OMP_CLAUSE_DEVICE_RESIDENT: + case OMP_CLAUSE_USE_DEVICE: + case OMP_CLAUSE_GANG: + case OMP_CLAUSE_WAIT: + case OMP_NO_CLAUSE_CACHE: + case OMP_CLAUSE_INDEPENDENT: + case OMP_CLAUSE_ASYNC: + case OMP_CLAUSE_WORKER: + case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_NUM_GANGS: + case OMP_CLAUSE_NUM_WORKERS: + case OMP_CLAUSE_VECTOR_LENGTH: default: gcc_unreachable (); } @@ -7988,6 +8028,19 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, ret = GS_ALL_DONE; break; + case OACC_KERNELS: + case OACC_DATA: + case OACC_HOST_DATA: + case OACC_DECLARE: + case OACC_UPDATE: + case OACC_ENTER_DATA: + case OACC_EXIT_DATA: + case OACC_WAIT: + case OACC_CACHE: + sorry ("directive not yet implemented"); + ret = GS_ALL_DONE; + break; + case OMP_PARALLEL: gimplify_omp_parallel (expr_p, pre_p); ret = GS_ALL_DONE; @@ -8396,6 +8449,15 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, && code != SWITCH_EXPR && code != TRY_FINALLY_EXPR && code != OACC_PARALLEL + && code != OACC_KERNELS + && code != OACC_DATA + && code != OACC_HOST_DATA + && code != OACC_DECLARE + && code != OACC_UPDATE + && code != OACC_ENTER_DATA + && code != OACC_EXIT_DATA + && code != OACC_WAIT + && code != OACC_CACHE && code != OMP_CRITICAL && code != OMP_FOR && code != OMP_MASTER diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 4bbe6d6..fd94e63 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -1491,6 +1491,18 @@ fixup_child_record_type (omp_context *ctx) TREE_TYPE (ctx->receiver_decl) = build_pointer_type (type); } +static bool +gimple_code_is_oacc (const_gimple g) +{ + switch (gimple_code (g)) + { + case GIMPLE_OACC_PARALLEL: + return true; + default: + return false; + } +} + /* Instantiate decls as necessary in CTX to satisfy the data sharing specified by CLAUSES. */ @@ -1552,8 +1564,13 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_FIRSTPRIVATE: case OMP_CLAUSE_REDUCTION: + if (gimple_code (ctx->stmt) == GIMPLE_OACC_PARALLEL) + { + sorry ("clause not supported yet"); + break; + } case OMP_CLAUSE_LINEAR: - gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); + gcc_assert (!gimple_code_is_oacc (ctx->stmt)); decl = OMP_CLAUSE_DECL (c); do_private: if (is_variable_sized (decl)) @@ -1591,7 +1608,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_COPYPRIVATE: case OMP_CLAUSE_COPYIN: - gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); + gcc_assert (!gimple_code_is_oacc (ctx->stmt)); decl = OMP_CLAUSE_DECL (c); by_ref = use_pointer_for_field (decl, NULL); install_var_field (decl, by_ref, 3, ctx); @@ -1602,8 +1619,13 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) ctx->default_kind = OMP_CLAUSE_DEFAULT_KIND (c); break; - case OMP_CLAUSE_FINAL: case OMP_CLAUSE_IF: + if (gimple_code (ctx->stmt) == GIMPLE_OACC_PARALLEL) + { + sorry ("clause not supported yet"); + break; + } + case OMP_CLAUSE_FINAL: case OMP_CLAUSE_NUM_THREADS: case OMP_CLAUSE_NUM_TEAMS: case OMP_CLAUSE_THREAD_LIMIT: @@ -1611,14 +1633,14 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_SCHEDULE: case OMP_CLAUSE_DIST_SCHEDULE: case OMP_CLAUSE_DEPEND: - gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); + gcc_assert (!gimple_code_is_oacc (ctx->stmt)); if (ctx->outer) scan_omp_op (&OMP_CLAUSE_OPERAND (c, 0), ctx->outer); break; case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: - gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); + gcc_assert (!gimple_code_is_oacc (ctx->stmt)); case OMP_CLAUSE_MAP: if (ctx->outer) scan_omp_op (&OMP_CLAUSE_SIZE (c), ctx->outer); @@ -1641,7 +1663,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) /* Ignore OMP_CLAUSE_MAP_POINTER kind for arrays in #pragma omp target data, there is nothing to map for those. */ - if (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL + if (!gimple_code_is_oacc (ctx->stmt) && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA && !POINTER_TYPE_P (TREE_TYPE (decl))) break; @@ -1710,17 +1732,34 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_MERGEABLE: case OMP_CLAUSE_PROC_BIND: case OMP_CLAUSE_SAFELEN: - gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); + gcc_assert (!gimple_code_is_oacc (ctx->stmt)); break; case OMP_CLAUSE_ALIGNED: - gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); + gcc_assert (!gimple_code_is_oacc (ctx->stmt)); decl = OMP_CLAUSE_DECL (c); if (is_global_var (decl) && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) install_var_local (decl, ctx); break; + case OMP_CLAUSE_HOST: + case OMP_CLAUSE_OACC_DEVICE: + case OMP_CLAUSE_DEVICE_RESIDENT: + case OMP_CLAUSE_USE_DEVICE: + case OMP_CLAUSE_GANG: + case OMP_CLAUSE_WAIT: + case OMP_NO_CLAUSE_CACHE: + case OMP_CLAUSE_INDEPENDENT: + case OMP_CLAUSE_ASYNC: + case OMP_CLAUSE_WORKER: + case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_NUM_GANGS: + case OMP_CLAUSE_NUM_WORKERS: + case OMP_CLAUSE_VECTOR_LENGTH: + sorry ("Clause not supported yet"); + break; + default: gcc_unreachable (); } @@ -1731,7 +1770,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) switch (OMP_CLAUSE_CODE (c)) { case OMP_CLAUSE_LASTPRIVATE: - gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); + gcc_assert (!gimple_code_is_oacc (ctx->stmt)); /* Let the corresponding firstprivate clause create the variable. */ if (OMP_CLAUSE_LASTPRIVATE_GIMPLE_SEQ (c)) @@ -1743,8 +1782,13 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_PRIVATE: case OMP_CLAUSE_FIRSTPRIVATE: case OMP_CLAUSE_REDUCTION: + if (gimple_code (ctx->stmt) == GIMPLE_OACC_PARALLEL) + { + sorry ("clause not supported yet"); + break; + } case OMP_CLAUSE_LINEAR: - gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); + gcc_assert (!gimple_code_is_oacc (ctx->stmt)); decl = OMP_CLAUSE_DECL (c); if (is_variable_sized (decl)) install_var_local (decl, ctx); @@ -1757,7 +1801,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) break; case OMP_CLAUSE_SHARED: - gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); + gcc_assert (!gimple_code_is_oacc (ctx->stmt)); /* Ignore shared directives in teams construct. */ if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS) break; @@ -1767,7 +1811,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) break; case OMP_CLAUSE_MAP: - if (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL + if (!gimple_code_is_oacc (ctx->stmt) && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA) break; decl = OMP_CLAUSE_DECL (c); @@ -1776,7 +1820,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) && lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl))) { - gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); + gcc_assert (!gimple_code_is_oacc (ctx->stmt)); break; } if (DECL_P (decl)) @@ -1804,10 +1848,15 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) } break; + case OMP_CLAUSE_IF: + if (gimple_code (ctx->stmt) == GIMPLE_OACC_PARALLEL) + { + sorry ("clause not supported yet"); + break; + } case OMP_CLAUSE_COPYPRIVATE: case OMP_CLAUSE_COPYIN: case OMP_CLAUSE_DEFAULT: - case OMP_CLAUSE_IF: case OMP_CLAUSE_NUM_THREADS: case OMP_CLAUSE_NUM_TEAMS: case OMP_CLAUSE_THREAD_LIMIT: @@ -1827,7 +1876,24 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE__LOOPTEMP_: case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: - gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); + gcc_assert (!gimple_code_is_oacc (ctx->stmt)); + break; + + case OMP_CLAUSE_HOST: + case OMP_CLAUSE_OACC_DEVICE: + case OMP_CLAUSE_DEVICE_RESIDENT: + case OMP_CLAUSE_USE_DEVICE: + case OMP_CLAUSE_GANG: + case OMP_CLAUSE_WAIT: + case OMP_NO_CLAUSE_CACHE: + case OMP_CLAUSE_INDEPENDENT: + case OMP_CLAUSE_ASYNC: + case OMP_CLAUSE_WORKER: + case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_NUM_GANGS: + case OMP_CLAUSE_NUM_WORKERS: + case OMP_CLAUSE_VECTOR_LENGTH: + sorry ("Clause not supported yet"); break; default: diff --git a/gcc/tree-core.h b/gcc/tree-core.h index 1b974f8..d7b4ef4 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -213,19 +213,19 @@ enum omp_clause_code { (c_parser_omp_variable_list). */ OMP_CLAUSE_ERROR = 0, - /* OpenMP clause: private (variable_list). */ + /* OpenMP/OpenACC clause: private (variable_list). */ OMP_CLAUSE_PRIVATE, /* OpenMP clause: shared (variable_list). */ OMP_CLAUSE_SHARED, - /* OpenMP clause: firstprivate (variable_list). */ + /* OpenMP/OpenACC clause: firstprivate (variable_list). */ OMP_CLAUSE_FIRSTPRIVATE, /* OpenMP clause: lastprivate (variable_list). */ OMP_CLAUSE_LASTPRIVATE, - /* OpenMP clause: reduction (operator:variable_list). + /* OpenMP/OpenACC clause: reduction (operator:variable_list). OMP_CLAUSE_REDUCTION_CODE: The tree_code of the operator. Operand 1: OMP_CLAUSE_REDUCTION_INIT: Stmt-list to initialize the var. Operand 2: OMP_CLAUSE_REDUCTION_MERGE: Stmt-list to merge private var @@ -265,10 +265,37 @@ enum omp_clause_code { OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */ OMP_CLAUSE_MAP, + /* OpenACC clause: host (variable_list). */ + OMP_CLAUSE_HOST, + + /* OpenACC clause: device (variable_list). */ + OMP_CLAUSE_OACC_DEVICE, + + /* OpenACC clause: device_resident (variable_list). */ + OMP_CLAUSE_DEVICE_RESIDENT, + + /* OpenACC clause: use_device (variable_list). */ + OMP_CLAUSE_USE_DEVICE, + + /* OpenACC clause: gang [(gang-argument-list)]. + Where + gang-argument-list: [gang-argument-list, ] gang-argument + gang-argument: [num:] integer-expression + | static: size-expression + size-expression: * | integer-expression. */ + OMP_CLAUSE_GANG, + + /* OpenACC clause/directive: wait [(integer-expression-list)]. */ + OMP_CLAUSE_WAIT, + + /* Internal structure to hold OpenACC cache directive's variable-list. + #pragma acc cache (variable-list). */ + OMP_NO_CLAUSE_CACHE, + /* Internal clause: temporary for combined loops expansion. */ OMP_CLAUSE__LOOPTEMP_, - /* OpenMP clause: if (scalar-expression). */ + /* OpenMP/OpenACC clause: if (scalar-expression). */ OMP_CLAUSE_IF, /* OpenMP clause: num_threads (integer-expression). */ @@ -281,12 +308,13 @@ enum omp_clause_code { OMP_CLAUSE_NOWAIT, /* OpenMP clause: ordered. */ + /* OpenACC clause: seq. */ OMP_CLAUSE_ORDERED, /* OpenMP clause: default. */ OMP_CLAUSE_DEFAULT, - /* OpenMP clause: collapse (constant-integer-expression). */ + /* OpenMP/OpenACC clause: collapse (constant-integer-expression). */ OMP_CLAUSE_COLLAPSE, /* OpenMP clause: untied. */ @@ -338,7 +366,28 @@ enum omp_clause_code { OMP_CLAUSE_TASKGROUP, /* Internally used only clause, holding SIMD uid. */ - OMP_CLAUSE__SIMDUID_ + OMP_CLAUSE__SIMDUID_, + + /* OpenACC clause: independent. */ + OMP_CLAUSE_INDEPENDENT, + + /* OpenACC clause: async [(integer-expression)]. */ + OMP_CLAUSE_ASYNC, + + /* OpenACC clause: worker [( [num:] integer-expression)]. */ + OMP_CLAUSE_WORKER, + + /* OpenACC clause: vector [( [length:] integer-expression)]. */ + OMP_CLAUSE_VECTOR, + + /* OpenACC clause: num_gangs (integer-expression). */ + OMP_CLAUSE_NUM_GANGS, + + /* OpenACC clause: num_workers (integer-expression). */ + OMP_CLAUSE_NUM_WORKERS, + + /* OpenACC clause: vector_length (integer-expression). */ + OMP_CLAUSE_VECTOR_LENGTH }; #undef DEFTREESTRUCT diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index 5a87728..5c9e249 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -326,6 +326,21 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags) case OMP_CLAUSE__LOOPTEMP_: name = "_looptemp_"; goto print_remap; + case OMP_CLAUSE_HOST: + name = "host"; + goto print_remap; + case OMP_CLAUSE_OACC_DEVICE: + name = "device"; + goto print_remap; + case OMP_CLAUSE_DEVICE_RESIDENT: + name = "device_resident"; + goto print_remap; + case OMP_CLAUSE_USE_DEVICE: + name = "use_device"; + goto print_remap; + case OMP_NO_CLAUSE_CACHE: + name = "_cache_"; + goto print_remap; print_remap: pp_string (buffer, name); pp_left_paren (buffer); @@ -634,6 +649,62 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags) pp_right_paren (buffer); break; + case OMP_CLAUSE_GANG: + pp_string (buffer, "gang("); + dump_generic_node (buffer, OMP_CLAUSE_DECL (clause), spc, flags, false); + pp_character(buffer, ')'); + break; + + case OMP_CLAUSE_WAIT: + pp_string (buffer, "wait("); + dump_generic_node (buffer, OMP_CLAUSE_DECL (clause), spc, flags, false); + pp_character(buffer, ')'); + break; + + case OMP_CLAUSE_ASYNC: + pp_string (buffer, "async"); + if (OMP_CLAUSE_DECL (clause)) + { + pp_character(buffer, '('); + dump_generic_node (buffer, OMP_CLAUSE_DECL (clause), + spc, flags, false); + pp_character(buffer, ')'); + } + break; + + case OMP_CLAUSE_WORKER: + pp_string (buffer, "worker("); + dump_generic_node (buffer, OMP_CLAUSE_DECL (clause), spc, flags, false); + pp_character(buffer, ')'); + break; + + case OMP_CLAUSE_VECTOR: + pp_string (buffer, "vector("); + dump_generic_node (buffer, OMP_CLAUSE_DECL (clause), spc, flags, false); + pp_character(buffer, ')'); + break; + + case OMP_CLAUSE_NUM_GANGS: + pp_string (buffer, "num_gangs("); + dump_generic_node (buffer, OMP_CLAUSE_NUM_GANGS_EXPR (clause), + spc, flags, false); + pp_character (buffer, ')'); + break; + + case OMP_CLAUSE_NUM_WORKERS: + pp_string (buffer, "num_workers("); + dump_generic_node (buffer, OMP_CLAUSE_NUM_WORKERS_EXPR (clause), + spc, flags, false); + pp_character (buffer, ')'); + break; + + case OMP_CLAUSE_VECTOR_LENGTH: + pp_string (buffer, "vector_length("); + dump_generic_node (buffer, OMP_CLAUSE_VECTOR_LENGTH_EXPR (clause), + spc, flags, false); + pp_character (buffer, ')'); + break; + case OMP_CLAUSE_INBRANCH: pp_string (buffer, "inbranch"); break; @@ -652,6 +723,9 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags) case OMP_CLAUSE_TASKGROUP: pp_string (buffer, "taskgroup"); break; + case OMP_CLAUSE_INDEPENDENT: + pp_string (buffer, "independent"); + break; default: /* Should never happen. */ @@ -2384,6 +2458,51 @@ dump_generic_node (pretty_printer *buffer, tree node, int spc, int flags, dump_omp_clauses (buffer, OACC_PARALLEL_CLAUSES (node), spc, flags); goto dump_omp_body; + case OACC_KERNELS: + pp_string (buffer, "#pragma acc kernels"); + dump_omp_clauses (buffer, OACC_KERNELS_CLAUSES (node), spc, flags); + goto dump_omp_body; + + case OACC_DATA: + pp_string (buffer, "#pragma acc data"); + dump_omp_clauses (buffer, OACC_DATA_CLAUSES (node), spc, flags); + goto dump_omp_body; + + case OACC_HOST_DATA: + pp_string (buffer, "#pragma acc host_data"); + dump_omp_clauses (buffer, OACC_HOST_DATA_CLAUSES (node), spc, flags); + goto dump_omp_body; + + case OACC_DECLARE: + pp_string (buffer, "#pragma acc declare"); + dump_omp_clauses (buffer, OACC_DECLARE_CLAUSES (node), spc, flags); + break; + + case OACC_UPDATE: + pp_string (buffer, "#pragma acc update"); + dump_omp_clauses (buffer, OACC_UPDATE_CLAUSES (node), spc, flags); + break; + + case OACC_ENTER_DATA: + pp_string (buffer, "#pragma acc enter data"); + dump_omp_clauses (buffer, OACC_ENTER_DATA_CLAUSES (node), spc, flags); + break; + + case OACC_EXIT_DATA: + pp_string (buffer, "#pragma acc exit data"); + dump_omp_clauses (buffer, OACC_EXIT_DATA_CLAUSES (node), spc, flags); + break; + + case OACC_WAIT: + pp_string (buffer, "#pragma acc wait"); + dump_omp_clauses (buffer, OACC_WAIT_CLAUSES (node), spc, flags); + break; + + case OACC_CACHE: + pp_string (buffer, "#pragma acc cache"); + dump_omp_clauses (buffer, OACC_CACHE_CLAUSES(node), spc, flags); + break; + case OMP_PARALLEL: pp_string (buffer, "#pragma omp parallel"); dump_omp_clauses (buffer, OMP_PARALLEL_CLAUSES (node), spc, flags); diff --git a/gcc/tree.c b/gcc/tree.c index 76e3efb..95df2a3 100644 --- a/gcc/tree.c +++ b/gcc/tree.c @@ -259,6 +259,13 @@ unsigned const char omp_clause_num_ops[] = 2, /* OMP_CLAUSE_FROM */ 2, /* OMP_CLAUSE_TO */ 2, /* OMP_CLAUSE_MAP */ + 1, /* OMP_CLAUSE_HOST */ + 1, /* OMP_CLAUSE_OACC_DEVICE */ + 1, /* OMP_CLAUSE_DEVICE_RESIDENT */ + 1, /* OMP_CLAUSE_USE_DEVICE */ + 1, /* OMP_CLAUSE_GANG */ + 1, /* OMP_CLAUSE_WAIT */ + 1, /* OMP_NO_CLAUSE_CACHE */ 1, /* OMP_CLAUSE__LOOPTEMP_ */ 1, /* OMP_CLAUSE_IF */ 1, /* OMP_CLAUSE_NUM_THREADS */ @@ -284,6 +291,13 @@ unsigned const char omp_clause_num_ops[] = 0, /* OMP_CLAUSE_SECTIONS */ 0, /* OMP_CLAUSE_TASKGROUP */ 1, /* OMP_CLAUSE__SIMDUID_ */ + 0, /* OMP_CLAUSE_INDEPENDENT */ + 1, /* OMP_CLAUSE_ASYNC */ + 1, /* OMP_CLAUSE_WORKER */ + 1, /* OMP_CLAUSE_VECTOR */ + 1, /* OMP_CLAUSE_NUM_GANGS */ + 1, /* OMP_CLAUSE_NUM_WORKERS */ + 1, /* OMP_CLAUSE_VECTOR_LENGTH */ }; const char * const omp_clause_code_name[] = @@ -303,6 +317,13 @@ const char * const omp_clause_code_name[] = "from", "to", "map", + "host", + "device", + "device_resident", + "use_device", + "gang", + "wait", + "_cache_", "_looptemp_", "if", "num_threads", @@ -327,7 +348,14 @@ const char * const omp_clause_code_name[] = "parallel", "sections", "taskgroup", - "_simduid_" + "_simduid_", + "independent", + "async", + "worker", + "vector", + "num_gangs", + "num_workers", + "vector_length" }; @@ -11034,6 +11062,19 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data, case OMP_CLAUSE: switch (OMP_CLAUSE_CODE (*tp)) { + case OMP_CLAUSE_HOST: + case OMP_CLAUSE_OACC_DEVICE: + case OMP_CLAUSE_DEVICE_RESIDENT: + case OMP_CLAUSE_USE_DEVICE: + case OMP_CLAUSE_GANG: + case OMP_CLAUSE_WAIT: + case OMP_NO_CLAUSE_CACHE: + case OMP_CLAUSE_ASYNC: + case OMP_CLAUSE_WORKER: + case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_NUM_GANGS: + case OMP_CLAUSE_NUM_WORKERS: + case OMP_CLAUSE_VECTOR_LENGTH: case OMP_CLAUSE_PRIVATE: case OMP_CLAUSE_SHARED: case OMP_CLAUSE_FIRSTPRIVATE: @@ -11056,6 +11097,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data, WALK_SUBTREE (OMP_CLAUSE_OPERAND (*tp, 0)); /* FALLTHRU */ + case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_NOWAIT: case OMP_CLAUSE_ORDERED: case OMP_CLAUSE_DEFAULT: diff --git a/gcc/tree.def b/gcc/tree.def index b921b12..623ebb0 100644 --- a/gcc/tree.def +++ b/gcc/tree.def @@ -1017,6 +1017,24 @@ DEFTREECODE (MEM_REF, "mem_ref", tcc_reference, 2) DEFTREECODE (OACC_PARALLEL, "oacc_parallel", tcc_statement, 2) +/* OpenACC - #pragma acc kernels [clause1 ... clauseN] + Operand 0: OACC_KERNELS_BODY: Sequence of kernels. + Operand 1: OACC_KERNELS_CLAUSES: List of clauses. */ + +DEFTREECODE (OACC_KERNELS, "oacc_kernels", tcc_statement, 2) + +/* OpenACC - #pragma acc data [clause1 ... clauseN] + Operand 0: OACC_DATA_BODY: Data construct body. + Operand 1: OACC_DATA_CLAUSES: List of clauses. */ + +DEFTREECODE (OACC_DATA, "oacc_data", tcc_statement, 2) + +/* OpenACC - #pragma acc host_data [clause1 ... clauseN] + Operand 0: OACC_HOST_DATA_BODY: Host_data construct body. + Operand 1: OACC_HOST_DATA_CLAUSES: List of clauses. */ + +DEFTREECODE (OACC_HOST_DATA, "oacc_host_data", tcc_statement, 2) + /* OpenMP - #pragma omp parallel [clause1 ... clauseN] Operand 0: OMP_PARALLEL_BODY: Code to be executed by all threads. Operand 1: OMP_PARALLEL_CLAUSES: List of clauses. */ @@ -1108,6 +1126,30 @@ DEFTREECODE (OMP_ORDERED, "omp_ordered", tcc_statement, 1) Operand 1: OMP_CRITICAL_NAME: Identifier for critical section. */ DEFTREECODE (OMP_CRITICAL, "omp_critical", tcc_statement, 2) +/* OpenACC - #pragma acc declare [clause1 ... clauseN] + Operand 0: OACC_DECLARE_CLAUSES: List of clauses. */ +DEFTREECODE (OACC_DECLARE, "oacc_declare", tcc_statement, 1) + +/* OpenACC - #pragma acc update [clause1 ... clauseN] + Operand 0: OACC_UPDATE_CLAUSES: List of clauses. */ +DEFTREECODE (OACC_UPDATE, "oacc_update", tcc_statement, 1) + +/* OpenACC - #pragma acc enter data [clause1 ... clauseN] + Operand 0: OACC_ENTER_DATA_CLAUSES: List of clauses. */ +DEFTREECODE (OACC_ENTER_DATA, "oacc_enter_data", tcc_statement, 1) + +/* OpenACC - #pragma acc exit data [clause1 ... clauseN] + Operand 0: OACC_EXIT_DATA_CLAUSES: List of clauses. */ +DEFTREECODE (OACC_EXIT_DATA, "oacc_exit_data", tcc_statement, 1) + +/* OpenACC - #pragma acc wait [clause1 ... clauseN] + Operand 0: OACC_WAIT_CLAUSES: List of clauses. */ +DEFTREECODE (OACC_WAIT, "oacc_wait", tcc_statement, 1) + +/* OpenACC - #pragma acc cache [clause1 ... clauseN] + Operand 0: OACC_CACHE_CLAUSES: List of clauses. */ +DEFTREECODE (OACC_CACHE, "oacc_cache", tcc_statement, 1) + /* OpenMP - #pragma omp target update [clause1 ... clauseN] Operand 0: OMP_TARGET_UPDATE_CLAUSES: List of clauses. */ DEFTREECODE (OMP_TARGET_UPDATE, "omp_target_update", tcc_statement, 1) diff --git a/gcc/tree.h b/gcc/tree.h index 202ad9e..dd90cde 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1155,7 +1155,7 @@ extern void protected_set_expr_location (tree, location_t); #define TRANSACTION_EXPR_RELAXED(NODE) \ (TRANSACTION_EXPR_CHECK (NODE)->base.public_flag) -/* OpenMP directive and clause accessors. */ +/* OpenMP and OpenACC directive and clause accessors. */ #define OMP_BODY(NODE) \ TREE_OPERAND (TREE_RANGE_CHECK (NODE, OACC_PARALLEL, OMP_CRITICAL), 0) @@ -1167,6 +1167,39 @@ extern void protected_set_expr_location (tree, location_t); #define OACC_PARALLEL_CLAUSES(NODE) \ TREE_OPERAND (OACC_PARALLEL_CHECK (NODE), 1) +#define OACC_KERNELS_BODY(NODE) \ + TREE_OPERAND (OACC_KERNELS_CHECK(NODE), 0) +#define OACC_KERNELS_CLAUSES(NODE) \ + TREE_OPERAND (OACC_KERNELS_CHECK(NODE), 1) + +#define OACC_DATA_BODY(NODE) \ + TREE_OPERAND (OACC_DATA_CHECK (NODE), 0) +#define OACC_DATA_CLAUSES(NODE) \ + TREE_OPERAND (OACC_DATA_CHECK (NODE), 1) + +#define OACC_HOST_DATA_BODY(NODE) \ + TREE_OPERAND (OACC_HOST_DATA_CHECK (NODE), 0) +#define OACC_HOST_DATA_CLAUSES(NODE) \ + TREE_OPERAND (OACC_HOST_DATA_CHECK (NODE), 1) + +#define OACC_DECLARE_CLAUSES(NODE) \ + TREE_OPERAND (OACC_DECLARE_CHECK (NODE), 0) + +#define OACC_ENTER_DATA_CLAUSES(NODE) \ + TREE_OPERAND (OACC_ENTER_DATA_CHECK (NODE), 0) + +#define OACC_EXIT_DATA_CLAUSES(NODE) \ + TREE_OPERAND (OACC_EXIT_DATA_CHECK (NODE), 0) + +#define OACC_UPDATE_CLAUSES(NODE) \ + TREE_OPERAND (OACC_UPDATE_CHECK (NODE), 0) + +#define OACC_WAIT_CLAUSES(NODE) \ + TREE_OPERAND (OACC_WAIT_CHECK (NODE), 0) + +#define OACC_CACHE_CLAUSES(NODE) \ + TREE_OPERAND (OACC_CACHE_CHECK (NODE), 0) + #define OMP_PARALLEL_BODY(NODE) TREE_OPERAND (OMP_PARALLEL_CHECK (NODE), 0) #define OMP_PARALLEL_CLAUSES(NODE) TREE_OPERAND (OMP_PARALLEL_CHECK (NODE), 1) @@ -1278,6 +1311,32 @@ extern void protected_set_expr_location (tree, location_t); #define OMP_CLAUSE_SCHEDULE_CHUNK_EXPR(NODE) \ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_SCHEDULE), 0) +/* OpenACC clause expressions */ +#define OMP_CLAUSE_GANG_EXPR(NODE) \ + OMP_CLAUSE_OPERAND ( \ + OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_GANG), 0) +#define OMP_WAIT_EXPR(NODE) \ + OMP_CLAUSE_OPERAND ( \ + OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_WAIT), 0) +#define OMP_CLAUSE_VECTOR_EXPR(NODE) \ + OMP_CLAUSE_OPERAND ( \ + OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_VECTOR_LENGTH), 0) +#define OMP_CLAUSE_ASYNC_EXPR(NODE) \ + OMP_CLAUSE_OPERAND ( \ + OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_ASYNC), 0) +#define OMP_CLAUSE_WORKER_EXPR(NODE) \ + OMP_CLAUSE_OPERAND ( \ + OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_WORKER), 0) +#define OMP_CLAUSE_NUM_GANGS_EXPR(NODE) \ + OMP_CLAUSE_OPERAND ( \ + OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_NUM_GANGS), 0) +#define OMP_CLAUSE_NUM_WORKERS_EXPR(NODE) \ + OMP_CLAUSE_OPERAND ( \ + OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_NUM_WORKERS), 0) +#define OMP_CLAUSE_VECTOR_LENGTH_EXPR(NODE) \ + OMP_CLAUSE_OPERAND ( \ + OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_VECTOR_LENGTH), 0) + #define OMP_CLAUSE_DEPEND_KIND(NODE) \ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_DEPEND)->omp_clause.subcode.depend_kind) -- 1.8.3.2 --------------010805090902090604010307 Content-Type: text/plain; charset=UTF-8; name="ChangeLog" Content-Transfer-Encoding: base64 Content-Disposition: attachment; filename="ChangeLog" Content-length: 2595 MTMtMDItMjAxNCAgSWxtaXIgVXNtYW5vdiAgPGkudXNtYW5vdkBzYW1zdW5n LmNvbT4KCglBZGQgT3BlbkFDQyAxLjAgc3VwcG9ydCB0byBHRU5FUklDLCBl eGNlcHQgbG9vcCBkaXJlY3RpdmUgYW5kIHN1YmFycmF5cy4KCglEbWl0cnkg Qm9jaGFybmlrb3YgPGRtaXRyeS5iQHNhbXN1bmcuY29tPgoJRXZnZW55IEdh dnJpbiA8ZS5nYXZyaW5Ac2Ftc3VuZy5jb20+CglJbG1pciBVc21hbm92IDxp LnVzbWFub3ZAc2Ftc3VuZy5jb20+CgoJZ2NjLwoJKiBnaW1wbGlmeS5jIChp c19naW1wbGVfc3RtdCk6IFN0dWIgT3BlbkFDQyBkaXJlY3RpdmVzIGFuZCBj bGF1c2VzLgoJKGdpbXBsaWZ5X3NjYW5fb21wX2NsYXVzZXMsIGdpbXBsaWZ5 X2FkanVzdF9vbXBfY2xhdXNlcyk6IExpa2V3aXNlLgoJKGdpbXBsaWZ5X2V4 cHIpOiBMaWtld2lzZS4KCSogb21wLWxvdy5jIChzY2FuX3NoYXJpbmdfY2xh dXNlcyk6IExpa2V3aXNlLgoJKGdpbXBsZV9jb2RlX2lzX29hY2MpOiBOZXcg aGVscGVyIGZ1bmN0aW9uLgoJKHNjYW5fc2hhcmluZ19jbGF1c2VzKTogVXNl IGl0LgoJKiB0cmVlLWNvcmUuaCAKCShPTVBfQ0xBVVNFX0hPU1QsIE9NUF9D TEFVU0VfT0FDQ19ERVZJQ0UsIE9NUF9DTEFVU0VfREVWSUNFX1JFU0lERU5U LAoJT01QX0NMQVVTRV9VU0VfREVWSUNFLCBPTVBfQ0xBVVNFX0dBTkcsIE9N UF9DTEFVU0VfV0FJVCwKCU9NUF9OT19DTEFVU0VfQ0FDSEUsIE9NUF9DTEFV U0VfSU5ERVBFTkRFTlQsIE9NUF9DTEFVU0VfQVNZTkMsCglPTVBfQ0xBVVNF X1dPUktFUiwgT01QX0NMQVVTRV9WRUNUT1IsIE9NUF9DTEFVU0VfTlVNX0dB TkdTLAoJT01QX0NMQVVTRV9OVU1fV09SS0VSUywgT01QX0NMQVVTRV9WRUNU T1JfTEVOR1RIKTogTmV3IGNsYXVzZXMuCgkqIHRyZWUtcHJldHR5LXByaW50 LmMgKGR1bXBfb21wX2NsYXVzZSk6IFByaW50IE9wZW5BQ0MgY2xhdXNlLgoJ KGR1bXBfZ2VuZXJpY19ub2RlKTogUHJpbnQgT3BlbkFDQyBkaXJlY3RpdmVz IGFuZCBpdHMgY2xhdXNlcy4KCSogdHJlZS5jIChvbXBfY2xhdXNlX251bV9v cHMpOiBBZGQgT3BlbkFDQyBjbGF1c2VzLgoJKG9tcF9jbGF1c2VfY29kZV9u YW1lKTogTGlrZXdpc2UuCgkod2Fsa190cmVlXzEpOiBMaWtld2lzZS4KCSog dHJlZS5kZWYgKE9BQ0NfS0VSTkVMUyk6IE5ldyB0cmVlIG5vZGUuCgkoT0FD Q19EQVRBLCBPQUNDX0hPU1RfREFUQSwgT0FDQ19ERUNMQVJFLCBPQUNDX1VQ REFURSk6IExpa2V3aXNlLgoJKE9BQ0NfRU5URVJfREFUQSwgT0FDQ19FWElU X0RBVEEsIE9BQ0NfV0FJVCwgT0FDQ19DQUNIRSk6IExpa2V3aXNlLgoJKiB0 cmVlLmggKE9BQ0NfS0VSTkVMU19CT0RZKTogTmV3IG1hY3Jvcy4KCShPQUND X0tFUk5FTFNfQ0xBVVNFUywgT0FDQ19DTEFVU0VfTlVNX0dBTkdTX0VYUFIs IAoJT01QX0NMQVVTRV9OVU1fV09SS0VSU19FWFBSLCBPTVBfQ0xBVVNFX1ZF Q1RPUl9MRU5HVEhfRVhQUiwgCglPTVBfQ0xBVVNFX1ZFQ1RPUl9FWFBSLCBP TVBfQ0xBVVNFX1dPUktFUl9FWFBSLCBPTVBfQ0xBVVNFX0dBTkdfRVhQUiwg CglPTVBfQ0xBVVNFX0FTWU5DX0VYUFIsIE9NUF9XQUlUX0VYUFIsIE9BQ0Nf REFUQV9CT0RZLCBPQUNDX0RBVEFfQ0xBVVNFUywKCU9BQ0NfREVDTEFSRV9D TEFVU0VTLCBPQUNDX1VQREFURV9DTEFVU0VTLCBPQUNDX1dBSVRfQ0xBVVNF UywgCglPQUNDX0NBQ0hFX0NMQVVTRVMsIE9BQ0NfSE9TVF9EQVRBX0JPRFks IAoJT0FDQ19IT1NUX0RBVEFfQ0xBVVNFUyk6IExpa2V3aXNlLgoJZ2NjL2Rv Yy8KCSogZ2VuZXJpYy50ZXhpOiBEb2N1bWVudCBPQUNDX0tFUk5FTFMsIE9B Q0NfREFUQSwgT0FDQ19IT1NUX0RBVEEsIAoJT0FDQ19ERUNMQVJFLCBPQUND X1VQREFURSwgT0FDQ19FTlRFUl9EQVRBLCBPQUNDX0VYSVRfREFUQSwgT0FD Q19XQUlULCAKCU9BQ0NfQ0FDSEUuCg== --------------010805090902090604010307--