LI81198: AN ILLEGAL MEMORY ACCESS WAS ENCOUNTERED ERROR.

APAR status

• Closed as program error.

Error description

• PROBLEM DESCRIPTION:
  
  Using default assignment operator in OpenMP target fails with
  "an illegal memory access was encountered"
  

Local fix

Problem summary

• USERS AFFECTED:
  Users who have test cases similar to the following are affected
  by this issue:
  #include <type_traits>
  #include <omp.h>
  #include <stdlib.h>
  #include <math.h>
  #include <stdio.h>
  
  #define N 1048576
  
  template <typename T, int LVL>
  struct simd_gpu
  {
  
    T d[LVL];
  
  #pragma omp declare target
  
  #ifdef USE_DEFAULT
    simd_gpu& operator=(const simd_gpu&)=default;
  #else
    simd_gpu& operator=(const simd_gpu& a){
  #pragma omp parallel for simd
      for(int i=0; i<LVL; i++)
      d[i] = a.d[i];
      return *this;
    }
  #endif
  
  simd_gpu operator*(const simd_gpu &a){
      simd_gpu temp;
  #pragma omp parallel for simd
      for(int i=0; i<LVL; i++)
      temp.d[i] = d[i] * a.d[i];
   return temp;
  }
  
    #pragma omp end declare target
  
  };
  
  void check (double *a, double *b, double *p)
  {
  #pragma omp parallel for
    for (int i = 0; i < N; i++)
    if (p[i] != a[i]*b[i]){
      printf("\n\n############################ error
  ############################\n\n");
      exit(1);
    }
    printf("\n\nSuccess !\n\n");
  }
  
  void vec_mult_gpu_simd (double *a, double *b, double *p) {
    omp_set_max_active_levels(3);
    printf("executing gpu simd\n");
    constexpr int LVL=32;
    typedef simd_gpu<double, LVL> simd_type;
    simd_type *as, *bs, *ps;
    as = reinterpret_cast<simd_type*>(a);
    bs = reinterpret_cast<simd_type*>(b);
    ps = reinterpret_cast<simd_type*>(p);
  
  #pragma omp target teams distribute parallel for
  map(ps[0:N/LVL], as[0:N/LVL], bs[0:N/LVL])
    for (int i = 0; i < N/LVL; i++)
      ps[i] = as[i] * bs[i]; //call operator* from simd_gpu and
  assign result by calling operator=
  }
  
  int main ()
  {
  double a1[N], b1[N], p1[N];
  #pragma omp parallel for
  for (int i = 0; i < N; i++){ //init
    a1[i] = ((i+1)%256 / N);
    b1[i] = i * M_PI;
  }
  
    vec_mult_gpu_simd (a1, b1, p1);
    check (a1, b1, p1);
  
    return 0;
  }
  
  
  PROBLEM DESCRIPTION:
  For the test case, the compiler backend doesn't generate the GPU
  kernel in SPMD mode.
  The reason for the compiler backend not doing that is because
  there is a "omp parallel do simd" construct inside the potential
  SPMD region.
  

Problem conclusion

• The limitation has been fixed in the compiler.
  

Temporary fix

Comments

APAR Information

• APAR is sysrouted FROM one or more of the following:

• APAR is sysrouted TO one or more of the following:

Fix information

• Fixed component name

  XL C/C++ LINUX

• Fixed component ID

  5725C7310

Applicable component levels