source: GPU/pipeline_xmlwf.cpp @ 2277

Last change on this file since 2277 was 1759, checked in by lindanl, 8 years ago

Pipelined Parabix on GPU

File size: 13.3 KB
RevLine 
[1759]1// System includes
2#include <stdio.h>
3#include <stdlib.h>
4#include <string.h>
5
6// OpenCL includes
7#include <CL/cl.h>
8
9// Project includes
10
11// Constants, globals
12const int ELEMENTS = 1024*128;   // elements in each vector
13const int WORK_GROUP_SIZE = 64;
14#define ENTRIES 8
15#define BitBlock unsigned long long
16  struct Basis_bits {
17  BitBlock bit_0;
18  BitBlock bit_1;
19  BitBlock bit_2;
20  BitBlock bit_3;
21  BitBlock bit_4;
22  BitBlock bit_5;
23  BitBlock bit_6;
24  BitBlock bit_7;
25};
26
27  struct Lex {
28  BitBlock CR;
29  BitBlock LF;
30  BitBlock HT;
31  BitBlock SP;
32  BitBlock CRLF;
33  BitBlock RefStart;
34  BitBlock Semicolon;
35  BitBlock Colon;
36  BitBlock LAngle;
37  BitBlock RAngle;
38  BitBlock LBracket;
39  BitBlock RBracket;
40  BitBlock Exclam;
41  BitBlock QMark;
42  BitBlock Hyphen;
43  BitBlock Equals;
44  BitBlock SQuote;
45  BitBlock DQuote;
46  BitBlock Slash;
47  BitBlock Hash;
48  BitBlock x;
49  BitBlock ASCII_name_start;
50  BitBlock ASCII_name_char;
51  BitBlock NameScan;
52  BitBlock Digit;
53  BitBlock Hex;
54  BitBlock WS;
55  BitBlock error;
56};
57
58struct Pdata{
59          int cur_stage[WORK_GROUP_SIZE];
60          struct Basis_bits basis_bits[WORK_GROUP_SIZE];
61          struct Lex lex[WORK_GROUP_SIZE];
62};
63
64// Signatures
65char* readSource(const char *sourceFilename); 
66
67#define BitBlock long long
68
69int main(int argc, char ** argv)
70{
71   printf("Running Transposition program\n\n");
72
73   size_t datasize = sizeof(BitBlock)*8*ELEMENTS;
74
75   BitBlock *S;   // Input array
76   BitBlock *P;   // Output array
77
78   // Allocate space for input/output data
79   S = (BitBlock *)malloc(datasize);
80   P = (BitBlock *)malloc(datasize/8);
81 
82   if(S == NULL || P == NULL) {
83      perror("malloc");
84      exit(-1);
85   }
86
87   // Initialize the input data
88   memset (S,0,datasize);
89   memset (P,0,datasize/8);
90
91   FILE *infile;
92   char * infilename = "test.xml";
93   infile = fopen(infilename, "rb");
94        if (!infile) {
95                fprintf(stderr, "Error: cannot open %s for input.\n", infilename);
96                exit(-1);
97        }
98   fread((void*)S,datasize,1,infile);
99
100   cl_int status;  // use as return value for most OpenCL functions
101
102   cl_uint numPlatforms = 0;
103   cl_platform_id *platforms;
104               
105   // Query for the number of recongnized platforms
106   status = clGetPlatformIDs(0, NULL, &numPlatforms);
107   if(status != CL_SUCCESS) {
108      printf("clGetPlatformIDs failed\n");
109      exit(-1);
110   }
111
112   // Make sure some platforms were found
113   if(numPlatforms == 0) {
114      printf("No platforms detected.\n");
115      exit(-1);
116   }
117
118   // Allocate enough space for each platform
119   platforms = (cl_platform_id*)malloc(numPlatforms*sizeof(cl_platform_id));
120   if(platforms == NULL) {
121      perror("malloc");
122      exit(-1);
123   }
124
125   // Fill in platforms
126   clGetPlatformIDs(numPlatforms, platforms, NULL);
127   if(status != CL_SUCCESS) {
128      printf("clGetPlatformIDs failed\n");
129      exit(-1);
130   }
131
132   // Print out some basic information about each platform
133   printf("%u platforms detected\n", numPlatforms);
134   for(unsigned int i = 0; i < numPlatforms; i++) {
135      char buf[100];
136      printf("Platform %u: \n", i);
137      status = clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR,
138                       sizeof(buf), buf, NULL);
139      printf("\tVendor: %s\n", buf);
140      status |= clGetPlatformInfo(platforms[i], CL_PLATFORM_NAME,
141                       sizeof(buf), buf, NULL);
142      printf("\tName: %s\n", buf);
143
144      if(status != CL_SUCCESS) {
145         printf("clGetPlatformInfo failed\n");
146         exit(-1);
147      }
148   }
149   printf("\n");
150
151   cl_uint numDevices = 0;
152   cl_device_id *devices;
153
154   // Retrive the number of devices present
155   status = clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_GPU, 0, NULL, 
156                           &numDevices);
157   
158  // clUtilCheckErrorVoid(status);
159   if(status != CL_SUCCESS) {
160      printf("clGetDeviceIDs failed\n");
161      exit(-1);
162   }
163
164   // Make sure some devices were found
165   if(numDevices == 0) {
166      printf("No devices detected.\n");
167      exit(-1);
168   }
169
170   // Allocate enough space for each device
171   devices = (cl_device_id*)malloc(numDevices*sizeof(cl_device_id));
172   if(devices == NULL) {
173      perror("malloc");
174      exit(-1);
175   }
176
177   // Fill in devices
178   status = clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_GPU, numDevices,
179                     devices, NULL);
180   if(status != CL_SUCCESS) {
181      printf("clGetDeviceIDs failed\n");
182      exit(-1);
183   }   
184
185   // Print out some basic information about each device
186   printf("%u devices detected\n", numDevices);
187   for(unsigned int i = 0; i < numDevices; i++) {
188      char buf[100];
189      printf("Device %u: \n", i);
190      status = clGetDeviceInfo(devices[i], CL_DEVICE_VENDOR,
191                       sizeof(buf), buf, NULL);
192      printf("\tDevice: %s\n", buf);
193      status |= clGetDeviceInfo(devices[i], CL_DEVICE_NAME,
194                       sizeof(buf), buf, NULL);
195      printf("\tName: %s\n", buf);
196
197      if(status != CL_SUCCESS) {
198         printf("clGetDeviceInfo failed\n");
199         exit(-1);
200      }
201   }
202   printf("\n");
203
204   cl_context context;
205
206   // Create a context and associate it with the devices
207   context = clCreateContext(NULL, numDevices, devices, NULL, NULL, &status);
208   if(status != CL_SUCCESS || context == NULL) {
209      printf("clCreateContext failed\n");
210      exit(-1);
211   }
212
213   cl_command_queue cmdQueue;
214
215   // Create a command queue and associate it with the device you
216   // want to execute on
217   cmdQueue = clCreateCommandQueue(context, devices[0],CL_QUEUE_PROFILING_ENABLE, &status);
218   if(status != CL_SUCCESS || cmdQueue == NULL) {
219      printf("clCreateCommandQueue failed\n");
220      exit(-1);
221   }
222
223   cl_mem d_S;  // Input buffers on device
224   cl_mem d_P;  // Output buffer on device
225   cl_mem d_D;
226
227   // Create a buffer object (d_S) that contains the data from the host ptr S
228   d_S = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,
229                   datasize, S, &status);
230   if(status != CL_SUCCESS || d_S == NULL) {
231      printf("clCreateBuffer failed\n");
232      exit(-1);
233   }
234
235   // Create a buffer object (d_P)
236   d_P  = clCreateBuffer(context, CL_MEM_READ_WRITE, 
237                   datasize/8, NULL, &status);
238   if(status != CL_SUCCESS || d_P == NULL) {
239      printf("clCreateBuffer failed\n");
240      exit(-1);
241   }
242
243   d_D  = clCreateBuffer(context, CL_MEM_READ_WRITE, 
244                   sizeof(struct Pdata)*ENTRIES, NULL, &status);
245   if(status != CL_SUCCESS || d_D == NULL) {
246      printf("clCreateBuffer failed\n");
247      exit(-1);
248   }
249   cl_program program;
250   
251   char *source;
252   const char *sourceFile = "tag_parsing.cl";
253   // This function reads in the source code of the program
254   source = readSource(sourceFile);
255
256   //printf("Program source is:\n%s\n", source);
257
258   // Create a program. The 'source' string is the code from the
259   // xmlwf.cl file.
260   program = clCreateProgramWithSource(context, 1, (const char**)&source, 
261                              NULL, &status);
262   if(status != CL_SUCCESS) {
263      printf("clCreateProgramWithSource failed\n");
264      exit(-1);
265   }
266
267   cl_int buildErr;
268   // Build (compile & link) the program for the devices.
269   // Save the return value in 'buildErr' (the following
270   // code will print any compilation errors to the screen)
271   buildErr = clBuildProgram(program, numDevices, devices, NULL, NULL, NULL);
272
273   // If there are build errors, print them to the screen
274   if(buildErr != CL_SUCCESS) {
275      printf("Program failed to build.\n");
276      cl_build_status buildStatus;
277      for(unsigned int i = 0; i < numDevices; i++) {
278         clGetProgramBuildInfo(program, devices[i], CL_PROGRAM_BUILD_STATUS,
279                          sizeof(cl_build_status), &buildStatus, NULL);
280         if(buildStatus == CL_SUCCESS) {
281            continue;
282         }
283
284         char *buildLog;
285         size_t buildLogSize;
286         clGetProgramBuildInfo(program, devices[i], CL_PROGRAM_BUILD_LOG,
287                          0, NULL, &buildLogSize);
288         buildLog = (char*)malloc(buildLogSize);
289         if(buildLog == NULL) {
290            perror("malloc");
291            exit(-1);
292         }
293         clGetProgramBuildInfo(program, devices[i], CL_PROGRAM_BUILD_LOG,
294                          buildLogSize, buildLog, NULL);
295         buildLog[buildLogSize-1] = '\0';
296         printf("Device %u Build Log:\n%s\n", i, buildLog);   
297         free(buildLog);
298      }
299      exit(0);
300   }
301   else {
302      printf("No build errors\n");
303   }
304
305
306   cl_kernel kernel;
307
308   // Create a kernel from the xmlwf function
309   kernel = clCreateKernel(program, "tag_parsing", &status);
310   if(status != CL_SUCCESS) {
311      printf("clCreateKernel failed\n");
312      exit(-1);
313   }
314
315   // Associate the input and output buffers with the kernel
316   status  = clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_S);
317   status |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_P);
318   status |= clSetKernelArg(kernel, 2, (WORK_GROUP_SIZE+1)*sizeof(BitBlock), NULL);
319   status |= clSetKernelArg(kernel, 3, WORK_GROUP_SIZE*sizeof(BitBlock), NULL);   
320   status |= clSetKernelArg(kernel, 4, sizeof(cl_mem), &d_D);   
321
322   if(status != CL_SUCCESS) {
323      printf("clSetKernelArg failed\n");
324      exit(-1);
325   }
326
327   cl_event ndrEvt;
328   cl_int eventStatus = CL_QUEUED;
329
330   // Define an index space (global work size) of threads for execution. 
331   // A workgroup size (local work size) is not required, but can be used.
332   size_t globalWorkSize[1];  // There are ELEMENTS threads
333   size_t localWorkSize[1];
334   globalWorkSize[0] = WORK_GROUP_SIZE*3;
335   localWorkSize[0] = WORK_GROUP_SIZE;
336
337   // Execute the kernel.
338   // 'globalWorkSize' is the 1D dimension of the work-items
339   status = clEnqueueNDRangeKernel(cmdQueue, kernel, 1, NULL, globalWorkSize, 
340                           localWorkSize, 0, NULL, &ndrEvt);
341   if(status != CL_SUCCESS) {
342      printf("clEnqueueNDRangeKernel failed\n");
343      exit(-1);
344   }
345    status = clFlush(cmdQueue);
346   if(status != CL_SUCCESS) {
347      printf("clFlush failed\n");
348      exit(-1);
349   }
350
351    /* wait for the kernel call to finish execution */
352   eventStatus = CL_QUEUED;
353   while(eventStatus != CL_COMPLETE)
354    {
355        status = clGetEventInfo(
356                        ndrEvt, 
357                        CL_EVENT_COMMAND_EXECUTION_STATUS, 
358                        sizeof(cl_int),
359                        &eventStatus,
360                        NULL);
361            if(status!=CL_SUCCESS){
362                                printf("clGetEventInfo failed.\n");
363                exit(-1);
364                        }
365    }
366
367
368   /* Calculate performance */
369    cl_ulong startTime=0;
370    cl_ulong endTime=0;
371       
372    /* Get kernel profiling info */
373    status = clGetEventProfilingInfo(ndrEvt,
374                                        CL_PROFILING_COMMAND_START,
375                                        sizeof(cl_ulong),
376                                        &startTime,
377                                        0);
378    if(status!=CL_SUCCESS){
379                        printf("clGetEventProfilingInfo failed.(startTime)\n");
380            exit(-1);
381        }
382
383    status = clGetEventProfilingInfo(ndrEvt,
384                                        CL_PROFILING_COMMAND_END,
385                                        sizeof(cl_ulong),
386                                        &endTime,
387                                        0);
388
389    if(status!=CL_SUCCESS){
390                        printf("clGetEventProfilingInfo failed.(endTime)\n");
391            exit(-1);
392        }
393
394    /* Print performance numbers */
395    unsigned long elapsed = (unsigned long) (endTime - startTime);
396        printf("Elapsed time is %i\n",elapsed);
397
398    clReleaseEvent(ndrEvt);
399       
400   // Read the OpenCL output buffer (d_C) to the host output array (C)
401   clEnqueueReadBuffer(cmdQueue, d_P, CL_TRUE, 0, datasize/8, P, 
402                  0, NULL, NULL);
403
404   // Verify correctness
405   
406   char * C = (char *)S;
407   for(int i=0;i<ELEMENTS;i++){
408           if(P[i]!=0){
409                        printf("error between position %i and %i : %llx\n",i*64,i*64+63,P[i]);
410                        for(int j=0;j<64;j++)
411                                printf("%c",C[i*64+j]);
412                        printf("\n");
413                        break;
414           }
415   }
416
417   clReleaseKernel(kernel);
418   clReleaseProgram(program);
419   clReleaseCommandQueue(cmdQueue);
420   clReleaseMemObject(d_S);
421   clReleaseMemObject(d_P);
422   clReleaseContext(context);
423
424   free(S);
425   free(P);
426   free(source);
427   free(platforms);
428   free(devices);
429
430}
431
432
433char* readSource(const char *sourceFilename) {
434
435   FILE *fp;
436   int err;
437   int size;
438
439   char *source;
440
441   fp = fopen(sourceFilename, "rb");
442   if(fp == NULL) {
443      printf("Could not open kernel file: %s\n", sourceFilename);
444      exit(-1);
445   }
446   
447   err = fseek(fp, 0, SEEK_END);
448   if(err != 0) {
449      printf("Error seeking to end of file\n");
450      exit(-1);
451   }
452
453   size = ftell(fp);
454   if(size < 0) {
455      printf("Error getting file position\n");
456      exit(-1);
457   }
458
459   err = fseek(fp, 0, SEEK_SET);
460   if(err != 0) {
461      printf("Error seeking to start of file\n");
462      exit(-1);
463   }
464
465   source = (char*)malloc(size+1);
466   if(source == NULL) {
467      printf("Error allocating %d bytes for the program source\n", size+1);
468      exit(-1);
469   }
470
471   err = fread(source, 1, size, fp);
472   if(err != size) {
473      printf("only read %d bytes\n", err);
474      exit(0);
475   }
476
477   source[size] = '\0';
478
479   return source;
480}
Note: See TracBrowser for help on using the repository browser.