[OpenMP] outlined functions for target regions in header files not always uniquely named

[mikerice1969]

The offloaded region is outlined to functions expecting DeviceID, FileID, ParentName, and LineNum to uniquely identify them. They are registered in CGOpenMPRuntime with:

OffloadEntriesTargetRegion[DeviceID][FileID][ParentName][LineNum] = Entry;

But some users use macros to generate different code across the include file like this:

#include <stdio.h>
int main() {
  int i = 0;
#define VALUE 1
#include "c.inc"

#undef VALUE
#define VALUE 2
#include "c.inc"

  printf("%d\n", i);
  return 0;
}

c.inc

#pragma omp target
{
  i = i + VALUE;
}

Two functions are created on the host, but only one is used in the device compile.

host:
define internal void @__omp_offloading_10309_859455e_foo_l1(i64 %i) #1
define internal void @__omp_offloading_10309_859455e_foo_l1.1(i64 %i) #1
device:
define weak void @__omp_offloading_10309_859455e_foo_l1(i64 %i) #0
!omp_offload.info = !{!0}
!0 = !{i32 0, i32 66313, i32 140068190, !"foo", i32 1, i32 0}

This would result in silently executing the wrong code (incrementing i by 1 twice in the above example).

[jdoerfert]

Would it be sufficient and doable to take the line number in the resulting file, not the original one (here the header)?

[mikerice1969]

If I understand the suggestion, it doesn't seem like that would work. For example if "c.inc" was included from another include "foo.h" which was included from the main source.

void foo1()
{
#define VER 1
#include "foo.h"
}
void foo2()
{
#undef VER
#define VER 2
#include "foo.h"
} 

It seems no matter what line number is used there would be a problem.
Maybe it would work if we encoded the whole include stack but I don't believe that is currently available in the AST.

Plus (I suspect) 99.9% of functions in include files would not be different, so by encoding this info we could be generating a ton of extra functions with the same text and different names.

[jdoerfert]

I don't get it. So, there is one main source file, let's call it main.c.
Somehow main.c includes c.inc twice, either directly or transitively.
In the pre-processed main.c the target regions from c.inc are now at different source lines.
Those lines are unique per construction.

How does the foo code snipped contradict that?

[mikerice1969]

Any include could have two or more targets, so if c.inc has two they would both be on the same source line in main. Sorry I didn't provide a full example.
f.c

 1  void foo1()
 2  {
 3    int i;
 4  #define VER 1
 5  #include "foo.h"
 6  }
 7  void foo2()
 8  {
 9    int i;
10  #undef VER
11  #define VER 2
12  #include "foo.h"
13  }

foo.h

1  #undef VALUE
2  #if VER==1
3  #define VALUE 1
4  #else
5  #define VALUE 2
6  #endif
7  #include "c.inc"
8
9  #undef VALUE
10  #if VER==1
11  #define VALUE 3
12  #else
13  #define VALUE 4
14  #endif
15  #include "c.inc"

c.inc

#pragma omp target
{
  i = i + VALUE;
}

This produces four kernels, none of which have a single single unique line number:
f.c:5 foo.h:7
f.c:5 foo.h.15
f.c:12 foo.h:7
f.c:12 foo.h:15

If you mean the line number in a preprocessed file would be unique, sure, but I don't think we actually have that information. If we do we'd still have the problem of matching exactly preprocessed files across the host and device compile.

[jdoerfert]

I get your point now, and yes, I meant preprocessed file. That said, I agree the matching is a problem. (One could argue target regions should be in both, but I guess that is not a great solution.)

What if we number them per line then? So foo1_5_1, foo1_5_2, ...
Basically what happened in your initial example on the host but only because we actually generated both outlined functions with the same name. Somehow we don't on the device, otherwise they'd match again, no?

[mikerice1969]

My impression is that any sort of numbering system could be broken by differences in macros between the host and device compile.

But if we think we can ignore that possibility, and we think we'll see the same target regions in the same order during host and device codegen, maybe we can get the same name in both.

I haven't looked at this code much so that's mostly speculation.

[jdoerfert]

Even without macros we can create mismatches. TBH, I think there is no "perfect" solution and the standard is not clear what should happen. We need to support the most common cases, so, if they match we need to get it right, but we can allow missing ones on either side.

[mikerice1969]

Proposed fix is here: https://reviews.llvm.org/D134816
@jdoerfert would you mind reviewing?