offload/test/offloading/info.c

   1 // RUN: %libomptarget-compile-generic \
   2 // RUN:     -gline-tables-only -fopenmp-extensions
   3 // RUN: env LIBOMPTARGET_INFO=63 %libomptarget-run-generic 2>&1 | \
   4 // RUN:   %fcheck-generic -allow-empty -check-prefixes=INFO
   5 // RUN: env LIBOMPTARGET_INFO=63 %libomptarget-run-amdgcn-amd-amdhsa 2>&1 | \
   6 // RUN:   %fcheck-amdgcn-amd-amdhsa -allow-empty -check-prefixes=INFO,AMDGPU
   7
   8 // FIXME: Fails due to optimized debugging in 'ptxas'.
   9 // UNSUPPORTED: nvptx64-nvidia-cuda-LTO
  10
  11 #include <omp.h>
  12 #include <stdio.h>
  13
  14 #define N 64
  15
  16 #pragma omp declare target
  17 int global;
  18 #pragma omp end declare target
  19
  20 extern void __tgt_set_info_flag(unsigned);
  21
  22 int main() {
  23   int A[N];
  24   int B[N];
  25   int C[N];
  26   int val = 1;
  27
  28 // clang-format off
  29 // INFO: info: Entering OpenMP data region with being_mapper at info.c:{{[0-9]+}}:{{[0-9]+}} with 3 arguments:
  30 // INFO: info: alloc(A[0:64])[256]
  31 // INFO: info: tofrom(B[0:64])[256]
  32 // INFO: info: to(C[0:64])[256]
  33 // INFO: info: Creating new map entry with HstPtrBase={{.*}}, HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, DynRefCount=1, HoldRefCount=0, Name=A[0:64]
  34 // INFO: info: Creating new map entry with HstPtrBase={{.*}}, HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, DynRefCount=0, HoldRefCount=1, Name=B[0:64]
  35 // INFO: info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=256, Name=B[0:64]
  36 // INFO: info: Creating new map entry with HstPtrBase={{.*}}, HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, DynRefCount=1, HoldRefCount=0, Name=C[0:64]
  37 // INFO: info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=256, Name=C[0:64]
  38 // INFO: info: OpenMP Host-Device pointer mappings after block at info.c:{{[0-9]+}}:{{[0-9]+}}:
  39 // INFO: info: Host Ptr           Target Ptr         Size (B) DynRefCount HoldRefCount Declaration
  40 // INFO: info: {{.*}}             {{.*}}             256      1           0            C[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
  41 // INFO: info: {{.*}}             {{.*}}             256      0           1            B[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
  42 // INFO: info: {{.*}}             {{.*}}             256      1           0            A[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
  43 // INFO: info: Entering OpenMP kernel at info.c:{{[0-9]+}}:{{[0-9]+}} with 1 arguments:
  44 // INFO: info: firstprivate(val)[4]
  45 // INFO: info: Launching kernel __omp_offloading_{{.*}}main{{.*}} with [{{[0-9]+}},1,1] blocks and [{{[0-9]+}},1,1] threads in Generic mode
  46 // AMDGPU: AMDGPU device {{[0-9]}} info: #Args: {{[0-9]}} Teams x Thrds: {{[0-9]+}}x {{[0-9]+}} (MaxFlatWorkGroupSize: {{[0-9]+}}) LDS Usage: {{[0-9]+}}B #SGPRs/VGPRs: {{[0-9]+}}/{{[0-9]+}} #SGPR/VGPR Spills: {{[0-9]+}}/{{[0-9]+}} Tripcount: {{[0-9]+}}
  47 // INFO: info: OpenMP Host-Device pointer mappings after block at info.c:{{[0-9]+}}:{{[0-9]+}}:
  48 // INFO: info: Host Ptr           Target Ptr         Size (B) DynRefCount HoldRefCount Declaration
  49 // INFO: info: {{.*}}             {{.*}}             256      1           0            C[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
  50 // INFO: info: {{.*}}             {{.*}}             256      0           1            B[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
  51 // INFO: info: {{.*}}             {{.*}}             256      1           0            A[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
  52 // INFO: info: Exiting OpenMP data region with end_mapper at info.c:{{[0-9]+}}:{{[0-9]+}} with 3 arguments:
  53 // INFO: info: alloc(A[0:64])[256]
  54 // INFO: info: tofrom(B[0:64])[256]
  55 // INFO: info: to(C[0:64])[256]
  56 // INFO: info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=256, Name=B[0:64]
  57 // INFO: info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=C[0:64]
  58 // INFO: info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=B[0:64]
  59 // INFO: info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=A[0:64]
  60 // INFO: info: OpenMP Host-Device pointer mappings after block at info.c:[[#%u,]]:[[#%u,]]:
  61 // INFO: info: Host Ptr  Target Ptr Size (B) DynRefCount HoldRefCount Declaration
  62 // INFO: info: [[#%#x,]] [[#%#x,]]  4        INF         0            global at unknown:0:0
  63 // clang-format on
  64 #pragma omp target data map(alloc : A[0 : N])                                  \
  65     map(ompx_hold, tofrom : B[0 : N]) map(to : C[0 : N])
  66 #pragma omp target firstprivate(val)
  67   { val = 1; }
  68
  69   __tgt_set_info_flag(0x0);
  70 // INFO-NOT: omptarget device 0 info: {{.*}}
  71 #pragma omp target
  72   {}
  73
  74   return 0;
  75 }