[PATCH] D103225: [AMDGPU] Replace non-kernel function uses of LDS globals by pointers.

[Dhruva Chakrabarti via Phabricator via llvm-commits]

dhruvachak added a comment.

I am seeing a run-time failure as well in addition to the compile-time failure reported earlier. Here's a test case, veccopy.c.

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

int main()
{

  int N = 100000;
  
  int a[N];
  int b[N];
  
  int i;
  
  for (i=0; i<N; i++){
    a[i]=0;
    b[i]=i;
  }

#pragma omp target teams distribute parallel for map(from: a[0:N]) map(to: b[0:N])

  {
    for (int j = 0; j< N; j++)
      a[j]=b[j];
  }
  
  int rc = 0;
  for (i=0; i<N; i++)
    if (a[i] != b[i] ) {
      rc++;
      printf ("Wrong value: a[%d]=%d\n", i, a[i]);
    }
  
  if (!rc){
    printf("Success\n");
    return EXIT_SUCCESS;
  } else{
    printf("Failure\n");
    return EXIT_FAILURE;
  }

}

Compile/run:
clang -O2 -target x86_64-pc-linux-gnu -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx906   veccopy.c -o veccopy
./veccopy 
[GPU Memory Error] Addr: 0x7f60a4210000 Reason: Page not present or supervisor privilege. 
Memory access fault by GPU node-1 (Agent handle: 0x168c5d0) on address 0x7f60a4210000. Reason: Page not present or supervisor privilege.
Aborted (core dumped)

Adding -mllvm -amdgpu-enable-lds-replace-with-pointer=false gets rid of the problem. 
Running with env var LIBOMPTARGET_KERNEL_TRACE=1
shows LDS went down from 16400B to 32B but it triggered the above problem.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D103225/new/

https://reviews.llvm.org/D103225