libgomp/testsuite/libgomp.c-c++-common/memcpyxd-bias-1.c

   1 /* { dg-do run } */
   2
   3 #include <stdlib.h>
   4 #include <stdint.h>
   5 #include <assert.h>
   6 #include <omp.h>
   7
   8 /* Say this is N rows and M columns.  */
   9 #define N 1024
  10 #define M 2048
  11
  12 #define row_offset 256
  13 #define row_length 512
  14 #define col_offset 128
  15 #define col_length 384
  16
  17 int
  18 main ()
  19 {
  20   int *arr2d = (int *) calloc (N * M, sizeof (int));
  21   uintptr_t dstptr;
  22   int hostdev = omp_get_initial_device ();
  23   int targdev;
  24
  25 #pragma omp target enter data map(to: arr2d[col_offset*M:col_length*M])
  26
  27 #pragma omp target map(from: targdev, dstptr) \
  28                    map(present, tofrom: arr2d[col_offset*M:col_length*M])
  29   {
  30     for (int j = col_offset; j < col_offset + col_length; j++)
  31       for (int i = row_offset; i < row_offset + row_length; i++)
  32         arr2d[j * M + i]++;
  33     targdev = omp_get_device_num ();
  34     dstptr = (uintptr_t) arr2d;
  35   }
  36
  37   /* Copy rectangular block back to the host.  */
  38   {
  39     size_t volume[2] = { col_length, row_length };
  40     size_t offsets[2] = { col_offset, row_offset };
  41     size_t dimensions[2] = { N, M };
  42     omp_target_memcpy_rect ((void *) arr2d, (const void *) dstptr,
  43                             sizeof (int), 2, &volume[0], &offsets[0],
  44                             &offsets[0], &dimensions[0], &dimensions[0],
  45                             hostdev, targdev);
  46   }
  47
  48 #pragma omp target exit data map(release: arr2d[col_offset*M:col_length*M])
  49
  50   for (int j = 0; j < N; j++)
  51     for (int i = 0; i < M; i++)
  52       if (i >= row_offset && i < row_offset + row_length
  53           && j >= col_offset && j < col_offset + col_length)
  54         assert (arr2d[j * M + i] == 1);
  55       else
  56         assert (arr2d[j * M + i] == 0);
  57
  58   free (arr2d);
  59
  60   return 0;
  61 }