/perf/kseta

To get this branch, use:
bzr branch http://darksoft.org/webbzr/perf/kseta

« back to all changes in this revision

Viewing changes to sources/mm/opencl7.c

  • Committer: Suren A. Chilingaryan
  • Date: 2013-09-30 06:47:09 UTC
  • Revision ID: csa@dside.dyndns.org-20130930064709-55cde0k5ci76t8z5
Simple matrix multiplication

Show diffs side-by-side

added added

removed removed

Lines of Context:
sources/mm/opencl7.c
 
1
#include <sys/types.h>
 
2
#include <stdio.h>
 
3
#include <string.h>
 
4
 
 
5
#include <CL/cl.h>
 
6
 
 
7
#define BLOCK_SIZE 16
 
8
 
 
9
#ifndef PPT
 
10
# define PPT 1
 
11
#endif 
 
12
 
 
13
#ifndef SHMEM
 
14
# define SHMEM 1
 
15
#endif
 
16
 
 
17
static cl_int err;
 
18
static cl_platform_id platform = 0;
 
19
static cl_device_id device = 0;
 
20
 
 
21
static cl_context_properties props[3] = { CL_CONTEXT_PLATFORM, 0, 0 };
 
22
static cl_context ctx = 0;
 
23
static cl_command_queue queues[2] = {0};
 
24
static cl_command_queue queue = 0;
 
25
 
 
26
static cl_mem dev_res[2], dev_a[2], dev_b[2];
 
27
static cl_mem pinned_mem_res, pinned_mem_a, pinned_mem_b;
 
28
float *pinned_res, *pinned_a, *pinned_b;
 
29
 
 
30
cl_program app;
 
31
cl_kernel kernel;
 
32
 
 
33
size_t matrix_size;
 
34
 
 
35
int exercise_required_alignment = PPT * BLOCK_SIZE;
 
36
 
 
37
int exercise_init(const char *name, size_t size) {
 
38
    int ret = 0;
 
39
    char param[1024];
 
40
 
 
41
    FILE *f;
 
42
    size_t len;
 
43
    char *source;
 
44
    
 
45
    int i;
 
46
 
 
47
    err = clGetPlatformIDs(1, &platform, NULL);
 
48
    if (err != CL_SUCCESS) {
 
49
        printf( "clGetPlatformIDs() failed with %d\n", err );
 
50
        return 1;
 
51
    }
 
52
 
 
53
    err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
 
54
    if (err != CL_SUCCESS) {
 
55
        printf( "clGetDeviceIDs() failed with %d\n", err );
 
56
        return 1;
 
57
    }
 
58
 
 
59
    err = clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(param), param, NULL);
 
60
    if (!err) {
 
61
        printf("Using device: %s\n", param);
 
62
    }
 
63
 
 
64
 
 
65
    props[1] = (cl_context_properties)platform;
 
66
    ctx = clCreateContext(props, 1, &device, NULL, NULL, &err);
 
67
    if (err != CL_SUCCESS) {
 
68
        printf( "clCreateContext() failed with %d\n", err );
 
69
        return 1;
 
70
    }
 
71
 
 
72
    for (i = 0; i < 2; i++) {
 
73
        queues[i] = clCreateCommandQueue(ctx, device, CL_QUEUE_PROFILING_ENABLE, &err);
 
74
        if (err != CL_SUCCESS) {
 
75
            printf( "clCreateCommandQueue() failed with %d\n", err );
 
76
            clReleaseContext(ctx);
 
77
            return 1;
 
78
        }
 
79
 
 
80
        dev_a[i] = clCreateBuffer(ctx, CL_MEM_READ_ONLY, size * size * sizeof(float), NULL, &err);
 
81
        if (err != CL_SUCCESS) {
 
82
            printf("clCreateBuffer() failed with %d\n", err);
 
83
            return 1;
 
84
        }
 
85
 
 
86
        dev_b[i] = clCreateBuffer(ctx, CL_MEM_READ_ONLY, size * size * sizeof(float), NULL, &err);
 
87
        if (err != CL_SUCCESS) {
 
88
            printf("clCreateBuffer() failed with %d\n", err);
 
89
            return 1;
 
90
        }
 
91
 
 
92
        dev_res[i] = clCreateBuffer(ctx, CL_MEM_READ_WRITE, size * size * sizeof(float), NULL, &err);
 
93
        if (err != CL_SUCCESS) {
 
94
            printf("clCreateBuffer() failed with %d\n", err);
 
95
            return 1;
 
96
        }
 
97
    }
 
98
    queue = queues[0];
 
99
 
 
100
    char *res = calloc(size * size, sizeof(float));
 
