clray

annotate src/ocl.cc @ 45:8047637961a2

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