161 lines · c
1// clang-format off2// RUN: %libomptarget-compile-run-and-check-generic3// REQUIRES: ompt4// REQUIRES: gpu5// clang-format on6 7/*8 * Example OpenMP program that registers EMI callbacks.9 * Explicitly testing for an initialized device num and10 * #pragma omp target [data enter / data exit / update]11 * The latter with the addition of a nowait clause.12 */13 14#include <omp.h>15#include <stdio.h>16 17#include "callbacks.h"18#include "register_emi.h"19 20#define N 10000021 22#pragma omp declare target23int c[N];24#pragma omp end declare target25 26int main() {27 int a[N];28 int b[N];29 30 int i;31 32 for (i = 0; i < N; i++)33 a[i] = 0;34 35 for (i = 0; i < N; i++)36 b[i] = i;37 38 for (i = 0; i < N; i++)39 c[i] = 0;40 41#pragma omp target enter data map(to : a)42#pragma omp target parallel for43 {44 for (int j = 0; j < N; j++)45 a[j] = b[j];46 }47#pragma omp target exit data map(from : a)48 49#pragma omp target parallel for map(alloc : c)50 {51 for (int j = 0; j < N; j++)52 c[j] = 2 * j + 1;53 }54#pragma omp target update from(c) nowait55#pragma omp barrier56 57 int rc = 0;58 for (i = 0; i < N; i++) {59 if (a[i] != i) {60 rc++;61 printf("Wrong value: a[%d]=%d\n", i, a[i]);62 }63 }64 65 for (i = 0; i < N; i++) {66 if (c[i] != 2 * i + 1) {67 rc++;68 printf("Wrong value: c[%d]=%d\n", i, c[i]);69 }70 }71 72 if (!rc)73 printf("Success\n");74 75 return rc;76}77 78// clang-format off79/// CHECK-NOT: Callback Target EMI:80/// CHECK-NOT: device_num=-181/// CHECK: Callback Init:82/// CHECK: Callback Load:83/// CHECK: Callback Target EMI: kind=ompt_target_enter_data endpoint=ompt_scope_begin84/// CHECK-NOT: device_num=-185/// CHECK-NOT: code=(nil)86/// CHECK: code=[[CODE1:.*]]87/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_alloc88/// CHECK: code=[[CODE1]]89/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_alloc90/// CHECK-NOT: dest=(nil)91/// CHECK: code=[[CODE1]]92/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_to_device93/// CHECK: code=[[CODE1]]94/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_to_device95/// CHECK: code=[[CODE1]]96/// CHECK: Callback Target EMI: kind=ompt_target_enter_data endpoint=ompt_scope_end97/// CHECK-NOT: device_num=-198/// CHECK: code=[[CODE1]]99/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_begin100/// CHECK-NOT: device_num=-1101/// CHECK-NOT: code=(nil)102/// CHECK: code=[[CODE2:.*]]103/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_alloc104/// CHECK: code=[[CODE2]]105/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_alloc106/// CHECK-NOT: dest=(nil)107/// CHECK: code=[[CODE2]]108/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_to_device109/// CHECK: code=[[CODE2]]110/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_to_device111/// CHECK: code=[[CODE2]]112/// CHECK: Callback Submit EMI: endpoint=ompt_scope_begin req_num_teams=1113/// CHECK: Callback Submit EMI: endpoint=ompt_scope_end req_num_teams=1114/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_from_device115/// CHECK: code=[[CODE2]]116/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_from_device117/// CHECK: code=[[CODE2]]118/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_delete119/// CHECK: code=[[CODE2]]120/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_delete121/// CHECK: code=[[CODE2]]122/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_end123/// CHECK-NOT: device_num=-1124/// CHECK: code=[[CODE2]]125/// CHECK: Callback Target EMI: kind=ompt_target_exit_data endpoint=ompt_scope_begin126/// CHECK-NOT: device_num=-1127/// CHECK-NOT: code=(nil)128/// CHECK: code=[[CODE3:.*]]129/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_from_device130/// CHECK: code=[[CODE3]]131/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_from_device132/// CHECK: code=[[CODE3]]133/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_delete134/// CHECK: code=[[CODE3]]135/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_delete136/// CHECK: code=[[CODE3]]137/// CHECK: Callback Target EMI: kind=ompt_target_exit_data endpoint=ompt_scope_end138/// CHECK-NOT: device_num=-1139/// CHECK: code=[[CODE3]]140/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_begin141/// CHECK-NOT: device_num=-1142/// CHECK-NOT: code=(nil)143/// CHECK: code=[[CODE4:.*]]144/// CHECK: Callback Submit EMI: endpoint=ompt_scope_begin req_num_teams=1145/// CHECK: Callback Submit EMI: endpoint=ompt_scope_end req_num_teams=1146/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_end147/// CHECK-NOT: device_num=-1148/// CHECK: code=[[CODE4]]149/// CHECK: Callback Target EMI: kind=ompt_target_update endpoint=ompt_scope_begin150/// CHECK-NOT: device_num=-1151/// CHECK-NOT: code=(nil)152/// CHECK: code=[[CODE5:.*]]153/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_from_device154/// CHECK: code=[[CODE5]]155/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_from_device156/// CHECK: code=[[CODE5]]157/// CHECK: Callback Target EMI: kind=ompt_target_update endpoint=ompt_scope_end158/// CHECK-NOT: device_num=-1159/// CHECK: code=[[CODE5]]160/// CHECK: Callback Fini:161