clray: 1169f3d04135 src/ocl.cc

clray

view src/ocl.cc @ 42:1169f3d04135

added CL/GL interop support for macosx (compiles) and windows (untested)

author	John Tsiombikas <nuclear@member.fsf.org>
date	Sat, 28 Aug 2010 02:01:16 +0100
parents	057b8575a1c1
children	f9eec11e5acc

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 }

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

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

136 {

137 int err;

138 cl_mem mem;

139 cl_mem_flags flags = rdwr | CL_MEM_ALLOC_HOST_PTR;

140

141 if(buf) {

142 flags |= CL_MEM_COPY_HOST_PTR;

143 }

144

145

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

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

148 return 0;

149 }

150

151 CLMemBuffer *mbuf = new CLMemBuffer;

152 mbuf->type = MEM_BUFFER;

153 mbuf->mem = mem;

154 mbuf->size = sz;

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

156 mbuf->ptr = 0;

157 mbuf->tex = 0;

158 return mbuf;

159 }

160

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

162 {

163 int err, pitch;

164 cl_mem mem;

165 cl_mem_flags flags = rdwr | CL_MEM_ALLOC_HOST_PTR;

166

167 if(pixels) {

168 flags |= CL_MEM_COPY_HOST_PTR;

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

170 } else {

171 pitch = 0;

172 }

173

174 cl_image_format fmt = {CL_RGBA, CL_FLOAT};

175

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

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

178 return 0;

179 }

180

181 CLMemBuffer *mbuf = new CLMemBuffer;

182 mbuf->type = IMAGE_BUFFER;

183 mbuf->mem = mem;

184 mbuf->size = ysz * pitch;

185 mbuf->xsz = xsz;

186 mbuf->ysz = ysz;

187 mbuf->ptr = 0;

188 mbuf->tex = 0;

189 return mbuf;

190 }

191

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

193 {

194 int err, xsz, ysz;

195 cl_mem mem;

196

197 glGetError(); // clear previous OpenGL errors

198

199 glPushAttrib(GL_TEXTURE_BIT);

200 glBindTexture(GL_TEXTURE_2D, tex);

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

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

203 glPopAttrib();

204

205 if(glGetError()) {

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

207 return 0;

208 }

209

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

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

212 return 0;

213 }

214

215 CLMemBuffer *mbuf = new CLMemBuffer;

216 mbuf->type = IMAGE_BUFFER;

217 mbuf->mem = mem;

218 mbuf->size = 0;

219 mbuf->xsz = xsz;

220 mbuf->ysz = ysz;

221 mbuf->ptr = 0;

222 mbuf->tex = tex;

223

224 return mbuf;

225 }

226

227 void destroy_mem_buffer(CLMemBuffer *mbuf)

228 {

229 if(mbuf) {

230 clReleaseMemObject(mbuf->mem);

231 delete mbuf;

232 }

233 }

234

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

236 {

237 if(!mbuf) return 0;

238

239 #ifndef NDEBUG

240 if(mbuf->ptr) {

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

242 }

243 #endif

244

245 int err;

246

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

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

249 if(!mbuf->ptr) {

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

251 return 0;

252 }

253 } else {

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

255

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

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

258 size_t pitch;

259

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

261 if(!mbuf->ptr) {

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

263 return 0;

264 }

265

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

267 }

268 return mbuf->ptr;

269 }

270

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

272 {

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

274

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

276 mbuf->ptr = 0;

277 }

278

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

280 {

281 if(!mbuf) return false;

282

283 int err;

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

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

286 return false;

287 }

288 return true;

289 }

290

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

292 {

293 if(!mbuf) return false;

294

295 int err;

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

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

298 return false;

299 }

300 return true;

301 }

302

303

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

305 {

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

307 return false;

308 }

309

310 int err;

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

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

313 return false;

314 }

315 return true;

316 }

317

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

319 {

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

321 return false;

322 }

323

324 int err;

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

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

327 return false;

328 }

329 return true;

330 }

331

332

333 CLArg::CLArg()

334 {

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

336 }

337

338

339 CLProgram::CLProgram(const char *kname)

340 {

341 prog = 0;

342 kernel = 0;

343 this->kname = kname;

344 args.resize(16);

345 built = false;

346

347 wait_event = last_event = 0;

348 }

349

350 CLProgram::~CLProgram()

351 {

352 if(wait_event) {

353 clReleaseEvent(wait_event);

354 }

355 if(last_event) {

356 clWaitForEvents(1, &last_event);

357 clReleaseEvent(last_event);

358 }

359

360 if(prog) {

361 clReleaseProgram(prog);

362 }

363 if(kernel) {

364 clReleaseKernel(kernel);

365 }

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

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

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

369 }

370 }

371 }

372

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

374 {

375 FILE *fp;

376 char *src;

377 struct stat st;

378

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

380

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

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

383 return false;

384 }

385

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

387

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

389

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

391 src[st.st_size] = 0;

392 fclose(fp);

393

394

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

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

397 delete [] src;

398 return false;

399 }

400 delete [] src;

401 return true;

402 }

403

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

405 {

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

407 args.resize(idx + 1);

408 }

409

410 CLArg *arg = &args[idx];

411 arg->type = ARGTYPE_INT;

412 arg->v.ival = val;

413 return true;

414 }

415

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

417 {

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

419 args.resize(idx + 1);

420 }

421

422 CLArg *arg = &args[idx];

423 arg->type = ARGTYPE_FLOAT;

