clray: 85fd61f374d9 src/ocl.cc

clray

view src/ocl.cc @ 12:85fd61f374d9

fixed the bloody intersection bug

author	John Tsiombikas <nuclear@member.fsf.org>
date	Tue, 03 Aug 2010 13:06:59 +0100
parents	d9a1bab1c3f5
children	407935b73af3

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 #ifndef _MSC_VER

9 #include <alloca.h>

10 #else

11 #include <malloc.h>

12 #endif

13 #include <sys/stat.h>

14 #include "ocl.h"

15 #include "ocl_errstr.h"

18 class InitCL {

19 public:

20 InitCL();

21 };

23 struct device_info {

24 cl_device_id id;

25 cl_device_type type;

26 unsigned int units;

27 unsigned int clock;

29 unsigned int dim;

30 size_t *work_item_sizes;

31 size_t work_group_size;

33 unsigned long mem_size;

34 };

36 static bool init_opencl(void);

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

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

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

40 static const char *devtypestr(cl_device_type type);

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

42 static const char *clstrerror(int err);

45 static InitCL initcl;

46 static cl_context ctx;

47 static cl_command_queue cmdq;

48 static device_info devinf;

50 InitCL::InitCL()

51 {

52 if(!init_opencl()) {

53 exit(0);

54 }

55 }

57 static bool init_opencl(void)

58 {

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

60 return false;

61 }

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

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

66 return false;

67 }

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

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

71 return false;

72 }

73 return true;

74 }

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

78 {

79 int err;

80 cl_mem mem;

81 cl_mem_flags flags = rdwr | CL_MEM_ALLOC_HOST_PTR;

83 if(buf) {

84 flags |= CL_MEM_COPY_HOST_PTR;

85 }

88 if(!(mem = clCreateBuffer(ctx, flags, sz, buf, &err))) {

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

90 return 0;

91 }

93 CLMemBuffer *mbuf = new CLMemBuffer;

94 mbuf->mem = mem;

95 mbuf->size = sz;

96 mbuf->ptr = 0;

97 return mbuf;

98 }

100 void destroy_mem_buffer(CLMemBuffer *mbuf)

101 {

102 if(mbuf) {

103 clReleaseMemObject(mbuf->mem);

104 delete mbuf;

105 }

106 }

107

108 void *map_mem_buffer(CLMemBuffer *mbuf, int rdwr)

109 {

110 if(!mbuf) return 0;

111

112 #ifndef NDEBUG

113 if(mbuf->ptr) {

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

115 }

116 #endif

117

118 int err;

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

120 if(!mbuf->ptr) {

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

122 return 0;

123 }

124 return mbuf->ptr;

125 }

126

127 void unmap_mem_buffer(CLMemBuffer *mbuf)

128 {

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

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

131 mbuf->ptr = 0;

132 }

133

134 bool write_mem_buffer(CLMemBuffer *mbuf, size_t sz, void *src)

135 {

136 if(!mbuf) return false;

137

138 int err;

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

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

141 return false;

142 }

143 return true;

144 }

145

146 bool read_mem_buffer(CLMemBuffer *mbuf, size_t sz, void *dest)

147 {

148 if(!mbuf) return false;

149

150 int err;

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

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

153 return false;

154 }

155 return true;

156 }

157

158

159 CLProgram::CLProgram(const char *kname)

160 {

161 prog = 0;

162 kernel = 0;

163 this->kname = kname;

164 args.resize(16);

165 built = false;

166 }

167

168 CLProgram::~CLProgram()

169 {

170 if(prog) {

171

172 clReleaseProgram(prog);

173 }

174 if(kernel) {

175

176 clReleaseKernel(kernel);

177 }

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

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

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

181 }

182 }

183 }

184

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

186 {

187 FILE *fp;

188 char *src;

189 struct stat st;

190

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

192

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

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

195 return false;

196 }

197

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

199

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

201

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

203 src[st.st_size] = 0;

204 fclose(fp);

205

206

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

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

209 delete [] src;

210 return false;

211 }

212 delete [] src;

213 return true;

214 }

215

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

217 {

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

219 args.resize(idx + 1);

220 }

221

222 CLArg *arg = &args[idx];

223 arg->type = ARGTYPE_INT;

224 arg->v.ival = val;

225 return true;

