https://gcc.gnu.org/bugzilla/show_bug.cgi?id=90861
Bug ID: 90861 Summary: OpenACC 'declare' not cleaning up for VLAs Product: gcc Version: 10.0 Status: UNCONFIRMED Keywords: openacc Severity: normal Priority: P3 Component: middle-end Assignee: unassigned at gcc dot gnu.org Reporter: tschwinge at gcc dot gnu.org Target Milestone: --- Given: #define N_f1 1000 int A_f1[N_f1]; #pragma acc declare copy(A_f1) ..., in '-fdump-tree-gimple' we get the expected: try { #pragma omp target oacc_declare map(to:A_f1 [len: 4000]) } finally { #pragma omp target oacc_declare map(from:A_f1) A_f1 = {CLOBBER}; } However, given: int N_f2 = 1000; int A_f2[N_f2]; #pragma acc declare copy(A_f2) ..., the cleanup ('from') is missing, so device memory gets allocated corresponding to local variable 'A_f2', but isn't deallocated and deassociated from the stack address of 'A_f2' when that one goes out of scope. As stack space will be re-used in different functions called consecutively, libgomp will then abort because of 'libgomp: Trying to map into device [0x7ffd86278870..0x7ffd86279810) object when [0x7ffd86278880..0x7ffd86279820) is already mapped', if running two such VLA functions consecutively, for example. This can be easily seen with the 'libgomp.oacc-c-c++-common/declare-vla.c' test case, when duplicating its 'main' function into two functions, which are then called consecutively. 'libgomp.oacc-c-c++-common/declare-vla.c' has been added in r246381 for PR80029; I don't know what the behavior would've been before that. In <http://mid.mail-archive.com/874ly0h7dc.fsf@euler.schwinge.homeip.net> there are unanswered questions, but I'm not sure if they're relevant to this problem here.