clray: e3b4457dc4d2 src/ocl.cc

clray

view src/ocl.cc @ 56:e3b4457dc4d2

added glFinish after swap-buffers to make the program absolutely correct in regards to mediating usage of the shared GL/CL texture image

author	John Tsiombikas <nuclear@member.fsf.org>
date	Mon, 06 Sep 2010 05:40:47 +0100
parents	f9eec11e5acc
children	8c858e1a89e8

line source

1 #define OCL_CC_

3 #include <stdio.h>

4 #include <stdlib.h>

5 #include <string.h>

6 #include <stdarg.h>

7 #include <errno.h>

8 #include <assert.h>

9 #ifndef _MSC_VER

10 #include <alloca.h>

11 #else

12 #include <malloc.h>

13 #endif

14 #include <sys/stat.h>

15 #include "ocl.h"

16 #include "ogl.h"

17 #include "ocl_errstr.h"

19 #if defined(unix) || defined(__unix__)

20 #include <X11/Xlib.h>

21 #include <GL/glx.h>

22 #endif

24 #ifdef __APPLE__

25 #include <OpenGL/CGLCurrent.h>

26 #endif

29 struct device_info {

30 cl_device_id id;

31 cl_device_type type;

32 unsigned int units;

33 unsigned int clock;

35 unsigned int dim;

36 size_t *work_item_sizes;

37 size_t work_group_size;

39 unsigned long mem_size;

41 char *extensions;

42 bool gl_sharing;

43 };

45 static int select_device(struct device_info *di, int (*devcmp)(struct device_info*, struct device_info*));

46 static int get_dev_info(cl_device_id dev, struct device_info *di);

47 static void destroy_dev_info(struct device_info *di);

48 static int devcmp(struct device_info *a, struct device_info *b);

49 static const char *devtypestr(cl_device_type type);

50 static void print_memsize(FILE *out, unsigned long memsz);

51 static const char *clstrerror(int err);

54 static cl_context ctx;

55 static cl_command_queue cmdq;

56 static device_info devinf;

58 bool init_opencl()

59 {

60 if(select_device(&devinf, devcmp) == -1) {

61 return false;

62 }

64 #ifndef CLGL_INTEROP

65 cl_context_properties *prop = 0;

66 #else

68 #if defined(__APPLE__)

69 CGLContextObj glctx = CGLGetCurrentContext();

70 CGLShareGroupObj sgrp = CGLGetShareGroup(glctx);

72 cl_context_properties prop[] = {

73 #ifdef CL_CONTEXT_PROPERTY_USE_CGL_SHAREGROUP_APPLE

74 CL_CONTEXT_PROPERTY_USE_CGL_SHAREGROUP_APPLE, (cl_context_properties)sgrp,

75 #else

76 CL_GL_CONTEXT_KHR, (cl_context_properties)glctx,

77 CL_CGL_SHAREGROUP_KHR, (cl_context_properties)sgrp,

78 #endif

79 0

80 };

81 #elif defined(unix) || defined(__unix__)

82 Display *dpy = glXGetCurrentDisplay();

83 GLXContext glctx = glXGetCurrentContext();

85 assert(dpy && glctx);

87 cl_context_properties prop[] = {

88 CL_GLX_DISPLAY_KHR, (cl_context_properties)dpy,

89 CL_GL_CONTEXT_KHR, (cl_context_properties)glctx,

90 0

91 };

92 #elif defined(WIN32) || defined(__WIN32__)

93 HGLRC glctx = wglGetCurrentContext();

94 HDC dc = wglGetCurrentDC();

96 cl_context_properties prop[] = {

97 CL_GL_CONTEXT_KHR, (cl_context_properties)glctx,

98 CL_WGL_HDC_KHR, (cl_context_properties)dc,

99 0

100 };

101 #else

102 #error "unknown or unsupported platform"

103 #endif

104

105 #endif /* CLGL_INTEROP */

106

107 if(!(ctx = clCreateContext(prop, 1, &devinf.id, 0, 0, 0))) {

108 fprintf(stderr, "failed to create opencl context\n");

109 return false;

110 }

111

112 if(!(cmdq = clCreateCommandQueue(ctx, devinf.id, 0, 0))) {

113 fprintf(stderr, "failed to create command queue\n");

114 return false;

115 }

116 return true;

117 }

