File: shared_update.c

package info (click to toggle)
llvm-toolchain-15 1%3A15.0.6-4
  • links: PTS, VCS
  • area: main
  • in suites: bookworm
  • size: 1,554,644 kB
  • sloc: cpp: 5,922,452; ansic: 1,012,136; asm: 674,362; python: 191,568; objc: 73,855; f90: 42,327; lisp: 31,913; pascal: 11,973; javascript: 10,144; sh: 9,421; perl: 7,447; ml: 5,527; awk: 3,523; makefile: 2,520; xml: 885; cs: 573; fortran: 567
file content (116 lines) | stat: -rw-r--r-- 2,805 bytes parent folder | download
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
// RUN: %libomptarget-compile-run-and-check-generic

// REQUIRES: unified_shared_memory

// amdgpu runtime crash
// UNSUPPORTED: amdgcn-amd-amdhsa

#include <stdio.h>
#include <omp.h>

// ---------------------------------------------------------------------------
// Various definitions copied from OpenMP RTL

extern void __tgt_register_requires(int64_t);

// End of definitions copied from OpenMP RTL.
// ---------------------------------------------------------------------------

#pragma omp requires unified_shared_memory

#define N 1024

int main(int argc, char *argv[]) {
  int fails;
  void *host_alloc, *device_alloc;
  void *host_data, *device_data;
  int *alloc = (int *)malloc(N * sizeof(int));
  int data[N];

  // Manual registration of requires flags for Clang versions
  // that do not support requires.
  __tgt_register_requires(8);

  for (int i = 0; i < N; ++i) {
    alloc[i] = 10;
    data[i] = 1;
  }

  host_data = &data[0];
  host_alloc = &alloc[0];

// implicit mapping of data
#pragma omp target map(tofrom : device_data, device_alloc)
  {
    device_data = &data[0];
    device_alloc = &alloc[0];

    for (int i = 0; i < N; i++) {
      alloc[i] += 1;
      data[i] += 1;
    }
  }

  // CHECK: Address of alloc on device matches host address.
  if (device_alloc == host_alloc)
    printf("Address of alloc on device matches host address.\n");

  // CHECK: Address of data on device matches host address.
  if (device_data == host_data)
    printf("Address of data on device matches host address.\n");

  // On the host, check that the arrays have been updated.
  // CHECK: Alloc device values updated: Succeeded
  fails = 0;
  for (int i = 0; i < N; i++) {
    if (alloc[i] != 11)
      fails++;
  }
  printf("Alloc device values updated: %s\n",
         (fails == 0) ? "Succeeded" : "Failed");

  // CHECK: Data device values updated: Succeeded
  fails = 0;
  for (int i = 0; i < N; i++) {
    if (data[i] != 2)
      fails++;
  }
  printf("Data device values updated: %s\n",
         (fails == 0) ? "Succeeded" : "Failed");

  //
  // Test that updates on the host snd on the device are both visible.
  //

  // Update on the host.
  for (int i = 0; i < N; ++i) {
    alloc[i] += 1;
    data[i] += 1;
  }

#pragma omp target
  {
    // CHECK: Alloc host values updated: Succeeded
    fails = 0;
    for (int i = 0; i < N; i++) {
      if (alloc[i] != 12)
        fails++;
    }
    printf("Alloc host values updated: %s\n",
           (fails == 0) ? "Succeeded" : "Failed");
    // CHECK: Data host values updated: Succeeded
    fails = 0;
    for (int i = 0; i < N; i++) {
      if (data[i] != 3)
        fails++;
    }
    printf("Data host values updated: %s\n",
           (fails == 0) ? "Succeeded" : "Failed");
  }

  free(alloc);

  printf("Done!\n");

  return 0;
}