I see valgrind memory leak starting from r230275: $ valgrind --leak-check=yes --trace-children=yes g++ /home/marxin/Programming/gcc/gcc/testsuite/c-c++-common/goacc/declare-1.c -fno-diagnostics-show-caret -fdiagnostics-color=never -nostdinc++ -I/home/marxin/Programming/gcc/objdir/x86_64-pc-linux-gnu/libstdc++-v3/include/x86_64-pc-linux-gnu -I/home/marxin/Programming/gcc/objdir/x86_64-pc-linux-gnu/libstdc++-v3/include -I/home/marxin/Programming/gcc/libstdc++-v3/libsupc++ -I/home/marxin/Programming/gcc/libstdc++-v3/include/backward -I/home/marxin/Programming/gcc/libstdc++-v3/testsuite/util -fmessage-length=0 -std=c++98 -fopenacc -S -o declare-1.s ... ==21322== 4,768 (56 direct, 4,712 indirect) bytes in 1 blocks are definitely lost in loss record 672 of 706 ==21322== at 0x4C2D0C5: calloc (in /usr/lib64/valgrind/vgpreload_memcheck-amd64-linux.so) ==21322== by 0x1C96570: xcalloc (xmalloc.c:162) ==21322== by 0xE24A72: new_omp_context(omp_region_type) (gimplify.c:400) ==21322== by 0xE3C4A0: gimplify_scan_omp_clauses(tree_node**, gimple**, omp_region_type, tree_code) (gimplify.c:7334) ==21322== by 0xE47532: gimplify_oacc_declare(tree_node**, gimple**) (gimplify.c:9238) ==21322== by 0xE52AC7: gimplify_expr(tree_node**, gimple**, gimple**, bool (*)(tree_node*), int) (gimplify.c:11683) ==21322== by 0xE39C71: gimplify_stmt(tree_node**, gimple**) (gimplify.c:6478) ==21322== by 0xE28AE5: gimplify_statement_list(tree_node**, gimple**) (gimplify.c:1716) ==21322== by 0xE528F3: gimplify_expr(tree_node**, gimple**, gimple**, bool (*)(tree_node*), int) (gimplify.c:11626) ==21322== by 0xE39C71: gimplify_stmt(tree_node**, gimple**) (gimplify.c:6478) ==21322== by 0xE2724F: gimplify_bind_expr(tree_node**, gimple**) (gimplify.c:1290) ==21322== by 0xE5196B: gimplify_expr(tree_node**, gimple**, gimple**, bool (*)(tree_node*), int) (gimplify.c:11398) ...
I think following change could cure it: --- gcc/gimplify.c.jj 2017-02-16 08:36:56.000000000 +0100 +++ gcc/gimplify.c 2017-03-14 17:30:02.040865055 +0100 @@ -9236,6 +9236,7 @@ gimplify_oacc_declare (tree *expr_p, gim clauses = OACC_DECLARE_CLAUSES (expr); gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA, OACC_DECLARE); + gimplify_adjust_omp_clauses (pre_p, NULL, &clauses, OACC_DECLARE); for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t)) { though can't test that right now.
That doesn't work, it is too early. With: --- gimplify.c.jj 2017-03-08 18:19:24.000000000 +0100 +++ gimplify.c 2017-03-20 07:22:05.587913531 +0100 @@ -9261,6 +9261,8 @@ gimplify_oacc_declare (tree *expr_p, gim omp_add_variable (gimplify_omp_ctxp, decl, GOVD_SEEN); } + gimplify_adjust_omp_clauses (pre_p, NULL, &clauses, OACC_DECLARE); + stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE, clauses); we don't ICE on it, but still no idea what is the right thing to do. From what I understand, #pragma acc declare is active for all code from the directive till end of function (unless the variables mentioned in its clauses go out of scope earlier), so in theory perhaps the right thing is to move the gimplify omp context structure up in the gimplify context tree so that they are just children of the function (what happens if you put #pragma acc declare inside the body of some other acc region?). Though on the other side, e.g. for #pragma omp target enter data we adjust the omp clauses right away, but that doesn't really affect the handling of variables later on. So, the important question is do you need the variables mentioned in acc declare clauses to be in the hash tables as GOMP_MAP after gimplify_oacc_declare returns or not? If not, then the above patch might DTRT, otherwise we need to figure out something different.
(In reply to Jakub Jelinek from comment #2) > That doesn't work, it is too early. > With: > --- gimplify.c.jj 2017-03-08 18:19:24.000000000 +0100 > +++ gimplify.c 2017-03-20 07:22:05.587913531 +0100 > @@ -9261,6 +9261,8 @@ gimplify_oacc_declare (tree *expr_p, gim > omp_add_variable (gimplify_omp_ctxp, decl, GOVD_SEEN); > } > > + gimplify_adjust_omp_clauses (pre_p, NULL, &clauses, OACC_DECLARE); > + > stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE, > clauses); > > we don't ICE on it, but still no idea what is the right thing to do. > From what I understand, #pragma acc declare is active for all code from the > directive till end of function (unless the variables mentioned in its > clauses go out of scope earlier), so in theory perhaps the right thing is to > move the gimplify omp context structure up in the gimplify context tree so > that they are just children of the function (what happens if you put #pragma > acc declare inside > the body of some other acc region?). You are correctly describing the behavior of declare in OpenACC. Declared variables live throughout the scope specified by the user. However, I believe the scope must be global or function, i.e. I don't think you can declare variables inside nested blocks in C/C++. > Though on the other side, e.g. for #pragma omp target enter data we adjust > the omp clauses right away, but that doesn't really affect the handling of > variables later on. So, the important question is do you need the variables > mentioned in acc declare clauses to be in the hash tables as GOMP_MAP after > gimplify_oacc_declare returns or not? If not, then the above patch might > DTRT, otherwise we need to figure out something different. I don't think this memory leak is present in gomp-4_0-branch. The patch that resolves this issue was posted back in September 2016 here <https://gcc.gnu.org/ml/gcc-patches/2016-09/msg01630.html>. I'll rebuild trunk with that patch to see if it resolves the memory leak there.
Author: cesar Date: Wed Mar 22 13:52:10 2017 New Revision: 246381 URL: https://gcc.gnu.org/viewcvs?rev=246381&root=gcc&view=rev Log: PR c++/80029 gcc/ * gimplify.c (is_oacc_declared): New function. (oacc_default_clause): Use it to set default flags for acc declared variables inside parallel regions. (gimplify_scan_omp_clauses): Strip firstprivate pointers for acc declared variables. (gimplify_oacc_declare): Gimplify the declare clauses. Add the declare attribute to any decl as necessary. libgomp/ * testsuite/libgomp.oacc-c-c++-common/declare-vla.c: New test. Added: trunk/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c Modified: trunk/gcc/ChangeLog trunk/gcc/gimplify.c trunk/libgomp/ChangeLog
Fixed in trunk.