118

119 void destroy_opencl()

120 {

121 if(cmdq) {

122 clReleaseCommandQueue(cmdq);

123 cmdq = 0;

124 }

125

126 if(ctx) {

127 clReleaseContext(ctx);

128 ctx = 0;

129 }

130 }

131

132

133 CLMemBuffer *create_mem_buffer(int rdwr, size_t sz, const void *buf)

134 {

135 int err;

136 cl_mem mem;

137 cl_mem_flags flags = rdwr | CL_MEM_ALLOC_HOST_PTR;

138

139 if(buf) {

140 flags |= CL_MEM_COPY_HOST_PTR;

141 }

142

143

144 if(!(mem = clCreateBuffer(ctx, flags, sz, (void*)buf, &err))) {

145 fprintf(stderr, "failed to create memory buffer: %s\n", clstrerror(err));

146 return 0;

147 }

148

149 CLMemBuffer *mbuf = new CLMemBuffer;

150 mbuf->type = MEM_BUFFER;

151 mbuf->mem = mem;

152 mbuf->size = sz;

153 mbuf->xsz = mbuf->ysz = 0;

154 mbuf->ptr = 0;

155 mbuf->tex = 0;

156 return mbuf;

157 }

158

159 CLMemBuffer *create_image_buffer(int rdwr, int xsz, int ysz, const void *pixels)

160 {

161 int err, pitch;

162 cl_mem mem;

163 cl_mem_flags flags = rdwr | CL_MEM_ALLOC_HOST_PTR;

164

165 if(pixels) {

166 flags |= CL_MEM_COPY_HOST_PTR;

167 pitch = xsz * 4 * sizeof(float);

168 } else {

169 pitch = 0;

170 }

171

172 cl_image_format fmt = {CL_RGBA, CL_FLOAT};

173

174 if(!(mem = clCreateImage2D(ctx, flags, &fmt, xsz, ysz, pitch, (void*)pixels, &err))) {

175 fprintf(stderr, "failed to create %dx%d image: %s\n", xsz, ysz, clstrerror(err));

176 return 0;

177 }

178

179 CLMemBuffer *mbuf = new CLMemBuffer;

180 mbuf->type = IMAGE_BUFFER;

181 mbuf->mem = mem;

182 mbuf->size = ysz * pitch;

183 mbuf->xsz = xsz;

184 mbuf->ysz = ysz;

185 mbuf->ptr = 0;

186 mbuf->tex = 0;

187 return mbuf;

188 }

189

190 CLMemBuffer *create_image_buffer(int rdwr, unsigned int tex)

191 {

192 int err, xsz, ysz;

193 cl_mem mem;

194

195 glGetError(); // clear previous OpenGL errors

196

197 glPushAttrib(GL_TEXTURE_BIT);

198 glBindTexture(GL_TEXTURE_2D, tex);

199 glGetTexLevelParameteriv(GL_TEXTURE_2D, 0, GL_TEXTURE_WIDTH, &xsz);

200 glGetTexLevelParameteriv(GL_TEXTURE_2D, 0, GL_TEXTURE_HEIGHT, &ysz);

201 glPopAttrib();

202

203 if(glGetError()) {

204 fprintf(stderr, "create_image_buffer: GL error while retreiving texture parameters for texture %u\n", tex);

205 return 0;

206 }

207

208 if(!(mem = clCreateFromGLTexture2D(ctx, rdwr, GL_TEXTURE_2D, 0, tex, &err))) {

209 fprintf(stderr, "failed to create memory buffer from GL texture %u: %s\n", tex, clstrerror(err));

210 return 0;

211 }

212

213 CLMemBuffer *mbuf = new CLMemBuffer;

214 mbuf->type = IMAGE_BUFFER;

215 mbuf->mem = mem;

216 mbuf->size = 0;

217 mbuf->xsz = xsz;

218 mbuf->ysz = ysz;

219 mbuf->ptr = 0;

220 mbuf->tex = tex;

221

222 return mbuf;

223 }

