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
17int global;
18#pragma omp end declare target
19
20extern void __tgt_set_info_flag(unsigned);
21
22int 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]+}} blocks and {{[0-9]+}} 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}
76

source code of offload/test/offloading/info.c