|
0
|
1 #include <stdio.h>
|
|
|
2 #include <stdlib.h>
|
|
|
3 #include <string.h>
|
|
|
4 #include <sys/stat.h>
|
|
|
5 #include <sys/mman.h>
|
|
|
6 #include <sys/types.h>
|
|
|
7 #include <fcntl.h>
|
|
|
8 #include <unistd.h>
|
|
|
9 #include <OpenCL/opencl.h>
|
|
|
10 #include <oclUtils.h>
|
|
|
11
|
|
1
|
12 #define OUT_PARAM_NUM 4
|
|
|
13 #define PRINT_PARAM_NUM 2
|
|
0
|
14
|
|
|
15 typedef struct {
|
|
|
16 caddr_t file_mmap;
|
|
|
17 off_t size;
|
|
|
18 } st_mmap_t;
|
|
|
19
|
|
1
|
20 typedef struct {
|
|
|
21 int size;
|
|
|
22 int remain_size;
|
|
|
23 } param_t;
|
|
0
|
24
|
|
|
25
|
|
|
26 static st_mmap_t
|
|
|
27 my_mmap(char *filename)
|
|
|
28 {
|
|
|
29
|
|
|
30 /*マッピングだよ!*/
|
|
|
31 int fd = -1;
|
|
|
32 int map = MAP_PRIVATE;
|
|
|
33 st_mmap_t st_mmap;
|
|
|
34 struct stat sb;
|
|
|
35
|
|
|
36 if ((fd=open(filename,O_RDONLY,0666))==0) {
|
|
|
37 fprintf(stderr,"can't open %s\n",filename);
|
|
|
38 }
|
|
|
39
|
|
|
40 if (fstat(fd,&sb)) {
|
|
|
41 fprintf(stderr,"can't fstat %s\n",filename);
|
|
|
42 }
|
|
|
43
|
|
1
|
44 st_mmap.size = sb.st_size;
|
|
0
|
45
|
|
1
|
46 printf("file size %d\n",(int)st_mmap.size);
|
|
0
|
47
|
|
|
48 st_mmap.file_mmap = (char*)mmap(NULL,st_mmap.size,PROT_READ,map,fd,(off_t)0);
|
|
|
49 if (st_mmap.file_mmap == (caddr_t)-1) {
|
|
|
50 fprintf(stderr,"Can't mmap file\n");
|
|
|
51 perror(NULL);
|
|
|
52 exit(0);
|
|
|
53 }
|
|
|
54
|
|
|
55 return st_mmap;
|
|
|
56
|
|
|
57 }
|
|
|
58
|
|
|
59
|
|
|
60
|
|
|
61 int main(int args, char *argv[])
|
|
|
62 {
|
|
|
63
|
|
|
64 char *filename = 0;
|
|
4
|
65 cl_device_type dev_type = CL_DEVICE_TYPE_DEFAULT;
|
|
0
|
66
|
|
|
67 for (int i = 1; argv[i]; ++i) {
|
|
|
68 if (strcmp(argv[i], "-file") == 0) {
|
|
|
69 filename = argv[i+1];
|
|
4
|
70 } else if (strcmp(argv[i], "-help") == 0) {
|
|
0
|
71 printf("Usage: ./word_count [-file filename]\n");
|
|
4
|
72 } else if (strcmp(argv[i], "-type") == 0) {
|
|
|
73
|
|
|
74 if (strcmp(argv[i+1], "gpu") == 0) {
|
|
|
75 dev_type = CL_DEVICE_TYPE_GPU;
|
|
|
76 } else if (strcmp(argv[i+1], "cpu") == 0) {
|
|
|
77 dev_type = CL_DEVICE_TYPE_CPU;
|
|
|
78 }
|
|
|
79
|
|
0
|
80 }
|
|
|
81 }
|
|
|
82
|
|
|
83 if (filename == 0) {
|
|
|
84 printf("Usage: ./word_count [-file filename]\n");
|
|
|
85 return 0;
|
|
|
86 }
|
|
|
87
|
|
|
88 //指定されたファイルをメモリにmap
|
|
|
89 st_mmap_t st_mmap = my_mmap(filename);
|
|
|
90
|
|
|
91 //kernelファイルの大きさ取得して、メモリ確保
|
|
|
92 int fd = -1;
|
|
2
|
93 const char *kernel_filename = "./kernel/word_count.cl";
|
|
0
|
94 if ((fd=open(kernel_filename,O_RDONLY,0666))==0) {
|
|
|
95 fprintf(stderr,"can't open %s\n",kernel_filename);
|
|
|
96 }
|
|
|
97
|
|
|
98 struct stat sb;
|
|
|
99 if (fstat(fd,&sb)) {
|
|
|
100 fprintf(stderr,"can't fstat %s\n",filename);
|
|
|
101 }
|
|
|
102
|
|
|
103 size_t source_size = sb.st_size;
|
|
|
104 void *source_str = malloc(source_size);
|
|
|
105
|
|
|
106 int err = read(fd, source_str, source_size);
|
|
|
107 if (err == -1) {
|
|
|
108 fprintf(stderr,"can't read %s\n",filename);
|
|
|
109 }
|
|
|
110
|
|
|
111 close(fd);
|
|
|
112
|
|
|
113 cl_platform_id platform_id = NULL;
|
|
|
114 cl_uint ret_num_platforms = NULL;
|
|
|
115 cl_int ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
|
|
|
116
|
|
|
117 cl_device_id device_id = NULL;
|
|
|
118 cl_uint ret_num_devices = NULL;
|
|
|
119 // CL_DEVICE_TYPE_CPU, CL_DEVICE_TYPE_GPU, と指定できる
|
|
|
120 // CL_DEVICE_TYPE_DEFAULT はどうなるのか
|
|
4
|
121 ret = clGetDeviceIDs( platform_id, dev_type,
|
|
0
|
122 1, &device_id, &ret_num_devices);
|
|
|
123
|
|
|
124 oclCheckError(ret, CL_SUCCESS);
|
|
|
125
|
|
|
126 //OpenCLコンテキストの作成
|
|
|
127 cl_context context = clCreateContext( NULL, 1, &device_id, NULL, NULL, &ret);
|
|
|
128 oclCheckError(ret, CL_SUCCESS);
|
|
|
129 cl_command_queue command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
|
|
|
130 oclCheckError(ret, CL_SUCCESS);
|
|
|
131 // カーネルプログラムを読み込む
|
|
|
132 cl_program program = clCreateProgramWithSource(context, 1, (const char **)&source_str,
|
|
|
133 (const size_t *)&source_size, &ret);
|
|
1
|
134
|
|
0
|
135 ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
|
|
|
136
|
|
1
|
137 if (ret != CL_SUCCESS) {
|
|
|
138 oclLogBuildInfo(program, device_id);
|
|
|
139 }
|
|
|
140
|
|
0
|
141 oclCheckError(ret, CL_SUCCESS);
|
|
|
142
|
|
|
143 //カーネルプログラムをビルド
|
|
|
144 //Task選択にあたる
|
|
|
145 cl_kernel kernel = clCreateKernel(program, "word_count", &ret);
|
|
|
146 oclCheckError(ret, CL_SUCCESS);
|
|
|
147
|
|
|
148 //カウントするテキストデータのメモリオブジェクト
|
|
|
149 cl_mem text_memobj = clCreateBuffer(context, CL_MEM_READ_ONLY,
|
|
2
|
150 st_mmap.size, NULL, &ret);
|
|
0
|
151
|
|
|
152 ret = clEnqueueWriteBuffer(command_queue, text_memobj,
|
|
|
153 CL_TRUE, 0, st_mmap.size, (char*)st_mmap.file_mmap,
|
|
|
154 0, NULL, NULL);
|
|
|
155
|
|
|
156 oclCheckError(ret, CL_SUCCESS);
|
|
|
157
|
|
|
158 /*
|
|
|
159 * 並列度の計算
|
|
|
160 * wcするファイルの大きさに合わせる
|
|
|
161 *
|
|
|
162 */
|
|
|
163
|
|
|
164 // このdivi_size はどうやって決めるよ
|
|
|
165 int divi_size = 1024;
|
|
2
|
166 //size_t global_work_size = (st_mmap.size + divi_size - 1) / divi_size;
|
|
|
167 size_t global_work_size = st_mmap.size / divi_size;
|
|
1
|
168 int remain_size = st_mmap.size - global_work_size * divi_size;
|
|
2
|
169 if (remain_size > 0) {
|
|
|
170 global_work_size += 1;
|
|
|
171 }
|
|
0
|
172
|
|
1
|
173 /*
|
|
|
174 * 行数、単語数, 分割地点のフラグを格納する配列
|
|
|
175 * word_num, line_num, head, tail
|
|
|
176 */
|
|
0
|
177
|
|
|
178 int out_size = sizeof(int) * OUT_PARAM_NUM * global_work_size;
|
|
|
179 cl_mem out_memobj = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
|
|
|
180 out_size * sizeof(char), NULL, &ret);
|
|
|
181
|
|
|
182 oclCheckError(ret, CL_SUCCESS);
|
|
|
183
|
|
2
|
184 // 必要なパラメータのオブジェクト
|
|
|
185 cl_mem param_memobj = clCreateBuffer(context, CL_MEM_READ_ONLY,
|
|
|
186 sizeof(param_t), NULL, &ret);
|
|
|
187
|
|
|
188 param_t param;
|
|
|
189 param.size = divi_size;
|
|
|
190 param.remain_size = remain_size;
|
|
|
191
|
|
|
192 ret = clEnqueueWriteBuffer(command_queue, param_memobj,
|
|
|
193 CL_TRUE, 0, sizeof(param_t), ¶m,
|
|
|
194 0, NULL, NULL);
|
|
|
195
|
|
|
196 oclCheckError(ret, CL_SUCCESS);
|
|
|
197
|
|
0
|
198 // 引数のSet
|
|
|
199 // memory object にしなくてもできるsetできるかも
|
|
|
200 ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&text_memobj);
|
|
|
201 ret |= clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)¶m_memobj);
|
|
|
202 ret |= clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&out_memobj);
|
|
|
203
|
|
|
204 oclCheckError(ret, CL_SUCCESS);
|
|
|
205
|
|
|
206 /*
|
|
|
207 * kernel実行
|
|
|
208 * 並列に処理せずに work-item ひとつで動かしたい場合は、clEnqueueNDRangeKernel の簡易版 clEnqueueTask が使える
|
|
|
209 *
|
|
1
|
210 */
|
|
0
|
211
|
|
|
212 //ret = clEnqueueTask(command_queue, kernel, 0, NULL, NULL);
|
|
|
213 // global_work_size は配列。並列動作させる時の次元数にあわせて、配列の次元数も決まるはず
|
|
|
214 ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global_work_size, NULL, 0, NULL, NULL);
|
|
|
215
|
|
4
|
216
|
|
0
|
217 oclCheckError(ret, CL_SUCCESS);
|
|
|
218
|
|
|
219 int *out_data = (int*)malloc(out_size);
|
|
|
220
|
|
4
|
221
|
|
|
222 #ifndef BLOKING
|
|
|
223
|
|
0
|
224 // 演算結果の読み込み
|
|
4
|
225 // CL_TRUE で bloking
|
|
0
|
226 ret = clEnqueueReadBuffer(command_queue, out_memobj, CL_TRUE, 0,
|
|
1
|
227 out_size, out_data, 0, NULL, NULL);
|
|
0
|
228
|
|
4
|
229 #else
|
|
|
230
|
|
|
231 cl_event ev;
|
|
|
232
|
|
|
233 // 演算結果の読み込み
|
|
|
234 // CL_FALSE で non-bloking
|
|
|
235 ret = clEnqueueReadBuffer(command_queue, out_memobj, CL_FALSE, 0,
|
|
|
236 out_size, out_data, 0, NULL, &ev);
|
|
|
237
|
|
|
238 // event object に関連する処理の完了を wait する
|
|
|
239 // これはeventひとつだけど、listにして、複数 wait できる
|
|
|
240 clWaitForEvents(1, &ev);
|
|
|
241
|
|
|
242 #endif
|
|
|
243
|
|
0
|
244 oclCheckError(ret, CL_SUCCESS);
|
|
|
245
|
|
1
|
246 printf("global_work_size %d\n", (int)global_work_size);
|
|
|
247
|
|
2
|
248 int word_line_num[PRINT_PARAM_NUM];
|
|
|
249
|
|
|
250 for (int i = 0; i < PRINT_PARAM_NUM; i++) {
|
|
|
251 word_line_num[i] = 0;
|
|
|
252 }
|
|
|
253
|
|
0
|
254 for (int i = 0; i < global_work_size; i++) {
|
|
2
|
255
|
|
1
|
256 for (int j = 0; j < PRINT_PARAM_NUM; j++) {
|
|
2
|
257
|
|
|
258 word_line_num[j] += out_data[i*OUT_PARAM_NUM+j];
|
|
|
259
|
|
0
|
260 }
|
|
2
|
261
|
|
3
|
262 //flagの判定
|
|
2
|
263 if ((i != global_work_size-1) &&
|
|
3
|
264 (out_data[i*OUT_PARAM_NUM+PRINT_PARAM_NUM+1] == 1) &&
|
|
|
265 (out_data[i*OUT_PARAM_NUM+PRINT_PARAM_NUM+4] == 0)) {
|
|
2
|
266
|
|
|
267 word_line_num[0] += 1;
|
|
|
268
|
|
|
269 }
|
|
0
|
270 }
|
|
|
271
|
|
4
|
272 printf("%d %d \n", word_line_num[0], word_line_num[1]);
|
|
2
|
273
|
|
0
|
274 clFlush(command_queue);
|
|
|
275 clFinish(command_queue);
|
|
|
276 clReleaseKernel(kernel);
|
|
|
277 clReleaseProgram(program);
|
|
|
278 clReleaseMemObject(text_memobj);
|
|
|
279 clReleaseMemObject(param_memobj);
|
|
|
280 clReleaseMemObject(out_memobj);
|
|
|
281 clReleaseCommandQueue(command_queue);
|
|
|
282 clReleaseContext(context);
|
|
|
283
|
|
|
284 free(source_str);
|
|
|
285 free(out_data);
|
|
|
286
|
|
|
287 return 0;
|
|
|
288
|
|
|
289 }
|
