| 1 | // RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu |
| 2 | // RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu |
| 3 | // RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu |
| 4 | // RUN: %libomptarget-compilexx-run-and-check-x86_64-unknown-linux-gnu |
| 5 | // RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda |
| 6 | |
| 7 | #include <assert.h> |
| 8 | #include <omp.h> |
| 9 | #include <stdio.h> |
| 10 | #include <stdlib.h> |
| 11 | |
| 12 | #define N 2 |
| 13 | |
| 14 | class MyObjectA { |
| 15 | public: |
| 16 | MyObjectA() { |
| 17 | data1 = 1; |
| 18 | data2 = 2; |
| 19 | } |
| 20 | void show() { |
| 21 | printf(format: "\t\tObject A Contents:\n" ); |
| 22 | printf(format: "\t\t\tdata1 = %d data2 = %d\n" , data1, data2); |
| 23 | } |
| 24 | void foo() { |
| 25 | data1 += 10; |
| 26 | data2 += 20; |
| 27 | } |
| 28 | int data1; |
| 29 | int data2; |
| 30 | }; |
| 31 | |
| 32 | class MyObjectB { |
| 33 | public: |
| 34 | MyObjectB() { |
| 35 | arr = new MyObjectA[N]; |
| 36 | len = N; |
| 37 | } |
| 38 | void show() { |
| 39 | printf(format: "\tObject B Contents:\n" ); |
| 40 | for (int i = 0; i < len; i++) |
| 41 | arr[i].show(); |
| 42 | } |
| 43 | void foo() { |
| 44 | for (int i = 0; i < len; i++) |
| 45 | arr[i].foo(); |
| 46 | } |
| 47 | MyObjectA *arr; |
| 48 | int len; |
| 49 | }; |
| 50 | #pragma omp declare mapper(MyObjectB obj) map(obj, obj.arr[ : obj.len]) |
| 51 | |
| 52 | class MyObjectC { |
| 53 | public: |
| 54 | MyObjectC() { |
| 55 | arr = new MyObjectB[N]; |
| 56 | len = N; |
| 57 | } |
| 58 | void show() { |
| 59 | printf(format: "Object C Contents:\n" ); |
| 60 | for (int i = 0; i < len; i++) |
| 61 | arr[i].show(); |
| 62 | } |
| 63 | void foo() { |
| 64 | for (int i = 0; i < len; i++) |
| 65 | arr[i].foo(); |
| 66 | } |
| 67 | MyObjectB *arr; |
| 68 | int len; |
| 69 | }; |
| 70 | #pragma omp declare mapper(MyObjectC obj) map(obj, obj.arr[ : obj.len]) |
| 71 | |
| 72 | int main(void) { |
| 73 | MyObjectC *outer = new MyObjectC[N]; |
| 74 | |
| 75 | printf(format: "Original data hierarchy:\n" ); |
| 76 | for (int i = 0; i < N; i++) |
| 77 | outer[i].show(); |
| 78 | |
| 79 | printf(format: "Sending data to device...\n" ); |
| 80 | #pragma omp target enter data map(to : outer[ : N]) |
| 81 | |
| 82 | printf(format: "Calling foo()...\n" ); |
| 83 | #pragma omp target teams distribute parallel for |
| 84 | for (int i = 0; i < N; i++) |
| 85 | outer[i].foo(); |
| 86 | |
| 87 | printf(format: "foo() complete!\n" ); |
| 88 | |
| 89 | printf(format: "Sending data back to host...\n" ); |
| 90 | #pragma omp target exit data map(from : outer[ : N]) |
| 91 | |
| 92 | printf(format: "Modified Data Hierarchy:\n" ); |
| 93 | for (int i = 0; i < N; i++) |
| 94 | outer[i].show(); |
| 95 | |
| 96 | printf(format: "Testing for correctness...\n" ); |
| 97 | for (int i = 0; i < N; ++i) |
| 98 | for (int j = 0; j < N; ++j) |
| 99 | for (int k = 0; k < N; ++k) { |
| 100 | printf(format: "outer[%d].arr[%d].arr[%d].data1 = %d.\n" , i, j, k, |
| 101 | outer[i].arr[j].arr[k].data1); |
| 102 | printf(format: "outer[%d].arr[%d].arr[%d].data2 = %d.\n" , i, j, k, |
| 103 | outer[i].arr[j].arr[k].data2); |
| 104 | assert(outer[i].arr[j].arr[k].data1 == 11 && |
| 105 | outer[i].arr[j].arr[k].data2 == 22); |
| 106 | } |
| 107 | // CHECK: outer[0].arr[0].arr[0].data1 = 11. |
| 108 | // CHECK: outer[0].arr[0].arr[0].data2 = 22. |
| 109 | // CHECK: outer[0].arr[0].arr[1].data1 = 11. |
| 110 | // CHECK: outer[0].arr[0].arr[1].data2 = 22. |
| 111 | // CHECK: outer[0].arr[1].arr[0].data1 = 11. |
| 112 | // CHECK: outer[0].arr[1].arr[0].data2 = 22. |
| 113 | // CHECK: outer[0].arr[1].arr[1].data1 = 11. |
| 114 | // CHECK: outer[0].arr[1].arr[1].data2 = 22. |
| 115 | // CHECK: outer[1].arr[0].arr[0].data1 = 11. |
| 116 | // CHECK: outer[1].arr[0].arr[0].data2 = 22. |
| 117 | // CHECK: outer[1].arr[0].arr[1].data1 = 11. |
| 118 | // CHECK: outer[1].arr[0].arr[1].data2 = 22. |
| 119 | // CHECK: outer[1].arr[1].arr[0].data1 = 11. |
| 120 | // CHECK: outer[1].arr[1].arr[0].data2 = 22. |
| 121 | // CHECK: outer[1].arr[1].arr[1].data1 = 11. |
| 122 | // CHECK: outer[1].arr[1].arr[1].data2 = 22. |
| 123 | assert(outer[1].arr[1].arr[0].data1 == 11 && |
| 124 | outer[1].arr[1].arr[0].data2 == 22 && |
| 125 | outer[1].arr[1].arr[1].data1 == 11 && |
| 126 | outer[1].arr[1].arr[1].data2 == 22); |
| 127 | |
| 128 | return 0; |
| 129 | } |
| 130 | |