/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 tutorials/1_transpose/transpose1.c

  • Committer: Suren A. Chilingaryan
  • Date: 2013-10-08 23:53:50 UTC
  • Revision ID: csa@dside.dyndns.org-20131008235350-hsu8oukzkh05gtcm
Add tutorials

Show diffs side-by-side

added added

removed removed

Lines of Context:
tutorials/1_transpose/transpose1.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 queue = 0;
 
24
 
 
25
static cl_mem dev_res, dev_a;
 
26
 
 
27
static cl_program app;
 
28
static cl_kernel kernel;
 
29
 
 
30
static size_t matrix_size;
 
31
 
 
32
int exercise_required_alignment = PPT * BLOCK_SIZE;
 
33
 
 
34
int exercise_init(const char *name, size_t size) {
 
35
    int ret = 0;
 
36
    char param[1024];
 
37
 
 
38
    FILE *f;
 
39
    size_t len;
 
40
    char *source;
 
41
 
 
42
    err = clGetPlatformIDs(1, &platform, NULL);
 
43
    if (err != CL_SUCCESS) {
 
44
        printf( "clGetPlatformIDs() failed with %d\n", err );
 
45
        return 1;
 
46
    }
 
47
 
 
48
    err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
 
49
    if (err != CL_SUCCESS) {
 
50
        printf( "clGetDeviceIDs() failed with %d\n", err );
 
51
        return 1;
 
52
    }
 
53
 
 
54
    err = clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(param), param, NULL);
 
55
    if (!err) {
 
56
        printf("Using device: %s\n", param);
 
57
    }
 
58
 
 
59
 
 
60
    props[1] = (cl_context_properties)platform;
 
61
    ctx = clCreateContext(props, 1, &device, NULL, NULL, &err);
 
62
    if (err != CL_SUCCESS) {
 
63
        printf( "clCreateContext() failed with %d\n", err );
 
64
        return 1;
 
65
    }
 
66
 
 
67
    queue = clCreateCommandQueue(ctx, device, CL_QUEUE_PROFILING_ENABLE, &err);
 
68
    if (err != CL_SUCCESS) {
 
69
        printf( "clCreateCommandQueue() failed with %d\n", err );
 
70
        clReleaseContext(ctx);
 
71
        return 1;
 
72
    }
 
73
 
 
74
    dev_a = clCreateBuffer(ctx, CL_MEM_READ_ONLY, size * size * sizeof(float), NULL, &err);
 
75
    if (err != CL_SUCCESS) {
 
76
        printf("clCreateBuffer() failed with %d\n", err);
 
77
        return 1;
 
78
    }
 
79
 
 
80
    dev_res = clCreateBuffer(ctx, CL_MEM_READ_WRITE, 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
    char *res = calloc(size * size, sizeof(float));
 
87
    if (res) {
 
88
        err = clEnqueueWriteBuffer(queue, dev_res, CL_TRUE, 0, size * size * sizeof(float), res, 0, NULL, NULL);
 
89
        if (err != CL_SUCCESS) {
 
90
            printf("clEnqueueWriteBuffer() failed with %d\n", err);
 
91
            return -1;
 
92
        }
 
93
        free(res);
 
94
 
 
95
        err = clFinish(queue);
 
96
        if (err != CL_SUCCESS) {
 
97
            printf("clFinish() failed with %d\n", err);
 
98
            return -1;
 
99
        }
 
100
    }
 
101
 
 
102
    sprintf(param, "%s.cl", name);
 
103
    f = fopen(param, "r");
 
104
    if (!f) {
 
105
        printf("Can't open file with OpenCL kernels\n");
 
106
        return 1;
 
107
    }
 
108
 
 
109
    fseek(f, 0, SEEK_END);
 
110
    len = ftell(f);
 
111
    fseek(f, 0, SEEK_SET);
 
112
    
 
113
    source = (char*)malloc(len + 128);
 
114
    if (!source) {
 
115
        printf("Can't allocate memory for OpenCL source\n");
 
116
        return 1;
 
117
    }
 
118
    
 
119
    sprintf(source, "#define BLOCK_SIZE %u\n#define PPT %u\n\n", BLOCK_SIZE, PPT);
 
120
    if (fread(source + strlen(source), 1, len, f) != len) {
 
121
        printf("Can't read OpenCL source\n");
 
122
        return 1;
 
123
    }
 
124
    fclose(f);
 
125
 
 
126
    len = strlen(source);
 
127
 
 
128
    app = clCreateProgramWithSource(ctx, 1, (const char**)&source, &len, &err);
 
129
    if (err != CL_SUCCESS) {
 
130
        printf("clCreateProgramWithSource() failed with %d\n", err);
 
131
        return 1;
 
132
    }
 
133
    
 
134
    err = clBuildProgram(app, 1, &device,  "", NULL, NULL);
 
135
    if (err != CL_SUCCESS) {
 
136
        printf("clBuildProgram() failed with %d\n", err);
 
137
        return 1;
 
138
    }
 
139
    
 
140
    free(source);
 
141
    
 
142
    kernel = clCreateKernel(app, "transpose", &err);
 
143
    if (err != CL_SUCCESS) {
 
144
        printf("clCreateKernel() failed with %d\n", err);
 
145
        return 1;
 
146
    }
 
147
 
 
148
    clSetKernelArg(kernel, 0, sizeof(cl_mem), &dev_res);
 
149
    clSetKernelArg(kernel, 1, sizeof(cl_mem), &dev_a);
 
150
    matrix_size = size;
 
151
    clSetKernelArg(kernel, 2, sizeof(unsigned long), &matrix_size);
 
152
    err = clSetKernelArg(kernel, 3, 2 * BLOCK_SIZE * BLOCK_SIZE  * sizeof(float), NULL);
 
153
 
 
154
    return 0;
 
155
}
 
