1 module oclcv.clcore; 2 3 import std.string : toStringz; 4 import std.conv : to; 5 import std.outbuffer : OutBuffer; 6 debug import std.stdio; 7 8 import core.stdc.stdio : printf, fread, fopen, fclose, FILE; 9 import core.stdc.stdlib : EXIT_FAILURE, exit, malloc, free; 10 import bc.string; 11 import bcaa; 12 import bindbc.opencl; 13 import dplug.core; 14 15 struct BlockDim { 16 int x = 1, y = 1, z = 1; 17 } 18 19 struct GridDim { 20 int x = 1, y = 1, z = 1; 21 } 22 23 24 alias MemFlag = int; 25 enum : MemFlag 26 { 27 MEM_FLAG_READ_WRITE = 1 << 0, 28 MEM_FLAG_WRITE_ONLY = 1 << 1, 29 MEM_FLAG_READ_ONLY = 1 << 2, 30 MEM_FLAG_USE_HOST_PTR = 1 << 3, // maybe needs to be removed, in cuda not trivial 31 MEM_FLAG_ALLOC_HOST_PTR = 1 << 4, 32 MEM_FLAG_COPY_HOST_PTR = 1 << 5 33 } 34 35 alias SyncMode = int; 36 enum : SyncMode 37 { 38 SYNC_MODE_ASYNC = 0, 39 SYNC_MODE_BLOCKING = 1 40 } 41 42 final class CLContext { 43 public: 44 @nogc nothrow: 45 46 this(int platform_id = 0, int device_id = 0, int num_streams = 1){ 47 48 loadDLib(); 49 50 cl_platform_id p_id; 51 cl_int err = 0; 52 cl_uint num_platforms, num_divices; 53 cl_platform_id* p_ids; 54 cl_device_id* d_ids; 55 56 clGetPlatformIDs(0, null, &num_platforms); 57 if(num_platforms > 0){ 58 p_ids = cast(cl_platform_id*)malloc(cl_platform_id.sizeof * num_platforms); 59 scope(exit) free(p_ids); 60 clGetPlatformIDs(num_platforms, p_ids, null); 61 if(platform_id < 0 || platform_id >= int(num_platforms)) { 62 printf("Incorrect platform id %d!\n", platform_id); 63 exit(EXIT_FAILURE); 64 } 65 p_id = p_ids[platform_id]; 66 }else { 67 printf("Not found any platforms\n"); 68 exit(EXIT_FAILURE); 69 } 70 71 clGetDeviceIDs(p_id, CL_DEVICE_TYPE_ALL, 0, null, &num_divices); 72 if(num_divices > 0){ 73 d_ids = cast(cl_device_id*)malloc(cl_device_id.sizeof * num_divices); 74 scope(exit) free(d_ids); 75 clGetDeviceIDs(p_id, CL_DEVICE_TYPE_ALL, num_divices, d_ids, null); 76 if(device_id < 0 || device_id >= int(num_divices)){ 77 printf("Incorrect device id %d!\n",device_id); 78 exit(EXIT_FAILURE); 79 } 80 cl_device_id_ = d_ids[device_id]; 81 } 82 else{ 83 printf("Not found any devices\n"); 84 exit(EXIT_FAILURE); 85 } 86 87 cl_context_properties[3] prop = [CL_CONTEXT_PLATFORM, cast(cl_context_properties)p_id, 0]; 88 cl_context_ = clCreateContext(prop.ptr, 1, &cl_device_id_, null, null, &err); 89 handleError(err, RCStringZ.from("creating context")); 90 printf("OpenCL context created! \n"); 91 92 cl_command_queues_ = (cast(cl_command_queue*)malloc(cl_command_queue.sizeof * num_streams))[0..num_streams]; 93 for(int i=0; i < num_streams; i++){ 94 cl_command_queues_[i] = clCreateCommandQueue(cl_context_, cl_device_id_, 95 CL_QUEUE_PROFILING_ENABLE, &err); 96 handleError(err, RCStringZ.from("creating ClCommandQueue")); 97 } 98 99 RCStringZ oss; 100 oss ~= nogcFormat!"Selected platform vendor: %s %s"(getPlatformInfo(p_id,CL_PLATFORM_VENDOR), 101 getPlatformInfo(p_id,CL_PLATFORM_VERSION)); 102 oss ~= nogcFormat!"Selected device name: %s"(getDevInfo(cl_device_id_, CL_DEVICE_NAME)); 103 oss ~= nogcFormat!"Selected device OpenCL device version: %s"( 104 getDevInfo(cl_device_id_, CL_DEVICE_VERSION)); 105 oss ~= nogcFormat!"Selected device OpenCL C device version: %s"( 106 getDevInfo(cl_device_id_, CL_DEVICE_OPENCL_C_VERSION)); 107 cl_info_ = oss; 108 } 109 110 ~this(){ 111 foreach(ref cq; cl_command_queues_) 112 clReleaseCommandQueue(cq); 113 free(cl_command_queues_.ptr); 114 clReleaseContext(cl_context_); 115 } 116 117 cl_context getCLContext() {return cl_context_;} 118 cl_device_id getDevId() {return cl_device_id_;} 119 120 121 cl_command_queue getCommandQueue(int id) { 122 return cl_command_queues_[id]; 123 } 124 125 void finish(int command_queue_id) { 126 cl_int err = clFinish(cl_command_queues_[command_queue_id]); 127 handleError(err, RCStringZ.from("finishing command queue")); 128 } 129 130 RCStringZ clInfo() {return cl_info_;} 131 132 private: 133 RCStringZ getPlatformInfo(cl_platform_id platform_id, int info_name){ 134 size_t info_size = 0; 135 clGetPlatformInfo(platform_id, info_name, 0, null, &info_size); 136 RCStringZ str; str.reserve(info_size); 137 138 clGetPlatformInfo(platform_id, info_name, info_size, str.data.ptr, null); 139 return str; 140 } 141 142 RCStringZ getDevInfo(cl_device_id dev_id, int info_name){ 143 size_t info_size = 0; 144 clGetDeviceInfo(dev_id, info_name, 0, null, &info_size); 145 RCStringZ str; str.reserve(info_size); 146 clGetDeviceInfo(dev_id, info_name, info_size, str.data.ptr, null); 147 return str; 148 } 149 150 void loadDLib(){ 151 if(CLContext.support_ == CLSupport.noLibrary){ 152 CLContext.support_ = loadOpenCL(); 153 debug writeln("Load CL: ", CLContext.support_); 154 if(CLContext.support_ == CLSupport.noLibrary || CLContext.support_ == CLSupport.badLibrary){ 155 debug _assert(0, "Problem loading opencl dynamic library"); 156 } 157 } 158 159 } 160 161 static CLSupport support_ = CLSupport.noLibrary; 162 cl_command_queue[] cl_command_queues_; 163 RCStringZ cl_info_; 164 cl_context cl_context_; 165 cl_device_id cl_device_id_; 166 } 167 168 alias DataType = int; 169 enum : DataType { 170 BYTE = 0, 171 UBYTE = 0, 172 SHORT = 1, 173 USHORT = 1, 174 INT = 2, 175 UINT = 2, 176 FLOAT = 3, 177 DOUBLE = 4, 178 LONG = 5, 179 ULONG = 5 180 } 181 @nogc nothrow: 182 183 size_t unitSize(int dtype) { 184 if (dtype == -1){ 185 const string f = __FILE__; 186 const int ln = __LINE__; 187 printf("dataType of BufferMeta must be set to a supported type %s:%d", f.ptr, ln); 188 exit(-1); 189 } else 190 if (dtype == 0){ 191 return byte.sizeof; 192 } else 193 if (dtype == 1){ 194 return short.sizeof; 195 } else 196 if (dtype == 2){ 197 return int.sizeof; 198 } else 199 if (dtype == 3){ 200 return float.sizeof; 201 } else 202 if (dtype == 4){ 203 return double.sizeof; 204 } else 205 if (dtype == 5){ 206 return long.sizeof; 207 } 208 209 return 0; 210 } 211 212 struct BufferMeta { 213 int dataType = -1; 214 size_t height; 215 size_t width; 216 size_t numberOfChannels = 1; 217 218 @nogc nothrow: 219 size_t memorySize(){ 220 return height * width * numberOfChannels * unitSize(dataType); 221 } 222 223 alias rows = height; 224 alias cols = width; 225 } 226 227 final class CLBuffer { 228 public: 229 @nogc nothrow: 230 231 this(CLContext ctx, BufferMeta buffer_meta, MemFlag flag = MEM_FLAG_READ_WRITE, 232 void[] host_data = null) 233 { 234 meta_data = buffer_meta; 235 236 this(ctx, flag, host_data); 237 } 238 239 private this(CLContext ctx, MemFlag flag = MEM_FLAG_READ_WRITE, 240 void[] host_data = null){ 241 context_ = ctx; 242 flag_ = flag; 243 244 cl_int err; 245 buffer_ = clCreateBuffer(context_.getCLContext(), getCLMemFlag(flag), 246 size_, null, &err); 247 handleError(err, RCStringZ.from("creating buffer")); 248 if(host_data){ 249 upload(host_data, SYNC_MODE_BLOCKING, 0); 250 } 251 } 252 253 ~this(){ 254 if(buffer_){ 255 cl_int err = clReleaseMemObject(buffer_) ; 256 handleError(err, RCStringZ.from("in releasing buffer")); 257 buffer_ = null; 258 } 259 } 260 261 @nogc nothrow: 262 // validate device memory for debug purpose 263 int validate(){ 264 //check if valid mem object, 265 cl_mem_object_type mem_type = 0; 266 clGetMemObjectInfo(buffer_, CL_MEM_TYPE, cl_mem_object_type.sizeof, &mem_type, null); 267 if (mem_type != CL_MEM_OBJECT_BUFFER) 268 { 269 debug writeln("CL_INVALID_MEM_OBJECT"); 270 return CL_INVALID_MEM_OBJECT; 271 } 272 //check if mem object have valid required size 273 if (size_ > 0) 274 { 275 size_t current_size; 276 clGetMemObjectInfo(buffer_, CL_MEM_SIZE, 277 size_t.sizeof, ¤t_size, null); 278 debug writefln("[CLBuffer.validate] Buffer size: %s bytes. Required size: %s", 279 current_size, size_); 280 281 if (current_size < size_) 282 return CL_INVALID_IMAGE_SIZE; 283 } 284 285 return CL_SUCCESS; 286 } 287 288 bool isNull(){return buffer_ == null;} 289 290 void upload(const(void)[] data, SyncMode block_queue = SYNC_MODE_BLOCKING, 291 int command_queue = 0){ 292 debug _assert(data.length == metaData().memorySize(), "Mismatch in source and destination memory sizes."); 293 upload(data.ptr, 0, size_, block_queue, command_queue); 294 } 295 296 void download(void[] data, SyncMode block_queue = SYNC_MODE_BLOCKING, 297 int command_queue = 0){ 298 debug _assert(data.length == metaData().memorySize(), "Mismatch in source and destination memory sizes."); 299 download(data.ptr, 0, size_, block_queue, command_queue); 300 } 301 302 cl_mem getCObject(){ 303 return buffer_; 304 } 305 306 BufferMeta metaData(){return meta_data;} 307 308 private: 309 310 void upload(const void* data, size_t offset, size_t size, 311 SyncMode block_queue,int command_queue){ 312 313 cl_bool b_Block = (block_queue == SYNC_MODE_BLOCKING) ? CL_TRUE : CL_FALSE; 314 cl_int err = clEnqueueWriteBuffer(context_.getCommandQueue(command_queue), 315 buffer_, b_Block, offset, size, data, 0, 316 null, null); 317 handleError(err, "enqueuing writing buffer".RCStringZ); 318 } 319 320 void download(void* data, size_t offset, size_t size, SyncMode block_queue, 321 int command_queue){ 322 cl_bool b_Block = (block_queue == SYNC_MODE_BLOCKING) ? CL_TRUE : CL_FALSE; 323 cl_int err = clEnqueueReadBuffer(context_.getCommandQueue(command_queue), 324 buffer_, b_Block, offset, size, data, 0, 325 null, null); 326 handleError(err, "enqueuing reading buffer".RCStringZ); 327 } 328 329 CLContext context_; 330 cl_mem buffer_; 331 MemFlag flag_; 332 333 BufferMeta meta_data; 334 size_t size_(){return meta_data.memorySize;} 335 336 } 337 338 final class CLKernel { 339 public: 340 @nogc nothrow: 341 342 this(CLContext context, cl_program program, string kernel_name){ 343 cl_int err = CL_SUCCESS; 344 context_ = context; 345 kernel_name_ = kernel_name; 346 kernel_ = clCreateKernel(program, RCStringZ.from(kernel_name_).data.ptr, &err); 347 handleError(err, RCStringZ.from("creating kernel: ", kernel_name)); 348 } 349 ~this(){ 350 cl_int err = clReleaseKernel(kernel_); 351 handleError(err, RCStringZ.from("releasing kernel objects")); 352 } 353 354 void launch(int queue_id, GridDim gd, BlockDim bd){ 355 size_t[3] global_w_offset = [0, 0, 0]; 356 size_t[3] global_w_size = [ 357 size_t(gd.x * bd.x), 358 size_t(gd.y * bd.y), 359 size_t(gd.z * bd.z)]; 360 size_t[3] local_w_size = [size_t(bd.x),size_t(bd.y),size_t(bd.z)]; 361 362 cl_int err = clEnqueueNDRangeKernel(context_.getCommandQueue(queue_id), 363 kernel_, 3, global_w_offset.ptr, global_w_size.ptr, 364 local_w_size.ptr, 0, null, null); 365 366 handleError(err, RCStringZ.from("enqueuing kernel")); 367 } 368 369 void launch(int queue_id, size_t* gwo, size_t* gws, size_t* lws){ 370 cl_int err = clEnqueueNDRangeKernel(context_.getCommandQueue(queue_id), 371 kernel_, 3, gwo, gws, lws, 0, null, null); 372 handleError(err, RCStringZ.from("enqueuing kernel")); 373 } 374 375 void setArgs(Args...)(Args args){ 376 import std.stdio; 377 cl_int err = CL_SUCCESS; 378 foreach(i, arg; args){ 379 380 static if(is(typeof(arg)==CLBuffer)){ 381 auto raw_mem = arg.getCObject(); 382 err = clSetKernelArg(kernel_, cl_uint(i), cl_mem.sizeof, cast(void*)&raw_mem); 383 } else { 384 err = clSetKernelArg(kernel_, cl_uint(i), typeof(arg).sizeof, cast(void*)&arg); 385 } 386 387 debug handleError(err, RCStringZ.from("setting kernel arguments of ", kernel_name_)); 388 } 389 } 390 391 private: 392 string kernel_name_; 393 cl_kernel kernel_; 394 CLContext context_; 395 } 396 397 final class CLProgram{ 398 public: 399 @nogc nothrow: 400 this(string source_path = "", CLContext context=null 401 , const(char)[] compilation_options="-I \"./\""){ 402 context_ = context; 403 404 enum MAX_SOURCE_SIZE = 0x100000; 405 406 FILE *fp; 407 char *source_str; 408 size_t source_size; 409 410 fp = fopen(RCStringZ.from(source_path).data.ptr, "r"); 411 if (!fp) { 412 printf("Failed to load kernel.\n"); 413 exit(1); 414 } 415 source_str = cast(char*)malloc(MAX_SOURCE_SIZE); 416 scope(exit) free(source_str); 417 source_size = fread( source_str, 1, MAX_SOURCE_SIZE, fp); 418 fclose( fp ); 419 420 createProgram(source_str, source_size, compilation_options); 421 createKernels(); 422 } 423 424 this(CTKernel ct_kernel, CLContext context=null 425 , const(char)[] compilation_options="-I \"./\""){ 426 context_ = context; 427 428 createProgram(ct_kernel.ptr, ct_kernel.length, compilation_options); 429 createKernels(); 430 } 431 432 ~this(){ 433 foreach(ref item; kernels_.byValue()){ 434 destroyFree(item); 435 item = null; 436 } 437 if(cl_program_ !is null){ 438 cl_int err = clReleaseProgram(cl_program_); 439 handleError(err, RCStringZ.from("releasing program")); 440 } 441 442 kernels_.free(); 443 } 444 bool createProgram(const char* source, size_t source_size, const(char)[] compilation_options = null){ 445 cl_int err; 446 RCStringZ cop; 447 if(compilation_options){ 448 //debug writeln(compilation_options); 449 cop = RCStringZ.from(compilation_options); 450 } 451 const size_t size_src = source_size; 452 453 cl_program_ = clCreateProgramWithSource(context_.getCLContext(), 1, 454 cast(const char **)&source, &size_src, &err); 455 handleError(err, RCStringZ.from("creating program with source data")); 456 cl_device_id dev_id = context_.getDevId(); 457 err = clBuildProgram(cl_program_, 1, &dev_id, cop.data.ptr, null, null); 458 459 cl_build_status build_status; 460 clGetProgramBuildInfo(cl_program_, context_.getDevId(), 461 CL_PROGRAM_BUILD_STATUS, cl_build_status.sizeof, &build_status, null); 462 463 if (build_status == CL_BUILD_PROGRAM_FAILURE) { 464 // Get the build log size 465 size_t _log_size; 466 clGetProgramBuildInfo(cl_program_, context_.getDevId(), CL_PROGRAM_BUILD_LOG, 0, null, &_log_size); 467 468 // Allocate memory for the log 469 char *log = cast(char*)malloc(_log_size); 470 471 // Get the build log 472 clGetProgramBuildInfo(cl_program_, context_.getDevId(), CL_PROGRAM_BUILD_LOG, _log_size, log, null); 473 474 // Print the build log 475 printf("Build log:\n%s\n", log); 476 477 // Free the allocated memory 478 free(log); 479 exit(-1); 480 } 481 482 /*handleError(err, RCStringZ.from(nogcFormat!"building program with source: %s\nUsing compilation options: %s\n"( 483 source, compilation_options)));*/ 484 return true; 485 } 486 487 bool createKernels(){ 488 cl_uint num_kernels = 0; 489 cl_int err; 490 err = clCreateKernelsInProgram(cl_program_, 0, null, &num_kernels); 491 if(num_kernels == 0) 492 err = CL_INVALID_BINARY; 493 if(err != CL_SUCCESS){ 494 char* build_log; 495 size_t log_size = 0; 496 497 clGetProgramBuildInfo(cl_program_, context_.getDevId(), 498 CL_PROGRAM_BUILD_LOG, 0, null, &log_size); 499 build_log = cast(char*)malloc(char.sizeof * log_size); 500 scope(exit) free(build_log); 501 clGetProgramBuildInfo(cl_program_, context_.getDevId(), 502 CL_PROGRAM_BUILD_LOG, log_size, build_log, null); 503 printf("%s \n", build_log); 504 handleError(err, RCStringZ.from("creating kernels")); 505 506 } 507 return true; 508 } 509 510 CLKernel getKernel(string kernel_name){ 511 CLKernel kernel; 512 513 if (auto kernptr = kernel_name in kernels_) 514 kernel = *kernptr; 515 else{ 516 kernel = mallocNew!CLKernel(context_, cl_program_, kernel_name); 517 kernels_[kernel_name] = kernel; 518 } 519 520 if(kernel is null){ 521 printf("kernel has been deleted or failed to create!\n"); 522 exit(EXIT_FAILURE); 523 } 524 525 return kernel; 526 } 527 528 void setCLContext(CLContext context) {context_ = context;} 529 530 private: 531 CLContext context_; 532 cl_program cl_program_; 533 Bcaa!(string, CLKernel) kernels_; 534 } 535 536 @nogc nothrow: 537 538 static void handleError()(cl_int err, RCStringZ msg, string f = __FILE__, int l = __LINE__){ 539 if(err != CL_SUCCESS){ 540 immutable(string) oerr = errorNumberToString(err); 541 printf("[OpenCL Error] in %s !: %s %s:%d\n", msg.data.ptr, oerr.ptr, f.ptr, l); 542 exit(EXIT_FAILURE); 543 } 544 } 545 546 immutable(string) errorNumberToString(cl_int errorNumber) 547 { 548 switch (errorNumber) 549 { 550 case CL_SUCCESS: 551 return "CL_SUCCESS"; 552 case CL_DEVICE_NOT_FOUND: 553 return "CL_DEVICE_NOT_FOUND"; 554 case CL_DEVICE_NOT_AVAILABLE: 555 return "CL_DEVICE_NOT_AVAILABLE"; 556 case CL_COMPILER_NOT_AVAILABLE: 557 return "CL_COMPILER_NOT_AVAILABLE"; 558 case CL_MEM_OBJECT_ALLOCATION_FAILURE: 559 return "CL_MEM_OBJECT_ALLOCATION_FAILURE"; 560 case CL_OUT_OF_RESOURCES: 561 return "CL_OUT_OF_RESOURCES"; 562 case CL_OUT_OF_HOST_MEMORY: 563 return "CL_OUT_OF_HOST_MEMORY"; 564 case CL_PROFILING_INFO_NOT_AVAILABLE: 565 return "CL_PROFILING_INFO_NOT_AVAILABLE"; 566 case CL_MEM_COPY_OVERLAP: 567 return "CL_MEM_COPY_OVERLAP"; 568 case CL_IMAGE_FORMAT_MISMATCH: 569 return "CL_IMAGE_FORMAT_MISMATCH"; 570 case CL_IMAGE_FORMAT_NOT_SUPPORTED: 571 return "CL_IMAGE_FORMAT_NOT_SUPPORTED"; 572 case CL_BUILD_PROGRAM_FAILURE: 573 return "CL_BUILD_PROGRAM_FAILURE"; 574 case CL_MAP_FAILURE: 575 return "CL_MAP_FAILURE"; 576 case CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST: 577 return "CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST"; 578 case CL_INVALID_VALUE: 579 return "CL_INVALID_VALUE"; 580 case CL_INVALID_DEVICE_TYPE: 581 return "CL_INVALID_DEVICE_TYPE"; 582 case CL_INVALID_PLATFORM: 583 return "CL_INVALID_PLATFORM"; 584 case CL_INVALID_DEVICE: 585 return "CL_INVALID_DEVICE"; 586 case CL_INVALID_CONTEXT: 587 return "CL_INVALID_CONTEXT"; 588 case CL_INVALID_QUEUE_PROPERTIES: 589 return "CL_INVALID_QUEUE_PROPERTIES"; 590 case CL_INVALID_COMMAND_QUEUE: 591 return "CL_INVALID_COMMAND_QUEUE"; 592 case CL_INVALID_HOST_PTR: 593 return "CL_INVALID_HOST_PTR"; 594 case CL_INVALID_MEM_OBJECT: 595 return "CL_INVALID_MEM_OBJECT"; 596 case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR: 597 return "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR"; 598 case CL_INVALID_IMAGE_SIZE: 599 return "CL_INVALID_IMAGE_SIZE"; 600 case CL_INVALID_SAMPLER: 601 return "CL_INVALID_SAMPLER"; 602 case CL_INVALID_BINARY: 603 return "CL_INVALID_BINARY"; 604 case CL_INVALID_BUILD_OPTIONS: 605 return "CL_INVALID_BUILD_OPTIONS"; 606 case CL_INVALID_PROGRAM: 607 return "CL_INVALID_PROGRAM"; 608 case CL_INVALID_PROGRAM_EXECUTABLE: 609 return "CL_INVALID_PROGRAM_EXECUTABLE"; 610 case CL_INVALID_KERNEL_NAME: 611 return "CL_INVALID_KERNEL_NAME"; 612 case CL_INVALID_KERNEL_DEFINITION: 613 return "CL_INVALID_KERNEL_DEFINITION"; 614 case CL_INVALID_KERNEL: 615 return "CL_INVALID_KERNEL"; 616 case CL_INVALID_ARG_INDEX: 617 return "CL_INVALID_ARG_INDEX"; 618 case CL_INVALID_ARG_VALUE: 619 return "CL_INVALID_ARG_VALUE"; 620 case CL_INVALID_ARG_SIZE: 621 return "CL_INVALID_ARG_SIZE"; 622 case CL_INVALID_KERNEL_ARGS: 623 return "CL_INVALID_KERNEL_ARGS"; 624 case CL_INVALID_WORK_DIMENSION: 625 return "CL_INVALID_WORK_DIMENSION"; 626 case CL_INVALID_WORK_GROUP_SIZE: 627 return "CL_INVALID_WORK_GROUP_SIZE"; 628 case CL_INVALID_WORK_ITEM_SIZE: 629 return "CL_INVALID_WORK_ITEM_SIZE"; 630 case CL_INVALID_GLOBAL_OFFSET: 631 return "CL_INVALID_GLOBAL_OFFSET"; 632 case CL_INVALID_EVENT_WAIT_LIST: 633 return "CL_INVALID_EVENT_WAIT_LIST"; 634 case CL_INVALID_EVENT: 635 return "CL_INVALID_EVENT"; 636 case CL_INVALID_OPERATION: 637 return "CL_INVALID_OPERATION"; 638 case CL_INVALID_GL_OBJECT: 639 return "CL_INVALID_GL_OBJECT"; 640 case CL_INVALID_BUFFER_SIZE: 641 return "CL_INVALID_BUFFER_SIZE"; 642 case CL_INVALID_MIP_LEVEL: 643 return "CL_INVALID_MIP_LEVEL"; 644 default: 645 return "Unknown error"; 646 } 647 } 648 649 static int getCLMemFlag(MemFlag mem_flag) 650 { 651 int ret = 0; 652 switch(mem_flag){ 653 case MEM_FLAG_READ_WRITE: 654 ret = ret | CL_MEM_READ_WRITE; break; 655 case MEM_FLAG_READ_ONLY: 656 ret = ret | CL_MEM_READ_ONLY; break; 657 case MEM_FLAG_WRITE_ONLY: 658 ret = ret | CL_MEM_WRITE_ONLY; break; 659 case MEM_FLAG_USE_HOST_PTR: 660 ret = ret | CL_MEM_USE_HOST_PTR; break; 661 case MEM_FLAG_ALLOC_HOST_PTR: 662 ret = ret | CL_MEM_ALLOC_HOST_PTR; break; 663 case MEM_FLAG_COPY_HOST_PTR: 664 ret = ret | CL_MEM_COPY_HOST_PTR; break; 665 default: break; 666 } 667 return ret; 668 } 669 670 void _assert(bool condition, string msg, string file = __FILE__, int line = __LINE__) 671 { 672 if(!condition) 673 { 674 printf("%s %s:%d\n", msg.ptr, file.ptr, line); 675 exit(EXIT_FAILURE); 676 } 677 } 678 679 // registering library kernels 680 enum CTKernel { 681 KGRAY = import("gray.cl"), 682 KYUV = import("yuv.cl"), 683 KHSV = import("hsv.cl"), 684 KSGM = import("sgm.cl"), 685 KINRANGE3 = import("inrange3.cl"), 686 KCOUNTNONZERO = import("countnonzero.cl"), 687 KMORPHED = import("morphed.cl"), 688 KRESIZE = import("resize.cl"), 689 KCONV = import("convolution.cl"), 690 KMEDIAN = import("medianfilteropencv.cl") 691 }