101
    if (res) {
 
102
        for (i = 0; i < 2; i++) {
 
103
            err = clEnqueueWriteBuffer(queue, dev_res[i], CL_TRUE, 0, size * size * sizeof(float), res, 0, NULL, NULL);
 
104
            if (err != CL_SUCCESS) {
 
105
                printf("clEnqueueWriteBuffer() failed with %d\n", err);
 
106
                return -1;
 
107
            }
 
108
        }
 
109
        free(res);
 
110
 
 
111
        err = clFinish(queue);
 
112
        if (err != CL_SUCCESS) {
 
113
            printf("clFinish() failed with %d\n", err);
 
114
            return -1;
 
115
        }
 
116
    }
 
117
 
 
118
    sprintf(param, "%s.cl", name);
 
119
    f = fopen(param, "r");
 
120
    if (!f) {
 
121
        printf("Can't open file with OpenCL kernels\n");
 
122
        return 1;
 
123
    }
 
124
 
 
125
    fseek(f, 0, SEEK_END);
 
126
    len = ftell(f);
 
127
    fseek(f, 0, SEEK_SET);
 
128
    
 
129
    source = (char*)malloc(len + 128);
 
130
    if (!source) {
 
131
        printf("Can't allocate memory for OpenCL source\n");
 
132
        return 1;
 
133
    }
 
134
    
 
135
    sprintf(source, "#define BLOCK_SIZE %lu\n#define PPT %lu\n\n", BLOCK_SIZE, PPT, PPT);
 
136
    if (fread(source + strlen(source), 1, len, f) != len) {
 
137
        printf("Can't read OpenCL source\n");
 
138
        return 1;
 
139
    }
 
140
    fclose(f);
 
141
 
 
142
    len = strlen(source);
 
143
 
 
144
    app = clCreateProgramWithSource(ctx, 1, (const char**)&source, &len, &err);
 
145
    if (err != CL_SUCCESS) {
 
146
        printf("clCreateProgramWithSource() failed with %d\n", err);
 
147
        return 1;
 
148
    }
 
149
    
 
150
    err = clBuildProgram(app, 1, &device,  "", NULL, NULL);
 
151
    if (err != CL_SUCCESS) {
 
152
        printf("clBuildProgram() failed with %d\n", err);
 
153
        return 1;
 
154
    }
 
155
    
 
156
    free(source);
 
157
    
 
158
    kernel = clCreateKernel(app, "multiply", &err);
 
159
    if (err != CL_SUCCESS) {
 
160
        printf("clCreateKernel() failed with %d\n", err);
 
161
        return 1;
 
162
    }
 
163
 
 
164
    matrix_size = size;
 
165
    clSetKernelArg(kernel, 3, sizeof(unsigned long), &matrix_size);
 
166
    err = clSetKernelArg(kernel, 4, 2 * SHMEM * BLOCK_SIZE * BLOCK_SIZE * sizeof(float), NULL);
 
167
 
 
168
    return 0;
 
169
}
 
170
 
 
171
int exercise_allocate(float **res, float **a, float **b, size_t size) {
 
172
    pinned_mem_a = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, size * size * sizeof(float), NULL, &err);
 
173
    if (err != CL_SUCCESS) {
 
174
        printf("clCreateBuffer() for pinned memory failed with %d\n", err);
 
175
        return 1;
 
176
    }
 
177
    
 
178
    pinned_a = (cl_float*)clEnqueueMapBuffer(queue, pinned_mem_a, CL_TRUE, CL_MAP_WRITE, 0, size * size * sizeof(float), 0, NULL, NULL, &err);
 
179
    if (err != CL_SUCCESS) {
 
180
        printf("clEnqueueMapBuffer() failed with %d\n", err);
 
181
        return 1;
 
182
    }
 
183
 
 
184
    pinned_mem_b = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, size * size * sizeof(float), NULL, &err);
 
185
    if (err != CL_SUCCESS) {
 
186
        printf("clCreateBuffer() for pinned memory failed with %d\n", err);
 
187
        return 1;
 
188
    }
 
189
    
 
190
    pinned_b = (cl_float*)clEnqueueMapBuffer(queue, pinned_mem_b, CL_TRUE, CL_MAP_WRITE, 0, size * size * sizeof(float), 0, NULL, NULL, &err);
 
191
    if (err != CL_SUCCESS) {
 
192
        printf("clEnqueueMapBuffer() failed with %d\n", err);
 
193
        return 1;
 
194
    }
 
195
 
 
196
    pinned_mem_res = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, size * size * sizeof(float), NULL, &err);
 
197
    if (err != CL_SUCCESS) {
 
198
        printf("clCreateBuffer() for pinned memory failed with %d\n", err);
 
199
        return 1;
 
200
    }
 
201
    
 
202
    pinned_res = (cl_float*)clEnqueueMapBuffer(queue, pinned_mem_res, CL_TRUE, CL_MAP_READ, 0, size * size * sizeof(float), 0, NULL, NULL, &err);
 
203
    if (err != CL_SUCCESS) {
 
204
        printf("clEnqueueMapBuffer() failed with %d\n", err);
 
205
        return 1;
 
206
    }
 
207
    
 
208
    *res = pinned_res;
 
209
    *a = pinned_a;
 
210
    *b = pinned_b;
 
