clray: 8c858e1a89e8 src/ocl.cc

clray

view src/ocl.cc @ 60:8c858e1a89e8

minor cleanups

author	John Tsiombikas <nuclear@member.fsf.org>
date	Mon, 28 Dec 2015 09:08:49 +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

106

107 #endif /* CLGL_INTEROP */

108

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

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

111 return false;

112 }

113

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 }

120

121 void destroy_opencl()

122 {

123 if(cmdq) {

124 clReleaseCommandQueue(cmdq);

125 cmdq = 0;

126 }

127

128 if(ctx) {

129 clReleaseContext(ctx);

130 ctx = 0;

131 }

132 }

133

134 void finish_opencl()

135 {

136 clFinish(cmdq);

137 }

138

139

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;

145

146 if(buf) {

147 flags |= CL_MEM_COPY_HOST_PTR;

148 }

149

150

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 }

155

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 }

165

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;

171

172 if(pixels) {

173 flags |= CL_MEM_COPY_HOST_PTR;

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

175 } else {

176 pitch = 0;

177 }

178

179 cl_image_format fmt = {CL_RGBA, CL_FLOAT};

180

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 }

185

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 }

196

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

198 {

199 int err, xsz, ysz;

200 cl_mem mem;

201

202 glGetError(); // clear previous OpenGL errors

203

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();

209

210 if(glGetError()) {

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

212 return 0;

213 }

214

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 }

219

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;

228

229 return mbuf;

230 }

231

232 void destroy_mem_buffer(CLMemBuffer *mbuf)

233 {

234 if(mbuf) {

235 clReleaseMemObject(mbuf->mem);

236 delete mbuf;

237 }

238 }

239

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

241 {

242 if(!mbuf) return 0;

243

244 #ifndef NDEBUG

245 if(mbuf->ptr) {

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

247 }

248 #endif

249

250 int err;

251

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);

260

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

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

263 size_t pitch;

264

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 }

270

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

272 }

273 return mbuf->ptr;

274 }

275

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

277 {

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

279

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

281 mbuf->ptr = 0;

282 }

283

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

285 {

286 if(!mbuf) return false;

287

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 }

295

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

297 {

298 if(!mbuf) return false;

299

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 }

307

308

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

310 {

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

312 return false;

313 }

314

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 }

322

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

324 {

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

326 return false;

327 }

328

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 }

336

337

338 CLArg::CLArg()

339 {

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

341 }

342

343

344 CLProgram::CLProgram(const char *kname)

345 {

346 prog = 0;

347 kernel = 0;

348 this->kname = kname;

349 args.resize(16);

350 built = false;

351

352 wait_event = last_event = 0;

353 }

354

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 }

364

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 }

377

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

379 {

380 FILE *fp;

381 char *src;

382 struct stat st;

383

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

385

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

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

388 return false;

389 }

390

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

392

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

394

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

396 src[st.st_size] = 0;

397 fclose(fp);

398

399

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 }

408

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

410 {

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

412 args.resize(idx + 1);

413 }

414

415 CLArg *arg = &args[idx];

416 arg->type = ARGTYPE_INT;

417 arg->v.ival = val;

418 return true;

419 }

420

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

422 {

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

424 args.resize(idx + 1);

425 }

426

427 CLArg *arg = &args[idx];

428 arg->type = ARGTYPE_FLOAT;

429 arg->v.fval = val;

430 return true;

431 }

432

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;

437

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 }

445

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 }

453

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;

458

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

460 return false;

461 }

462

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 }

470

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;

475

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

477 return false;

478 }

479

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 }

487

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 }

495

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 }

506

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);

513

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);

517

518 clReleaseProgram(prog);

519 prog = 0;

520 return false;

521 }

522

523

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 }

530

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

532 int err;

533

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

535 break;

536 }

537

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;

545

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;

552

553 case ARGTYPE_MEM_BUF:

554 {

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

556

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;

563

564 default:

565 break;

566 }

567 }

568

569 built = true;

570 return true;

571

572 fail:

573 clReleaseProgram(prog);

574 clReleaseKernel(kernel);

575 prog = 0;

576 kernel = 0;

577 return false;

578 }

579

580 bool CLProgram::run() const

581 {

582 return run(1, 1);

583 }

584

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

586 {

587 va_list ap;

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

589

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);

595

596 if(last_event) {

597 clReleaseEvent(last_event);

598 }

599

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 }

606

607 if(wait_event) {

608 clReleaseEvent(wait_event);

609 wait_event = 0;

610 }

611 return true;

612 }

613

614 void CLProgram::set_wait_event(cl_event ev)

615 {

616 if(wait_event) {

617 clReleaseEvent(wait_event);

618 }

619 wait_event = ev;

620 }

621

622 cl_event CLProgram::get_last_event() const

623 {

624 return last_event;

625 }

626

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];

632

633 dev_inf->work_item_sizes = 0;

634

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 }

643

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

645 char buf[512];

646

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 }

654

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);

660

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

662 struct device_info di;

663

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

665 destroy_dev_info(&di);

666 return -1;

667 }

668

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);

673

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');

682

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');

687

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

689

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 }

696

697 if(num_dev) {

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

699 return 0;

700 }

701

702 return -1;

703 }

704

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

706 {

707 di->id = dev;

708

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);

713

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

715

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);

719

720 size_t ext_str_len;

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

722

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;

726

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 }

732

733 return 0;

734 }

735

736 static void destroy_dev_info(struct device_info *di)

737 {

738 delete [] di->work_item_sizes;

739 delete [] di->extensions;

740 }

741

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;

746

747 return aval - bval;

748 }

749

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 }

764

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};

770

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 }

779

780 memsz /= 1024;

781 }

782 }

783

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 }