224

225 void destroy_mem_buffer(CLMemBuffer *mbuf)

226 {

227 if(mbuf) {

228 clReleaseMemObject(mbuf->mem);

229 delete mbuf;

230 }

231 }

232

233 void *map_mem_buffer(CLMemBuffer *mbuf, int rdwr, cl_event *ev)

234 {

235 if(!mbuf) return 0;

236

237 #ifndef NDEBUG

238 if(mbuf->ptr) {

239 fprintf(stderr, "WARNING: map_mem_buffer called on already mapped buffer\n");

240 }

241 #endif

242

243 int err;

244

245 if(mbuf->type == MEM_BUFFER) {

246 mbuf->ptr = clEnqueueMapBuffer(cmdq, mbuf->mem, 1, rdwr, 0, mbuf->size, 0, 0, ev, &err);

247 if(!mbuf->ptr) {

248 fprintf(stderr, "failed to map buffer: %s\n", clstrerror(err));

249 return 0;

250 }

251 } else {

252 assert(mbuf->type == IMAGE_BUFFER);

253

254 size_t orig[] = {0, 0, 0};

255 size_t rgn[] = {mbuf->xsz, mbuf->ysz, 1};

256 size_t pitch;

257

258 mbuf->ptr = clEnqueueMapImage(cmdq, mbuf->mem, 1, rdwr, orig, rgn, &pitch, 0, 0, 0, ev, &err);

259 if(!mbuf->ptr) {

260 fprintf(stderr, "failed to map image: %s\n", clstrerror(err));

261 return 0;

262 }

263

264 assert(pitch == mbuf->xsz * 4 * sizeof(float));

265 }

266 return mbuf->ptr;

267 }

268

269 void unmap_mem_buffer(CLMemBuffer *mbuf, cl_event *ev)

270 {

271 if(!mbuf || !mbuf->ptr) return;

272

273 clEnqueueUnmapMemObject(cmdq, mbuf->mem, mbuf->ptr, 0, 0, ev);

274 mbuf->ptr = 0;

275 }

276

277 bool write_mem_buffer(CLMemBuffer *mbuf, size_t sz, const void *src, cl_event *ev)

278 {

279 if(!mbuf) return false;

280

281 int err;

282 if((err = clEnqueueWriteBuffer(cmdq, mbuf->mem, 1, 0, sz, src, 0, 0, ev)) != 0) {

283 fprintf(stderr, "failed to write buffer: %s\n", clstrerror(err));

284 return false;

285 }

286 return true;

287 }

288

289 bool read_mem_buffer(CLMemBuffer *mbuf, size_t sz, void *dest, cl_event *ev)

290 {

291 if(!mbuf) return false;

292

293 int err;

294 if((err = clEnqueueReadBuffer(cmdq, mbuf->mem, 1, 0, sz, dest, 0, 0, ev)) != 0) {

295 fprintf(stderr, "failed to read buffer: %s\n", clstrerror(err));

296 return false;

297 }

298 return true;

299 }

300

301

302 bool acquire_gl_object(CLMemBuffer *mbuf, cl_event *ev)

303 {

304 if(!mbuf || !mbuf->tex) {

305 return false;

306 }

307

308 int err;

309 if((err = clEnqueueAcquireGLObjects(cmdq, 1, &mbuf->mem, 0, 0, ev)) != 0) {

310 fprintf(stderr, "failed to acquire gl object: %s\n", clstrerror(err));

311 return false;

312 }

313 return true;

314 }

