161 lines · c
1// RUN: %libomptarget-compile-generic2// RUN: env HSA_XNACK=1 \3// RUN: %libomptarget-run-generic | %fcheck-generic4// XFAIL: nvptx64-nvidia-cuda5// XFAIL: nvptx64-nvidia-cuda-LTO6 7// REQUIRES: unified_shared_memory8 9#include <omp.h>10#include <stdio.h>11 12// End of definitions copied from OpenMP RTL.13// ---------------------------------------------------------------------------14 15#pragma omp requires unified_shared_memory16 17#define N 102418 19void init(int A[], int B[], int C[]) {20 for (int i = 0; i < N; ++i) {21 A[i] = 0;22 B[i] = 1;23 C[i] = i;24 }25}26 27int main(int argc, char *argv[]) {28 const int device = omp_get_default_device();29 30 // CHECK: Initial device: [[INITIAL_DEVICE:[0-9]+]]31 printf("Initial device: %d\n", omp_get_initial_device());32 // CHECK: Num devices: [[INITIAL_DEVICE]]33 printf("Num devices: %d\n", omp_get_num_devices());34 35 //36 // Target alloc & target memcpy37 //38 int A[N], B[N], C[N];39 40 // Init41 init(A, B, C);42 43 int *pA, *pB, *pC;44 45 // map ptrs46 pA = &A[0];47 pB = &B[0];48 pC = &C[0];49 50 int *d_A = (int *)omp_target_alloc(N * sizeof(int), device);51 int *d_B = (int *)omp_target_alloc(N * sizeof(int), device);52 int *d_C = (int *)omp_target_alloc(N * sizeof(int), device);53 54 // CHECK: omp_target_alloc succeeded55 printf("omp_target_alloc %s\n", d_A && d_B && d_C ? "succeeded" : "failed");56 57 omp_target_memcpy(d_B, pB, N * sizeof(int), 0, 0, device,58 omp_get_initial_device());59 omp_target_memcpy(d_C, pC, N * sizeof(int), 0, 0, device,60 omp_get_initial_device());61 62#pragma omp target is_device_ptr(d_A, d_B, d_C) device(device)63 {64#pragma omp parallel for schedule(static, 1)65 for (int i = 0; i < N; i++) {66 d_A[i] = d_B[i] + d_C[i] + 1;67 }68 }69 70 omp_target_memcpy(pA, d_A, N * sizeof(int), 0, 0, omp_get_initial_device(),71 device);72 73 // CHECK: Test omp_target_memcpy: Succeeded74 int fail = 0;75 for (int i = 0; i < N; ++i) {76 if (A[i] != i + 2)77 fail++;78 }79 if (fail) {80 printf("Test omp_target_memcpy: Failed\n");81 } else {82 printf("Test omp_target_memcpy: Succeeded\n");83 }84 85 //86 // target_is_present and target_associate/disassociate_ptr87 //88 init(A, B, C);89 90 // CHECK: B is not present, associating it...91 // CHECK: omp_target_associate_ptr B succeeded92 if (!omp_target_is_present(B, device)) {93 printf("B is not present, associating it...\n");94 int rc = omp_target_associate_ptr(B, d_B, N * sizeof(int), 0, device);95 printf("omp_target_associate_ptr B %s\n", !rc ? "succeeded" : "failed");96 }97 98 // CHECK: C is not present, associating it...99 // CHECK: omp_target_associate_ptr C succeeded100 if (!omp_target_is_present(C, device)) {101 printf("C is not present, associating it...\n");102 int rc = omp_target_associate_ptr(C, d_C, N * sizeof(int), 0, device);103 printf("omp_target_associate_ptr C %s\n", !rc ? "succeeded" : "failed");104 }105 106// CHECK: Inside target data: A is not present107// CHECK: Inside target data: B is present108// CHECK: Inside target data: C is present109#pragma omp target data map(from : B, C) device(device)110 {111 printf("Inside target data: A is%s present\n",112 omp_target_is_present(A, device) ? "" : " not");113 printf("Inside target data: B is%s present\n",114 omp_target_is_present(B, device) ? "" : " not");115 printf("Inside target data: C is%s present\n",116 omp_target_is_present(C, device) ? "" : " not");117 118#pragma omp target map(from : A) device(device)119 {120#pragma omp parallel for schedule(static, 1)121 for (int i = 0; i < N; i++)122 A[i] = B[i] + C[i] + 1;123 }124 }125 126 // CHECK: B is present, disassociating it...127 // CHECK: omp_target_disassociate_ptr B succeeded128 // CHECK: C is present, disassociating it...129 // CHECK: omp_target_disassociate_ptr C succeeded130 if (omp_target_is_present(B, device)) {131 printf("B is present, disassociating it...\n");132 int rc = omp_target_disassociate_ptr(B, device);133 printf("omp_target_disassociate_ptr B %s\n", !rc ? "succeeded" : "failed");134 }135 if (omp_target_is_present(C, device)) {136 printf("C is present, disassociating it...\n");137 int rc = omp_target_disassociate_ptr(C, device);138 printf("omp_target_disassociate_ptr C %s\n", !rc ? "succeeded" : "failed");139 }140 141 // CHECK: Test omp_target_associate_ptr: Succeeded142 fail = 0;143 for (int i = 0; i < N; ++i) {144 if (A[i] != i + 2)145 fail++;146 }147 if (fail) {148 printf("Test omp_target_associate_ptr: Failed\n");149 } else {150 printf("Test omp_target_associate_ptr: Succeeded\n");151 }152 153 omp_target_free(d_A, device);154 omp_target_free(d_B, device);155 omp_target_free(d_C, device);156 157 printf("Done!\n");158 159 return 0;160}161