target-29.c 4.1 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112
  1. #include <omp.h>
  2. #include <stdlib.h>
  3. struct S { char p[64]; int a; int b[2]; long c[4]; int *d; char q[64]; };
  4. __attribute__((noinline, noclone)) void
  5. foo (struct S s)
  6. {
  7. int d = omp_get_default_device ();
  8. int id = omp_get_initial_device ();
  9. int sep = 1;
  10. if (d < 0 || d >= omp_get_num_devices ())
  11. d = id;
  12. int err;
  13. #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) map(to: sep) map(from: err)
  14. {
  15. err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
  16. err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
  17. s.a = 35; s.b[0] = 36; s.b[1] = 37;
  18. s.c[1] = 38; s.c[2] = 39; s.d[-2] = 40; s.d[-1] = 41; s.d[0] = 42;
  19. sep = 0;
  20. }
  21. if (err) abort ();
  22. err = s.a != 35 || s.b[0] != 36 || s.b[1] != 37;
  23. err |= s.c[1] != 38 || s.c[2] != 39 || s.d[-2] != 40 || s.d[-1] != 41 || s.d[0] != 42;
  24. if (err) abort ();
  25. s.a = 50; s.b[0] = 49; s.b[1] = 48;
  26. s.c[1] = 47; s.c[2] = 46; s.d[-2] = 45; s.d[-1] = 44; s.d[0] = 43;
  27. if (sep
  28. && (omp_target_is_present (&s.a, d)
  29. || omp_target_is_present (s.b, d)
  30. || omp_target_is_present (&s.c[1], d)
  31. || omp_target_is_present (s.d, d)
  32. || omp_target_is_present (&s.d[-2], d)))
  33. abort ();
  34. #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
  35. {
  36. if (!omp_target_is_present (&s.a, d)
  37. || !omp_target_is_present (s.b, d)
  38. || !omp_target_is_present (&s.c[1], d)
  39. || !omp_target_is_present (s.d, d)
  40. || !omp_target_is_present (&s.d[-2], d))
  41. abort ();
  42. #pragma omp target update to(s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
  43. #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) map(from: err)
  44. {
  45. err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
  46. err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
  47. s.a = 17; s.b[0] = 18; s.b[1] = 19;
  48. s.c[1] = 20; s.c[2] = 21; s.d[-2] = 22; s.d[-1] = 23; s.d[0] = 24;
  49. }
  50. #pragma omp target update from(s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
  51. }
  52. if (sep
  53. && (omp_target_is_present (&s.a, d)
  54. || omp_target_is_present (s.b, d)
  55. || omp_target_is_present (&s.c[1], d)
  56. || omp_target_is_present (s.d, d)
  57. || omp_target_is_present (&s.d[-2], d)))
  58. abort ();
  59. if (err) abort ();
  60. err = s.a != 17 || s.b[0] != 18 || s.b[1] != 19;
  61. err |= s.c[1] != 20 || s.c[2] != 21 || s.d[-2] != 22 || s.d[-1] != 23 || s.d[0] != 24;
  62. if (err) abort ();
  63. s.a = 33; s.b[0] = 34; s.b[1] = 35;
  64. s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
  65. #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
  66. if (!omp_target_is_present (&s.a, d)
  67. || !omp_target_is_present (s.b, d)
  68. || !omp_target_is_present (&s.c[1], d)
  69. || !omp_target_is_present (s.d, d)
  70. || !omp_target_is_present (&s.d[-2], d))
  71. abort ();
  72. #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
  73. #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) map(from: err)
  74. {
  75. err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
  76. err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
  77. s.a = 49; s.b[0] = 48; s.b[1] = 47;
  78. s.c[1] = 46; s.c[2] = 45; s.d[-2] = 44; s.d[-1] = 43; s.d[0] = 42;
  79. }
  80. #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
  81. if (!omp_target_is_present (&s.a, d)
  82. || !omp_target_is_present (s.b, d)
  83. || !omp_target_is_present (&s.c[1], d)
  84. || !omp_target_is_present (s.d, d)
  85. || !omp_target_is_present (&s.d[-2], d))
  86. abort ();
  87. #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
  88. if (sep
  89. && (omp_target_is_present (&s.a, d)
  90. || omp_target_is_present (s.b, d)
  91. || omp_target_is_present (&s.c[1], d)
  92. || omp_target_is_present (s.d, d)
  93. || omp_target_is_present (&s.d[-2], d)))
  94. abort ();
  95. if (err) abort ();
  96. err = s.a != 49 || s.b[0] != 48 || s.b[1] != 47;
  97. err |= s.c[1] != 46 || s.c[2] != 45 || s.d[-2] != 44 || s.d[-1] != 43 || s.d[0] != 42;
  98. if (err) abort ();
  99. }
  100. int
  101. main ()
  102. {
  103. int d[3] = { 18, 19, 20 };
  104. struct S s = { {}, 11, { 12, 13 }, { 14, 15, 16, 17 }, d + 2, {} };
  105. foo (s);
  106. return 0;
  107. }