315

316 bool release_gl_object(CLMemBuffer *mbuf, cl_event *ev)

317 {

318 if(!mbuf || !mbuf->tex) {

319 return false;

320 }

321

322 int err;

323 if((err = clEnqueueReleaseGLObjects(cmdq, 1, &mbuf->mem, 0, 0, ev)) != 0) {

324 fprintf(stderr, "failed to release gl object: %s\n", clstrerror(err));

325 return false;

326 }

327 return true;

328 }

329

330

331 CLArg::CLArg()

332 {

333 memset(this, 0, sizeof *this);

334 }

335

336

337 CLProgram::CLProgram(const char *kname)

338 {

339 prog = 0;

340 kernel = 0;

341 this->kname = kname;

342 args.resize(16);

343 built = false;

344

345 wait_event = last_event = 0;

346 }

347

348 CLProgram::~CLProgram()

349 {

350 if(wait_event) {

351 clReleaseEvent(wait_event);

352 }

353 if(last_event) {

354 clWaitForEvents(1, &last_event);

355 clReleaseEvent(last_event);

356 }

357

358 if(prog) {

359 clReleaseProgram(prog);

360 }

361 if(kernel) {

362 clReleaseKernel(kernel);

363 }

364 for(size_t i=0; i<args.size(); i++) {

365 if(args[i].type == ARGTYPE_MEM_BUF) {

366 destroy_mem_buffer(args[i].v.mbuf);

367 }

368 }

369 }

370

371 bool CLProgram::load(const char *fname)

372 {

373 FILE *fp;

374 char *src;

375 struct stat st;

376

377 printf("loading opencl program (%s)\n", fname);

378

379 if(!(fp = fopen(fname, "rb"))) {

380 fprintf(stderr, "failed to open %s: %s\n", fname, strerror(errno));

381 return false;

382 }

383

384 fstat(fileno(fp), &st);

385

386 src = new char[st.st_size + 1];

387

388 fread(src, 1, st.st_size, fp);

389 src[st.st_size] = 0;

390 fclose(fp);

391

392

393 if(!(prog = clCreateProgramWithSource(ctx, 1, (const char**)&src, 0, 0))) {

394 fprintf(stderr, "error creating program object: %s\n", fname);

395 delete [] src;

396 return false;

397 }

398 delete [] src;

399 return true;

400 }

401

402 bool CLProgram::set_argi(int idx, int val)

403 {

404 if((int)args.size() <= idx) {

405 args.resize(idx + 1);

406 }

407

408 CLArg *arg = &args[idx];

409 arg->type = ARGTYPE_INT;

410 arg->v.ival = val;

411 return true;

412 }

413

414 bool CLProgram::set_argf(int idx, float val)

415 {

416 if((int)args.size() <= idx) {

417 args.resize(idx + 1);

418 }

419

420 CLArg *arg = &args[idx];

421 arg->type = ARGTYPE_FLOAT;

422 arg->v.fval = val;

423 return true;

424 }

425

426 bool CLProgram::set_arg_buffer(int idx, int rdwr, size_t sz, const void *ptr)

427 {

428 printf("create argument %d buffer: %d bytes\n", idx, (int)sz);

429 CLMemBuffer *buf;

430

431 if(sz <= 0) {

432 fprintf(stderr, "invalid size while creating argument buffer %d: %d bytes\n", idx, (int)sz);

433 return false;

434 }

435 if(!(buf = create_mem_buffer(rdwr, sz, ptr))) {

436 return false;

437 }

438

439 if((int)args.size() <= idx) {

440 args.resize(idx + 1);

441 }

442 args[idx].type = ARGTYPE_MEM_BUF;

443 args[idx].v.mbuf = buf;

444 return true;

445 }

446

