source: GPU/xmlwf.cpp @ 5817

Last change on this file since 5817 was 1668, checked in by lindanl, 7 years ago

Parabix on GPU : start tag parsing

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