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
|
/* { dg-do run } */
/* { dg-require-effective-target offload_device } */
#include <stdlib.h>
#include <omp.h>
#define EPS 0.000001
#define THRESHOLD 1000
const int MAX = 1800;
void check (float *a, float *b, int N)
{
int i;
for (i = 0; i < N; i++)
if (a[i] - b[i] > EPS || b[i] - a[i] > EPS)
abort ();
}
void init (float *a1, float *a2, int N)
{
float s = -1;
int i;
for (i = 0; i < N; i++)
{
a1[i] = s;
a2[i] = i;
s = -s;
}
}
void init_again (float *a1, float *a2, int N)
{
float s = -1;
int i;
for (i = 0; i < N; i++)
{
a1[i] = s * 10;
a2[i] = i;
s = -s;
}
}
void vec_mult_ref (float *p, float *v1, float *v2, int N)
{
int i;
init (v1, v2, N);
for (i = 0; i < N; i++)
p[i] = v1[i] * v2[i];
init_again (v1, v2, N);
for (i = 0; i < N; i++)
p[i] = p[i] + (v1[i] * v2[i]);
}
void vec_mult (float *p, float *v1, float *v2, int N)
{
int i;
init (v1, v2, N);
#pragma omp target data if(N > THRESHOLD) map(from: p[0:N])
{
#pragma omp target if (N > THRESHOLD) map(to: v1[:N], v2[:N])
{
if (omp_is_initial_device ())
abort;
#pragma omp parallel for
for (i = 0; i < N; i++)
p[i] = v1[i] * v2[i];
}
init_again (v1, v2, N);
#pragma omp target if (N > THRESHOLD) map(to: v1[:N], v2[:N])
{
if (omp_is_initial_device ())
abort ();
#pragma omp parallel for
for (i = 0; i < N; i++)
p[i] = p[i] + (v1[i] * v2[i]);
}
}
}
int main ()
{
float *p1 = (float *) malloc (MAX * sizeof (float));
float *p2 = (float *) malloc (MAX * sizeof (float));
float *v1 = (float *) malloc (MAX * sizeof (float));
float *v2 = (float *) malloc (MAX * sizeof (float));
vec_mult_ref (p1, v1, v2, MAX);
vec_mult (p2, v1, v2, MAX);
check (p1, p2, MAX);
free (p1);
free (p2);
free (v1);
free (v2);
return 0;
}
|
