$ clang++ -fopenmp --offload-arch=sm_80 main.cpp -foffload-lto -O3 -ffast-math -v
...
nvlink info : Function properties for '__omp_offloading_10307_18c6188_main_l9':
nvlink info : used 28 registers, 8 stack, 40 bytes smem, 360 bytes cmem[0], 0 bytes lmem
Once I comment the second offload region
nvlink info : Function properties for '__omp_offloading_10307_18c606c_main_l9':
nvlink info : used 16 registers, 0 stack, 12 bytes smem, 360 bytes cmem[0], 0 bytes lmem
register count, stack, smen all changed.
The first offload kernel generation should not be affected by other kernels. There are undesired entanglement.
reproducer main.cpp
#include <unistd.h>
#include <cstdio>
#define N 100
int main()
{
int a[N];
#pragma omp target teams distribute parallel for map(a)
for(int i=0; i<N; i++)
{ a[i]=i; }
int b = 0;
#pragma omp target teams distribute parallel for map(a) reduction(+:b)
for(int i=0; i<N; i++)
{ a[i]+=1; b+=a[i]; }
}
Using the reproducer
Once I comment the second offload region
register count, stack, smen all changed. The first offload kernel generation should not be affected by other kernels. There are undesired entanglement.
reproducer
main.cpp