clray: 407935b73af3 src/ocl.cc

clray

view src/ocl.cc @ 13:407935b73af3

bollocks

author	John Tsiombikas <nuclear@member.fsf.org>
date	Wed, 04 Aug 2010 04:51:06 +0100
parents	85fd61f374d9
children	29f9330cfa4b

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 printf("create argument %d buffer: %d bytes\n", idx, (int)sz);

243 CLMemBuffer *buf;

244

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

246 return false;

247 }

248

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

250 args.resize(idx + 1);

251 }

252 args[idx].type = ARGTYPE_MEM_BUF;

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

254 return true;

255 }

256

257 CLMemBuffer *CLProgram::get_arg_buffer(int arg)

258 {

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

260 return 0;

261 }

262 return args[arg].v.mbuf;

263 }

264

265 bool CLProgram::build()

266 {

267 int err;

268

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

270 size_t sz;

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

272

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

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

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

276

277 clReleaseProgram(prog);

278 prog = 0;

279 return false;

280 }

281

282

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

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

285 clReleaseProgram(prog);

286 prog = 0;

287 return false;

288 }

289

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

291 int err;

292

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

294 break;

295 }

296

297 switch(args[i].type) {

298 case ARGTYPE_INT:

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

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

301 goto fail;

302 }

303 break;

304

305 case ARGTYPE_FLOAT:

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

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

308 goto fail;

309 }

310 break;

311

312 case ARGTYPE_MEM_BUF:

313 {

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

315

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

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

318 goto fail;

319 }

320 }

321 break;

322

323 default:

324 break;

325 }

326 }

327

328 built = true;

329 return true;

330

331 fail:

332 clReleaseProgram(prog);

333 clReleaseKernel(kernel);

334 prog = 0;

335 kernel = 0;

336 return false;

337 }

338

339 bool CLProgram::run() const

340 {

341 return run(1, 1);

342 }

343

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

345 {

346 if(!built) {

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

348 return false;

349 }

350 }

351

352 va_list ap;

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

354

355 va_start(ap, dim);

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

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

358 }

359 va_end(ap);

360

361 int err;

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

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

364 return false;

365 }

366 return true;

367 }

368

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

370 {

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

372 cl_device_id dev[32];

373 cl_platform_id plat[32];

374

375 dev_inf->work_item_sizes = 0;

376

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

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

379 return -1;

380 }

381 if(!num_plat) {

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

383 return -1;

384 }

385

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

387 char buf[512];

388

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

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

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

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

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

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

395 }

396

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

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

399 return -1;

400 }

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

402

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

404 struct device_info di;

405

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

407 free(dev_inf->work_item_sizes);

408 return -1;

409 }

410

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

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

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

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

415

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

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

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

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

420 printf(", ");

421 }

422 }

423 putchar('\n');

424

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

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

427 print_memsize(stdout, di.mem_size);

428 putchar('\n');

429

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

431 free(dev_inf->work_item_sizes);

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

433 sel = i;

434 }

435 }

436

437 if(num_dev) {

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

439 return 0;

440 }

441

442 return -1;

443 }

444

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

446 {

447 di->id = dev;

448

449

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

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

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

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

454

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

456

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

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

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

460

461 return 0;

462 }

463

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

465 {

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

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

468

469 return aval - bval;

470 }

471

472 static const char *devtypestr(cl_device_type type)

473 {

474 switch(type) {

475 case CL_DEVICE_TYPE_CPU:

476 return "cpu";

477 case CL_DEVICE_TYPE_GPU:

478 return "gpu";

479 case CL_DEVICE_TYPE_ACCELERATOR:

480 return "accelerator";

481 default:

482 break;

483 }

484 return "unknown";

485 }

486

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

488 {

489 int i;

490 unsigned long memsz = bytes;

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

492

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

494 if(memsz < 1024) {

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

496 if(i > 0) {

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

498 }

499 return;

500 }

501

502 memsz /= 1024;

503 }

504 }

505

506 static const char *clstrerror(int err)

507 {

508 if(err > 0) {

509 return "<invalid error code>";

510 }

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

512 return "<unknown error>";

513 }

514 return ocl_errstr[-err];

515 }