clray: 980bc07be868 src/ocl.cc

clray

view src/ocl.cc @ 39:980bc07be868

Implemented OpenGL/OpenCL interop, and removed the texture copy

author	John Tsiombikas <nuclear@member.fsf.org>
date	Fri, 27 Aug 2010 18:30:09 +0100
parents	4cf4919c3812
children	1bcbb53b3505

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

25 struct device_info {

26 cl_device_id id;

27 cl_device_type type;

28 unsigned int units;

29 unsigned int clock;

31 unsigned int dim;

32 size_t *work_item_sizes;

33 size_t work_group_size;

35 unsigned long mem_size;

36 };

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

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

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

41 static const char *devtypestr(cl_device_type type);

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

43 static const char *clstrerror(int err);

46 static cl_context ctx;

47 static cl_command_queue cmdq;

48 static device_info devinf;

50 bool init_opencl()

51 {

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

53 return false;

54 }

56 #if defined(__APPLE__)

57 #error "CL/GL context sharing not implemented on MacOSX yet"

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

59 Display *dpy = glXGetCurrentDisplay();

60 GLXContext glctx = glXGetCurrentContext();

62 assert(dpy && glctx);

64 cl_context_properties prop[] = {

65 CL_GLX_DISPLAY_KHR, (cl_context_properties)dpy,

66 CL_GL_CONTEXT_KHR, (cl_context_properties)glctx,

67 0

68 };

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

70 #error "CL/GL context sharing not implemented on windows yet"

71 #else

72 #error "unknown or unsupported platform"

73 #endif

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

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

77 return false;

78 }

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

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

82 return false;

83 }

84 return true;

85 }

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

89 {

90 int err;

91 cl_mem mem;

92 cl_mem_flags flags = rdwr | CL_MEM_ALLOC_HOST_PTR;

94 if(buf) {

95 flags |= CL_MEM_COPY_HOST_PTR;

96 }

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

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

101 return 0;

102 }

103

104 CLMemBuffer *mbuf = new CLMemBuffer;

105 mbuf->mem = mem;

106 mbuf->size = sz;

107 mbuf->ptr = 0;

108 mbuf->tex = 0;

109 return mbuf;

110 }

111

112 CLMemBuffer *create_mem_buffer(int rdwr, unsigned int tex)

113 {

114 int err;

115 cl_mem mem;

116

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

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

119 return 0;

120 }

121

122 CLMemBuffer *mbuf = new CLMemBuffer;

123 mbuf->mem = mem;

124 mbuf->size = 0;

125 mbuf->ptr = 0;

126 mbuf->tex = tex;

127 return mbuf;

128 }

129

130 void destroy_mem_buffer(CLMemBuffer *mbuf)

131 {

132 if(mbuf) {

133 clReleaseMemObject(mbuf->mem);

134 delete mbuf;

135 }

136 }

137

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

139 {

140 if(!mbuf) return 0;

141

142 #ifndef NDEBUG

143 if(mbuf->ptr) {

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

145 }

146 #endif

147

148 int err;

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

150 if(!mbuf->ptr) {

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

152 return 0;

153 }

154 return mbuf->ptr;

155 }

156

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

158 {

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

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

161 mbuf->ptr = 0;

162 }

163

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

165 {

166 if(!mbuf) return false;

167

168 int err;

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

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

171 return false;

172 }

173 return true;

174 }

175

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

177 {

178 if(!mbuf) return false;

179

180 int err;

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

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

183 return false;

184 }

185 return true;

186 }

187

188

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

190 {

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

192 return false;

193 }

194

195 int err;

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

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

198 return false;

199 }

200 return true;

201 }

202

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

204 {

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

206 return false;

207 }

208

209 int err;

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

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

212 return false;

213 }

214 return true;

215 }

216

217

218 CLArg::CLArg()

219 {

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

221 }

222

223

224 CLProgram::CLProgram(const char *kname)

225 {

226 prog = 0;

227 kernel = 0;

228 this->kname = kname;

229 args.resize(16);

230 built = false;

231

232 wait_event = last_event = 0;

233 }

234

235 CLProgram::~CLProgram()

236 {

237 if(wait_event) {

238 clReleaseEvent(wait_event);

239 }

240 if(last_event) {

241 clReleaseEvent(last_event);

242 }

243

244 if(prog) {

245

246 clReleaseProgram(prog);

247 }

248 if(kernel) {

249

250 clReleaseKernel(kernel);

251 }

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

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

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

255 }

256 }

257 }

258

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

260 {

261 FILE *fp;

262 char *src;

263 struct stat st;

264

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

266

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

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

269 return false;

270 }

271

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

273

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

275

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

277 src[st.st_size] = 0;

278 fclose(fp);

279

280

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

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

283 delete [] src;

284 return false;

285 }

286 delete [] src;

287 return true;

288 }

289

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

291 {

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

293 args.resize(idx + 1);

294 }

295

296 CLArg *arg = &args[idx];

297 arg->type = ARGTYPE_INT;

298 arg->v.ival = val;

299 return true;

300 }

301

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

