Index: libomptarget/test/offloading/target_data/test_target_data_if.c =================================================================== --- libomptarget/test/offloading/target_data/test_target_data_if.c +++ libomptarget/test/offloading/target_data/test_target_data_if.c @@ -3,56 +3,53 @@ // RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu // RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu -#include #include +#include /*Test for OpenMP 4.5 target data with if*/ -int main () -{ - int a[1024]; - int b[1024]; - int c[1024]; - int size,i=0,success=0; +int main() { + int a[1024]; + int b[1024]; + int c[1024]; + int size, i = 0, success = 0; + + for (i = 0; i < 1024; i++) { + a[i] = 1; + b[i] = i; + c[i] = 0; + } - for (i = 0; i < 1024; i++) + for (size = 250; size < 1024; size += 256) { +#pragma omp target data if (size > 256) map(to : size) \ + map(tofrom : c[0 : size]) \ + map(to : a[0 : size], b[0 : size], i) { - a[i] = 1; - b[i] = i; - c[i] = 0; +#pragma omp target if (size > 256) + { + if (!omp_is_initial_device()) { + for (i = 0; i < size; i++) + c[i] = a[i] + b[i]; + } /*end-if*/ + } } - for(size=250; size <1024; size+=256){ - #pragma omp target data if(size > 256) map(to:size) map(tofrom: c[0:size]) map(to:a[0:size],b[0:size],i) - { - #pragma omp target if(size > 256) - { - if(!omp_is_initial_device()){ - for (i = 0; i < size; i++) - c[i] = a[i] + b[i]; - }/*end-if*/ + for (i = 0; i < size; i++) { + if (size < 256) { + if (c[i] != 0) { + success = 1; } - } + } else { + if (c[i] != i + 1) { + success = 1; + } + } /*end-else*/ + } + } /*end-for*/ - for (i = 0; i < size; i++){ - if(size < 256){ - if(c[i] != 0) - { - success=1; - } - } else{ - if(c[i] != i+1) - { - success=1; - } - }/*end-else*/ - } - }/*end-for*/ - - if(success) - printf("Test failed\n"); - else - printf("Test passed\n"); + if (success) + printf("Test failed\n"); + else + printf("Test passed\n"); - - return success; + return success; } Index: libomptarget/test/offloading/target_data/test_target_data_map.c =================================================================== --- libomptarget/test/offloading/target_data/test_target_data_map.c +++ libomptarget/test/offloading/target_data/test_target_data_map.c @@ -3,218 +3,234 @@ // RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu // RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu +#include +#include #include #include -#include -#include #define N 1000 int test_map_from() { - printf("test_map_from\n"); - - int sum = 0, sum2 = 0, errors = 0; - - // host arrays: heap and stack - int *h_array_h = (int*) malloc(N * sizeof(int)); - int h_array_s[N]; - - #pragma omp target data map(from:h_array_h[0:N]) map(from:h_array_s[0:N]) - { - #pragma omp target - if (!omp_is_initial_device()) - for (int i = 0; i < N; ++i) { - h_array_h[i] = 1; - h_array_s[i] = 2; - } - } + printf("test_map_from\n"); - // checking results - for(int i = 0; i < N; ++i) { - sum += h_array_h[i]; - sum2 += h_array_s[i]; - } + int sum = 0, sum2 = 0, errors = 0; - errors = (N != sum) || (2*N != sum2); - if(!errors) - printf("Test passed sum=%d, sum2=%d, N=%d\n", sum, sum2, N); - else - printf("Test failed: sum=%d, sum2=%d, N=%d\n", sum, sum2, N); + // host arrays: heap and stack + int *h_array_h = (int *)malloc(N * sizeof(int)); + int h_array_s[N]; - return errors; +#pragma omp target data map(from : h_array_h[0 : N]) \ + map(from : h_array_s[0 : N]) + { +#pragma omp target + if (!omp_is_initial_device()) + for (int i = 0; i < N; ++i) { + h_array_h[i] = 1; + h_array_s[i] = 2; + } + } + + // checking results + for (int i = 0; i < N; ++i) { + sum += h_array_h[i]; + sum2 += h_array_s[i]; + } + + free(h_array_h); + errors = (N != sum) || (2 * N != sum2); + if (!errors) + printf("Test passed sum=%d, sum2=%d, N=%d\n", sum, sum2, N); + else + printf("Test failed: sum=%d, sum2=%d, N=%d\n", sum, sum2, N); + + return errors; } int test_map_tofrom() { - - printf("test_map_tofrom\n"); - - int sum = 0, sum2 = 0, errors = 0; - - // host arrays: heap and stack - int *h_array_h = (int*) malloc(N * sizeof(int)); - int h_array_s[N]; - - for(int i = 0; i < N; ++i) { - h_array_h[i] = 0; - h_array_s[i] = 0; - } - #pragma omp target data map(tofrom:h_array_h[0:N]) map(tofrom:h_array_s[0:N]) - { - #pragma omp target - if (!omp_is_initial_device()) - for(int i = 0; i < N; ++i) { - h_array_h[i] = 1; - h_array_s[i] = 1; - } - } // end of data target region - - // checking errors - for(int i = 0; i < N; ++i) { - sum += h_array_h[i]; - sum2 += h_array_s[i]; - } + printf("test_map_tofrom\n"); - errors = (N != sum) || (N != sum2); - if(!errors) - printf("Test passed sum=%d, sum2=%d, N=%d\n", sum, sum2, N); - else - printf("Test failed: sum=%d, sum2=%d, N=%d\n", sum, sum2, N); + int sum = 0, sum2 = 0, errors = 0; - return errors; -} + // host arrays: heap and stack + int *h_array_h = (int *)malloc(N * sizeof(int)); + int h_array_s[N]; -int test_map_to() { + for (int i = 0; i < N; ++i) { + h_array_h[i] = 0; + h_array_s[i] = 0; + } - printf("test_map_to\n"); - - int sum = 0, sum2 = 0, errors = 0; - // host arrays: heap and stack - int *h_array_h = (int*) malloc(N * sizeof(int)); - int *h_array2_h = (int*) malloc(N*sizeof(int)); - int h_array_s[N]; - int h_array2_s[N]; - - // initializing arrays - for (int i = 0; i < N; ++i) { +#pragma omp target data map(tofrom : h_array_h[0 : N]) \ + map(tofrom : h_array_s[0 : N]) + { +#pragma omp target + if (!omp_is_initial_device()) + for (int i = 0; i < N; ++i) { h_array_h[i] = 1; h_array_s[i] = 1; - h_array2_h[i] = 0; - h_array2_s[i] = 0; - } - - // device arrays to get the data from the device - int *d_array = (int*) omp_target_alloc(N*sizeof(int), omp_get_default_device()); - int *d_array2 = (int*) omp_target_alloc(N*sizeof(int), omp_get_default_device()); - - #pragma omp target data map(to:h_array_h[0:N]) map(to:h_array_s[0:N]) - { - #pragma omp target is_device_ptr(d_array, d_array2) - if (!omp_is_initial_device()) - for (int i = 0; i < N; ++i) { - d_array[i] = h_array_h[i]; - d_array2[i] = h_array_s[i]; - } - } - - // copy from d to h - omp_target_memcpy(h_array2_h, d_array, N * sizeof(int), 0, 0, omp_get_initial_device(), omp_get_default_device()); - omp_target_memcpy(h_array2_s, d_array2, N * sizeof(int), 0, 0, omp_get_initial_device(), omp_get_default_device()); - // deallocating device arrays - omp_target_free(d_array, omp_get_default_device()); - omp_target_free(d_array2, omp_get_default_device()); - - // checking errors - for(int i = 0; i < N; ++i) { - sum += h_array2_h[i]; - sum2 += h_array2_s[i]; - } - - errors = (N != sum) || (N != sum2); - if(!errors) - printf("Test passed sum=%d, sum2=%d, N=%d\n", sum, sum2, N); - else - printf("Test failed: sum=%d, sum2=%d, N=%d\n", sum, sum2, N); - - return errors; + } + } // end of data target region + + // checking errors + for (int i = 0; i < N; ++i) { + sum += h_array_h[i]; + sum2 += h_array_s[i]; + } + + free(h_array_h); + errors = (N != sum) || (N != sum2); + if (!errors) + printf("Test passed sum=%d, sum2=%d, N=%d\n", sum, sum2, N); + else + printf("Test failed: sum=%d, sum2=%d, N=%d\n", sum, sum2, N); + + return errors; } -int test_map_to_from() { +int test_map_to() { - printf("test_map_to_from\n"); - - int sum = 0, errors = 0; - int *h_array_h = (int*) malloc(N*sizeof(int)); - int *h_array2_h = (int*) malloc(N*sizeof(int)); + printf("test_map_to\n"); + + int sum = 0, sum2 = 0, errors = 0; + // host arrays: heap and stack + int *h_array_h = (int *)malloc(N * sizeof(int)); + int *h_array2_h = (int *)malloc(N * sizeof(int)); + int h_array_s[N]; + int h_array2_s[N]; + + // initializing arrays + for (int i = 0; i < N; ++i) { + h_array_h[i] = 1; + h_array_s[i] = 1; + h_array2_h[i] = 0; + h_array2_s[i] = 0; + } + + // device arrays to get the data from the device + int *d_array = + (int *)omp_target_alloc(N * sizeof(int), omp_get_default_device()); + int *d_array2 = + (int *)omp_target_alloc(N * sizeof(int), omp_get_default_device()); + +#pragma omp target data map(to : h_array_h[0 : N]) map(to : h_array_s[0 : N]) + { +#pragma omp target is_device_ptr(d_array, d_array2) + if (!omp_is_initial_device()) + for (int i = 0; i < N; ++i) { + d_array[i] = h_array_h[i]; + d_array2[i] = h_array_s[i]; + } + } + + // copy from d to h + omp_target_memcpy(h_array2_h, d_array, N * sizeof(int), 0, 0, + omp_get_initial_device(), omp_get_default_device()); + omp_target_memcpy(h_array2_s, d_array2, N * sizeof(int), 0, 0, + omp_get_initial_device(), omp_get_default_device()); + // deallocating device arrays + omp_target_free(d_array, omp_get_default_device()); + omp_target_free(d_array2, omp_get_default_device()); + + // checking errors + for (int i = 0; i < N; ++i) { + sum += h_array2_h[i]; + sum2 += h_array2_s[i]; + } + + free(h_array_h); + free(h_array2_h); + errors = (N != sum) || (N != sum2); + if (!errors) + printf("Test passed sum=%d, sum2=%d, N=%d\n", sum, sum2, N); + else + printf("Test failed: sum=%d, sum2=%d, N=%d\n", sum, sum2, N); + + return errors; +} - for (int i = 0; i < N; ++i) { - h_array_h[i] = 1; - h_array2_h[i] = 0; - } +int test_map_to_from() { - #pragma omp target data map(to:h_array_h[0:N]) map(from:h_array2_h[0:N]) - { - #pragma omp target - if (!omp_is_initial_device()) - for (int i = 0; i < N; ++i) - h_array2_h[i] = h_array_h[i]; - } - - // cheking errors - for(int i = 0; i < N; ++i) - sum += h_array2_h[i]; - - errors = N - sum; - if(!errors) - printf("Test passed sum=%d, N=%d\n", sum, N); - else - printf("Test failed sum=%d, N=%d\n", sum, N); - - return errors; + printf("test_map_to_from\n"); + + int sum = 0, errors = 0; + int *h_array_h = (int *)malloc(N * sizeof(int)); + int *h_array2_h = (int *)malloc(N * sizeof(int)); + + for (int i = 0; i < N; ++i) { + h_array_h[i] = 1; + h_array2_h[i] = 0; + } + +#pragma omp target data map(to : h_array_h[0 : N]) map(from : h_array2_h[0 : N]) + { +#pragma omp target + if (!omp_is_initial_device()) + for (int i = 0; i < N; ++i) + h_array2_h[i] = h_array_h[i]; + } + + // cheking errors + for (int i = 0; i < N; ++i) + sum += h_array2_h[i]; + + free(h_array_h); + free(h_array2_h); + errors = N - sum; + if (!errors) + printf("Test passed sum=%d, N=%d\n", sum, N); + else + printf("Test failed sum=%d, N=%d\n", sum, N); + + return errors; } int test_map_alloc() { - puts("test_map_alloc"); - - int sum = 0, errors = 0; - int *h_array_h = (int*) malloc(N*sizeof(int)); - int *d_sum = (int*) omp_target_alloc(sizeof(int), omp_get_default_device()); - #pragma omp target data map(alloc:h_array_h[0:N]) - { - #pragma omp target is_device_ptr(d_sum) - if (!omp_is_initial_device()) { - for (int i = 0; i < N; ++i) { - h_array_h[i] = 1; - } - // cheking errors - d_sum[0] = 0; // pointer arithmetic is not supported on the devices for the device address returned by omp_target_alloc - for (int i = 0; i < N; ++i) - d_sum[0] += h_array_h[i]; - } - assert(0 == omp_target_memcpy(&sum, d_sum, sizeof(int), 0, 0, omp_get_initial_device(), omp_get_default_device())); - } - omp_target_free(d_sum, omp_get_default_device()); - - errors = N - sum; - if(!errors) - printf("Test passed sum=%d, N=%d\n", sum, N); - else - printf("Test failed sum=%d, N=%d\n", sum, N); + puts("test_map_alloc"); - return errors; + int sum = 0, errors = 0; + int *h_array_h = (int *)malloc(N * sizeof(int)); + int *d_sum = (int *)omp_target_alloc(sizeof(int), omp_get_default_device()); +#pragma omp target data map(alloc : h_array_h[0 : N]) + { +#pragma omp target is_device_ptr(d_sum) + if (!omp_is_initial_device()) { + for (int i = 0; i < N; ++i) { + h_array_h[i] = 1; + } + // cheking errors + d_sum[0] = 0; // pointer arithmetic is not supported on the devices for + // the device address returned by omp_target_alloc + for (int i = 0; i < N; ++i) + d_sum[0] += h_array_h[i]; + } + assert(0 == omp_target_memcpy(&sum, d_sum, sizeof(int), 0, 0, + omp_get_initial_device(), + omp_get_default_device())); + } + omp_target_free(d_sum, omp_get_default_device()); + + free(h_array_h); + errors = N - sum; + if (!errors) + printf("Test passed sum=%d, N=%d\n", sum, N); + else + printf("Test failed sum=%d, N=%d\n", sum, N); + + return errors; } int main() { - int errors = 0; + int errors = 0; - errors += test_map_from(); - errors += test_map_tofrom(); - errors += test_map_to(); - errors += test_map_to_from(); - errors += test_map_alloc(); + errors += test_map_from(); + errors += test_map_tofrom(); + errors += test_map_to(); + errors += test_map_to_from(); + errors += test_map_alloc(); - return errors; + return errors; } Index: libomptarget/test/offloading/target_data/test_target_data_map_array_sections.c =================================================================== --- libomptarget/test/offloading/target_data/test_target_data_map_array_sections.c +++ libomptarget/test/offloading/target_data/test_target_data_map_array_sections.c @@ -3,9 +3,9 @@ // RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu // RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu +#include #include #include -#include #define N 1000 @@ -14,372 +14,388 @@ void init_3d(int a[N][2][2]); int test_lower_length_1d() { - // array sections of the form a[lower:length] - puts("test_lower_length_1d"); - // If a list item is an array section, it must specify contiguous storage. - - int errors = 0; - - // stack - int a1d [N]; - init_1d(a1d); - - #pragma omp target data map(from:a1d[1:N-2]) - { - #pragma omp target - if (!omp_is_initial_device()) - for (int i = 1; i < N-1; ++i) - a1d[i] = 1; - } - - // checking errors - for (int i = 0; i < N; ++i) - if (i == 0 || i == N-1) errors += a1d[i] == 0? 0 : 1; - else errors += a1d[i] == 1 ? 0 : 1; - - if(!errors) - printf("Test passed\n"); + // array sections of the form a[lower:length] + puts("test_lower_length_1d"); + // If a list item is an array section, it must specify contiguous storage. + + int errors = 0; + + // stack + int a1d[N]; + init_1d(a1d); + +#pragma omp target data map(from : a1d[1 : N - 2]) + { +#pragma omp target + if (!omp_is_initial_device()) + for (int i = 1; i < N - 1; ++i) + a1d[i] = 1; + } + + // checking errors + for (int i = 0; i < N; ++i) + if (i == 0 || i == N - 1) + errors += a1d[i] == 0 ? 0 : 1; else - printf("Test failed\n"); + errors += a1d[i] == 1 ? 0 : 1; + + if (!errors) + printf("Test passed\n"); + else + printf("Test failed\n"); - return errors; + return errors; } int test_lower_length_2d() { - // array sections of the form a[lower:length] - puts("test_lower_length_2d"); - // If a list item is an array section, it must specify contiguous storage. - - int errors = 0; - - // stack - int a2d [N][2]; - init_2d(a2d); - - #pragma omp target data map(from:a2d[1:N-2][0:2]) - { - #pragma omp target - if (!omp_is_initial_device()) - for (int i = 1; i < N-1; ++i) { - a2d[i][0] = 1; - a2d[i][1] = 1; - } - } - - // checking errors - for (int i = 0; i < N; ++i) - if (i == 0 || i == N-1) errors += a2d[i][0] == 0 && a2d[i][1] == 0? 0 : 1; - else errors += a2d[i][0] == 1 && a2d[i][1] == 1? 0 : 1; - - if(!errors) - printf("Test passed\n"); + // array sections of the form a[lower:length] + puts("test_lower_length_2d"); + // If a list item is an array section, it must specify contiguous storage. + + int errors = 0; + + // stack + int a2d[N][2]; + init_2d(a2d); + +#pragma omp target data map(from : a2d[1 : N - 2][0 : 2]) + { +#pragma omp target + if (!omp_is_initial_device()) + for (int i = 1; i < N - 1; ++i) { + a2d[i][0] = 1; + a2d[i][1] = 1; + } + } + + // checking errors + for (int i = 0; i < N; ++i) + if (i == 0 || i == N - 1) + errors += a2d[i][0] == 0 && a2d[i][1] == 0 ? 0 : 1; else - printf("Test failed\n"); + errors += a2d[i][0] == 1 && a2d[i][1] == 1 ? 0 : 1; + + if (!errors) + printf("Test passed\n"); + else + printf("Test failed\n"); - return errors; + return errors; } int test_lower_length_3d() { - // array sections of the form a[lower:length] - puts("test_lower_length_3d"); - // If a list item is an array section, it must specify contiguous storage. - - int errors = 0; - - // stack - int a3d [N][2][2]; - init_3d(a3d); - int a3d2 [N][2][2]; - init_3d(a3d2); - - #pragma omp target data map(from:a3d[1:N-2][0:2][0:2], a3d2[0:N][0:2][0:2]) - { - #pragma omp target - if (!omp_is_initial_device()) - for (int i = 0; i < N; ++i) - for (int j = 0; j < 2; ++j) { - if (i > 0 && i < N-1) { - a3d[i][j][0] = 1; - a3d[i][j][1] = 1; - } - a3d2[i][j][0] = 1; - a3d2[i][j][1] = 1; - } - } - - // checking errors - for (int i = 0; i < N; ++i) + // array sections of the form a[lower:length] + puts("test_lower_length_3d"); + // If a list item is an array section, it must specify contiguous storage. + + int errors = 0; + + // stack + int a3d[N][2][2]; + init_3d(a3d); + int a3d2[N][2][2]; + init_3d(a3d2); + +#pragma omp target data map( \ + from : a3d[1 : N - 2][0 : 2][0 : 2], a3d2[0 : N][0 : 2][0 : 2]) + { +#pragma omp target + if (!omp_is_initial_device()) + for (int i = 0; i < N; ++i) for (int j = 0; j < 2; ++j) { - // a3d - if (i == 0 || i == N-1) errors += a3d[i][j][0] == 0 && a3d[i][j][1] == 0? 0 : 1; - else errors += a3d[i][j][0] == 1 && a3d[i][j][1] == 1? 0 : 1; - // a3d2 - errors += a3d2[i][j][0] == 1 && a3d2[i][j][1] == 1? 0 : 1; + if (i > 0 && i < N - 1) { + a3d[i][j][0] = 1; + a3d[i][j][1] = 1; + } + a3d2[i][j][0] = 1; + a3d2[i][j][1] = 1; } + } + + // checking errors + for (int i = 0; i < N; ++i) + for (int j = 0; j < 2; ++j) { + // a3d + if (i == 0 || i == N - 1) + errors += a3d[i][j][0] == 0 && a3d[i][j][1] == 0 ? 0 : 1; + else + errors += a3d[i][j][0] == 1 && a3d[i][j][1] == 1 ? 0 : 1; + // a3d2 + errors += a3d2[i][j][0] == 1 && a3d2[i][j][1] == 1 ? 0 : 1; + } - if(!errors) - printf("Test passed\n"); - else - printf("Test failed\n"); + if (!errors) + printf("Test passed\n"); + else + printf("Test failed\n"); - return errors; + return errors; } int test_length_1d() { - // array sections of the form a[:length] - puts("test_length_1d"); - - int errors = 0; - - // stack - int a1d [N]; - init_1d(a1d); - - #pragma omp target data map(from:a1d[:N-2]) - { - #pragma omp target - if (!omp_is_initial_device()) - for (int i = 0; i < N-2; ++i) - a1d[i] = 1; - } - - // checking errors - for (int i = 0; i < N-2; ++i) - errors += a1d[i] == 1 ? 0 : 1; - // N-2 - errors += a1d[N-2] == 0? 0 : 1; - // N-1 - errors += a1d[N-1] == 0? 0 : 1; - - if(!errors) - printf("Test passed\n"); - else - printf("Test failed\n"); - - return errors; + // array sections of the form a[:length] + puts("test_length_1d"); + + int errors = 0; + + // stack + int a1d[N]; + init_1d(a1d); + +#pragma omp target data map(from : a1d[ : N - 2]) + { +#pragma omp target + if (!omp_is_initial_device()) + for (int i = 0; i < N - 2; ++i) + a1d[i] = 1; + } + + // checking errors + for (int i = 0; i < N - 2; ++i) + errors += a1d[i] == 1 ? 0 : 1; + // N-2 + errors += a1d[N - 2] == 0 ? 0 : 1; + // N-1 + errors += a1d[N - 1] == 0 ? 0 : 1; + + if (!errors) + printf("Test passed\n"); + else + printf("Test failed\n"); + + return errors; } int test_length_2d() { - // array sections of the form a[:length] - puts("test_length_2d"); - - int errors = 0; - - // stack - int a2d [N][2]; - init_2d(a2d); - - #pragma omp target data map(from:a2d[:N-2][:2]) - { - #pragma omp target - if (!omp_is_initial_device()) - for (int i = 0; i < N-2; ++i) { - a2d[i][0] = 1; - a2d[i][1] = 1; - } - } - - // checking errors - for (int i = 0; i < N-2; ++i) - errors += a2d[i][0] == 1 && a2d[i][1] == 1? 0 : 1; - errors += a2d[N-2][0] == 0 && a2d[N-2][1] == 0? 0 : 1; - errors += a2d[N-1][0] == 0 && a2d[N-1][1] == 0? 0 : 1; - - if(!errors) - printf("Test passed\n"); - else - printf("Test failed\n"); - - return errors; + // array sections of the form a[:length] + puts("test_length_2d"); + + int errors = 0; + + // stack + int a2d[N][2]; + init_2d(a2d); + +#pragma omp target data map(from : a2d[ : N - 2][ : 2]) + { +#pragma omp target + if (!omp_is_initial_device()) + for (int i = 0; i < N - 2; ++i) { + a2d[i][0] = 1; + a2d[i][1] = 1; + } + } + + // checking errors + for (int i = 0; i < N - 2; ++i) + errors += a2d[i][0] == 1 && a2d[i][1] == 1 ? 0 : 1; + errors += a2d[N - 2][0] == 0 && a2d[N - 2][1] == 0 ? 0 : 1; + errors += a2d[N - 1][0] == 0 && a2d[N - 1][1] == 0 ? 0 : 1; + + if (!errors) + printf("Test passed\n"); + else + printf("Test failed\n"); + + return errors; } int test_length_3d() { - // array sections of the form a[:length] - puts("test_length_3d"); - - int errors = 0; - - // stack - int a3d [N][2][2]; - init_3d(a3d); - int a3d2 [N][2][2]; - init_3d(a3d2); - - #pragma omp target data map(from:a3d[:N-2][:2][:2], a3d2[:N][:2][:2]) - { - #pragma omp target - if (!omp_is_initial_device()) - for (int i = 0; i < N; ++i) - for (int j = 0; j < 2; ++j) { - if (i < N-2) { - a3d[i][j][0] = 1; - a3d[i][j][1] = 1; - } - a3d2[i][j][0] = 1; - a3d2[i][j][1] = 1; - } - } - - // checking errors - for (int i = 0; i < N; ++i) + // array sections of the form a[:length] + puts("test_length_3d"); + + int errors = 0; + + // stack + int a3d[N][2][2]; + init_3d(a3d); + int a3d2[N][2][2]; + init_3d(a3d2); + +#pragma omp target data map( \ + from : a3d[ : N - 2][ : 2][ : 2], a3d2[ : N][ : 2][ : 2]) + { +#pragma omp target + if (!omp_is_initial_device()) + for (int i = 0; i < N; ++i) for (int j = 0; j < 2; ++j) { - if (i >= N-2) errors += a3d[i][j][0] == 0 && a3d[i][j][1] == 0? 0 : 1; - else errors += a3d[i][j][0] == 1 && a3d[i][j][1] == 1? 0 : 1; - // a3d2 - errors += a3d2[i][j][0] == 1 && a3d2[i][j][1] == 1? 0 : 1; + if (i < N - 2) { + a3d[i][j][0] = 1; + a3d[i][j][1] = 1; + } + a3d2[i][j][0] = 1; + a3d2[i][j][1] = 1; } + } + + // checking errors + for (int i = 0; i < N; ++i) + for (int j = 0; j < 2; ++j) { + if (i >= N - 2) + errors += a3d[i][j][0] == 0 && a3d[i][j][1] == 0 ? 0 : 1; + else + errors += a3d[i][j][0] == 1 && a3d[i][j][1] == 1 ? 0 : 1; + // a3d2 + errors += a3d2[i][j][0] == 1 && a3d2[i][j][1] == 1 ? 0 : 1; + } - if(!errors) - printf("Test passed\n"); - else - printf("Test failed\n"); + if (!errors) + printf("Test passed\n"); + else + printf("Test failed\n"); - return errors; + return errors; } int test_lower_1d() { - // array sections of the form a[lower:] - puts("test_lower_1d"); - - int errors = 0; - - // stack - int a1d [N]; - init_1d(a1d); - - #pragma omp target data map(from:a1d[1:]) - { - #pragma omp target - if (!omp_is_initial_device()) - for (int i = 1; i < N; ++i) - a1d[i] = 1; - } - - // checking errors - for (int i = 0; i < N; ++i) - if (i == 0) errors += a1d[i] == 0? 0 : 1; - else errors += a1d[i] == 1 ? 0 : 1; - - if(!errors) - printf("Test passed\n"); + // array sections of the form a[lower:] + puts("test_lower_1d"); + + int errors = 0; + + // stack + int a1d[N]; + init_1d(a1d); + +#pragma omp target data map(from : a1d[1 : ]) + { +#pragma omp target + if (!omp_is_initial_device()) + for (int i = 1; i < N; ++i) + a1d[i] = 1; + } + + // checking errors + for (int i = 0; i < N; ++i) + if (i == 0) + errors += a1d[i] == 0 ? 0 : 1; else - printf("Test failed\n"); + errors += a1d[i] == 1 ? 0 : 1; + + if (!errors) + printf("Test passed\n"); + else + printf("Test failed\n"); - return errors; + return errors; } int test_lower_2d() { - // array sections of the form a[lower:] - puts("test_lower_2d"); - - int errors = 0; - - // stack - int a2d [N][2]; - init_2d(a2d); - - #pragma omp target data map(from:a2d[1:][0:]) - { - #pragma omp target - if (!omp_is_initial_device()) - for (int i = 1; i < N; ++i) { - a2d[i][0] = 1; - a2d[i][1] = 1; - } - } - - // checking errors - for (int i = 0; i < N; ++i) - if (i == 0) errors += a2d[i][0] == 0 && a2d[i][1] == 0? 0 : 1; - else errors += a2d[i][0] == 1 && a2d[i][1] == 1? 0 : 1; - - if(!errors) - printf("Test passed\n"); + // array sections of the form a[lower:] + puts("test_lower_2d"); + + int errors = 0; + + // stack + int a2d[N][2]; + init_2d(a2d); + +#pragma omp target data map(from : a2d[1 : ][0 : ]) + { +#pragma omp target + if (!omp_is_initial_device()) + for (int i = 1; i < N; ++i) { + a2d[i][0] = 1; + a2d[i][1] = 1; + } + } + + // checking errors + for (int i = 0; i < N; ++i) + if (i == 0) + errors += a2d[i][0] == 0 && a2d[i][1] == 0 ? 0 : 1; else - printf("Test failed\n"); + errors += a2d[i][0] == 1 && a2d[i][1] == 1 ? 0 : 1; - return errors; + if (!errors) + printf("Test passed\n"); + else + printf("Test failed\n"); + + return errors; } int test_lower_3d() { - // array sections of the form a[lower:] - puts("test_lower_3d"); - - int errors = 0; - - // stack - int a3d [N][2][2]; - init_3d(a3d); - int a3d2 [N][2][2]; - init_3d(a3d2); - - #pragma omp target data map(from:a3d[1:][0:][0:], a3d2[0:][0:][0:]) - { - #pragma omp target - if (!omp_is_initial_device()) - for (int i = 0; i < N; ++i) - for (int j = 0; j < 2; ++j) { - if (i > 0) { - a3d[i][j][0] = 1; - a3d[i][j][1] = 1; - } - a3d2[i][j][0] = 1; - a3d2[i][j][1] = 1; - } - } - - // checking errors - for (int i = 0; i < N; ++i) + // array sections of the form a[lower:] + puts("test_lower_3d"); + + int errors = 0; + + // stack + int a3d[N][2][2]; + init_3d(a3d); + int a3d2[N][2][2]; + init_3d(a3d2); + +#pragma omp target data map(from : a3d[1 : ][0 : ][0 : ], \ + a3d2[0 : ][0 : ][0 : ]) + { +#pragma omp target + if (!omp_is_initial_device()) + for (int i = 0; i < N; ++i) for (int j = 0; j < 2; ++j) { - // a3d - if (i == 0) errors += a3d[i][j][0] == 0 && a3d[i][j][1] == 0? 0 : 1; - else errors += a3d[i][j][0] == 1 && a3d[i][j][1] == 1? 0 : 1; - // a3d2 - errors += a3d2[i][j][0] == 1 && a3d2[i][j][1] == 1? 0 : 1; + if (i > 0) { + a3d[i][j][0] = 1; + a3d[i][j][1] = 1; + } + a3d2[i][j][0] = 1; + a3d2[i][j][1] = 1; } + } + + // checking errors + for (int i = 0; i < N; ++i) + for (int j = 0; j < 2; ++j) { + // a3d + if (i == 0) + errors += a3d[i][j][0] == 0 && a3d[i][j][1] == 0 ? 0 : 1; + else + errors += a3d[i][j][0] == 1 && a3d[i][j][1] == 1 ? 0 : 1; + // a3d2 + errors += a3d2[i][j][0] == 1 && a3d2[i][j][1] == 1 ? 0 : 1; + } - if(!errors) - printf("Test passed\n"); - else - printf("Test failed\n"); + if (!errors) + printf("Test passed\n"); + else + printf("Test failed\n"); - return errors; + return errors; } int main() { - int errors = 0; + int errors = 0; - errors += test_lower_length_1d(); - errors += test_lower_length_2d(); - errors += test_lower_length_3d(); - errors += test_length_1d(); - errors += test_length_2d(); - errors += test_length_3d(); - errors += test_lower_1d(); - errors += test_lower_2d(); - errors += test_lower_3d(); + errors += test_lower_length_1d(); + errors += test_lower_length_2d(); + errors += test_lower_length_3d(); + errors += test_length_1d(); + errors += test_length_2d(); + errors += test_length_3d(); + errors += test_lower_1d(); + errors += test_lower_2d(); + errors += test_lower_3d(); - return errors; + return errors; } void init_1d(int a[N]) { - for (int i = 0; i < N; ++i) - a[i] = 0; + for (int i = 0; i < N; ++i) + a[i] = 0; } void init_2d(int a[N][2]) { - for (int i = 0; i < N; ++i) { - a[i][0] = 0; - a[i][1] = 0; - } + for (int i = 0; i < N; ++i) { + a[i][0] = 0; + a[i][1] = 0; + } } void init_3d(int a[N][2][2]) { - for (int i = 0; i < N; ++i) - for (int j = 0; j < 2; ++j) { - a[i][j][0] = 0; - a[i][j][1] = 0; - } + for (int i = 0; i < N; ++i) + for (int j = 0; j < 2; ++j) { + a[i][j][0] = 0; + a[i][j][1] = 0; + } } - Index: libomptarget/test/offloading/target_data/test_target_data_map_classes.cpp =================================================================== --- libomptarget/test/offloading/target_data/test_target_data_map_classes.cpp +++ libomptarget/test/offloading/target_data/test_target_data_map_classes.cpp @@ -11,74 +11,69 @@ #define N 1000 class A { - - public: - // breaks encapsulation - int *h_array; - int size; - int sum; - - A(int *array, const int s) : h_array(array), size(s) { - sum = 0; - } - - // TODO: Add virtual once supported - ~A() { - h_array = NULL; - } + +public: + // breaks encapsulation + int *h_array; + int size; + int sum; + + A(int *array, const int s) : h_array(array), size(s) { sum = 0; } + + // TODO: Add virtual once supported + ~A() { h_array = NULL; } }; int test_map_tofrom_class() { - cout << "test_map_tofrom_class" << endl; - - int sum = 0, errors = 0; - - int *array = new int[N]; - A* obj = new A(array, N); - - // mapping an object + array - #pragma omp target data map(from:array[0:N]) map(tofrom:obj[0:1]) - { - #pragma omp target - if (!omp_is_initial_device()) { - // assign device array ptr to device obj - int* tmp_h_array = obj->h_array; - obj->h_array = array; - int tmp = 0; - for (int i = 0; i < N; ++i) { - obj->h_array[i] = 1; - tmp += 1; - } - // swap array device ptr to host ptr - obj->h_array = tmp_h_array; - - obj->sum = tmp; - } + cout << "test_map_tofrom_class" << endl; + + int sum = 0, errors = 0; + + int *array = new int[N]; + A *obj = new A(array, N); + +// mapping an object + array +#pragma omp target data map(from : array[0 : N]) map(tofrom : obj[0 : 1]) + { +#pragma omp target + if (!omp_is_initial_device()) { + // assign device array ptr to device obj + int *tmp_h_array = obj->h_array; + obj->h_array = array; + int tmp = 0; + for (int i = 0; i < N; ++i) { + obj->h_array[i] = 1; + tmp += 1; + } + // swap array device ptr to host ptr + obj->h_array = tmp_h_array; + + obj->sum = tmp; } + } - // checking results - for(int i = 0; i < N; ++i) - sum += obj->h_array[i]; + // checking results + for (int i = 0; i < N; ++i) + sum += obj->h_array[i]; - errors = (N != sum) || (N != obj->sum); - if(!errors) - cout << "Test passed: sum=" << sum << ", N=" << N << endl; - else - cout << "Test failed: sum=" << sum << ", N=" << N << endl; - - delete obj; - delete[] array; + errors = (N != sum) || (N != obj->sum); + if (!errors) + cout << "Test passed: sum=" << sum << ", N=" << N << endl; + else + cout << "Test failed: sum=" << sum << ", N=" << N << endl; - return errors; + delete obj; + delete[] array; + + return errors; } int main() { - int errors = 0; + int errors = 0; - errors += test_map_tofrom_class(); + errors += test_map_tofrom_class(); - return errors; + return errors; } - Index: libomptarget/test/offloading/target_data/test_target_data_map_devices.c =================================================================== --- libomptarget/test/offloading/target_data/test_target_data_map_devices.c +++ libomptarget/test/offloading/target_data/test_target_data_map_devices.c @@ -3,166 +3,171 @@ // RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu // RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu +#include +#include #include #include -#include -#include #define N 1000 int test_map_set_default_dev() { - printf("test_map_set_default_dev\n"); - - // Get number of devices - int num_dev = omp_get_num_devices(); - printf("num_devices: %d\n", num_dev); - - int def_dev = omp_get_default_device(); - printf("initial device: %d\n", omp_get_initial_device()); - printf("default device: %d\n", def_dev); - - int sum[num_dev] , errors = 0; - int h_matrix[num_dev][N]; - - for (int dev = 0; dev < num_dev; ++dev) { - omp_set_default_device(dev); - #pragma omp target data map(from:h_matrix[dev][0:N]) - { - assert(dev == omp_get_default_device()); - #pragma omp target map(from:h_matrix[dev][0:N]) - if (!omp_is_initial_device()) - for (int i = 0; i < N; ++i) - h_matrix[dev][i] = dev; - } - } + printf("test_map_set_default_dev\n"); + + // Get number of devices + int num_dev = omp_get_num_devices(); + printf("num_devices: %d\n", num_dev); + + int def_dev = omp_get_default_device(); + printf("initial device: %d\n", omp_get_initial_device()); + printf("default device: %d\n", def_dev); - // checking results - errors = 0; - for (int dev = 0; dev < num_dev; ++dev) { - sum[dev] = h_matrix[dev][0]; - for(int i = 1; i < N; ++i) - sum[dev] += h_matrix[dev][i]; - errors |= (dev*N != sum[dev]); + int sum[num_dev], errors = 0; + int h_matrix[num_dev][N]; + + for (int dev = 0; dev < num_dev; ++dev) { + omp_set_default_device(dev); +#pragma omp target data map(from : h_matrix[dev][0 : N]) + { + assert(dev == omp_get_default_device()); +#pragma omp target map(from : h_matrix[dev][0 : N]) + if (!omp_is_initial_device()) + for (int i = 0; i < N; ++i) + h_matrix[dev][i] = dev; } + } + + // checking results + errors = 0; + for (int dev = 0; dev < num_dev; ++dev) { + sum[dev] = h_matrix[dev][0]; + for (int i = 1; i < N; ++i) + sum[dev] += h_matrix[dev][i]; + errors |= (dev * N != sum[dev]); + } - if(!errors) - printf("Test passed: num_devices = %d\n", num_dev); - else - printf("Test failed: num_devices = %d\n", num_dev); + if (!errors) + printf("Test passed: num_devices = %d\n", num_dev); + else + printf("Test failed: num_devices = %d\n", num_dev); - omp_set_default_device(def_dev); + omp_set_default_device(def_dev); - return errors; + return errors; } int test_map_device() { - printf("test_map_device\n"); - - // Get number of devices - int num_dev = omp_get_num_devices(); - printf("num_devices: %d\n", num_dev); - - printf("initial device: %d\n", omp_get_initial_device()); - printf("default device: %d\n", omp_get_default_device()); - - int sum[num_dev] , errors = 0; - int h_matrix[num_dev][N]; - - for (int dev = 0; dev < num_dev; ++dev) { - #pragma omp target data map(from:h_matrix[dev][0:N]) device(dev) - { - #pragma omp target map(from:h_matrix[dev][0:N]) device(dev) - if (!omp_is_initial_device()) - for (int i = 0; i < N; ++i) - h_matrix[dev][i] = dev; - } - } + printf("test_map_device\n"); - // checking results - errors = 0; - for (int dev = 0; dev < num_dev; ++dev) { - sum[dev] = h_matrix[dev][0]; - for(int i = 1; i < N; ++i) - sum[dev] += h_matrix[dev][i]; - errors |= (dev*N != sum[dev]); - } + // Get number of devices + int num_dev = omp_get_num_devices(); + printf("num_devices: %d\n", num_dev); - if(!errors) - printf("Test passed: num_devices = %d\n", num_dev); - else - printf("Test failed: num_devices = %d\n", num_dev); + printf("initial device: %d\n", omp_get_initial_device()); + printf("default device: %d\n", omp_get_default_device()); - return errors; + int sum[num_dev], errors = 0; + int h_matrix[num_dev][N]; + + for (int dev = 0; dev < num_dev; ++dev) { +#pragma omp target data map(from : h_matrix[dev][0 : N]) device(dev) + { +#pragma omp target map(from : h_matrix[dev][0 : N]) device(dev) + if (!omp_is_initial_device()) + for (int i = 0; i < N; ++i) + h_matrix[dev][i] = dev; + } + } + + // checking results + errors = 0; + for (int dev = 0; dev < num_dev; ++dev) { + sum[dev] = h_matrix[dev][0]; + for (int i = 1; i < N; ++i) + sum[dev] += h_matrix[dev][i]; + errors |= (dev * N != sum[dev]); + } + + if (!errors) + printf("Test passed: num_devices = %d\n", num_dev); + else + printf("Test failed: num_devices = %d\n", num_dev); + + return errors; } int test_more_devices() { - printf("test_more_devices\n"); - - // Get number of devices - int real_num_dev = omp_get_num_devices(); - int num_dev = real_num_dev + 2; - printf("num_devices: %d, real_num_devices: %d\n", num_dev, real_num_dev); - - int def_dev = omp_get_default_device(); - printf("initial device: %d\n", omp_get_initial_device()); - printf("default device: %d\n", def_dev); - - int sum[num_dev], errors = 0, extra_on_host = 0; - int h_matrix[num_dev][N]; - - // omp_set_default_device is implementation dependent - for (int dev = 0; dev < num_dev; ++dev) { - omp_set_default_device(dev); - int dev_tmp = omp_get_default_device(); - #pragma omp target data map(from:h_matrix[dev_tmp][0:N]) - { - assert(dev == dev_tmp); - #pragma omp target map(from:h_matrix[dev_tmp][0:N]) map(tofrom:extra_on_host) - { - if (!omp_is_initial_device()) - for (int i = 0; i < N; ++i) - h_matrix[dev_tmp][i] = dev; - else if (dev >= real_num_dev) { - // allows implementations that map extra devices to the host - extra_on_host = 1; - for (int i = 0; i < N; ++i) - h_matrix[dev_tmp][i] = dev; - } - } + printf("test_more_devices\n"); + + // Get number of devices + int real_num_dev = omp_get_num_devices(); + int num_dev = real_num_dev + 2; + printf("num_devices: %d, real_num_devices: %d\n", num_dev, real_num_dev); + + int def_dev = omp_get_default_device(); + printf("initial device: %d\n", omp_get_initial_device()); + printf("default device: %d\n", def_dev); + + int sum[num_dev], errors = 0, extra_on_host = 0; + int h_matrix[num_dev][N]; + + // omp_set_default_device is implementation dependent + for (int dev = 0; dev < num_dev; ++dev) { + omp_set_default_device(dev); + int dev_tmp = omp_get_default_device(); +#pragma omp target data map(from : h_matrix[dev_tmp][0 : N]) + { + assert(dev == dev_tmp); +#pragma omp target map(from : h_matrix[dev_tmp][0 : N]) \ + map(tofrom : extra_on_host) + { + if (!omp_is_initial_device()) + for (int i = 0; i < N; ++i) + h_matrix[dev_tmp][i] = dev; + else if (dev >= real_num_dev) { + // allows implementations that map extra devices to the host + extra_on_host = 1; + for (int i = 0; i < N; ++i) + h_matrix[dev_tmp][i] = dev; } + } } - - // checking results - errors = 0; - for (int dev = 0; dev < num_dev; ++dev) { - sum[dev] = h_matrix[dev][0]; - for(int i = 1; i < N; ++i) - sum[dev] += h_matrix[dev][i]; - errors |= (dev*N != sum[dev]); - if (errors) - printf("Error at dev=%d\n", dev); - } - - if(!errors) - printf("Test passed: num_devices = %d, real_num_dev = %d, extra_on_host = %d\n", num_dev, real_num_dev, extra_on_host); - else - printf("Test failed: num_devices = %d, real_num_dev = %d, extra_on_host = %d\n", num_dev, real_num_dev, extra_on_host); - - omp_set_default_device(def_dev); - - return errors; + } + + // checking results + errors = 0; + for (int dev = 0; dev < num_dev; ++dev) { + sum[dev] = h_matrix[dev][0]; + for (int i = 1; i < N; ++i) + sum[dev] += h_matrix[dev][i]; + errors |= (dev * N != sum[dev]); + if (errors) + printf("Error at dev=%d\n", dev); + } + + if (!errors) + printf("Test passed: num_devices = %d, real_num_dev = %d, extra_on_host = " + "%d\n", + num_dev, real_num_dev, extra_on_host); + else + printf("Test failed: num_devices = %d, real_num_dev = %d, extra_on_host = " + "%d\n", + num_dev, real_num_dev, extra_on_host); + + omp_set_default_device(def_dev); + + return errors; } int main() { - int errors = 0; + int errors = 0; - errors += test_map_set_default_dev(); - errors += test_map_device(); - errors += test_more_devices(); + errors += test_map_set_default_dev(); + errors += test_map_device(); + errors += test_more_devices(); - return errors; + return errors; } Index: libomptarget/test/offloading/target_data/test_target_data_use_device_ptr.c =================================================================== --- libomptarget/test/offloading/target_data/test_target_data_use_device_ptr.c +++ libomptarget/test/offloading/target_data/test_target_data_use_device_ptr.c @@ -3,56 +3,50 @@ // RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu // RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu - -#include #include +#include /*Test for OpenMP 4.5 target data with use_device_ptr*/ -int main () -{ +int main() { int success = 1; int len = 10000; int *array_device = NULL; int *array_host = NULL; - int i,j; - - array_device = (int*) malloc(len*sizeof(int)); - array_host = (int*) malloc(len*sizeof(int)); + int i, j; + array_device = (int *)malloc(len * sizeof(int)); + array_host = (int *)malloc(len * sizeof(int)); - for(i=0; i < len; i++) - array_host[i] = i; + for (i = 0; i < len; i++) + array_host[i] = i; - #pragma omp target data map(tofrom: array_device[0:len]) use_device_ptr(array_device) +#pragma omp target data map(tofrom : array_device[0 : len]) \ + use_device_ptr(array_device) { - #pragma omp target is_device_ptr(array_device) map(tofrom: array_host[0:len]) +#pragma omp target is_device_ptr(array_device) map(tofrom : array_host[0 : len]) { - if(!omp_is_initial_device()){ /*execute computation only on device*/ - for(i=0; i < len; i++) - { - *(array_device+i) = i; - array_host[i] += *(array_device+i); - }/*end-for*/ - }/*end-if*/ + if (!omp_is_initial_device()) { /*execute computation only on device*/ + for (i = 0; i < len; i++) { + *(array_device + i) = i; + array_host[i] += *(array_device + i); + } /*end-for*/ + } /*end-if*/ } } - for(i=0; i < len; i++) { - if(array_host[i]!= 2*i) - success = 0; + for (i = 0; i < len; i++) { + if (array_host[i] != 2 * i) + success = 0; } - + free(array_device); free(array_host); - if(success) { - printf("\nSuccess\n"); - } - else { + if (success) { + printf("\nSuccess\n"); + } else { printf("\nFailed\n"); } return success; - } -