Miscompiling LDS variable in openmp

[JonChesterfield]

Test case from aomp smoke (teams_nest.c) reproduced inline,

#include <stdio.h>
#include <omp.h>


int main(void) {

  int fail = 0;

  //                                                                                                                                                                                                               
  // Test: num_teams and omp_get_team_num()                                                                                                                                                                        
  #pragma omp target
 {
   printf("Num_teams=%d\n", omp_get_num_teams());
 }
  #pragma omp target
 {
    #pragma omp teams
    {
     	if (omp_get_team_num() == 0)
        printf("Num_teams=%d\n", omp_get_num_teams());
      #pragma omp distribute
      for (int i=0; i< 10; i++)
        printf("team %d thread %d\n", omp_get_team_num(), omp_get_thread_num());
    }
  }

  return fail;
}

Compile with

clang -O2 -fopenmp --offload-arch=gfx1010  teams_nest.c -o teams_nest -mllvm -openmp-opt-disable=true

Fails with

LLVM ERROR: Absolute address LDS variable inconsistent with variable alignment

Save temps suggests LDS variables have had addresses allocated before codegen. I guess that's related to the thin-lto work, though I thought that was only in rocm and this is reproducing in trunk.

The LDS lowering makes a very definite "allocating whole code object at a time" assumption which means it's not safe to run repeatedly. The thinlto work was supposed to run it on pieces and then recombine them without running the pass a second time, but it looks like the pass is running a second time and decides to increase the alignment of some already allocated variables, i.e. miscompile them.

This specific error can probably be fixed by not over-aligning any already allocated variables, but it seems the writing is on the wall for the closed world assumption LDS lowering currently relies on.

[JonChesterfield]

Side note, the asan-only prepass AMDGPUSwLowerLDS expects to be able to create and allocate LDS variables itself. That definitely conflicts with the invariants of the lowering pass, the mixed some-variables-allocated-and-some-not is wholly unsupported. It would have been a hard error if the thin-lto work hadn't slacked off some of the diagnostics.

It's likely a good idea to make the LDS lowering robust to arbitrary allocation and interleaving of object construction given there are now two in tree use cases for breaching the whole module allocation model, with "ISA level linking" a risk in the background as a third.

[JonChesterfield]

Related ccb3a8f - code has moved around a bit (pulling private methods out to call from some asan code, which probably breaks invariants in the lowering pass), but looks like we should still be fatal-error when some variables are allocated and others aren't.

Root error of this bug is probably unconditionally calling superAlignLDSGlobals, even when some variables are already allocated, and then ignoring that they are already allocated when setting their alignment. Point fix for that is straightforward.

With test cases tbd, this will unblock the current failure mode and expose other ones

     // in AMDGPULowerModuleLDS::superAlignLDSGlobals
      if (GV.isAbsoluteSymbolRef()) {
        continue;
      }

Point fix at #106421

[llvmbot]

@llvm/issue-subscribers-openmp

Author: Jon Chesterfield (JonChesterfield)

Test case from aomp smoke (teams_nest.c) reproduced inline,

#include <stdio.h>
#include <omp.h>


int main(void) {

  int fail = 0;

  //                                                                                                                                                                                                               
  // Test: num_teams and omp_get_team_num()                                                                                                                                                                        
  #pragma omp target
 {
   printf("Num_teams=%d\n", omp_get_num_teams());
 }
  #pragma omp target
 {
    #pragma omp teams
    {
     	if (omp_get_team_num() == 0)
        printf("Num_teams=%d\n", omp_get_num_teams());
      #pragma omp distribute
      for (int i=0; i< 10; i++)
        printf("team %d thread %d\n", omp_get_team_num(), omp_get_thread_num());
    }
  }

  return fail;
}

Compile with

clang -O2 -fopenmp --offload-arch=gfx1010  teams_nest.c -o teams_nest -mllvm -openmp-opt-disable=true

Fails with

LLVM ERROR: Absolute address LDS variable inconsistent with variable alignment

Save temps suggests LDS variables have had addresses allocated before codegen. I guess that's related to the thin-lto work, though I thought that was only in rocm and this is reproducing in trunk.

The LDS lowering makes a very definite "allocating whole code object at a time" assumption which means it's not safe to run repeatedly. The thinlto work was supposed to run it on pieces and then recombine them without running the pass a second time, but it looks like the pass is running a second time and decides to increase the alignment of some already allocated variables, i.e. miscompile them.

This specific error can probably be fixed by not over-aligning any already allocated variables, but it seems the writing is on the wall for the closed world assumption LDS lowering currently relies on.