clray

view src/ocl.cc @ 61:0f174fd60f19

added readme file and license
author John Tsiombikas <nuclear@member.fsf.org>
date Mon, 28 Dec 2015 10:27:23 +0200
parents 8047637961a2
children
line source
1 #define OCL_CC_
3 #define CL_USE_DEPRECATED_OPENCL_1_1_APIS
5 #include <stdio.h>
6 #include <stdlib.h>
7 #include <string.h>
8 #include <stdarg.h>
9 #include <errno.h>
10 #include <assert.h>
11 #ifndef _MSC_VER
12 #include <alloca.h>
13 #else
14 #include <malloc.h>
15 #endif
16 #include <sys/stat.h>
17 #include "ocl.h"
18 #include "ogl.h"
19 #include "ocl_errstr.h"
21 #if defined(unix) || defined(__unix__)
22 #include <X11/Xlib.h>
23 #include <GL/glx.h>
24 #endif
26 #ifdef __APPLE__
27 #include <OpenGL/CGLCurrent.h>
28 #endif
31 struct device_info {
32 cl_device_id id;
33 cl_device_type type;
34 unsigned int units;
35 unsigned int clock;
37 unsigned int dim;
38 size_t *work_item_sizes;
39 size_t work_group_size;
41 unsigned long mem_size;
43 char *extensions;
44 bool gl_sharing;
45 };
47 static int select_device(struct device_info *di, int (*devcmp)(struct device_info*, struct device_info*));
48 static int get_dev_info(cl_device_id dev, struct device_info *di);
49 static void destroy_dev_info(struct device_info *di);
50 static int devcmp(struct device_info *a, struct device_info *b);
51 static const char *devtypestr(cl_device_type type);
52 static void print_memsize(FILE *out, unsigned long memsz);
53 static const char *clstrerror(int err);
56 static cl_context ctx;
57 static cl_command_queue cmdq;
58 static device_info devinf;
60 bool init_opencl()
61 {
62 if(select_device(&devinf, devcmp) == -1) {
63 return false;
64 }
66 #ifndef CLGL_INTEROP
67 cl_context_properties *prop = 0;
68 #else
70 #if defined(__APPLE__)
71 CGLContextObj glctx = CGLGetCurrentContext();
72 CGLShareGroupObj sgrp = CGLGetShareGroup(glctx);
74 cl_context_properties prop[] = {
75 #ifdef CL_CONTEXT_PROPERTY_USE_CGL_SHAREGROUP_APPLE
76 CL_CONTEXT_PROPERTY_USE_CGL_SHAREGROUP_APPLE, (cl_context_properties)sgrp,
77 #else
78 CL_GL_CONTEXT_KHR, (cl_context_properties)glctx,
79 CL_CGL_SHAREGROUP_KHR, (cl_context_properties)sgrp,
80 #endif
81 0
82 };
83 #elif defined(unix) || defined(__unix__)
84 Display *dpy = glXGetCurrentDisplay();
85 GLXContext glctx = glXGetCurrentContext();
87 assert(dpy && glctx);
89 cl_context_properties prop[] = {
90 CL_GLX_DISPLAY_KHR, (cl_context_properties)dpy,
91 CL_GL_CONTEXT_KHR, (cl_context_properties)glctx,
92 0
93 };
94 #elif defined(WIN32) || defined(__WIN32__)
95 HGLRC glctx = wglGetCurrentContext();
96 HDC dc = wglGetCurrentDC();
98 cl_context_properties prop[] = {
99 CL_GL_CONTEXT_KHR, (cl_context_properties)glctx,
100 CL_WGL_HDC_KHR, (cl_context_properties)dc,
101 0
102 };
103 #else
104 #error "unknown or unsupported platform"
105 #endif
107 #endif /* CLGL_INTEROP */
109 if(!(ctx = clCreateContext(prop, 1, &devinf.id, 0, 0, 0))) {
110 fprintf(stderr, "failed to create opencl context\n");
111 return false;
112 }
114 if(!(cmdq = clCreateCommandQueue(ctx, devinf.id, 0, 0))) {
115 fprintf(stderr, "failed to create command queue\n");
116 return false;
117 }
118 return true;
119 }
121 void destroy_opencl()
122 {
123 if(cmdq) {
124 clReleaseCommandQueue(cmdq);
125 cmdq = 0;
126 }
128 if(ctx) {
129 clReleaseContext(ctx);
130 ctx = 0;
131 }
132 }
134 void finish_opencl()
135 {
136 clFinish(cmdq);
137 }
140 CLMemBuffer *create_mem_buffer(int rdwr, size_t sz, const void *buf)
141 {
142 int err;
143 cl_mem mem;
144 cl_mem_flags flags = rdwr | CL_MEM_ALLOC_HOST_PTR;
146 if(buf) {
147 flags |= CL_MEM_COPY_HOST_PTR;
148 }
151 if(!(mem = clCreateBuffer(ctx, flags, sz, (void*)buf, &err))) {
152 fprintf(stderr, "failed to create memory buffer: %s\n", clstrerror(err));
153 return 0;
154 }
156 CLMemBuffer *mbuf = new CLMemBuffer;
157 mbuf->type = MEM_BUFFER;
158 mbuf->mem = mem;
159 mbuf->size = sz;
160 mbuf->xsz = mbuf->ysz = 0;
161 mbuf->ptr = 0;
162 mbuf->tex = 0;
163 return mbuf;
164 }
166 CLMemBuffer *create_image_buffer(int rdwr, int xsz, int ysz, const void *pixels)
167 {
168 int err, pitch;
169 cl_mem mem;
170 cl_mem_flags flags = rdwr | CL_MEM_ALLOC_HOST_PTR;
172 if(pixels) {
173 flags |= CL_MEM_COPY_HOST_PTR;
174 pitch = xsz * 4 * sizeof(float);
175 } else {
176 pitch = 0;
177 }
179 cl_image_format fmt = {CL_RGBA, CL_FLOAT};
181 if(!(mem = clCreateImage2D(ctx, flags, &fmt, xsz, ysz, pitch, (void*)pixels, &err))) {
182 fprintf(stderr, "failed to create %dx%d image: %s\n", xsz, ysz, clstrerror(err));
183 return 0;
184 }
186 CLMemBuffer *mbuf = new CLMemBuffer;
187 mbuf->type = IMAGE_BUFFER;
188 mbuf->mem = mem;
189 mbuf->size = ysz * pitch;
190 mbuf->xsz = xsz;
191 mbuf->ysz = ysz;
192 mbuf->ptr = 0;
193 mbuf->tex = 0;
194 return mbuf;
195 }
197 CLMemBuffer *create_image_buffer(int rdwr, unsigned int tex)
198 {
199 int err, xsz, ysz;
200 cl_mem mem;
202 glGetError(); // clear previous OpenGL errors
204 glPushAttrib(GL_TEXTURE_BIT);
205 glBindTexture(GL_TEXTURE_2D, tex);
206 glGetTexLevelParameteriv(GL_TEXTURE_2D, 0, GL_TEXTURE_WIDTH, &xsz);
207 glGetTexLevelParameteriv(GL_TEXTURE_2D, 0, GL_TEXTURE_HEIGHT, &ysz);
208 glPopAttrib();
210 if(glGetError()) {
211 fprintf(stderr, "create_image_buffer: GL error while retreiving texture parameters for texture %u\n", tex);
212 return 0;
213 }
215 if(!(mem = clCreateFromGLTexture2D(ctx, rdwr, GL_TEXTURE_2D, 0, tex, &err))) {
216 fprintf(stderr, "failed to create memory buffer from GL texture %u: %s\n", tex, clstrerror(err));
217 return 0;
218 }
220 CLMemBuffer *mbuf = new CLMemBuffer;
221 mbuf->type = IMAGE_BUFFER;
222 mbuf->mem = mem;
223 mbuf->size = 0;
224 mbuf->xsz = xsz;
225 mbuf->ysz = ysz;
226 mbuf->ptr = 0;
227 mbuf->tex = tex;
229 return mbuf;
230 }
232 void destroy_mem_buffer(CLMemBuffer *mbuf)
233 {
234 if(mbuf) {
235 clReleaseMemObject(mbuf->mem);
236 delete mbuf;
237 }
238 }
240 void *map_mem_buffer(CLMemBuffer *mbuf, int rdwr, cl_event *ev)
241 {
242 if(!mbuf) return 0;
244 #ifndef NDEBUG
245 if(mbuf->ptr) {
246 fprintf(stderr, "WARNING: map_mem_buffer called on already mapped buffer\n");
247 }
248 #endif
250 int err;
252 if(mbuf->type == MEM_BUFFER) {
253 mbuf->ptr = clEnqueueMapBuffer(cmdq, mbuf->mem, 1, rdwr, 0, mbuf->size, 0, 0, ev, &err);
254 if(!mbuf->ptr) {
255 fprintf(stderr, "failed to map buffer: %s\n", clstrerror(err));
256 return 0;
257 }
258 } else {
259 assert(mbuf->type == IMAGE_BUFFER);
261 size_t orig[] = {0, 0, 0};
262 size_t rgn[] = {mbuf->xsz, mbuf->ysz, 1};
263 size_t pitch;
265 mbuf->ptr = clEnqueueMapImage(cmdq, mbuf->mem, 1, rdwr, orig, rgn, &pitch, 0, 0, 0, ev, &err);
266 if(!mbuf->ptr) {
267 fprintf(stderr, "failed to map image: %s\n", clstrerror(err));
268 return 0;
269 }
271 assert(pitch == mbuf->xsz * 4 * sizeof(float));
272 }
273 return mbuf->ptr;
274 }
276 void unmap_mem_buffer(CLMemBuffer *mbuf, cl_event *ev)
277 {
278 if(!mbuf || !mbuf->ptr) return;
280 clEnqueueUnmapMemObject(cmdq, mbuf->mem, mbuf->ptr, 0, 0, ev);
281 mbuf->ptr = 0;
282 }
284 bool write_mem_buffer(CLMemBuffer *mbuf, size_t sz, const void *src, cl_event *ev)
285 {
286 if(!mbuf) return false;
288 int err;
289 if((err = clEnqueueWriteBuffer(cmdq, mbuf->mem, 1, 0, sz, src, 0, 0, ev)) != 0) {
290 fprintf(stderr, "failed to write buffer: %s\n", clstrerror(err));
291 return false;
292 }
293 return true;
294 }
296 bool read_mem_buffer(CLMemBuffer *mbuf, size_t sz, void *dest, cl_event *ev)
297 {
298 if(!mbuf) return false;
300 int err;
301 if((err = clEnqueueReadBuffer(cmdq, mbuf->mem, 1, 0, sz, dest, 0, 0, ev)) != 0) {
302 fprintf(stderr, "failed to read buffer: %s\n", clstrerror(err));
303 return false;
304 }
305 return true;
306 }
309 bool acquire_gl_object(CLMemBuffer *mbuf, cl_event *ev)
310 {
311 if(!mbuf || !mbuf->tex) {
312 return false;
313 }
315 int err;
316 if((err = clEnqueueAcquireGLObjects(cmdq, 1, &mbuf->mem, 0, 0, ev)) != 0) {
317 fprintf(stderr, "failed to acquire gl object: %s\n", clstrerror(err));
318 return false;
319 }
320 return true;
321 }
323 bool release_gl_object(CLMemBuffer *mbuf, cl_event *ev)
324 {
325 if(!mbuf || !mbuf->tex) {
326 return false;
327 }
329 int err;
330 if((err = clEnqueueReleaseGLObjects(cmdq, 1, &mbuf->mem, 0, 0, ev)) != 0) {
331 fprintf(stderr, "failed to release gl object: %s\n", clstrerror(err));
332 return false;
333 }
334 return true;
335 }
338 CLArg::CLArg()
339 {
340 memset(this, 0, sizeof *this);
341 }
344 CLProgram::CLProgram(const char *kname)
345 {
346 prog = 0;
347 kernel = 0;
348 this->kname = kname;
349 args.resize(16);
350 built = false;
352 wait_event = last_event = 0;
353 }
355 CLProgram::~CLProgram()
356 {
357 if(wait_event) {
358 clReleaseEvent(wait_event);
359 }
360 if(last_event) {
361 clWaitForEvents(1, &last_event);
362 clReleaseEvent(last_event);
363 }
365 if(prog) {
366 clReleaseProgram(prog);
367 }
368 if(kernel) {
369 clReleaseKernel(kernel);
370 }
371 for(size_t i=0; i<args.size(); i++) {
372 if(args[i].type == ARGTYPE_MEM_BUF) {
373 destroy_mem_buffer(args[i].v.mbuf);
374 }
375 }
376 }
378 bool CLProgram::load(const char *fname)
379 {
380 FILE *fp;
381 char *src;
382 struct stat st;
384 printf("loading opencl program (%s)\n", fname);
386 if(!(fp = fopen(fname, "rb"))) {
387 fprintf(stderr, "failed to open %s: %s\n", fname, strerror(errno));
388 return false;
389 }
391 fstat(fileno(fp), &st);
393 src = new char[st.st_size + 1];
395 fread(src, 1, st.st_size, fp);
396 src[st.st_size] = 0;
397 fclose(fp);
400 if(!(prog = clCreateProgramWithSource(ctx, 1, (const char**)&src, 0, 0))) {
401 fprintf(stderr, "error creating program object: %s\n", fname);
402 delete [] src;
403 return false;
404 }
405 delete [] src;
406 return true;
407 }
409 bool CLProgram::set_argi(int idx, int val)
410 {
411 if((int)args.size() <= idx) {
412 args.resize(idx + 1);
413 }
415 CLArg *arg = &args[idx];
416 arg->type = ARGTYPE_INT;
417 arg->v.ival = val;
418 return true;
419 }
421 bool CLProgram::set_argf(int idx, float val)
422 {
423 if((int)args.size() <= idx) {
424 args.resize(idx + 1);
425 }
427 CLArg *arg = &args[idx];
428 arg->type = ARGTYPE_FLOAT;
429 arg->v.fval = val;
430 return true;
431 }
433 bool CLProgram::set_arg_buffer(int idx, int rdwr, size_t sz, const void *ptr)
434 {
435 printf("create argument %d buffer: %d bytes\n", idx, (int)sz);
436 CLMemBuffer *buf;
438 if(sz <= 0) {
439 fprintf(stderr, "invalid size while creating argument buffer %d: %d bytes\n", idx, (int)sz);
440 return false;
441 }
442 if(!(buf = create_mem_buffer(rdwr, sz, ptr))) {
443 return false;
444 }
446 if((int)args.size() <= idx) {
447 args.resize(idx + 1);
448 }
449 args[idx].type = ARGTYPE_MEM_BUF;
450 args[idx].v.mbuf = buf;
451 return true;
452 }
454 bool CLProgram::set_arg_image(int idx, int rdwr, int xsz, int ysz, const void *pix)
455 {
456 printf("create argument %d from %dx%d image\n", idx, xsz, ysz);
457 CLMemBuffer *buf;
459 if(!(buf = create_image_buffer(rdwr, xsz, ysz, pix))) {
460 return false;
461 }
463 if((int)args.size() <= idx) {
464 args.resize(idx + 1);
465 }
466 args[idx].type = ARGTYPE_MEM_BUF;
467 args[idx].v.mbuf = buf;
468 return true;
469 }
471 bool CLProgram::set_arg_texture(int idx, int rdwr, unsigned int tex)
472 {
473 printf("create argument %d from texture %u\n", idx, tex);
474 CLMemBuffer *buf;
476 if(!(buf = create_image_buffer(rdwr, tex))) {
477 return false;
478 }
480 if((int)args.size() <= idx) {
481 args.resize(idx + 1);
482 }
483 args[idx].type = ARGTYPE_MEM_BUF;
484 args[idx].v.mbuf = buf;
485 return true;
486 }
488 CLMemBuffer *CLProgram::get_arg_buffer(int arg)
489 {
490 if(arg < 0 || arg >= (int)args.size() || args[arg].type != ARGTYPE_MEM_BUF) {
491 return 0;
492 }
493 return args[arg].v.mbuf;
494 }
496 int CLProgram::get_num_args() const
497 {
498 int num_args = 0;
499 for(size_t i=0; i<args.size(); i++) {
500 if(args[i].type != ARGTYPE_NONE) {
501 num_args++;
502 }
503 }
504 return num_args;
505 }
507 bool CLProgram::build(const char *opt)
508 {
509 int err;
510 if((err = clBuildProgram(prog, 0, 0, opt, 0, 0)) != 0) {
511 size_t sz;
512 clGetProgramBuildInfo(prog, devinf.id, CL_PROGRAM_BUILD_LOG, 0, 0, &sz);
514 char *errlog = (char*)alloca(sz + 1);
515 clGetProgramBuildInfo(prog, devinf.id, CL_PROGRAM_BUILD_LOG, sz, errlog, 0);
516 fprintf(stderr, "failed to build program: %s\n%s\n", clstrerror(err), errlog);
518 clReleaseProgram(prog);
519 prog = 0;
520 return false;
521 }
524 if(!(kernel = clCreateKernel(prog, kname.c_str(), 0))) {
525 fprintf(stderr, "failed to create kernel: %s\n", kname.c_str());
526 clReleaseProgram(prog);
527 prog = 0;
528 return false;
529 }
531 for(size_t i=0; i<args.size(); i++) {
532 int err;
534 if(args[i].type == ARGTYPE_NONE) {
535 break;
536 }
538 switch(args[i].type) {
539 case ARGTYPE_INT:
540 if((err = clSetKernelArg(kernel, i, sizeof(int), &args[i].v.ival)) != 0) {
541 fprintf(stderr, "failed to bind kernel argument %d: %s\n", (int)i, clstrerror(err));
542 goto fail;
543 }
544 break;
546 case ARGTYPE_FLOAT:
547 if((err = clSetKernelArg(kernel, i, sizeof(float), &args[i].v.fval)) != 0) {
548 fprintf(stderr, "failed to bind kernel argument %d: %s\n", (int)i, clstrerror(err));
549 goto fail;
550 }
551 break;
553 case ARGTYPE_MEM_BUF:
554 {
555 CLMemBuffer *mbuf = args[i].v.mbuf;
557 if((err = clSetKernelArg(kernel, i, sizeof mbuf->mem, &mbuf->mem)) != 0) {
558 fprintf(stderr, "failed to bind kernel argument %d: %s\n", (int)i, clstrerror(err));
559 goto fail;
560 }
561 }
562 break;
564 default:
565 break;
566 }
567 }
569 built = true;
570 return true;
572 fail:
573 clReleaseProgram(prog);
574 clReleaseKernel(kernel);
575 prog = 0;
576 kernel = 0;
577 return false;
578 }
580 bool CLProgram::run() const
581 {
582 return run(1, 1);
583 }
585 bool CLProgram::run(int dim, ...) const
586 {
587 va_list ap;
588 size_t *global_size = (size_t*)alloca(dim * sizeof *global_size);
590 va_start(ap, dim);
591 for(int i=0; i<dim; i++) {
592 global_size[i] = va_arg(ap, int);
593 }
594 va_end(ap);
596 if(last_event) {
597 clReleaseEvent(last_event);
598 }
600 int err;
601 if((err = clEnqueueNDRangeKernel(cmdq, kernel, dim, 0, global_size, 0,
602 wait_event ? 1 : 0, wait_event ? &wait_event : 0, &last_event)) != 0) {
603 fprintf(stderr, "error executing kernel: %s\n", clstrerror(err));
604 return false;
605 }
607 if(wait_event) {
608 clReleaseEvent(wait_event);
609 wait_event = 0;
610 }
611 return true;
612 }
614 void CLProgram::set_wait_event(cl_event ev)
615 {
616 if(wait_event) {
617 clReleaseEvent(wait_event);
618 }
619 wait_event = ev;
620 }
622 cl_event CLProgram::get_last_event() const
623 {
624 return last_event;
625 }
627 static int select_device(struct device_info *dev_inf, int (*devcmp)(struct device_info*, struct device_info*))
628 {
629 unsigned int i, j, num_dev, num_plat, sel, ret;
630 cl_device_id dev[32];
631 cl_platform_id plat[32];
633 dev_inf->work_item_sizes = 0;
635 if((ret = clGetPlatformIDs(32, plat, &num_plat)) != 0) {
636 fprintf(stderr, "clGetPlatformIDs failed: %s\n", clstrerror(ret));
637 return -1;
638 }
639 if(!num_plat) {
640 fprintf(stderr, "OpenCL not available!\n");
641 return -1;
642 }
644 for(i=0; i<num_plat; i++) {
645 char buf[512];
647 clGetPlatformInfo(plat[i], CL_PLATFORM_NAME, sizeof buf, buf, 0);
648 printf("[%d]: %s", i, buf);
649 clGetPlatformInfo(plat[i], CL_PLATFORM_VENDOR, sizeof buf, buf, 0);
650 printf(", %s", buf);
651 clGetPlatformInfo(plat[i], CL_PLATFORM_VERSION, sizeof buf, buf, 0);
652 printf(" (%s)\n", buf);
653 }
655 if((ret = clGetDeviceIDs(plat[0], CL_DEVICE_TYPE_ALL, 32, dev, &num_dev)) != 0) {
656 fprintf(stderr, "clGetDeviceIDs failed: %s\n", clstrerror(ret));
657 return -1;
658 }
659 printf("found %d cl devices.\n", num_dev);
661 for(i=0; i<num_dev; i++) {
662 struct device_info di;
664 if(get_dev_info(dev[i], &di) == -1) {
665 destroy_dev_info(&di);
666 return -1;
667 }
669 printf("--> device %u (%s)\n", i, devtypestr(di.type));
670 printf("max compute units: %u\n", di.units);
671 printf("max clock frequency: %u\n", di.clock);
672 printf("max work item dimensions: %u\n", di.dim);
674 printf("max work item sizes: ");
675 for(j=0; j<di.dim; j++) {
676 printf("%u", (unsigned int)di.work_item_sizes[j]);
677 if(di.dim - j > 1) {
678 printf(", ");
679 }
680 }
681 putchar('\n');
683 printf("max work group size: %u\n", (unsigned int)di.work_group_size);
684 printf("max object allocation size: ");
685 print_memsize(stdout, di.mem_size);
686 putchar('\n');
688 printf("extensions: %s\n", di.extensions);
690 if(devcmp(&di, dev_inf) > 0) {
691 free(dev_inf->work_item_sizes);
692 memcpy(dev_inf, &di, sizeof di);
693 sel = i;
694 }
695 }
697 if(num_dev) {
698 printf("\nusing device: %d\n", sel);
699 return 0;
700 }
702 return -1;
703 }
705 static int get_dev_info(cl_device_id dev, struct device_info *di)
706 {
707 di->id = dev;
709 clGetDeviceInfo(dev, CL_DEVICE_TYPE, sizeof di->type, &di->type, 0);
710 clGetDeviceInfo(dev, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof di->units, &di->units, 0);
711 clGetDeviceInfo(dev, CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof di->clock, &di->clock, 0);
712 clGetDeviceInfo(dev, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, sizeof di->dim, &di->dim, 0);
714 di->work_item_sizes = new size_t[di->dim];
716 clGetDeviceInfo(dev, CL_DEVICE_MAX_WORK_ITEM_SIZES, di->dim * sizeof *di->work_item_sizes, di->work_item_sizes, 0);
717 clGetDeviceInfo(dev, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof di->work_group_size, &di->work_group_size, 0);
718 clGetDeviceInfo(dev, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof di->mem_size, &di->mem_size, 0);
720 size_t ext_str_len;
721 clGetDeviceInfo(dev, CL_DEVICE_EXTENSIONS, 0, 0, &ext_str_len);
723 di->extensions = new char[ext_str_len + 1];
724 clGetDeviceInfo(dev, CL_DEVICE_EXTENSIONS, ext_str_len, di->extensions, 0);
725 di->extensions[ext_str_len] = 0;
727 if(strstr(di->extensions, "cl_khr_gl_sharing") || strstr(di->extensions, "cl_APPLE_gl_sharing")) {
728 di->gl_sharing = true;
729 } else {
730 di->gl_sharing = false;
731 }
733 return 0;
734 }
736 static void destroy_dev_info(struct device_info *di)
737 {
738 delete [] di->work_item_sizes;
739 delete [] di->extensions;
740 }
742 static int devcmp(struct device_info *a, struct device_info *b)
743 {
744 unsigned int aval = a->units * a->clock;
745 unsigned int bval = b->units * b->clock;
747 return aval - bval;
748 }
750 static const char *devtypestr(cl_device_type type)
751 {
752 switch(type) {
753 case CL_DEVICE_TYPE_CPU:
754 return "cpu";
755 case CL_DEVICE_TYPE_GPU:
756 return "gpu";
757 case CL_DEVICE_TYPE_ACCELERATOR:
758 return "accelerator";
759 default:
760 break;
761 }
762 return "unknown";
763 }
765 static void print_memsize(FILE *out, unsigned long bytes)
766 {
767 int i;
768 unsigned long memsz = bytes;
769 const char *suffix[] = {"bytes", "kb", "mb", "gb", "tb", "pb", 0};
771 for(i=0; suffix[i]; i++) {
772 if(memsz < 1024) {
773 fprintf(out, "%lu %s", memsz, suffix[i]);
774 if(i > 0) {
775 fprintf(out, " (%lu bytes)", bytes);
776 }
777 return;
778 }
780 memsz /= 1024;
781 }
782 }
784 static const char *clstrerror(int err)
785 {
786 if(err > 0) {
787 return "<invalid error code>";
788 }
789 if(err <= -(int)(sizeof ocl_errstr / sizeof *ocl_errstr)) {
790 return "<unknown error>";
791 }
792 return ocl_errstr[-err];
793 }