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
|
/* { dg-do run } */
/* { dg-require-effective-target offload_device } */
#include <omp.h>
#include <stdlib.h>
#define EPS 0.00001
#define N 10000
#pragma omp declare target
void init (float *a, float *b, int n)
{
int i;
for (i = 0; i < n; i++)
{
a[i] = 0.1 * i;
b[i] = 0.01 * i * i;
}
}
#pragma omp end declare target
void vec_mult_ref (float *p, float *v1, float *v2, int n)
{
int i;
v1 = (float *) malloc (n * sizeof (float));
v2 = (float *) malloc (n * sizeof (float));
init (v1, v2, n);
for (i = 0; i < n; i++)
p[i] = v1[i] * v2[i];
free (v1);
free (v2);
}
void vec_mult (float *p, float *v1, float *v2, int n)
{
int i;
#pragma omp task shared(v1, v2) depend(out: v1, v2)
#pragma omp target map(v1, v2)
{
if (omp_is_initial_device ())
abort ();
v1 = (float *) malloc (n * sizeof (float));
v2 = (float *) malloc (n * sizeof (float));
init (v1, v2, n);
}
#pragma omp task shared(v1, v2) depend(in: v1, v2)
#pragma omp target map(to: v1, v2) map(from: p[0:n])
{
if (omp_is_initial_device ())
abort ();
#pragma omp parallel for
for (i = 0; i < n; i++)
p[i] = v1[i] * v2[i];
free (v1);
free (v2);
}
}
void check (float *a, float *b, int n)
{
int i;
for (i = 0 ; i < n ; i++)
{
float err = (a[i] == 0.0) ? b[i] : (b[i] - a[i]) / a[i];
if (((err > 0) ? err : -err) > EPS)
abort ();
}
}
int main ()
{
float *p1 = (float *) malloc (N * sizeof (float));
float *p2 = (float *) malloc (N * sizeof (float));
float *v1, *v2;
vec_mult_ref (p1, v1, v2, N);
vec_mult (p2, v1, v2, N);
check (p1, p2, N);
free (p1);
free (p2);
return 0;
}
|
