152 lines · c
1// Check that specifying device as omp_get_initial_device():2// - Doesn't cause the runtime to fail.3// - Offloads code to the host.4// - Doesn't transfer data. In this case, just check that neither host data nor5// default device data are affected by the specified transfers.6// - Works whether it's specified directly or as the default device.7 8// RUN: %libomptarget-compile-run-and-check-generic9 10#include <omp.h>11#include <stdio.h>12 13static void check(char *X, int Dev) {14 printf(" host X = %c\n", *X);15 char DV = -1;16#pragma omp target device(Dev) map(from : DV)17 DV = *X;18 printf("device X = %c\n", DV);19}20 21#define CHECK_DATA() check(&X, DevDefault)22 23int main(void) {24 int DevDefault = omp_get_default_device();25 int DevInit = omp_get_initial_device();26 27 //--------------------------------------------------28 // Initialize data on the host and default device.29 //--------------------------------------------------30 31 // CHECK: host X = h32 // CHECK-NEXT: device X = d33 char X = 'd';34#pragma omp target enter data map(to : X)35 X = 'h';36 CHECK_DATA();37 38//--------------------------------------------------39// Check behavior when specifying host directly.40//--------------------------------------------------41 42// CHECK-NEXT: omp_is_initial_device() = 143// CHECK-NEXT: host X = h44// CHECK-NEXT: device X = d45#pragma omp target device(DevInit) map(always, tofrom : X)46 printf("omp_is_initial_device() = %d\n", omp_is_initial_device());47 CHECK_DATA();48 49// CHECK-NEXT: omp_is_initial_device() = 150// CHECK-NEXT: host X = h51// CHECK-NEXT: device X = d52#pragma omp target teams device(DevInit) num_teams(1) map(always, tofrom : X)53 printf("omp_is_initial_device() = %d\n", omp_is_initial_device());54 CHECK_DATA();55 56// Check that __kmpc_push_target_tripcount_mapper doesn't fail. I'm not sure57// how to check that it actually pushes to the initial device.58#pragma omp target teams device(DevInit) num_teams(1)59#pragma omp distribute60 for (int i = 0; i < 2; ++i)61 ;62 63// CHECK-NEXT: host X = h64// CHECK-NEXT: device X = d65#pragma omp target data device(DevInit) map(always, tofrom : X)66 ;67 CHECK_DATA();68 69// CHECK-NEXT: host X = h70// CHECK-NEXT: device X = d71#pragma omp target enter data device(DevInit) map(always, to : X)72 ;73 CHECK_DATA();74 75// CHECK-NEXT: host X = h76// CHECK-NEXT: device X = d77#pragma omp target exit data device(DevInit) map(always, from : X)78 ;79 CHECK_DATA();80 81// CHECK-NEXT: host X = h82// CHECK-NEXT: device X = d83#pragma omp target update device(DevInit) to(X)84 ;85 CHECK_DATA();86 87// CHECK-NEXT: host X = h88// CHECK-NEXT: device X = d89#pragma omp target update device(DevInit) from(X)90 ;91 CHECK_DATA();92 93 //--------------------------------------------------94 // Check behavior when device defaults to host.95 //--------------------------------------------------96 97 omp_set_default_device(DevInit);98 99// CHECK-NEXT: omp_is_initial_device() = 1100// CHECK-NEXT: host X = h101// CHECK-NEXT: device X = d102#pragma omp target map(always, tofrom : X)103 printf("omp_is_initial_device() = %d\n", omp_is_initial_device());104 CHECK_DATA();105 106// CHECK-NEXT: omp_is_initial_device() = 1107// CHECK-NEXT: host X = h108// CHECK-NEXT: device X = d109#pragma omp target teams num_teams(1) map(always, tofrom : X)110 printf("omp_is_initial_device() = %d\n", omp_is_initial_device());111 CHECK_DATA();112 113// Check that __kmpc_push_target_tripcount_mapper doesn't fail. I'm not sure114// how to check that it actually pushes to the initial device.115#pragma omp target teams num_teams(1)116#pragma omp distribute117 for (int i = 0; i < 2; ++i)118 ;119 120// CHECK-NEXT: host X = h121// CHECK-NEXT: device X = d122#pragma omp target data map(always, tofrom : X)123 ;124 CHECK_DATA();125 126// CHECK-NEXT: host X = h127// CHECK-NEXT: device X = d128#pragma omp target enter data map(always, to : X)129 ;130 CHECK_DATA();131 132// CHECK-NEXT: host X = h133// CHECK-NEXT: device X = d134#pragma omp target exit data map(always, from : X)135 ;136 CHECK_DATA();137 138// CHECK-NEXT: host X = h139// CHECK-NEXT: device X = d140#pragma omp target update to(X)141 ;142 CHECK_DATA();143 144// CHECK-NEXT: host X = h145// CHECK-NEXT: device X = d146#pragma omp target update from(X)147 ;148 CHECK_DATA();149 150 return 0;151}152