226 }

227

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

229 {

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

231 args.resize(idx + 1);

232 }

233

234 CLArg *arg = &args[idx];

235 arg->type = ARGTYPE_FLOAT;

236 arg->v.fval = val;

237 return true;

238 }

239

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

241 {

242 CLMemBuffer *buf;

243

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

245 return false;

246 }

247

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

249 args.resize(idx + 1);

250 }

251 args[idx].type = ARGTYPE_MEM_BUF;

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

253 return true;

254 }

255

256 CLMemBuffer *CLProgram::get_arg_buffer(int arg)

257 {

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

259 return 0;

260 }

261 return args[arg].v.mbuf;

262 }

263

264 bool CLProgram::build()

265 {

266 int err;

267

268 if((err = clBuildProgram(prog, 0, 0, 0, 0, 0)) != 0) {

269 size_t sz;

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

271

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

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

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

275

276 clReleaseProgram(prog);

277 prog = 0;

278 return false;

279 }

280

281

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

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

284 clReleaseProgram(prog);

285 prog = 0;

286 return false;

287 }

288

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

290 int err;

291

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

293 break;

294 }

295

296 switch(args[i].type) {

297 case ARGTYPE_INT:

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

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

300 goto fail;

301 }

302 break;

303

304 case ARGTYPE_FLOAT:

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

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

307 goto fail;

308 }

309 break;

310

311 case ARGTYPE_MEM_BUF:

312 {

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

314

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

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

317 goto fail;

318 }

319 }

320 break;

321

322 default:

323 break;

324 }

325 }

326

327 built = true;

328 return true;

329

330 fail:

331 clReleaseProgram(prog);

332 clReleaseKernel(kernel);

333 prog = 0;

334 kernel = 0;

335 return false;

336 }

337

338 bool CLProgram::run() const

339 {

340 return run(1, 1);

341 }

342

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

344 {

345 if(!built) {

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

347 return false;

348 }

349 }

350

351 va_list ap;

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

353

354 va_start(ap, dim);

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

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

357 }

358 va_end(ap);

359

360 int err;

361 if((err = clEnqueueNDRangeKernel(cmdq, kernel, dim, 0, global_size, 0, 0, 0, 0)) != 0) {

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

363 return false;

364 }

365 return true;

366 }

367

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

369 {

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

371 cl_device_id dev[32];

372 cl_platform_id plat[32];

373

374 dev_inf->work_item_sizes = 0;

375

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

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

378 return -1;

379 }

380 if(!num_plat) {

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

382 return -1;

383 }

384

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

386 char buf[512];

387

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

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

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

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

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

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

394 }

395

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

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

398 return -1;

399 }

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

401

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

403 struct device_info di;

404

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

406 free(dev_inf->work_item_sizes);

407 return -1;

408 }

409

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

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

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

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

414

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

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

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

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

419 printf(", ");

420 }

421 }

422 putchar('\n');

423

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

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

426 print_memsize(stdout, di.mem_size);

427 putchar('\n');

428

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

430 free(dev_inf->work_item_sizes);

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

432 sel = i;

433 }

434 }

435

436 if(num_dev) {

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

438 return 0;

439 }

440

441 return -1;

442 }

443

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

445 {

446 di->id = dev;

447

448

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

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

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

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

453

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

455

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

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

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

459

460 return 0;

461 }

462

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

464 {

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

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

467

468 return aval - bval;

469 }

470

471 static const char *devtypestr(cl_device_type type)

472 {

473 switch(type) {

474 case CL_DEVICE_TYPE_CPU:

475 return "cpu";

476 case CL_DEVICE_TYPE_GPU:

477 return "gpu";

478 case CL_DEVICE_TYPE_ACCELERATOR:

479 return "accelerator";

480 default:

481 break;

482 }

483 return "unknown";

484 }

485

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

487 {

488 int i;

489 unsigned long memsz = bytes;

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

491

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

493 if(memsz < 1024) {

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

495 if(i > 0) {

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

497 }

498 return;

499 }

500

501 memsz /= 1024;

502 }

503 }

504

505 static const char *clstrerror(int err)

506 {

507 if(err > 0) {

508 return "<invalid error code>";

509 }

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

511 return "<unknown error>";

512 }

513 return ocl_errstr[-err];

514 }