File: target_array_extension.c

package info (click to toggle)
llvm-toolchain-16 1%3A16.0.6-15~deb12u1
  • links: PTS, VCS
  • area: main
  • in suites: bookworm
  • size: 1,634,792 kB
  • sloc: cpp: 6,179,261; ansic: 1,216,205; asm: 741,319; python: 196,614; objc: 75,325; f90: 49,640; lisp: 32,396; pascal: 12,286; sh: 9,394; perl: 7,442; ml: 5,494; awk: 3,523; makefile: 2,723; javascript: 1,206; xml: 886; fortran: 581; cs: 573
file content (89 lines) | stat: -rw-r--r-- 2,906 bytes parent folder | download | duplicates (2)
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
// --------------------------------------------------
// Check extends before
// --------------------------------------------------

// RUN: %libomptarget-compile-generic \
// RUN:   -fopenmp-version=51 -DEXTENDS=BEFORE
// RUN: %libomptarget-run-fail-generic 2>&1 \
// RUN: | %fcheck-generic

// --------------------------------------------------
// Check extends after
// --------------------------------------------------

// RUN: %libomptarget-compile-generic \
// RUN:   -fopenmp-version=51 -DEXTENDS=AFTER
// RUN: %libomptarget-run-fail-generic 2>&1 \
// RUN: | %fcheck-generic

// END.

#include <stdio.h>

#define BEFORE 0
#define AFTER 1

#define SIZE 100

#if EXTENDS == BEFORE
#define SMALL_BEG (SIZE - 2)
#define SMALL_END SIZE
#define LARGE_BEG 0
#define LARGE_END SIZE
#elif EXTENDS == AFTER
#define SMALL_BEG 0
#define SMALL_END 2
#define LARGE_BEG 0
#define LARGE_END SIZE
#else
#error EXTENDS undefined
#endif

#define SMALL_SIZE (SMALL_END - SMALL_BEG)
#define LARGE_SIZE (LARGE_END - LARGE_BEG)

#define SMALL                                                                  \
  SMALL_BEG:                                                                   \
  SMALL_SIZE
#define LARGE                                                                  \
  LARGE_BEG:                                                                   \
  LARGE_SIZE

int main() {
  int arr[SIZE];

  // CHECK: addr=0x[[#%x,SMALL_ADDR:]], size=[[#%u,SMALL_BYTES:]]
  fprintf(stderr, "addr=%p, size=%ld\n", &arr[SMALL_BEG],
          SMALL_SIZE * sizeof arr[0]);

  // CHECK: addr=0x[[#%x,LARGE_ADDR:]], size=[[#%u,LARGE_BYTES:]]
  fprintf(stderr, "addr=%p, size=%ld\n", &arr[LARGE_BEG],
          LARGE_SIZE * sizeof arr[0]);

  // CHECK-NOT: Libomptarget
#pragma omp target data map(alloc : arr[LARGE])
  {
#pragma omp target map(present, tofrom : arr[SMALL])
    ;
  }

  // CHECK: arr is present
  fprintf(stderr, "arr is present\n");

  // CHECK: Libomptarget message: explicit extension not allowed: host address specified is 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes), but device allocation maps to host at 0x{{0*}}[[#SMALL_ADDR]] ([[#SMALL_BYTES]] bytes)
  // CHECK: Libomptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes)
  // CHECK: Libomptarget error: Call to getTargetPointer returned null pointer ('present' map type modifier).
  // CHECK: Libomptarget error: Call to targetDataBegin failed, abort target.
  // CHECK: Libomptarget error: Failed to process data before launching the kernel.
  // CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
#pragma omp target data map(alloc : arr[SMALL])
  {
#pragma omp target map(present, tofrom : arr[LARGE])
    ;
  }

  // CHECK-NOT: arr is present
  fprintf(stderr, "arr is present\n");

  return 0;
}