OpenMP clause ordering (GPU offload) and is_device_ptr strange behavior

[Christos_Kotsalos]

Hello,

I have this very simple example: 

#include <stdio.h>

int main() {
  double *ptr = reinterpret_cast<double*>(0xdeadbeef);
  printf("ptr=%p\n", ptr);
  #pragma omp target parallel for simd is_device_ptr(ptr) if(true)
  for(int i = 0; i < 1; ++i) {
    printf("ptr=%p\n", ptr);
  }
  #pragma omp target parallel for simd if(true) is_device_ptr(ptr)
  for(int i = 0; i < 1; ++i) {
    printf("ptr=%p\n", ptr);
  }
  return 100;
}

For which I follow these steps in Intel DevCloud:

1. qsub -I -l nodes=1:gpu:ppn=2 -d .
2. icpx -qopenmp -fopenmp-targets=spir64 openmp_target_offload_clause_ordering.cpp
3. export OMP_TARGET_OFFLOAD=MANDATORY
4. ./a.out

and while it compiles it fails in runtime.

I have tested it with NVIDIA compilers (v22.2) and I had no issue (before the v22.2 there was an issue with the clause ordering -between if and is_device_ptr- but it was a confirmed bug).

I would really apreciate any feedback.

Thank you very much

Best,

Christos  

[ShivaniK_Intel]

Hi,

Thank you for posting in the Intel forums

We are able to reproduce the issue at our end by following the steps you have provided. We are working on it and will get back to you.

Thanks & Regards

Shivani

[Christos_Kotsalos]

Hi Shivani,

Thank you very much for your quick response.

I am looking forward to hearing from you regarding this compiler issue.

Best,

Christos

[ShivaniK_Intel · Accepted solution]

Hi,

From the code snippet you have shared with us we have observed that the pointer you are passing is not a valid device pointer.

Could you please try with the below-changed code snippet and let us know if you further face any issues?

#include <stdio.h>
#include <omp.h>
int main() {
 double *ptr = reinterpret_cast<double*>(omp_target_alloc(size, 0)); 
 printf("ptr=%p\n", ptr);
 #pragma omp target parallel for simd is_device_ptr(ptr) if(true)
 for(int i = 0; i < 1; ++i) {
  printf("ptr=%p\n", ptr);
 }
 #pragma omp target parallel for simd if(true) is_device_ptr(ptr)
 for(int i = 0; i < 1; ++i) {
 printf("ptr=%p\n", ptr);
 }
 return 100;
} 

Replace:

double *ptr = reinterpret_cast<double*>(omp_target_alloc(size, 0)); 

instead of:

 double *ptr = reinterpret_cast<double*>(0xdeadbeef);

Thanks & Regards

Shivani

[Christos_Kotsalos]

Hi Shivani,

Thank you very much for your prompt response.

Indeed, after your suggestion the code works without any issue.

Thanks for resolving it!

Best,

Christos

[ShivaniK_Intel]

Hi,

Thanks for accepting our solution. If you need any additional information, please post a new question as this thread will no longer be monitored by Intel.

Thanks & Regards

Shivani