1 // RUN: %libomptarget-compile-nvptx64-nvidia-cuda \
2 // RUN: -gline-tables-only -fopenmp-extensions
3 // RUN: env LIBOMPTARGET_INFO=63 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | \
4 // RUN: %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=INFO
5 // REQUIRES: nvptx64-nvidia-cuda
6
7 #include <stdio.h>
8 #include <omp.h>
9
10 #define N 64
11
12 #pragma omp declare target
13 int global;
14 #pragma omp end declare target
15
16 extern void __tgt_set_info_flag(unsigned);
17
main()18 int main() {
19 int A[N];
20 int B[N];
21 int C[N];
22 int val = 1;
23
24 // INFO: CUDA device 0 info: Device supports up to {{[0-9]+}} CUDA blocks and {{[0-9]+}} threads with a warp size of {{[0-9]+}}
25 // INFO: Libomptarget device 0 info: Entering OpenMP data region at info.c:{{[0-9]+}}:{{[0-9]+}} with 3 arguments:
26 // INFO: Libomptarget device 0 info: alloc(A[0:64])[256]
27 // INFO: Libomptarget device 0 info: tofrom(B[0:64])[256]
28 // INFO: Libomptarget device 0 info: to(C[0:64])[256]
29 // INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBase={{.*}}, HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, DynRefCount=1, HoldRefCount=0, Name=A[0:64]
30 // INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBase={{.*}}, HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, DynRefCount=0, HoldRefCount=1, Name=B[0:64]
31 // INFO: Libomptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=256, Name=B[0:64]
32 // INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBase={{.*}}, HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, DynRefCount=1, HoldRefCount=0, Name=C[0:64]
33 // INFO: Libomptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=256, Name=C[0:64]
34 // INFO: Libomptarget device 0 info: OpenMP Host-Device pointer mappings after block at info.c:{{[0-9]+}}:{{[0-9]+}}:
35 // INFO: Libomptarget device 0 info: Host Ptr Target Ptr Size (B) DynRefCount HoldRefCount Declaration
36 // INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 0 C[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
37 // INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 0 1 B[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
38 // INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 0 A[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
39 // INFO: Libomptarget device 0 info: Entering OpenMP kernel at info.c:{{[0-9]+}}:{{[0-9]+}} with 1 arguments:
40 // INFO: Libomptarget device 0 info: firstprivate(val)[4]
41 // INFO: CUDA device 0 info: Launching kernel __omp_offloading_{{.*}}main{{.*}} with {{[0-9]+}} blocks and {{[0-9]+}} threads in Generic mode
42 // INFO: Libomptarget device 0 info: OpenMP Host-Device pointer mappings after block at info.c:{{[0-9]+}}:{{[0-9]+}}:
43 // INFO: Libomptarget device 0 info: Host Ptr Target Ptr Size (B) DynRefCount HoldRefCount Declaration
44 // INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 0 C[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
45 // INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 0 1 B[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
46 // INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 0 A[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
47 // INFO: Libomptarget device 0 info: Exiting OpenMP data region at info.c:{{[0-9]+}}:{{[0-9]+}} with 3 arguments:
48 // INFO: Libomptarget device 0 info: alloc(A[0:64])[256]
49 // INFO: Libomptarget device 0 info: tofrom(B[0:64])[256]
50 // INFO: Libomptarget device 0 info: to(C[0:64])[256]
51 // INFO: Libomptarget device 0 info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=256, Name=B[0:64]
52 // INFO: Libomptarget device 0 info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=C[0:64]
53 // INFO: Libomptarget device 0 info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=B[0:64]
54 // INFO: Libomptarget device 0 info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=A[0:64]
55 // INFO: Libomptarget device 0 info: OpenMP Host-Device pointer mappings after block at info.c:[[#%u,]]:[[#%u,]]:
56 // INFO: Libomptarget device 0 info: Host Ptr Target Ptr Size (B) DynRefCount HoldRefCount Declaration
57 // INFO: Libomptarget device 0 info: [[#%#x,]] [[#%#x,]] 4 INF 0 global at unknown:0:0
58 #pragma omp target data map(alloc:A[0:N]) map(ompx_hold,tofrom:B[0:N]) map(to:C[0:N])
59 #pragma omp target firstprivate(val)
60 { val = 1; }
61
62 __tgt_set_info_flag(0x0);
63 // INFO-NOT: Libomptarget device 0 info: {{.*}}
64 #pragma omp target
65 { }
66
67 return 0;
68 }
69