303 {

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

305 args.resize(idx + 1);

306 }

307

308 CLArg *arg = &args[idx];

309 arg->type = ARGTYPE_FLOAT;

310 arg->v.fval = val;

311 return true;

312 }

313

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

315 {

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

317 CLMemBuffer *buf;

318

319 if(sz <= 0) {

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

321 return false;

322 }

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

324 return false;

325 }

326

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

328 args.resize(idx + 1);

329 }

330 args[idx].type = ARGTYPE_MEM_BUF;

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

332 return true;

333 }

334

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

336 {

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

338 CLMemBuffer *buf;

339

340 if(!(buf = create_mem_buffer(rdwr, tex))) {

341 return false;

342 }

343

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

345 args.resize(idx + 1);

346 }

347 args[idx].type = ARGTYPE_MEM_BUF;

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

349 return true;

350 }

351

352 CLMemBuffer *CLProgram::get_arg_buffer(int arg)

353 {

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

355 return 0;

356 }

357 return args[arg].v.mbuf;

358 }

359

360 int CLProgram::get_num_args() const

361 {

362 int num_args = 0;

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

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

365 num_args++;

366 }

367 }

368 return num_args;

369 }

370

371 bool CLProgram::build()

372 {

373 int err;

374

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

376 size_t sz;

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

378

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

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

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

382

383 clReleaseProgram(prog);

384 prog = 0;

385 return false;

386 }

387

388

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

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

391 clReleaseProgram(prog);

392 prog = 0;

393 return false;

394 }

395

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

397 int err;

398

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

400 break;

401 }

402

403 switch(args[i].type) {

404 case ARGTYPE_INT:

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

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

407 goto fail;

408 }

409 break;

410

411 case ARGTYPE_FLOAT:

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

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

414 goto fail;

415 }

416 break;

417

418 case ARGTYPE_MEM_BUF:

419 {

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

421

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

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

424 goto fail;

425 }

426 }

427 break;

428

429 default:

430 break;

431 }

432 }

433

434 built = true;

435 return true;

436

437 fail:

438 clReleaseProgram(prog);

439 clReleaseKernel(kernel);

440 prog = 0;

441 kernel = 0;

442 return false;

443 }

444

445 bool CLProgram::run() const

446 {

447 return run(1, 1);

448 }

449

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

451 {

452 if(!built) {

453 if(!((CLProgram*)this)->build()) {

454 return false;

455 }

456 }

457

458 va_list ap;

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

460

461 va_start(ap, dim);

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

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

464 }

465 va_end(ap);

466

467 if(last_event) {

468 clReleaseEvent(last_event);

469 }

470

471 int err;

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

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

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

475 return false;

476 }

477

478 if(wait_event) {

479 clReleaseEvent(wait_event);

480 wait_event = 0;

481 }

482 return true;

483 }

484

485 void CLProgram::set_wait_event(cl_event ev)

486 {

487 if(wait_event) {

488 clReleaseEvent(wait_event);

489 }

490 wait_event = ev;

491 }

492

493 cl_event CLProgram::get_last_event() const

494 {

495 return last_event;

496 }

497

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

499 {

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

501 cl_device_id dev[32];

502 cl_platform_id plat[32];

503

504 dev_inf->work_item_sizes = 0;

505

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

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

508 return -1;

509 }

510 if(!num_plat) {

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

512 return -1;

513 }

514

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

516 char buf[512];

517

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

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

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

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

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

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

524 }

525

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

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

528 return -1;

529 }

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

531

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

533 struct device_info di;

534

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

536 free(dev_inf->work_item_sizes);

537 return -1;

538 }

539

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

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

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

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

544

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

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

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

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

549 printf(", ");

550 }

551 }

552 putchar('\n');

553

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

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

556 print_memsize(stdout, di.mem_size);

557 putchar('\n');

558

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

560 free(dev_inf->work_item_sizes);

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

562 sel = i;

563 }

564 }

565

566 if(num_dev) {

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

568 return 0;

569 }

570

571 return -1;

572 }

573

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

575 {

576 di->id = dev;

577

578

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

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

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

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

583

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

585

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

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

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

589

590 return 0;

591 }

592

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

594 {

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

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

597

598 return aval - bval;

599 }

600

601 static const char *devtypestr(cl_device_type type)

602 {

603 switch(type) {

604 case CL_DEVICE_TYPE_CPU:

605 return "cpu";

606 case CL_DEVICE_TYPE_GPU:

607 return "gpu";

608 case CL_DEVICE_TYPE_ACCELERATOR:

609 return "accelerator";

610 default:

611 break;

612 }

613 return "unknown";

614 }

615

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

617 {

618 int i;

619 unsigned long memsz = bytes;

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

621

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

623 if(memsz < 1024) {

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

625 if(i > 0) {

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

627 }

628 return;

629 }

630

631 memsz /= 1024;

632 }

633 }

634

635 static const char *clstrerror(int err)

636 {

637 if(err > 0) {

638 return "<invalid error code>";

639 }

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

641 return "<unknown error>";

642 }

643 return ocl_errstr[-err];

644 }