156
 
 
157
int exercise_allocate(float **res, float **a, size_t size) {
 
158
    return 0;
 
159
}
 
160
 
 
161
 
 
162
void exercise_free() {
 
163
    clReleaseKernel(kernel);
 
164
    clReleaseProgram(app);
 
165
    
 
166
    clReleaseCommandQueue(queue);
 
167
    clReleaseContext(ctx);
 
168
 
 
169
    clReleaseMemObject(dev_res);
 
170
    clReleaseMemObject(dev_a);
 
171
}
 
172
 
 
173
size_t exercise(float *res, float *a, size_t size, int iterations) {
 
174
    int i;
 
175
    
 
176
    cl_event event = NULL;
 
177
    size_t runtime = 0;
 
178
    cl_ulong start, end;
 
179
 
 
180
    size_t local_size[] = {BLOCK_SIZE, BLOCK_SIZE};
 
181
    size_t global_size[] = {size/PPT, size/PPT};
 
182
 
 
183
    for (i = 0; i < iterations; i++) {
 
184
        err = clEnqueueWriteBuffer(queue, dev_a, CL_TRUE, 0, size * size * sizeof(float), a, 0, NULL, NULL);
 
185
        if (err != CL_SUCCESS) {
 
186
            printf("clEnqueueWriteBuffer() failed with %d\n", err);
 
187
            return -11;
 
188
        }
 
189
 
 
190
        err = clEnqueueNDRangeKernel(queue, kernel, 2, 0, global_size, local_size, 0, NULL, &event);
 
191
        if (err != CL_SUCCESS) {
 
192
            printf("clEnqueueNDRangeKernel() failed with %d\n", err);
 
193
            return -1;
 
194
        }
 
195
 
 
196
        err = clEnqueueReadBuffer(queue, dev_res, CL_TRUE, 0, size * size * sizeof(float), res, 0, NULL, NULL);
 
197
        if (err != CL_SUCCESS) {
 
198
            printf("clEnqueueReadBuffer() failed with %d\n", err);
 
199
            return -1;
 
200
        }
 
201
 
 
202
        err = clFinish(queue);
 
203
        if (err != CL_SUCCESS) {
 
204
            printf("clFinish() failed with %d\n", err);
 
205
            return -1;
 
206
        }
 
207
        
 
208
        err = clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL);
 
209
        if (err != CL_SUCCESS) {
 
210
            printf("clGetEventProfilingInfo() failed with %d\n", err);
 
211
            return -1;
 
212
        }
 
213
        
 
214
        err = clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL);
 
215
        if (err != CL_SUCCESS) {
 
216
            printf("clGetEventProfilingInfo() failed with %d\n", err);
 
217
            return -1;
 
218
        }
 
219
        
 
220
        runtime += end - start;
 
221
    }
 
222
 
 
223
    return runtime;
 
224
}
 
225