447 bool CLProgram::set_arg_image(int idx, int rdwr, int xsz, int ysz, const void *pix)

448 {

449 printf("create argument %d from %dx%d image\n", idx, xsz, ysz);

450 CLMemBuffer *buf;

451

452 if(!(buf = create_image_buffer(rdwr, xsz, ysz, pix))) {

453 return false;

454 }

455

456 if((int)args.size() <= idx) {

457 args.resize(idx + 1);

458 }

459 args[idx].type = ARGTYPE_MEM_BUF;

460 args[idx].v.mbuf = buf;

461 return true;

462 }

463

464 bool CLProgram::set_arg_texture(int idx, int rdwr, unsigned int tex)

465 {

466 printf("create argument %d from texture %u\n", idx, tex);

467 CLMemBuffer *buf;

468

469 if(!(buf = create_image_buffer(rdwr, tex))) {

470 return false;

471 }

472

473 if((int)args.size() <= idx) {

474 args.resize(idx + 1);

475 }

476 args[idx].type = ARGTYPE_MEM_BUF;

477 args[idx].v.mbuf = buf;

478 return true;

479 }

480

481 CLMemBuffer *CLProgram::get_arg_buffer(int arg)

482 {

483 if(arg < 0 || arg >= (int)args.size() || args[arg].type != ARGTYPE_MEM_BUF) {

484 return 0;

485 }

486 return args[arg].v.mbuf;

487 }

488

489 int CLProgram::get_num_args() const

490 {

491 int num_args = 0;

492 for(size_t i=0; i<args.size(); i++) {

493 if(args[i].type != ARGTYPE_NONE) {

494 num_args++;

495 }

496 }

497 return num_args;

498 }

499

500 bool CLProgram::build(const char *opt)

501 {

502 int err;

503 if((err = clBuildProgram(prog, 0, 0, opt, 0, 0)) != 0) {

504 size_t sz;

505 clGetProgramBuildInfo(prog, devinf.id, CL_PROGRAM_BUILD_LOG, 0, 0, &sz);

506

507 char *errlog = (char*)alloca(sz + 1);

508 clGetProgramBuildInfo(prog, devinf.id, CL_PROGRAM_BUILD_LOG, sz, errlog, 0);

509 fprintf(stderr, "failed to build program: %s\n%s\n", clstrerror(err), errlog);

510

511 clReleaseProgram(prog);

512 prog = 0;

513 return false;

514 }

515

516

517 if(!(kernel = clCreateKernel(prog, kname.c_str(), 0))) {

518 fprintf(stderr, "failed to create kernel: %s\n", kname.c_str());

519 clReleaseProgram(prog);

520 prog = 0;

521 return false;

522 }

523

524 for(size_t i=0; i<args.size(); i++) {

525 int err;

526

527 if(args[i].type == ARGTYPE_NONE) {

528 break;

529 }

530

531 switch(args[i].type) {

532 case ARGTYPE_INT:

533 if((err = clSetKernelArg(kernel, i, sizeof(int), &args[i].v.ival)) != 0) {

534 fprintf(stderr, "failed to bind kernel argument %d: %s\n", (int)i, clstrerror(err));

535 goto fail;

536 }

537 break;

538

539 case ARGTYPE_FLOAT:

540 if((err = clSetKernelArg(kernel, i, sizeof(float), &args[i].v.fval)) != 0) {

541 fprintf(stderr, "failed to bind kernel argument %d: %s\n", (int)i, clstrerror(err));

542 goto fail;

543 }

544 break;

545

546 case ARGTYPE_MEM_BUF:

547 {

548 CLMemBuffer *mbuf = args[i].v.mbuf;

549

550 if((err = clSetKernelArg(kernel, i, sizeof mbuf->mem, &mbuf->mem)) != 0) {

551 fprintf(stderr, "failed to bind kernel argument %d: %s\n", (int)i, clstrerror(err));

552 goto fail;

553 }

554 }

555 break;

556

557 default:

558 break;

559 }

560 }

561

562 built = true;

563 return true;

564

565 fail:

566 clReleaseProgram(prog);

567 clReleaseKernel(kernel);

568 prog = 0;

569 kernel = 0;

570 return false;

571 }