424 arg->v.fval = val;

425 return true;

426 }

427

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

429 {

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

431 CLMemBuffer *buf;

432

433 if(sz <= 0) {

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

435 return false;

436 }

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

438 return false;

439 }

440

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

442 args.resize(idx + 1);

443 }

444 args[idx].type = ARGTYPE_MEM_BUF;

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

446 return true;

447 }

448

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

450 {

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

452 CLMemBuffer *buf;

453

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

455 return false;

456 }

457

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

459 args.resize(idx + 1);

460 }

461 args[idx].type = ARGTYPE_MEM_BUF;

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

463 return true;

464 }

465

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

467 {

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

469 CLMemBuffer *buf;

470

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

472 return false;

473 }

474

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

476 args.resize(idx + 1);

477 }

478 args[idx].type = ARGTYPE_MEM_BUF;

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

480 return true;

481 }

482

483 CLMemBuffer *CLProgram::get_arg_buffer(int arg)

484 {

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

486 return 0;

487 }

488 return args[arg].v.mbuf;

489 }

490

491 int CLProgram::get_num_args() const

492 {

493 int num_args = 0;

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

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

496 num_args++;

497 }

498 }

499 return num_args;

500 }

501

502 bool CLProgram::build()

503 {

504 int err;

505

506 if((err = clBuildProgram(prog, 0, 0, "-cl-mad-enable", 0, 0)) != 0) {

507 size_t sz;

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

509

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

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

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

513

514 clReleaseProgram(prog);

515 prog = 0;

516 return false;

517 }

518

519

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

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

522 clReleaseProgram(prog);

523 prog = 0;

524 return false;

525 }

526

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

528 int err;

529

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

531 break;

532 }

533

534 switch(args[i].type) {

535 case ARGTYPE_INT:

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

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

538 goto fail;

539 }

540 break;

541

542 case ARGTYPE_FLOAT:

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

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

545 goto fail;

546 }

547 break;

548

549 case ARGTYPE_MEM_BUF:

550 {

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

552

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

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

555 goto fail;

556 }

557 }

558 break;

559

560 default:

561 break;

562 }

563 }

564

565 built = true;

566 return true;

567

568 fail:

569 clReleaseProgram(prog);

570 clReleaseKernel(kernel);

571 prog = 0;

572 kernel = 0;

573 return false;

574 }

575

576 bool CLProgram::run() const

577 {

578 return run(1, 1);

579 }

580

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

582 {

583 va_list ap;

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

585

586 va_start(ap, dim);

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

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

589 }

590 va_end(ap);

591

592 if(last_event) {

593 clReleaseEvent(last_event);

594 }

595

596 int err;

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

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

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

600 return false;

601 }

602

603 if(wait_event) {

604 clReleaseEvent(wait_event);

605 wait_event = 0;

606 }

607 return true;

608 }

609

610 void CLProgram::set_wait_event(cl_event ev)

611 {

612 if(wait_event) {

613 clReleaseEvent(wait_event);

614 }

615 wait_event = ev;

616 }

617

618 cl_event CLProgram::get_last_event() const

619 {

620 return last_event;

621 }

622

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

624 {

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

626 cl_device_id dev[32];

627 cl_platform_id plat[32];

628

629 dev_inf->work_item_sizes = 0;

630

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

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

633 return -1;

634 }

635 if(!num_plat) {

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

637 return -1;

638 }

639

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

641 char buf[512];

642

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

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

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

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

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

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

649 }

650

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

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

653 return -1;

654 }

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

656

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

658 struct device_info di;

659

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

661 destroy_dev_info(&di);

662 return -1;

663 }

664

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

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

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

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

669

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

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

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

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

674 printf(", ");

675 }

676 }

677 putchar('\n');

678

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

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

681 print_memsize(stdout, di.mem_size);

682 putchar('\n');

683

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

685

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

687 free(dev_inf->work_item_sizes);

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

689 sel = i;

690 }

691 }

692

693 if(num_dev) {

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

695 return 0;

696 }

697

698 return -1;

699 }

700

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

702 {

703 di->id = dev;

704

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

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

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

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

709

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

711

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

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

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

715

716 size_t ext_str_len;

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

718

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

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

721 di->extensions[ext_str_len] = 0;

722

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

724 di->gl_sharing = true;

725 } else {

726 di->gl_sharing = false;

727 }

728

729 return 0;

730 }

731

732 static void destroy_dev_info(struct device_info *di)

733 {

734 delete [] di->work_item_sizes;

735 delete [] di->extensions;

736 }

737

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

739 {

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

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

742

743 return aval - bval;

744 }

745

746 static const char *devtypestr(cl_device_type type)

747 {

748 switch(type) {

749 case CL_DEVICE_TYPE_CPU:

750 return "cpu";

751 case CL_DEVICE_TYPE_GPU:

752 return "gpu";

753 case CL_DEVICE_TYPE_ACCELERATOR:

754 return "accelerator";

755 default:

756 break;

757 }

758 return "unknown";

759 }

760

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

762 {

763 int i;

764 unsigned long memsz = bytes;

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

766

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

768 if(memsz < 1024) {

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

770 if(i > 0) {

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

772 }

773 return;

774 }

775

776 memsz /= 1024;

777 }

778 }

779

780 static const char *clstrerror(int err)

781 {

782 if(err > 0) {

783 return "<invalid error code>";

784 }

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

786 return "<unknown error>";

787 }

788 return ocl_errstr[-err];

789 }