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 } 76