211
 
 
212
    return 0;
 
213
}
 
214
 
 
215
 
 
216
void exercise_free() {
 
217
    int i;
 
218
    
 
219
    clReleaseKernel(kernel);
 
220
    clReleaseProgram(app);
 
221
 
 
222
    clEnqueueUnmapMemObject(queue, pinned_mem_a, pinned_a, 0, NULL, NULL);
 
223
    clEnqueueUnmapMemObject(queue, pinned_mem_b, pinned_b, 0, NULL, NULL);
 
224
    clEnqueueUnmapMemObject(queue, pinned_mem_res, pinned_res, 0, NULL, NULL);
 
225
 
 
226
    clReleaseMemObject(pinned_mem_a);
 
227
    clReleaseMemObject(pinned_mem_b);
 
228
    clReleaseMemObject(pinned_mem_res);
 
229
 
 
230
    clReleaseCommandQueue(queue);
 
231
    clReleaseContext(ctx);
 
232
 
 
233
    for (i = 0; i < 2; i++) {
 
234
        clReleaseMemObject(dev_res[i]);
 
235
        clReleaseMemObject(dev_b[i]);
 
236
        clReleaseMemObject(dev_a[i]);
 
237
    }
 
238
}
 
239
 
 
240
size_t exercise(float *res, float *a, float *b, size_t size, int iterations) {
 
241
    int i, pos;
 
242
    
 
243
    cl_event events[iterations];
 
244
    size_t runtime = 0;
 
245
    cl_ulong start, end;
 
246
 
 
247
    size_t local_size[] = {BLOCK_SIZE, BLOCK_SIZE};
 
248
    size_t global_size[] = {size/PPT, size/PPT};
 
249
 
 
250
    for (i = 0; i <= iterations; i++) {
 
251
        if (i < iterations) {
 
252
            pos = i%2;
 
253
            
 
254
            err = clEnqueueWriteBuffer(queues[pos], dev_a[pos], CL_FALSE, 0, size * size * sizeof(float), a, 0, NULL, NULL);
 
255
            if (err != CL_SUCCESS) {
 
256
                printf("clEnqueueWriteBuffer() failed with %d\n", err);
 
257
                return -11;
 
258
            }
 
259
 
 
260
            err = clEnqueueWriteBuffer(queues[pos], dev_b[pos], CL_FALSE, 0, size * size * sizeof(float), b, 0, NULL, NULL);
 
261
            if (err != CL_SUCCESS) {
 
262
                printf("clEnqueueWriteBuffer() failed with %d\n", err);
 
263
                return -1;
 
264
            }
 
265
            
 
266
            if (!i) continue;
 
267
        }
 
268
        
 
269
        pos = (i + 1)%2;
 
270
 
 
271
        clSetKernelArg(kernel, 0, sizeof(cl_mem), &dev_res[pos]);
 
272
        clSetKernelArg(kernel, 1, sizeof(cl_mem), &dev_a[pos]);
 
273
        clSetKernelArg(kernel, 2, sizeof(cl_mem), &dev_b[pos]);
 
274
 
 
275
        err = clEnqueueNDRangeKernel(queues[pos], kernel, 2, 0, global_size, local_size, 0, NULL, &events[i-1]);
 
276
        if (err != CL_SUCCESS) {
 
277
            printf("clEnqueueNDRangeKernel() failed with %d\n", err);
 
278
            return -1;
 
279
        }
 
280
 
 
281
        err = clEnqueueReadBuffer(queues[pos], dev_res[pos], CL_FALSE, 0, size * size * sizeof(float), res, 0, NULL, NULL);
 
282
        if (err != CL_SUCCESS) {
 
283
            printf("clEnqueueReadBuffer() failed with %d\n", err);
 
284
            return -1;
 
285
        }
 
286
    }
 
287
 
 
288
    err = clFinish(queues[0]);
 
289
    if (err != CL_SUCCESS) {
 
290
        printf("clFinish() failed with %d\n", err);
 
291
        return -1;
 
292
    }
 
293
 
 
294
    err = clFinish(queues[1]);
 
295
    if (err != CL_SUCCESS) {
 
296
        printf("clFinish() failed with %d\n", err);
 
297
        return -1;
 
298
    }
 
299
 
 
300
        
 
301
    for (i = 0; i < iterations; i++) {
 
302
        err = clGetEventProfilingInfo(events[i], CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL);
 
303
        if (err != CL_SUCCESS) {
 
304
            printf("clGetEventProfilingInfo() failed with %d\n", err);
 
305
            return -1;
 
306
        }
 
307
        
 
308
        err = clGetEventProfilingInfo(events[i], CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL);
 
309
        if (err != CL_SUCCESS) {
 
310
            printf("clGetEventProfilingInfo() failed with %d\n", err);
 
311
            return -1;
 
312
        }
 
313
        
 
314
        runtime += end - start;
 
315
    }
 
316
 
 
317
 
 
318
    return runtime;
 
319
}
 
320