[Bug libgomp/114445] New: [OpenMP] indirect - potential race when creating reverse-offload hash

[Bug libgomp/114445] New: [OpenMP] indirect - potential race when creating reverse-offload hash

[burnus at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=114445

            Bug ID: 114445
           Summary: [OpenMP] indirect - potential race when creating
                    reverse-offload hash
           Product: gcc
           Version: 14.0
            Status: UNCONFIRMED
          Keywords: openmp, wrong-code
          Severity: normal
          Priority: P3
         Component: libgomp
          Assignee: unassigned at gcc dot gnu.org
          Reporter: burnus at gcc dot gnu.org
                CC: jakub at gcc dot gnu.org
  Target Milestone: ---

It looks as if when starting two kernels very quickly after another as in

  #pragma omp target nowait
    ...
  #pragma omp target
    ...

Two device threads might concurrently create the hash in
libgomp/config/accel/target-indirect.c's build_indirect_map

  if (!indirect_htab)
    {
      /* Count the number of entries in the NULL-terminated address map.  */
      for (map_entry = GOMP_INDIRECT_ADDR_MAP; *map_entry;
           map_entry += 2, num_ind_funcs++);

      indirect_htab = htab_create (num_ind_funcs);

The issue only occurs if the assignment of the allocated memory done in the
first thread occurs after the second kernel has checked for indirect_htab ==
NULL, which seems to be rather unlikely to happen in the real world, but IMHO
cannot be ruled out.

Thus, I was wondering (see email) whether a tmp variable + CAS should be used
for indirect_htab; if the swap failed, the memory could just be freed as
another process was faster - and has also completed by then the creation of the
htab.

See also

  r14-9629-g637e76b90e8b045c5e25206a41e3be55deace8d5
  openmp: Change to using a hashtab to lookup offload target addresses for
indirect function calls

and review email at

  https://gcc.gnu.org/pipermail/gcc-patches/2024-March/647755.html