572

573 bool CLProgram::run() const

574 {

575 return run(1, 1);

576 }

577

578 bool CLProgram::run(int dim, ...) const

579 {

580 va_list ap;

581 size_t *global_size = (size_t*)alloca(dim * sizeof *global_size);

582

583 va_start(ap, dim);

584 for(int i=0; i<dim; i++) {

585 global_size[i] = va_arg(ap, int);

586 }

587 va_end(ap);

588

589 if(last_event) {

590 clReleaseEvent(last_event);

591 }

592

593 int err;

594 if((err = clEnqueueNDRangeKernel(cmdq, kernel, dim, 0, global_size, 0,

595 wait_event ? 1 : 0, wait_event ? &wait_event : 0, &last_event)) != 0) {

596 fprintf(stderr, "error executing kernel: %s\n", clstrerror(err));

597 return false;

598 }

599

600 if(wait_event) {

601 clReleaseEvent(wait_event);

602 wait_event = 0;

603 }

604 return true;

605 }

606

607 void CLProgram::set_wait_event(cl_event ev)

608 {

609 if(wait_event) {

610 clReleaseEvent(wait_event);

611 }

612 wait_event = ev;

613 }

614

615 cl_event CLProgram::get_last_event() const

616 {

617 return last_event;

618 }

619

620 static int select_device(struct device_info *dev_inf, int (*devcmp)(struct device_info*, struct device_info*))

621 {

622 unsigned int i, j, num_dev, num_plat, sel, ret;

623 cl_device_id dev[32];

624 cl_platform_id plat[32];

625

626 dev_inf->work_item_sizes = 0;

627

628 if((ret = clGetPlatformIDs(32, plat, &num_plat)) != 0) {

629 fprintf(stderr, "clGetPlatformIDs failed: %s\n", clstrerror(ret));

630 return -1;

631 }

632 if(!num_plat) {

633 fprintf(stderr, "OpenCL not available!\n");

634 return -1;

635 }

636

637 for(i=0; i<num_plat; i++) {

638 char buf[512];

639

640 clGetPlatformInfo(plat[i], CL_PLATFORM_NAME, sizeof buf, buf, 0);

641 printf("[%d]: %s", i, buf);

642 clGetPlatformInfo(plat[i], CL_PLATFORM_VENDOR, sizeof buf, buf, 0);

643 printf(", %s", buf);

644 clGetPlatformInfo(plat[i], CL_PLATFORM_VERSION, sizeof buf, buf, 0);

645 printf(" (%s)\n", buf);

646 }

647

648 if((ret = clGetDeviceIDs(plat[0], CL_DEVICE_TYPE_ALL, 32, dev, &num_dev)) != 0) {

649 fprintf(stderr, "clGetDeviceIDs failed: %s\n", clstrerror(ret));

650 return -1;

651 }

652 printf("found %d cl devices.\n", num_dev);

653

654 for(i=0; i<num_dev; i++) {

655 struct device_info di;

656

657 if(get_dev_info(dev[i], &di) == -1) {

658 destroy_dev_info(&di);

659 return -1;

660 }

661

662 printf("--> device %u (%s)\n", i, devtypestr(di.type));

663 printf("max compute units: %u\n", di.units);

664 printf("max clock frequency: %u\n", di.clock);

665 printf("max work item dimensions: %u\n", di.dim);

666

667 printf("max work item sizes: ");

668 for(j=0; j<di.dim; j++) {

669 printf("%u", (unsigned int)di.work_item_sizes[j]);

670 if(di.dim - j > 1) {

671 printf(", ");

672 }

673 }

674 putchar('\n');

675

676 printf("max work group size: %u\n", (unsigned int)di.work_group_size);

677 printf("max object allocation size: ");

678 print_memsize(stdout, di.mem_size);

679 putchar('\n');

680

681 printf("extensions: %s\n", di.extensions);

682

683 if(devcmp(&di, dev_inf) > 0) {

684 free(dev_inf->work_item_sizes);

685 memcpy(dev_inf, &di, sizeof di);

686 sel = i;

687 }

688 }

689

690 if(num_dev) {

691 printf("\nusing device: %d\n", sel);

692 return 0;

693 }

694

695 return -1;

696 }

