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
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
| #include "mini_blas.h"
|
| void axpy_cpu(int N, float ALPHA, float *X, int INCX, float *Y, int INCY)
| {
| int i;
| for(i = 0; i < N; ++i) Y[i*INCY] += ALPHA*X[i*INCX];
| }
|
| void scal_cpu(int N, float ALPHA, float *X, int INCX)
| {
| int i;
| for(i = 0; i < N; ++i) X[i*INCX] *= ALPHA;
| }
|
| void copy_cpu(int N, float *X, int INCX, float *Y, int INCY)
| {
| int i;
| for(i = 0; i < N; ++i) Y[i*INCY] = X[i*INCX];
| }
|
| float dot_cpu(int N, float *X, int INCX, float *Y, int INCY)
| {
| int i;
| float dot = 0;
| for(i = 0; i < N; ++i) dot += X[i*INCX] * Y[i*INCY];
| return dot;
| }
|
| #ifdef GPU
| #include "opencl.h"
|
| cl_kernel get_axpy_kernel()
| {
| static int init = 0;
| static cl_kernel kernel;
| if(!init){
| kernel = get_kernel("src/axpy.cl", "axpy", 0);
| init = 1;
| }
| return kernel;
| }
|
| cl_kernel get_copy_kernel()
| {
| static int init = 0;
| static cl_kernel kernel;
| if(!init){
| kernel = get_kernel("src/axpy.cl", "copy", 0);
| init = 1;
| }
| return kernel;
| }
|
| cl_kernel get_scal_kernel()
| {
| static int init = 0;
| static cl_kernel kernel;
| if(!init){
| kernel = get_kernel("src/axpy.cl", "scal", 0);
| init = 1;
| }
| return kernel;
| }
|
|
| void axpy_ongpu(int N, float ALPHA, cl_mem X, int INCX, cl_mem Y, int INCY)
| {
| axpy_ongpu_offset(N,ALPHA,X,0,INCX,Y,0,INCY);
| }
|
| void axpy_ongpu_offset(int N, float ALPHA, cl_mem X, int OFFX, int INCX, cl_mem Y, int OFFY, int INCY)
| {
| cl_setup();
| cl_kernel kernel = get_axpy_kernel();
| cl_command_queue queue = cl.queue;
|
| cl_uint i = 0;
| cl.error = clSetKernelArg(kernel, i++, sizeof(N), (void*) &N);
| cl.error = clSetKernelArg(kernel, i++, sizeof(ALPHA), (void*) &ALPHA);
| cl.error = clSetKernelArg(kernel, i++, sizeof(X), (void*) &X);
| cl.error = clSetKernelArg(kernel, i++, sizeof(OFFX), (void*) &OFFX);
| cl.error = clSetKernelArg(kernel, i++, sizeof(INCX), (void*) &INCX);
| cl.error = clSetKernelArg(kernel, i++, sizeof(Y), (void*) &Y);
| cl.error = clSetKernelArg(kernel, i++, sizeof(OFFY), (void*) &OFFY);
| cl.error = clSetKernelArg(kernel, i++, sizeof(INCY), (void*) &INCY);
| check_error(cl);
|
| const size_t global_size[] = {N};
|
| clEnqueueNDRangeKernel(queue, kernel, 1, 0, global_size, 0, 0, 0, 0);
| check_error(cl);
|
| }
| void copy_ongpu(int N, cl_mem X, int INCX, cl_mem Y, int INCY)
| {
| copy_ongpu_offset(N,X,0,INCX,Y,0,INCY);
| }
| void copy_ongpu_offset(int N, cl_mem X, int OFFX, int INCX, cl_mem Y, int OFFY, int INCY)
| {
| cl_setup();
| cl_kernel kernel = get_copy_kernel();
| cl_command_queue queue = cl.queue;
|
| cl_uint i = 0;
| cl.error = clSetKernelArg(kernel, i++, sizeof(N), (void*) &N);
| cl.error = clSetKernelArg(kernel, i++, sizeof(X), (void*) &X);
| cl.error = clSetKernelArg(kernel, i++, sizeof(OFFX), (void*) &OFFX);
| cl.error = clSetKernelArg(kernel, i++, sizeof(INCX), (void*) &INCX);
| cl.error = clSetKernelArg(kernel, i++, sizeof(Y), (void*) &Y);
| cl.error = clSetKernelArg(kernel, i++, sizeof(OFFY), (void*) &OFFY);
| cl.error = clSetKernelArg(kernel, i++, sizeof(INCY), (void*) &INCY);
| check_error(cl);
|
| const size_t global_size[] = {N};
|
| clEnqueueNDRangeKernel(queue, kernel, 1, 0, global_size, 0, 0, 0, 0);
| check_error(cl);
| }
| void scal_ongpu(int N, float ALPHA, cl_mem X, int INCX)
| {
| cl_setup();
| cl_kernel kernel = get_scal_kernel();
| cl_command_queue queue = cl.queue;
|
| cl_uint i = 0;
| cl.error = clSetKernelArg(kernel, i++, sizeof(N), (void*) &N);
| cl.error = clSetKernelArg(kernel, i++, sizeof(ALPHA), (void*) &ALPHA);
| cl.error = clSetKernelArg(kernel, i++, sizeof(X), (void*) &X);
| cl.error = clSetKernelArg(kernel, i++, sizeof(INCX), (void*) &INCX);
| check_error(cl);
|
| const size_t global_size[] = {N};
|
| clEnqueueNDRangeKernel(queue, kernel, 1, 0, global_size, 0, 0, 0, 0);
| check_error(cl);
| }
| #endif
|
|
