231 lines · c
1// RUN: %libomptarget-compile-generic -fopenmp-extensions2// RUN: %libomptarget-run-generic | %fcheck-generic -strict-whitespace3 4#include <omp.h>5#include <stdio.h>6 7#pragma omp begin declare target8#define MAX_NAME_SIZE 1009char N1[MAX_NAME_SIZE], N2[MAX_NAME_SIZE];10int V1, V2;11void copy_name(char *dst, char *src) {12 int i;13 for (i = 0; i < MAX_NAME_SIZE - 1 && src[i]; ++i)14 dst[i] = src[i];15 dst[i] = 0;16}17#pragma omp end declare target18 19#define CHECK_PRESENCE(Var1, Var2, Var3) \20 printf(" presence of %s, %s, %s: %d, %d, %d\n", #Var1, #Var2, #Var3, \21 omp_target_is_present(&(Var1), omp_get_default_device()), \22 omp_target_is_present(&(Var2), omp_get_default_device()), \23 omp_target_is_present(&(Var3), omp_get_default_device()))24 25#define CHECK_VALUES_HELPER(N1, N2, Var1, Var2) \26 printf(" values of %s, %s: %d, %d\n", N1, N2, (Var1), (Var2))27 28#define CHECK_VALUES_DELAYED(Var1, Var2) \29 copy_name(N1, #Var1); \30 copy_name(N2, #Var2); \31 V1 = (Var1); \32 V2 = (Var2);33 34#define CHECK_DELAYED_VALUES() \35 _Pragma("omp target update from(N1, N2, V1, V2)") \36 CHECK_VALUES_HELPER(N1, N2, V1, V2)37 38#define CHECK_VALUES(Var1, Var2) \39 CHECK_VALUES_HELPER(#Var1, #Var2, (Var1), (Var2))40 41int main() {42 struct S {43 int i;44 int j;45 } s;46 // CHECK: presence of s, s.i, s.j: 0, 0, 047 CHECK_PRESENCE(s, s.i, s.j);48 49 // =======================================================================50 // Check that ompx_hold keeps entire struct present.51 52 // -----------------------------------------------------------------------53 // CHECK-LABEL: check:{{.*}}54 printf("check: ompx_hold only on first member\n");55 s.i = 20;56 s.j = 30;57#pragma omp target data map(tofrom : s) map(ompx_hold, tofrom : s.i) \58 map(tofrom : s.j)59 {60 // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 161 CHECK_PRESENCE(s, s.i, s.j);62#pragma omp target map(tofrom : s)63 {64 s.i = 21;65 s.j = 31;66 }67#pragma omp target exit data map(delete : s, s.i)68 // ompx_hold on s.i applies to all of s.69 // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 170 // CHECK-NEXT: values of s.i, s.j: 20, 3071 CHECK_PRESENCE(s, s.i, s.j);72 CHECK_VALUES(s.i, s.j);73 }74 // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 075 // CHECK-NEXT: values of s.i, s.j: 21, 3176 CHECK_PRESENCE(s, s.i, s.j);77 CHECK_VALUES(s.i, s.j);78 79 // -----------------------------------------------------------------------80 // CHECK-LABEL: check:{{.*}}81 printf("check: ompx_hold only on last member\n");82 s.i = 20;83 s.j = 30;84#pragma omp target data map(tofrom : s) map(tofrom : s.i) \85 map(ompx_hold, tofrom : s.j)86 {87 // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 188 CHECK_PRESENCE(s, s.i, s.j);89#pragma omp target map(tofrom : s)90 {91 s.i = 21;92 s.j = 31;93 }94#pragma omp target exit data map(delete : s, s.i)95 // ompx_hold on s.j applies to all of s.96 // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 197 // CHECK-NEXT: values of s.i, s.j: 20, 3098 CHECK_PRESENCE(s, s.i, s.j);99 CHECK_VALUES(s.i, s.j);100 }101 // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0102 // CHECK-NEXT: values of s.i, s.j: 21, 31103 CHECK_PRESENCE(s, s.i, s.j);104 CHECK_VALUES(s.i, s.j);105 106 // -----------------------------------------------------------------------107 // CHECK-LABEL: check:{{.*}}108 printf("check: ompx_hold only on struct\n");109 s.i = 20;110 s.j = 30;111#pragma omp target data map(ompx_hold, tofrom : s) map(tofrom : s.i) \112 map(tofrom : s.j)113 {114 // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1115 CHECK_PRESENCE(s, s.i, s.j);116#pragma omp target map(tofrom : s)117 {118 s.i = 21;119 s.j = 31;120 }121#pragma omp target exit data map(delete : s, s.i)122 // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1123 // CHECK-NEXT: values of s.i, s.j: 20, 30124 CHECK_PRESENCE(s, s.i, s.j);125 CHECK_VALUES(s.i, s.j);126 }127 // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0128 // CHECK-NEXT: values of s.i, s.j: 21, 31129 CHECK_PRESENCE(s, s.i, s.j);130 CHECK_VALUES(s.i, s.j);131 132 // =======================================================================133 // Check that transfer to/from host checks reference count correctly.134 135 // -----------------------------------------------------------------------136 // CHECK-LABEL: check:{{.*}}137 printf("check: parent DynRefCount=1 is not sufficient for transfer\n");138 s.i = 20;139 s.j = 30;140#pragma omp target data map(ompx_hold, tofrom : s)141#pragma omp target data map(ompx_hold, tofrom : s)142 {143 // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1144 CHECK_PRESENCE(s, s.i, s.j);145#pragma omp target map(from : s.i, s.j)146 {147 s.i = 21;148 s.j = 31;149 } // No transfer here even though parent's DynRefCount=1.150 // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1151 // CHECK-NEXT: values of s.i, s.j: 20, 30152 CHECK_PRESENCE(s, s.i, s.j);153 CHECK_VALUES(s.i, s.j);154#pragma omp target map(to : s.i, s.j)155 { // No transfer here even though parent's DynRefCount=1.156 // CHECK-NEXT: values of s.i, s.j: 21, 31157 CHECK_VALUES_DELAYED(s.i, s.j);158 }159 CHECK_DELAYED_VALUES();160 }161 // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0162 // CHECK-NEXT: values of s.i, s.j: 21, 31163 CHECK_PRESENCE(s, s.i, s.j);164 CHECK_VALUES(s.i, s.j);165 166 // -----------------------------------------------------------------------167 // CHECK-LABEL: check:{{.*}}168 printf("check: parent HoldRefCount=1 is not sufficient for transfer\n");169 s.i = 20;170 s.j = 30;171#pragma omp target data map(tofrom : s)172#pragma omp target data map(tofrom : s)173 {174 // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1175 CHECK_PRESENCE(s, s.i, s.j);176#pragma omp target map(ompx_hold, from : s.i, s.j)177 {178 s.i = 21;179 s.j = 31;180 } // No transfer here even though parent's HoldRefCount=1.181 // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1182 // CHECK-NEXT: values of s.i, s.j: 20, 30183 CHECK_PRESENCE(s, s.i, s.j);184 CHECK_VALUES(s.i, s.j);185#pragma omp target map(ompx_hold, to : s.i, s.j)186 { // No transfer here even though parent's HoldRefCount=1.187 // CHECK-NEXT: values of s.i, s.j: 21, 31188 CHECK_VALUES_DELAYED(s.i, s.j);189 }190 CHECK_DELAYED_VALUES();191 }192 // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0193 // CHECK-NEXT: values of s.i, s.j: 21, 31194 CHECK_PRESENCE(s, s.i, s.j);195 CHECK_VALUES(s.i, s.j);196 197 // -----------------------------------------------------------------------198 // CHECK-LABEL: check:{{.*}}199 //200 // At the beginning of a region, if the parent's TotalRefCount=1, then the201 // transfer should happen.202 //203 // At the end of a region, it also must be true that the reference count being204 // decremented is the reference count that is 1.205 printf("check: parent TotalRefCount=1 is not sufficient for transfer\n");206 s.i = 20;207 s.j = 30;208#pragma omp target data map(ompx_hold, tofrom : s)209 {210 // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1211 CHECK_PRESENCE(s, s.i, s.j);212#pragma omp target map(ompx_hold, tofrom : s.i, s.j)213 {214 s.i = 21;215 s.j = 31;216 }217#pragma omp target exit data map(from : s.i, s.j)218 // No transfer here even though parent's TotalRefCount=1.219 // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1220 // CHECK-NEXT: values of s.i, s.j: 20, 30221 CHECK_PRESENCE(s, s.i, s.j);222 CHECK_VALUES(s.i, s.j);223 }224 // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0225 // CHECK-NEXT: values of s.i, s.j: 21, 31226 CHECK_PRESENCE(s, s.i, s.j);227 CHECK_VALUES(s.i, s.j);228 229 return 0;230}231