697

698 static int get_dev_info(cl_device_id dev, struct device_info *di)

699 {

700 di->id = dev;

701

702 clGetDeviceInfo(dev, CL_DEVICE_TYPE, sizeof di->type, &di->type, 0);

703 clGetDeviceInfo(dev, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof di->units, &di->units, 0);

704 clGetDeviceInfo(dev, CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof di->clock, &di->clock, 0);

705 clGetDeviceInfo(dev, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, sizeof di->dim, &di->dim, 0);

706

707 di->work_item_sizes = new size_t[di->dim];

708

709 clGetDeviceInfo(dev, CL_DEVICE_MAX_WORK_ITEM_SIZES, di->dim * sizeof *di->work_item_sizes, di->work_item_sizes, 0);

710 clGetDeviceInfo(dev, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof di->work_group_size, &di->work_group_size, 0);

711 clGetDeviceInfo(dev, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof di->mem_size, &di->mem_size, 0);

712

713 size_t ext_str_len;

714 clGetDeviceInfo(dev, CL_DEVICE_EXTENSIONS, 0, 0, &ext_str_len);

715

716 di->extensions = new char[ext_str_len + 1];

717 clGetDeviceInfo(dev, CL_DEVICE_EXTENSIONS, ext_str_len, di->extensions, 0);

718 di->extensions[ext_str_len] = 0;

719

720 if(strstr(di->extensions, "cl_khr_gl_sharing") || strstr(di->extensions, "cl_APPLE_gl_sharing")) {

721 di->gl_sharing = true;

722 } else {

723 di->gl_sharing = false;

724 }

725

726 return 0;

727 }

728

729 static void destroy_dev_info(struct device_info *di)

730 {

731 delete [] di->work_item_sizes;

732 delete [] di->extensions;

733 }

734

735 static int devcmp(struct device_info *a, struct device_info *b)

736 {

737 unsigned int aval = a->units * a->clock;

738 unsigned int bval = b->units * b->clock;

739

740 return aval - bval;

741 }

742

743 static const char *devtypestr(cl_device_type type)

744 {

745 switch(type) {

746 case CL_DEVICE_TYPE_CPU:

747 return "cpu";

748 case CL_DEVICE_TYPE_GPU:

749 return "gpu";

750 case CL_DEVICE_TYPE_ACCELERATOR:

751 return "accelerator";

752 default:

753 break;

754 }

755 return "unknown";

756 }

757

758 static void print_memsize(FILE *out, unsigned long bytes)

759 {

760 int i;

761 unsigned long memsz = bytes;

762 const char *suffix[] = {"bytes", "kb", "mb", "gb", "tb", "pb", 0};

763

764 for(i=0; suffix[i]; i++) {

765 if(memsz < 1024) {

766 fprintf(out, "%lu %s", memsz, suffix[i]);

767 if(i > 0) {

768 fprintf(out, " (%lu bytes)", bytes);

769 }

770 return;

771 }

772

773 memsz /= 1024;

774 }

775 }

776

777 static const char *clstrerror(int err)

778 {

779 if(err > 0) {

780 return "<invalid error code>";

781 }

782 if(err <= -(int)(sizeof ocl_errstr / sizeof *ocl_errstr)) {

783 return "<unknown error>";

784 }

785 return ocl_errstr[-err];

786 }