ViewVC Help
View File | Revision Log | Show Annotations | Download File
/cvs/OpenCL/OpenCL.xs
(Generate patch)

Comparing OpenCL/OpenCL.xs (file contents):
Revision 1.52 by root, Tue Apr 24 23:53:12 2012 UTC vs.
Revision 1.62 by root, Tue May 1 22:04:56 2012 UTC

57typedef cl_program OpenCL__Program; 57typedef cl_program OpenCL__Program;
58typedef cl_kernel OpenCL__Kernel; 58typedef cl_kernel OpenCL__Kernel;
59typedef cl_event OpenCL__Event; 59typedef cl_event OpenCL__Event;
60typedef cl_event OpenCL__UserEvent; 60typedef cl_event OpenCL__UserEvent;
61 61
62typedef struct mapped * OpenCL__Mapped;
63
62typedef SV *FUTURE; 64typedef SV *FUTURE;
65
66static HV
67 *stash_platform,
68 *stash_device,
69 *stash_context,
70 *stash_queue,
71 *stash_program,
72 *stash_kernel,
73 *stash_sampler,
74 *stash_event,
75 *stash_userevent,
76 *stash_memory,
77 *stash_buffer,
78 *stash_bufferobj,
79 *stash_image,
80 *stash_image1d,
81 *stash_image1darray,
82 *stash_image1dbuffer,
83 *stash_image2d,
84 *stash_image2darray,
85 *stash_image3d,
86 *stash_mapped,
87 *stash_mappedbuffer,
88 *stash_mappedimage;
63 89
64/*****************************************************************************/ 90/*****************************************************************************/
65 91
66// name must include a leading underscore 92// name must include a leading underscore
67// all of this horrors would be unneceesary if somebody wrote a proper OpenGL module 93// all of this horrors would be unneceesary if somebody wrote a proper OpenGL module
260 croak ("%s: %s is not a property list (either undef or [type => value, ...])", func, svname); 286 croak ("%s: %s is not a property list (either undef or [type => value, ...])", func, svname);
261} 287}
262 288
263/*****************************************************************************/ 289/*****************************************************************************/
264 290
265#define NEW_CLOBJ(class,ptr) sv_setref_pv (sv_newmortal (), class, ptr) 291static SV *
266#define PUSH_CLOBJ(class,ptr) PUSHs (NEW_CLOBJ (class, ptr)) 292new_clobj (HV *stash, IV id)
267#define XPUSH_CLOBJ(class,ptr) XPUSHs (NEW_CLOBJ (class, ptr)) 293{
294 return sv_2mortal (sv_bless (newRV_noinc (newSViv (id)), stash));
295}
296
297#define PUSH_CLOBJ(stash,id) PUSHs (new_clobj ((stash), (IV)(id)))
298#define XPUSH_CLOBJ(stash,id) XPUSHs (new_clobj ((stash), (IV)(id)))
268 299
269/* cl objects are either \$iv, or [$iv, ...] */ 300/* cl objects are either \$iv, or [$iv, ...] */
270/* they can be upgraded at runtime to the array form */ 301/* they can be upgraded at runtime to the array form */
271static void * 302static void *
272SvCLOBJ (const char *func, const char *svname, SV *sv, const char *pkg) 303SvCLOBJ (const char *func, const char *svname, SV *sv, const char *pkg)
273{ 304{
274 if (SvROK (sv) && sv_derived_from (sv, pkg)) 305 if (SvROK (sv) && sv_derived_from (sv, pkg))
275 {
276 SV *rv = SvRV (sv);
277
278#if CLOBJ_PUSH
279 if (SvTYPE (rv) == SVt_PVAV)
280 rv = AvARRAY (rv)[0];
281#endif
282
283 return (void *)SvIV (rv); 306 return (void *)SvIV (SvRV (sv));
284 }
285 307
286 croak ("%s: %s is not of type %s", func, svname, pkg); 308 croak ("%s: %s is not of type %s", func, svname, pkg);
287} 309}
288
289#if CLOBJ_PUSH
290
291static void
292CLOBJ_push (SV *self, SV *data)
293{
294 SV *rv = SvRV (self);
295
296 if (SvTYPE (rv) != SVt_PVAV)
297 {
298 AV *av = newAV ();
299 av_push (av, rv);
300 rv = (SV *)av;
301 SvRV_set (self, rv);
302 }
303
304 av_push ((AV *)rv, data);
305}
306
307static SV *
308sv_struct (STRLEN size)
309{
310 SV *sv = newSV (size);
311 SvPOK_only (sv);
312 return sv;
313}
314
315static void *
316CLOBJ_push_struct (SV *self, STRLEN size)
317{
318 SV *sv = sv_struct (size);
319 CLOBJ_push (self, sv);
320 return SvPVX (sv);
321}
322
323#endif
324 310
325/*****************************************************************************/ 311/*****************************************************************************/
326/* callback stuff */ 312/* callback stuff */
327 313
328/* default context callback, log to stderr */ 314/* default context callback, log to stderr */
386 372
387 X_UNLOCK (eq_lock); 373 X_UNLOCK (eq_lock);
388 374
389 return res; 375 return res;
390} 376}
391
392#if 0
393static void
394mem_free (pTHX_ void *p)
395{
396 free (p);
397}
398//SAVEDESTRUCTOR_X (mem_free, item);
399#endif
400 377
401static void 378static void
402eq_poll (void) 379eq_poll (void)
403{ 380{
404 eq_item *item; 381 eq_item *item;
470 } 447 }
471 448
472static SV * 449static SV *
473new_clobj_context (cl_context ctx, void *user_data) 450new_clobj_context (cl_context ctx, void *user_data)
474{ 451{
475 SV *sv = NEW_CLOBJ ("OpenCL::Context", ctx); 452 SV *sv = new_clobj (stash_context, (IV)ctx);
476 453
477 if (user_data) 454 if (user_data)
478 sv_magicext (SvRV (sv), user_data, PERL_MAGIC_ext, 0, 0, 0); 455 sv_magicext (SvRV (sv), user_data, PERL_MAGIC_ext, 0, 0, 0);
479 456
480 return sv; 457 return sv;
487 464
488static void 465static void
489eq_program_push (void *data1, void *data2, void *data3) 466eq_program_push (void *data1, void *data2, void *data3)
490{ 467{
491 dSP; 468 dSP;
492 PUSH_CLOBJ ("OpenCL::Program", data1); 469 PUSH_CLOBJ (stash_program, data1);
493 PUTBACK; 470 PUTBACK;
494} 471}
495 472
496static eq_vtbl eq_program_vtbl = { 1, eq_program_push }; 473static eq_vtbl eq_program_vtbl = { 1, eq_program_push };
497 474
544 521
545static void 522static void
546eq_event_push (void *data1, void *data2, void *data3) 523eq_event_push (void *data1, void *data2, void *data3)
547{ 524{
548 dSP; 525 dSP;
549 PUSH_CLOBJ ("OpenCL::Event", data1); 526 PUSH_CLOBJ (stash_event, data1);
550 PUSHs (sv_2mortal (newSViv ((IV)data2))); 527 PUSHs (sv_2mortal (newSViv ((IV)data2)));
551 PUTBACK; 528 PUTBACK;
552} 529}
553 530
554static eq_vtbl eq_event_vtbl = { 1, eq_event_push }; 531static eq_vtbl eq_event_vtbl = { 1, eq_event_push };
559 clRetainEvent (event); 536 clRetainEvent (event);
560 eq_enq (&eq_event_vtbl, user_data, (void *)event, (void *)(IV)event_command_exec_status, 0); 537 eq_enq (&eq_event_vtbl, user_data, (void *)event, (void *)(IV)event_command_exec_status, 0);
561} 538}
562 539
563/*****************************************************************************/ 540/*****************************************************************************/
541/* utilities for XS code */
564 542
565static size_t 543static size_t
566img_row_pitch (cl_mem img) 544img_row_pitch (cl_mem img)
567{ 545{
568 size_t res; 546 size_t res;
569 clGetImageInfo (img, CL_IMAGE_ROW_PITCH, sizeof (res), &res, 0); 547 clGetImageInfo (img, CL_IMAGE_ROW_PITCH, sizeof (res), &res, 0);
570 return res; 548 return res;
571} 549}
572 550
573static cl_event * 551static cl_event *
574event_list (SV **items, cl_uint *rcount) 552event_list (SV **items, cl_uint *rcount, cl_event extra)
575{ 553{
576 cl_uint count = *rcount; 554 cl_uint count = *rcount;
577 555
578 if (!count) 556 if (count > 0x7fffffff) // yeha, it's a hack - the caller might have underflowed
557 count = 0;
558
559 if (!count && !extra)
579 return 0; 560 return 0;
580 561
581 cl_event *list = tmpbuf (sizeof (cl_event) * count); 562 cl_event *list = tmpbuf (sizeof (cl_event) * (count + 1));
582 int i = 0; 563 int i = 0;
583 564
584 do 565 while (count--)
585 {
586 --count;
587 if (SvOK (items [count])) 566 if (SvOK (items [count]))
588 list [i++] = SvCLOBJ ("clEnqueue", "wait_events", items [count], "OpenCL::Event"); 567 list [i++] = SvCLOBJ ("clEnqueue", "wait_events", items [count], "OpenCL::Event");
589 } 568
590 while (count); 569 if (extra)
570 list [i++] = extra;
591 571
592 *rcount = i; 572 *rcount = i;
593 573
594 return i ? list : 0; 574 return i ? list : 0;
595} 575}
596 576
597#define EVENT_LIST(items,count) \ 577#define EVENT_LIST(skip) \
598 cl_uint event_list_count = (count); \ 578 cl_uint event_list_count = items - (skip); \
599 cl_event *event_list_ptr = event_list (&ST (items), &event_list_count) 579 cl_event *event_list_ptr = event_list (&ST (skip), &event_list_count, 0)
600 580
601#define INFO(class) \ 581#define INFO(class) \
602{ \ 582{ \
603 size_t size; \ 583 size_t size; \
604 NEED_SUCCESS (Get ## class ## Info, (self, name, 0, 0, &size)); \ 584 NEED_SUCCESS (Get ## class ## Info, (self, name, 0, 0, &size)); \
607 SvPOK_only (sv); \ 587 SvPOK_only (sv); \
608 SvCUR_set (sv, size); \ 588 SvCUR_set (sv, size); \
609 NEED_SUCCESS (Get ## class ## Info, (self, name, size, SvPVX (sv), 0)); \ 589 NEED_SUCCESS (Get ## class ## Info, (self, name, size, SvPVX (sv), 0)); \
610 XPUSHs (sv); \ 590 XPUSHs (sv); \
611} 591}
592
593/*****************************************************************************/
594/* mapped_xxx */
595
596static OpenCL__Mapped
597SvMAPPED (SV *sv)
598{
599 // no typechecking atm., keep your fingers crossed
600 return (OpenCL__Mapped)SvMAGIC (SvRV (sv))->mg_ptr;
601}
602
603struct mapped
604{
605 cl_command_queue queue;
606 cl_mem memobj;
607 void *ptr;
608 size_t cb;
609 cl_event event;
610 size_t row_pitch;
611 size_t slice_pitch;
612};
613
614static SV *
615mapped_new (HV *stash, cl_command_queue queue, cl_mem memobj, cl_map_flags flags, void *ptr, size_t cb, cl_event ev, size_t row_pitch, size_t slice_pitch)
616{
617 SV *data = newSV (0);
618 SvUPGRADE (data, SVt_PVMG);
619
620 OpenCL__Mapped mapped;
621 New (0, mapped, 1, struct mapped);
622
623 clRetainCommandQueue (queue);
624
625 mapped->queue = queue;
626 mapped->memobj = memobj;
627 mapped->ptr = ptr;
628 mapped->cb = cb;
629 mapped->event = ev;
630 mapped->row_pitch = row_pitch;
631 mapped->slice_pitch = slice_pitch;
632
633 sv_magicext (data, 0, PERL_MAGIC_ext, 0, (char *)mapped, 0);
634
635 if (SvLEN (data))
636 Safefree (data);
637
638 SvPVX (data) = (char *)ptr;
639 SvCUR_set (data, cb);
640 SvLEN_set (data, 0);
641 SvPOK_only (data);
642
643 SV *obj = sv_2mortal (sv_bless (newRV_noinc (data), stash));
644
645 if (!(flags & CL_MAP_WRITE))
646 SvREADONLY_on (data);
647
648 return obj;
649}
650
651static void
652mapped_detach (SV *sv, OpenCL__Mapped mapped)
653{
654 SV *data = SvRV (sv);
655
656 // the next check checks both before AND after detach, where SvPVX should STILL be 0
657 if (SvPVX (data) != (char *)mapped->ptr)
658 warn ("FATAL: OpenCL memory mapped scalar changed location, detected");
659 else
660 {
661 SvREADONLY_off (data);
662 SvCUR_set (data, 0);
663 SvPVX (data) = 0;
664 SvOK_off (data);
665 }
666
667 mapped->ptr = 0;
668}
669
670static void
671mapped_unmap (SV *self, OpenCL__Mapped mapped, cl_command_queue queue, SV **wait_list, cl_uint event_list_count)
672{
673 cl_event *event_list_ptr = event_list (wait_list, &event_list_count, mapped->event);
674 cl_event ev;
675
676 NEED_SUCCESS (EnqueueUnmapMemObject, (queue, mapped->memobj, mapped->ptr, event_list_count, event_list_ptr, &ev));
677
678 clReleaseEvent (mapped->event);
679 mapped->event = ev;
680
681 mapped_detach (self, mapped);
682}
683
684/*****************************************************************************/
612 685
613MODULE = OpenCL PACKAGE = OpenCL 686MODULE = OpenCL PACKAGE = OpenCL
614 687
615PROTOTYPES: ENABLE 688PROTOTYPES: ENABLE
616 689
644 }; 717 };
645 718
646 for (civ = const_iv + sizeof (const_iv) / sizeof (const_iv [0]); civ > const_iv; civ--) 719 for (civ = const_iv + sizeof (const_iv) / sizeof (const_iv [0]); civ > const_iv; civ--)
647 newCONSTSUB (stash, (char *)civ[-1].name, newSViv (civ[-1].iv)); 720 newCONSTSUB (stash, (char *)civ[-1].name, newSViv (civ[-1].iv));
648 721
722 stash_platform = gv_stashpv ("OpenCL::Platform", GV_ADD);
723 stash_device = gv_stashpv ("OpenCL::Device", GV_ADD);
724 stash_context = gv_stashpv ("OpenCL::Context", GV_ADD);
725 stash_queue = gv_stashpv ("OpenCL::Queue", GV_ADD);
726 stash_program = gv_stashpv ("OpenCL::Program", GV_ADD);
727 stash_kernel = gv_stashpv ("OpenCL::Kernel", GV_ADD);
728 stash_sampler = gv_stashpv ("OpenCL::Sampler", GV_ADD);
729 stash_event = gv_stashpv ("OpenCL::Event", GV_ADD);
730 stash_userevent = gv_stashpv ("OpenCL::UserEvent", GV_ADD);
731 stash_memory = gv_stashpv ("OpenCL::Memory", GV_ADD);
732 stash_buffer = gv_stashpv ("OpenCL::Buffer", GV_ADD);
733 stash_bufferobj = gv_stashpv ("OpenCL::BufferObj", GV_ADD);
734 stash_image = gv_stashpv ("OpenCL::Image", GV_ADD);
735 stash_image1d = gv_stashpv ("OpenCL::Image1D", GV_ADD);
736 stash_image1darray = gv_stashpv ("OpenCL::Image1DArray", GV_ADD);
737 stash_image1dbuffer = gv_stashpv ("OpenCL::Image1DBuffer", GV_ADD);
738 stash_image2d = gv_stashpv ("OpenCL::Image2D", GV_ADD);
739 stash_image2darray = gv_stashpv ("OpenCL::Image2DArray", GV_ADD);
740 stash_image3d = gv_stashpv ("OpenCL::Image3D", GV_ADD);
741 stash_mapped = gv_stashpv ("OpenCL::Mapped", GV_ADD);
742 stash_mappedbuffer = gv_stashpv ("OpenCL::MappedBuffer", GV_ADD);
743 stash_mappedimage = gv_stashpv ("OpenCL::MappedImage", GV_ADD);
744
649 sv_setiv (perl_get_sv ("OpenCL::POLL_FUNC", TRUE), (IV)eq_poll_interrupt); 745 sv_setiv (perl_get_sv ("OpenCL::POLL_FUNC", TRUE), (IV)eq_poll_interrupt);
650} 746}
651 747
652cl_int 748cl_int
653errno () 749errno ()
655 RETVAL = res; 751 RETVAL = res;
656 OUTPUT: 752 OUTPUT:
657 RETVAL 753 RETVAL
658 754
659const char * 755const char *
660err2str (cl_int err) 756err2str (cl_int err = res)
661 757
662const char * 758const char *
663enum2str (cl_uint value) 759enum2str (cl_uint value)
664 760
665void 761void
673 list = tmpbuf (sizeof (*list) * count); 769 list = tmpbuf (sizeof (*list) * count);
674 NEED_SUCCESS (GetPlatformIDs, (count, list, 0)); 770 NEED_SUCCESS (GetPlatformIDs, (count, list, 0));
675 771
676 EXTEND (SP, count); 772 EXTEND (SP, count);
677 for (i = 0; i < count; ++i) 773 for (i = 0; i < count; ++i)
678 PUSH_CLOBJ ("OpenCL::Platform", list [i]); 774 PUSH_CLOBJ (stash_platform, (IV)list [i]);
679 775
680void 776void
681context_from_type (cl_context_properties *properties = 0, cl_device_type type = CL_DEVICE_TYPE_DEFAULT, SV *notify = &PL_sv_undef) 777context_from_type (cl_context_properties *properties = 0, cl_device_type type = CL_DEVICE_TYPE_DEFAULT, SV *notify = &PL_sv_undef)
682 PPCODE: 778 PPCODE:
683 CONTEXT_NOTIFY_CALLBACK; 779 CONTEXT_NOTIFY_CALLBACK;
690 /* der Gipfel der Kunst */ 786 /* der Gipfel der Kunst */
691 787
692void 788void
693wait_for_events (...) 789wait_for_events (...)
694 CODE: 790 CODE:
695 EVENT_LIST (0, items); 791 EVENT_LIST (0);
696 NEED_SUCCESS (WaitForEvents, (event_list_count, event_list_ptr)); 792 NEED_SUCCESS (WaitForEvents, (event_list_count, event_list_ptr));
697 793
698PROTOTYPES: DISABLE 794PROTOTYPES: DISABLE
699 795
700MODULE = OpenCL PACKAGE = OpenCL::Platform 796MODULE = OpenCL PACKAGE = OpenCL::Platform
743 list = tmpbuf (sizeof (*list) * count); 839 list = tmpbuf (sizeof (*list) * count);
744 NEED_SUCCESS (GetDeviceIDs, (self, type, count, list, 0)); 840 NEED_SUCCESS (GetDeviceIDs, (self, type, count, list, 0));
745 841
746 EXTEND (SP, count); 842 EXTEND (SP, count);
747 for (i = 0; i < count; ++i) 843 for (i = 0; i < count; ++i)
748 PUSHs (sv_setref_pv (sv_newmortal (), "OpenCL::Device", list [i])); 844 PUSH_CLOBJ (stash_device, list [i]);
749 845
750void 846void
751context (OpenCL::Platform self, cl_context_properties *properties, SV *devices, SV *notify = &PL_sv_undef) 847context (OpenCL::Platform self, cl_context_properties *properties, SV *devices, SV *notify = &PL_sv_undef)
752 PPCODE: 848 PPCODE:
753 if (!SvROK (devices) || SvTYPE (SvRV (devices)) != SVt_PVAV) 849 if (!SvROK (devices) || SvTYPE (SvRV (devices)) != SVt_PVAV)
789 PPCODE: 885 PPCODE:
790 cl_device_type value [1]; 886 cl_device_type value [1];
791 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_TYPE, sizeof (value), value, 0)); 887 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_TYPE, sizeof (value), value, 0));
792 EXTEND (SP, 1); 888 EXTEND (SP, 1);
793 const int i = 0; 889 const int i = 0;
794 PUSHs (sv_2mortal (newSViv (value [i]))); 890 PUSHs (sv_2mortal (newSVuv (value [i])));
795 891
796void 892void
797vendor_id (OpenCL::Device self) 893vendor_id (OpenCL::Device self)
798 ALIAS: 894 ALIAS:
799 vendor_id = CL_DEVICE_VENDOR_ID 895 vendor_id = CL_DEVICE_VENDOR_ID
944 PPCODE: 1040 PPCODE:
945 cl_command_queue_properties value [1]; 1041 cl_command_queue_properties value [1];
946 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_QUEUE_PROPERTIES, sizeof (value), value, 0)); 1042 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_QUEUE_PROPERTIES, sizeof (value), value, 0));
947 EXTEND (SP, 1); 1043 EXTEND (SP, 1);
948 const int i = 0; 1044 const int i = 0;
949 PUSHs (sv_2mortal (newSViv (value [i]))); 1045 PUSHs (sv_2mortal (newSVuv (value [i])));
950 1046
951void 1047void
952platform (OpenCL::Device self) 1048platform (OpenCL::Device self)
953 PPCODE: 1049 PPCODE:
954 cl_platform_id value [1]; 1050 cl_platform_id value [1];
955 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_PLATFORM, sizeof (value), value, 0)); 1051 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_PLATFORM, sizeof (value), value, 0));
956 EXTEND (SP, 1); 1052 EXTEND (SP, 1);
957 const int i = 0; 1053 const int i = 0;
958 { 1054 {
959 PUSH_CLOBJ ("OpenCL::Platform", value [i]); 1055 PUSH_CLOBJ (stash_platform, value [i]);
960 } 1056 }
961 1057
962void 1058void
963name (OpenCL::Device self) 1059name (OpenCL::Device self)
964 ALIAS: 1060 ALIAS:
983 cl_device_id value [1]; 1079 cl_device_id value [1];
984 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_PARENT_DEVICE_EXT, sizeof (value), value, 0)); 1080 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_PARENT_DEVICE_EXT, sizeof (value), value, 0));
985 EXTEND (SP, 1); 1081 EXTEND (SP, 1);
986 const int i = 0; 1082 const int i = 0;
987 { 1083 {
988 PUSH_CLOBJ ("OpenCL::Device", value [i]); 1084 PUSH_CLOBJ (stash_device, value [i]);
989 } 1085 }
990 1086
991void 1087void
992partition_types_ext (OpenCL::Device self) 1088partition_types_ext (OpenCL::Device self)
993 ALIAS: 1089 ALIAS:
1020 1116
1021void 1117void
1022queue (OpenCL::Context self, OpenCL::Device device, cl_command_queue_properties properties = 0) 1118queue (OpenCL::Context self, OpenCL::Device device, cl_command_queue_properties properties = 0)
1023 PPCODE: 1119 PPCODE:
1024 NEED_SUCCESS_ARG (cl_command_queue queue, CreateCommandQueue, (self, device, properties, &res)); 1120 NEED_SUCCESS_ARG (cl_command_queue queue, CreateCommandQueue, (self, device, properties, &res));
1025 XPUSH_CLOBJ ("OpenCL::Queue", queue); 1121 XPUSH_CLOBJ (stash_queue, queue);
1026 1122
1027void 1123void
1028user_event (OpenCL::Context self) 1124user_event (OpenCL::Context self)
1029 PPCODE: 1125 PPCODE:
1030 NEED_SUCCESS_ARG (cl_event ev, CreateUserEvent, (self, &res)); 1126 NEED_SUCCESS_ARG (cl_event ev, CreateUserEvent, (self, &res));
1031 XPUSH_CLOBJ ("OpenCL::UserEvent", ev); 1127 XPUSH_CLOBJ (stash_userevent, ev);
1032 1128
1033void 1129void
1034buffer (OpenCL::Context self, cl_mem_flags flags, size_t len) 1130buffer (OpenCL::Context self, cl_mem_flags flags, size_t len)
1035 PPCODE: 1131 PPCODE:
1036 if (flags & (CL_MEM_USE_HOST_PTR | CL_MEM_COPY_HOST_PTR)) 1132 if (flags & (CL_MEM_USE_HOST_PTR | CL_MEM_COPY_HOST_PTR))
1037 croak ("OpenCL::Context::buffer: cannot use/copy host ptr when no data is given, use $context->buffer_sv instead?"); 1133 croak ("OpenCL::Context::buffer: cannot use/copy host ptr when no data is given, use $context->buffer_sv instead?");
1038 1134
1039 NEED_SUCCESS_ARG (cl_mem mem, CreateBuffer, (self, flags, len, 0, &res)); 1135 NEED_SUCCESS_ARG (cl_mem mem, CreateBuffer, (self, flags, len, 0, &res));
1040 XPUSH_CLOBJ ("OpenCL::BufferObj", mem); 1136 XPUSH_CLOBJ (stash_bufferobj, mem);
1041 1137
1042void 1138void
1043buffer_sv (OpenCL::Context self, cl_mem_flags flags, SV *data) 1139buffer_sv (OpenCL::Context self, cl_mem_flags flags, SV *data)
1044 PPCODE: 1140 PPCODE:
1045 STRLEN len; 1141 STRLEN len;
1046 char *ptr = SvOK (data) ? SvPVbyte (data, len) : 0; 1142 char *ptr = SvOK (data) ? SvPVbyte (data, len) : 0;
1047 if (!(flags & (CL_MEM_USE_HOST_PTR | CL_MEM_COPY_HOST_PTR))) 1143 if (!(flags & (CL_MEM_USE_HOST_PTR | CL_MEM_COPY_HOST_PTR)))
1048 croak ("OpenCL::Context::buffer_sv: you have to specify use or copy host ptr when buffer data is given, use $context->buffer instead?"); 1144 croak ("OpenCL::Context::buffer_sv: you have to specify use or copy host ptr when buffer data is given, use $context->buffer instead?");
1049 NEED_SUCCESS_ARG (cl_mem mem, CreateBuffer, (self, flags, len, ptr, &res)); 1145 NEED_SUCCESS_ARG (cl_mem mem, CreateBuffer, (self, flags, len, ptr, &res));
1050 XPUSH_CLOBJ ("OpenCL::BufferObj", mem); 1146 XPUSH_CLOBJ (stash_bufferobj, mem);
1051 1147
1052#if CL_VERSION_1_2 1148#if CL_VERSION_1_2
1053 1149
1054void 1150void
1055image (OpenCL::Context self, cl_mem_flags flags, cl_channel_order channel_order, cl_channel_type channel_type, cl_mem_object_type type, size_t width, size_t height, size_t depth, size_t array_size = 0, size_t row_pitch = 0, size_t slice_pitch = 0, cl_uint num_mip_level = 0, cl_uint num_samples = 0, SV *data = &PL_sv_undef) 1151image (OpenCL::Context self, cl_mem_flags flags, cl_channel_order channel_order, cl_channel_type channel_type, cl_mem_object_type type, size_t width, size_t height, size_t depth = 0, size_t array_size = 0, size_t row_pitch = 0, size_t slice_pitch = 0, cl_uint num_mip_level = 0, cl_uint num_samples = 0, SV *data = &PL_sv_undef)
1056 PPCODE: 1152 PPCODE:
1057 STRLEN len; 1153 STRLEN len;
1058 char *ptr = SvOK (data) ? SvPVbyte (data, len) : 0; 1154 char *ptr = SvOK (data) ? SvPVbyte (data, len) : 0;
1059 const cl_image_format format = { channel_order, channel_type }; 1155 const cl_image_format format = { channel_order, channel_type };
1060 const cl_image_desc desc = { 1156 const cl_image_desc desc = {
1063 array_size, row_pitch, slice_pitch, 1159 array_size, row_pitch, slice_pitch,
1064 num_mip_level, num_samples, 1160 num_mip_level, num_samples,
1065 type == CL_MEM_OBJECT_IMAGE1D_BUFFER ? (cl_mem)SvCLOBJ ("OpenCL::Context::Image", "data", data, "OpenCL::Buffer") : 0 1161 type == CL_MEM_OBJECT_IMAGE1D_BUFFER ? (cl_mem)SvCLOBJ ("OpenCL::Context::Image", "data", data, "OpenCL::Buffer") : 0
1066 }; 1162 };
1067 NEED_SUCCESS_ARG (cl_mem mem, CreateImage, (self, flags, &format, &desc, ptr, &res)); 1163 NEED_SUCCESS_ARG (cl_mem mem, CreateImage, (self, flags, &format, &desc, ptr, &res));
1068 char *klass = "OpenCL::Image"; 1164 HV *stash = stash_image;
1069 switch (type) 1165 switch (type)
1070 { 1166 {
1071 case CL_MEM_OBJECT_IMAGE1D_BUFFER: klass = "OpenCL::Image1DBuffer"; break; 1167 case CL_MEM_OBJECT_IMAGE1D_BUFFER: stash = stash_image1dbuffer; break;
1072 case CL_MEM_OBJECT_IMAGE1D: klass = "OpenCL::Image1D"; break; 1168 case CL_MEM_OBJECT_IMAGE1D: stash = stash_image1d; break;
1073 case CL_MEM_OBJECT_IMAGE1D_ARRAY: klass = "OpenCL::Image2DArray"; break; 1169 case CL_MEM_OBJECT_IMAGE1D_ARRAY: stash = stash_image2darray; break;
1074 case CL_MEM_OBJECT_IMAGE2D: klass = "OpenCL::Image2D"; break; 1170 case CL_MEM_OBJECT_IMAGE2D: stash = stash_image2d; break;
1075 case CL_MEM_OBJECT_IMAGE2D_ARRAY: klass = "OpenCL::Image2DArray"; break; 1171 case CL_MEM_OBJECT_IMAGE2D_ARRAY: stash = stash_image2darray; break;
1076 case CL_MEM_OBJECT_IMAGE3D: klass = "OpenCL::Image3D"; break; 1172 case CL_MEM_OBJECT_IMAGE3D: stash = stash_image3d; break;
1077 } 1173 }
1078 XPUSH_CLOBJ (klass, mem); 1174 XPUSH_CLOBJ (stash, mem);
1079 1175
1080#endif 1176#endif
1081 1177
1082void 1178void
1083image2d (OpenCL::Context self, cl_mem_flags flags, cl_channel_order channel_order, cl_channel_type channel_type, size_t width, size_t height, size_t row_pitch = 0, SV *data = &PL_sv_undef) 1179image2d (OpenCL::Context self, cl_mem_flags flags, cl_channel_order channel_order, cl_channel_type channel_type, size_t width, size_t height, size_t row_pitch = 0, SV *data = &PL_sv_undef)
1089 NEED_SUCCESS_ARG (cl_mem mem, CreateImage2D, (self, flags, &format, width, height, row_pitch, ptr, &res)); 1185 NEED_SUCCESS_ARG (cl_mem mem, CreateImage2D, (self, flags, &format, width, height, row_pitch, ptr, &res));
1090#else 1186#else
1091 const cl_image_desc desc = { CL_MEM_OBJECT_IMAGE2D, width, height, 0, 0, row_pitch, 0, 0, 0, 0 }; 1187 const cl_image_desc desc = { CL_MEM_OBJECT_IMAGE2D, width, height, 0, 0, row_pitch, 0, 0, 0, 0 };
1092 NEED_SUCCESS_ARG (cl_mem mem, CreateImage, (self, flags, &format, &desc, ptr, &res)); 1188 NEED_SUCCESS_ARG (cl_mem mem, CreateImage, (self, flags, &format, &desc, ptr, &res));
1093#endif 1189#endif
1094 XPUSH_CLOBJ ("OpenCL::Image2D", mem); 1190 XPUSH_CLOBJ (stash_image2d, mem);
1095 1191
1096void 1192void
1097image3d (OpenCL::Context self, cl_mem_flags flags, cl_channel_order channel_order, cl_channel_type channel_type, size_t width, size_t height, size_t depth, size_t row_pitch = 0, size_t slice_pitch = 0, SV *data = &PL_sv_undef) 1193image3d (OpenCL::Context self, cl_mem_flags flags, cl_channel_order channel_order, cl_channel_type channel_type, size_t width, size_t height, size_t depth, size_t row_pitch = 0, size_t slice_pitch = 0, SV *data = &PL_sv_undef)
1098 PPCODE: 1194 PPCODE:
1099 STRLEN len; 1195 STRLEN len;
1103 NEED_SUCCESS_ARG (cl_mem mem, CreateImage3D, (self, flags, &format, width, height, depth, row_pitch, slice_pitch, ptr, &res)); 1199 NEED_SUCCESS_ARG (cl_mem mem, CreateImage3D, (self, flags, &format, width, height, depth, row_pitch, slice_pitch, ptr, &res));
1104#else 1200#else
1105 const cl_image_desc desc = { CL_MEM_OBJECT_IMAGE3D, width, height, depth, 0, row_pitch, slice_pitch, 0, 0, 0 }; 1201 const cl_image_desc desc = { CL_MEM_OBJECT_IMAGE3D, width, height, depth, 0, row_pitch, slice_pitch, 0, 0, 0 };
1106 NEED_SUCCESS_ARG (cl_mem mem, CreateImage, (self, flags, &format, &desc, ptr, &res)); 1202 NEED_SUCCESS_ARG (cl_mem mem, CreateImage, (self, flags, &format, &desc, ptr, &res));
1107#endif 1203#endif
1108 XPUSH_CLOBJ ("OpenCL::Image3D", mem); 1204 XPUSH_CLOBJ (stash_image3d, mem);
1109 1205
1110#if cl_apple_gl_sharing || cl_khr_gl_sharing 1206#if cl_apple_gl_sharing || cl_khr_gl_sharing
1111 1207
1112void 1208void
1113gl_buffer (OpenCL::Context self, cl_mem_flags flags, cl_GLuint bufobj) 1209gl_buffer (OpenCL::Context self, cl_mem_flags flags, cl_GLuint bufobj)
1114 PPCODE: 1210 PPCODE:
1115 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLBuffer, (self, flags, bufobj, &res)); 1211 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLBuffer, (self, flags, bufobj, &res));
1116 XPUSH_CLOBJ ("OpenCL::BufferObj", mem); 1212 XPUSH_CLOBJ (stash_bufferobj, mem);
1117 1213
1118void 1214void
1119gl_renderbuffer (OpenCL::Context self, cl_mem_flags flags, cl_GLuint renderbuffer) 1215gl_renderbuffer (OpenCL::Context self, cl_mem_flags flags, cl_GLuint renderbuffer)
1120 PPCODE: 1216 PPCODE:
1121 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLRenderbuffer, (self, flags, renderbuffer, &res)); 1217 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLRenderbuffer, (self, flags, renderbuffer, &res));
1122 XPUSH_CLOBJ ("OpenCL::Image2D", mem); 1218 XPUSH_CLOBJ (stash_image2d, mem);
1123 1219
1124#if CL_VERSION_1_2 1220#if CL_VERSION_1_2
1125 1221
1126void 1222void
1127gl_texture (OpenCL::Context self, cl_mem_flags flags, cl_GLenum target, cl_GLint miplevel, cl_GLuint texture) 1223gl_texture (OpenCL::Context self, cl_mem_flags flags, cl_GLenum target, cl_GLint miplevel, cl_GLuint texture)
1128 ALIAS: 1224 ALIAS:
1129 PPCODE: 1225 PPCODE:
1130 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture, (self, flags, target, miplevel, texture, &res)); 1226 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture, (self, flags, target, miplevel, texture, &res));
1131 cl_gl_object_type type; 1227 cl_gl_object_type type;
1132 NEED_SUCCESS (GetGLObjectInfo, (mem, &type, 0)); // TODO: use target instead? 1228 NEED_SUCCESS (GetGLObjectInfo, (mem, &type, 0)); // TODO: use target instead?
1133 char *klass = "OpenCL::Memory"; 1229 HV *stash = stash_memory;
1134 switch (type) 1230 switch (type)
1135 { 1231 {
1136 case CL_GL_OBJECT_TEXTURE_BUFFER: klass = "OpenCL::Image1DBuffer"; break; 1232 case CL_GL_OBJECT_TEXTURE_BUFFER: stash = stash_image1dbuffer; break;
1137 case CL_GL_OBJECT_TEXTURE1D: klass = "OpenCL::Image1D"; break; 1233 case CL_GL_OBJECT_TEXTURE1D: stash = stash_image1d; break;
1138 case CL_GL_OBJECT_TEXTURE1D_ARRAY: klass = "OpenCL::Image2DArray"; break; 1234 case CL_GL_OBJECT_TEXTURE1D_ARRAY: stash = stash_image2darray; break;
1139 case CL_GL_OBJECT_TEXTURE2D: klass = "OpenCL::Image2D"; break; 1235 case CL_GL_OBJECT_TEXTURE2D: stash = stash_image2d; break;
1140 case CL_GL_OBJECT_TEXTURE2D_ARRAY: klass = "OpenCL::Image2DArray"; break; 1236 case CL_GL_OBJECT_TEXTURE2D_ARRAY: stash = stash_image2darray; break;
1141 case CL_GL_OBJECT_TEXTURE3D: klass = "OpenCL::Image3D"; break; 1237 case CL_GL_OBJECT_TEXTURE3D: stash = stash_image3d; break;
1142 } 1238 }
1143 XPUSH_CLOBJ (klass, mem); 1239 XPUSH_CLOBJ (stash, mem);
1144 1240
1145#endif 1241#endif
1146 1242
1147void 1243void
1148gl_texture2d (OpenCL::Context self, cl_mem_flags flags, cl_GLenum target, cl_GLint miplevel, cl_GLuint texture) 1244gl_texture2d (OpenCL::Context self, cl_mem_flags flags, cl_GLenum target, cl_GLint miplevel, cl_GLuint texture)
1150#if PREFER_1_1 1246#if PREFER_1_1
1151 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture2D, (self, flags, target, miplevel, texture, &res)); 1247 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture2D, (self, flags, target, miplevel, texture, &res));
1152#else 1248#else
1153 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture , (self, flags, target, miplevel, texture, &res)); 1249 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture , (self, flags, target, miplevel, texture, &res));
1154#endif 1250#endif
1155 XPUSH_CLOBJ ("OpenCL::Image2D", mem); 1251 XPUSH_CLOBJ (stash_image2d, mem);
1156 1252
1157void 1253void
1158gl_texture3d (OpenCL::Context self, cl_mem_flags flags, cl_GLenum target, cl_GLint miplevel, cl_GLuint texture) 1254gl_texture3d (OpenCL::Context self, cl_mem_flags flags, cl_GLenum target, cl_GLint miplevel, cl_GLuint texture)
1159 PPCODE: 1255 PPCODE:
1160#if PREFER_1_1 1256#if PREFER_1_1
1161 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture3D, (self, flags, target, miplevel, texture, &res)); 1257 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture3D, (self, flags, target, miplevel, texture, &res));
1162#else 1258#else
1163 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture , (self, flags, target, miplevel, texture, &res)); 1259 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture , (self, flags, target, miplevel, texture, &res));
1164#endif 1260#endif
1165 XPUSH_CLOBJ ("OpenCL::Image3D", mem); 1261 XPUSH_CLOBJ (stash_image3d, mem);
1166 1262
1167#endif 1263#endif
1168 1264
1169void 1265void
1170supported_image_formats (OpenCL::Context self, cl_mem_flags flags, cl_mem_object_type image_type) 1266supported_image_formats (OpenCL::Context self, cl_mem_flags flags, cl_mem_object_type image_type)
1190 1286
1191void 1287void
1192sampler (OpenCL::Context self, cl_bool normalized_coords, cl_addressing_mode addressing_mode, cl_filter_mode filter_mode) 1288sampler (OpenCL::Context self, cl_bool normalized_coords, cl_addressing_mode addressing_mode, cl_filter_mode filter_mode)
1193 PPCODE: 1289 PPCODE:
1194 NEED_SUCCESS_ARG (cl_sampler sampler, CreateSampler, (self, normalized_coords, addressing_mode, filter_mode, &res)); 1290 NEED_SUCCESS_ARG (cl_sampler sampler, CreateSampler, (self, normalized_coords, addressing_mode, filter_mode, &res));
1195 XPUSH_CLOBJ ("OpenCL::Sampler", sampler); 1291 XPUSH_CLOBJ (stash_sampler, sampler);
1196 1292
1197void 1293void
1198program_with_source (OpenCL::Context self, SV *program) 1294program_with_source (OpenCL::Context self, SV *program)
1199 PPCODE: 1295 PPCODE:
1200 STRLEN len; 1296 STRLEN len;
1201 size_t len2; 1297 size_t len2;
1202 const char *ptr = SvPVbyte (program, len); 1298 const char *ptr = SvPVbyte (program, len);
1203 1299
1204 len2 = len; 1300 len2 = len;
1205 NEED_SUCCESS_ARG (cl_program prog, CreateProgramWithSource, (self, 1, &ptr, &len2, &res)); 1301 NEED_SUCCESS_ARG (cl_program prog, CreateProgramWithSource, (self, 1, &ptr, &len2, &res));
1206 XPUSH_CLOBJ ("OpenCL::Program", prog); 1302 XPUSH_CLOBJ (stash_program, prog);
1207 1303
1208#BEGIN:context 1304#BEGIN:context
1209 1305
1210void 1306void
1211reference_count (OpenCL::Context self) 1307reference_count (OpenCL::Context self)
1228 NEED_SUCCESS (GetContextInfo, (self, CL_CONTEXT_DEVICES, size, value, 0)); 1324 NEED_SUCCESS (GetContextInfo, (self, CL_CONTEXT_DEVICES, size, value, 0));
1229 int i, n = size / sizeof (*value); 1325 int i, n = size / sizeof (*value);
1230 EXTEND (SP, n); 1326 EXTEND (SP, n);
1231 for (i = 0; i < n; ++i) 1327 for (i = 0; i < n; ++i)
1232 { 1328 {
1233 PUSH_CLOBJ ("OpenCL::Device", value [i]); 1329 PUSH_CLOBJ (stash_device, value [i]);
1234 } 1330 }
1235 1331
1236void 1332void
1237properties (OpenCL::Context self) 1333properties (OpenCL::Context self)
1238 PPCODE: 1334 PPCODE:
1253DESTROY (OpenCL::Queue self) 1349DESTROY (OpenCL::Queue self)
1254 CODE: 1350 CODE:
1255 clReleaseCommandQueue (self); 1351 clReleaseCommandQueue (self);
1256 1352
1257void 1353void
1258enqueue_read_buffer (OpenCL::Queue self, OpenCL::Buffer mem, cl_bool blocking, size_t offset, size_t len, SV *data, ...) 1354read_buffer (OpenCL::Queue self, OpenCL::Buffer mem, cl_bool blocking, size_t offset, size_t len, SV *data, ...)
1355 ALIAS:
1356 enqueue_read_buffer = 0
1259 PPCODE: 1357 PPCODE:
1260 cl_event ev = 0; 1358 cl_event ev = 0;
1261 EVENT_LIST (6, items - 6); 1359 EVENT_LIST (6);
1262 1360
1263 SvUPGRADE (data, SVt_PV); 1361 SvUPGRADE (data, SVt_PV);
1264 SvGROW (data, len); 1362 SvGROW (data, len);
1265 SvPOK_only (data); 1363 SvPOK_only (data);
1266 SvCUR_set (data, len); 1364 SvCUR_set (data, len);
1267 NEED_SUCCESS (EnqueueReadBuffer, (self, mem, blocking, offset, len, SvPVX (data), event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1365 NEED_SUCCESS (EnqueueReadBuffer, (self, mem, blocking, offset, len, SvPVX (data), event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1268 1366
1269 if (ev) 1367 if (ev)
1270 XPUSH_CLOBJ ("OpenCL::Event", ev); 1368 XPUSH_CLOBJ (stash_event, ev);
1271 1369
1272void 1370void
1273enqueue_write_buffer (OpenCL::Queue self, OpenCL::Buffer mem, cl_bool blocking, size_t offset, SV *data, ...) 1371write_buffer (OpenCL::Queue self, OpenCL::Buffer mem, cl_bool blocking, size_t offset, SV *data, ...)
1372 ALIAS:
1373 enqueue_write_buffer = 0
1274 PPCODE: 1374 PPCODE:
1275 cl_event ev = 0; 1375 cl_event ev = 0;
1276 STRLEN len; 1376 STRLEN len;
1277 char *ptr = SvPVbyte (data, len); 1377 char *ptr = SvPVbyte (data, len);
1278 EVENT_LIST (5, items - 5); 1378 EVENT_LIST (5);
1279 1379
1280 NEED_SUCCESS (EnqueueWriteBuffer, (self, mem, blocking, offset, len, ptr, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1380 NEED_SUCCESS (EnqueueWriteBuffer, (self, mem, blocking, offset, len, ptr, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1281 1381
1282 if (ev) 1382 if (ev)
1283 XPUSH_CLOBJ ("OpenCL::Event", ev); 1383 XPUSH_CLOBJ (stash_event, ev);
1284 1384
1285#if CL_VERSION_1_2 1385#if CL_VERSION_1_2
1286 1386
1287void 1387void
1288enqueue_fill_buffer (OpenCL::Queue self, OpenCL::Buffer mem, SV *data, size_t offset, size_t size, ...) 1388fill_buffer (OpenCL::Queue self, OpenCL::Buffer mem, SV *data, size_t offset, size_t size, ...)
1389 ALIAS:
1390 enqueue_fill_buffer = 0
1289 PPCODE: 1391 PPCODE:
1290 cl_event ev = 0; 1392 cl_event ev = 0;
1291 STRLEN len; 1393 STRLEN len;
1292 char *ptr = SvPVbyte (data, len); 1394 char *ptr = SvPVbyte (data, len);
1293 EVENT_LIST (5, items - 5); 1395 EVENT_LIST (5);
1294 1396
1295 NEED_SUCCESS (EnqueueFillBuffer, (self, mem, ptr, len, offset, size, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1397 NEED_SUCCESS (EnqueueFillBuffer, (self, mem, ptr, len, offset, size, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1296 1398
1297 if (ev) 1399 if (ev)
1298 XPUSH_CLOBJ ("OpenCL::Event", ev); 1400 XPUSH_CLOBJ (stash_event, ev);
1299 1401
1300void 1402void
1301enqueue_fill_image (OpenCL::Queue self, OpenCL::Image img, NV r, NV g, NV b, NV a, size_t x, size_t y, size_t z, size_t width, size_t height, size_t depth, ...) 1403fill_image (OpenCL::Queue self, OpenCL::Image img, NV r, NV g, NV b, NV a, size_t x, size_t y, size_t z, size_t width, size_t height, size_t depth, ...)
1404 ALIAS:
1405 enqueue_fill_image = 0
1302 PPCODE: 1406 PPCODE:
1303 cl_event ev = 0; 1407 cl_event ev = 0;
1304 STRLEN len; 1408 STRLEN len;
1305 const size_t origin [3] = { x, y, z }; 1409 const size_t origin [3] = { x, y, z };
1306 const size_t region [3] = { width, height, depth }; 1410 const size_t region [3] = { width, height, depth };
1307 EVENT_LIST (12, items - 12); 1411 EVENT_LIST (12);
1308 1412
1309 const cl_float c_f [4] = { r, g, b, a }; 1413 const cl_float c_f [4] = { r, g, b, a };
1310 const cl_uint c_u [4] = { r, g, b, a }; 1414 const cl_uint c_u [4] = { r, g, b, a };
1311 const cl_int c_s [4] = { r, g, b, a }; 1415 const cl_int c_s [4] = { r, g, b, a };
1312 const void *c_fus [3] = { &c_f, &c_u, &c_s }; 1416 const void *c_fus [3] = { &c_f, &c_u, &c_s };
1315 NEED_SUCCESS (GetImageInfo, (img, CL_IMAGE_FORMAT, sizeof (format), &format, 0)); 1419 NEED_SUCCESS (GetImageInfo, (img, CL_IMAGE_FORMAT, sizeof (format), &format, 0));
1316 assert (sizeof (fus) == CL_FLOAT + 1 - CL_SNORM_INT8); 1420 assert (sizeof (fus) == CL_FLOAT + 1 - CL_SNORM_INT8);
1317 if (format.image_channel_data_type < CL_SNORM_INT8 || CL_FLOAT < format.image_channel_data_type) 1421 if (format.image_channel_data_type < CL_SNORM_INT8 || CL_FLOAT < format.image_channel_data_type)
1318 croak ("enqueue_fill_image: image has unsupported channel type, only opencl 1.2 channel types supported."); 1422 croak ("enqueue_fill_image: image has unsupported channel type, only opencl 1.2 channel types supported.");
1319 1423
1320 NEED_SUCCESS (EnqueueFillImage, (self, img, c_fus [fus [format.image_channel_data_type]], 1424 NEED_SUCCESS (EnqueueFillImage, (self, img, c_fus [fus [format.image_channel_data_type - CL_SNORM_INT8]],
1321 origin, region, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1425 origin, region, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1322 1426
1323 if (ev) 1427 if (ev)
1324 XPUSH_CLOBJ ("OpenCL::Event", ev); 1428 XPUSH_CLOBJ (stash_event, ev);
1325 1429
1326#endif 1430#endif
1327 1431
1328void 1432void
1329enqueue_copy_buffer (OpenCL::Queue self, OpenCL::Buffer src, OpenCL::Buffer dst, size_t src_offset, size_t dst_offset, size_t len, ...) 1433copy_buffer (OpenCL::Queue self, OpenCL::Buffer src, OpenCL::Buffer dst, size_t src_offset, size_t dst_offset, size_t len, ...)
1434 ALIAS:
1435 enqueue_copy_buffer = 0
1330 PPCODE: 1436 PPCODE:
1331 cl_event ev = 0; 1437 cl_event ev = 0;
1332 EVENT_LIST (6, items - 6); 1438 EVENT_LIST (6);
1333 1439
1334 NEED_SUCCESS (EnqueueCopyBuffer, (self, src, dst, src_offset, dst_offset, len, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1440 NEED_SUCCESS (EnqueueCopyBuffer, (self, src, dst, src_offset, dst_offset, len, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1335 1441
1336 if (ev) 1442 if (ev)
1337 XPUSH_CLOBJ ("OpenCL::Event", ev); 1443 XPUSH_CLOBJ (stash_event, ev);
1338 1444
1339void 1445void
1340enqueue_read_buffer_rect (OpenCL::Queue self, OpenCL::Memory buf, cl_bool blocking, size_t buf_x, size_t buf_y, size_t buf_z, size_t host_x, size_t host_y, size_t host_z, size_t width, size_t height, size_t depth, size_t buf_row_pitch, size_t buf_slice_pitch, size_t host_row_pitch, size_t host_slice_pitch, SV *data, ...) 1446read_buffer_rect (OpenCL::Queue self, OpenCL::Memory buf, cl_bool blocking, size_t buf_x, size_t buf_y, size_t buf_z, size_t host_x, size_t host_y, size_t host_z, size_t width, size_t height, size_t depth, size_t buf_row_pitch, size_t buf_slice_pitch, size_t host_row_pitch, size_t host_slice_pitch, SV *data, ...)
1447 ALIAS:
1448 enqueue_read_buffer_rect = 0
1341 PPCODE: 1449 PPCODE:
1342 cl_event ev = 0; 1450 cl_event ev = 0;
1343 const size_t buf_origin [3] = { buf_x , buf_y , buf_z }; 1451 const size_t buf_origin [3] = { buf_x , buf_y , buf_z };
1344 const size_t host_origin[3] = { host_x, host_y, host_z }; 1452 const size_t host_origin[3] = { host_x, host_y, host_z };
1345 const size_t region[3] = { width, height, depth }; 1453 const size_t region[3] = { width, height, depth };
1346 EVENT_LIST (17, items - 17); 1454 EVENT_LIST (17);
1347 1455
1348 if (!buf_row_pitch) 1456 if (!buf_row_pitch)
1349 buf_row_pitch = region [0]; 1457 buf_row_pitch = region [0];
1350 1458
1351 if (!buf_slice_pitch) 1459 if (!buf_slice_pitch)
1364 SvPOK_only (data); 1472 SvPOK_only (data);
1365 SvCUR_set (data, len); 1473 SvCUR_set (data, len);
1366 NEED_SUCCESS (EnqueueReadBufferRect, (self, buf, blocking, buf_origin, host_origin, region, buf_row_pitch, buf_slice_pitch, host_row_pitch, host_slice_pitch, SvPVX (data), event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1474 NEED_SUCCESS (EnqueueReadBufferRect, (self, buf, blocking, buf_origin, host_origin, region, buf_row_pitch, buf_slice_pitch, host_row_pitch, host_slice_pitch, SvPVX (data), event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1367 1475
1368 if (ev) 1476 if (ev)
1369 XPUSH_CLOBJ ("OpenCL::Event", ev); 1477 XPUSH_CLOBJ (stash_event, ev);
1370 1478
1371void 1479void
1372enqueue_write_buffer_rect (OpenCL::Queue self, OpenCL::Memory buf, cl_bool blocking, size_t buf_x, size_t buf_y, size_t buf_z, size_t host_x, size_t host_y, size_t host_z, size_t width, size_t height, size_t depth, size_t buf_row_pitch, size_t buf_slice_pitch, size_t host_row_pitch, size_t host_slice_pitch, SV *data, ...) 1480write_buffer_rect (OpenCL::Queue self, OpenCL::Memory buf, cl_bool blocking, size_t buf_x, size_t buf_y, size_t buf_z, size_t host_x, size_t host_y, size_t host_z, size_t width, size_t height, size_t depth, size_t buf_row_pitch, size_t buf_slice_pitch, size_t host_row_pitch, size_t host_slice_pitch, SV *data, ...)
1481 ALIAS:
1482 enqueue_write_buffer_rect = 0
1373 PPCODE: 1483 PPCODE:
1374 cl_event ev = 0; 1484 cl_event ev = 0;
1375 const size_t buf_origin [3] = { buf_x , buf_y , buf_z }; 1485 const size_t buf_origin [3] = { buf_x , buf_y , buf_z };
1376 const size_t host_origin[3] = { host_x, host_y, host_z }; 1486 const size_t host_origin[3] = { host_x, host_y, host_z };
1377 const size_t region[3] = { width, height, depth }; 1487 const size_t region[3] = { width, height, depth };
1378 STRLEN len; 1488 STRLEN len;
1379 char *ptr = SvPVbyte (data, len); 1489 char *ptr = SvPVbyte (data, len);
1380 EVENT_LIST (17, items - 17); 1490 EVENT_LIST (17);
1381 1491
1382 if (!buf_row_pitch) 1492 if (!buf_row_pitch)
1383 buf_row_pitch = region [0]; 1493 buf_row_pitch = region [0];
1384 1494
1385 if (!buf_slice_pitch) 1495 if (!buf_slice_pitch)
1397 croak ("clEnqueueWriteImage: data string is shorter than what would be transferred"); 1507 croak ("clEnqueueWriteImage: data string is shorter than what would be transferred");
1398 1508
1399 NEED_SUCCESS (EnqueueWriteBufferRect, (self, buf, blocking, buf_origin, host_origin, region, buf_row_pitch, buf_slice_pitch, host_row_pitch, host_slice_pitch, ptr, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1509 NEED_SUCCESS (EnqueueWriteBufferRect, (self, buf, blocking, buf_origin, host_origin, region, buf_row_pitch, buf_slice_pitch, host_row_pitch, host_slice_pitch, ptr, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1400 1510
1401 if (ev) 1511 if (ev)
1402 XPUSH_CLOBJ ("OpenCL::Event", ev); 1512 XPUSH_CLOBJ (stash_event, ev);
1403 1513
1404void 1514void
1405enqueue_copy_buffer_rect (OpenCL::Queue self, OpenCL::Buffer src, OpenCL::Buffer dst, size_t src_x, size_t src_y, size_t src_z, size_t dst_x, size_t dst_y, size_t dst_z, size_t width, size_t height, size_t depth, size_t src_row_pitch, size_t src_slice_pitch, size_t dst_row_pitch, size_t dst_slice_pitch, ...) 1515copy_buffer_rect (OpenCL::Queue self, OpenCL::Buffer src, OpenCL::Buffer dst, size_t src_x, size_t src_y, size_t src_z, size_t dst_x, size_t dst_y, size_t dst_z, size_t width, size_t height, size_t depth, size_t src_row_pitch, size_t src_slice_pitch, size_t dst_row_pitch, size_t dst_slice_pitch, ...)
1516 ALIAS:
1517 enqueue_copy_buffer_rect = 0
1406 PPCODE: 1518 PPCODE:
1407 cl_event ev = 0; 1519 cl_event ev = 0;
1408 const size_t src_origin[3] = { src_x, src_y, src_z }; 1520 const size_t src_origin[3] = { src_x, src_y, src_z };
1409 const size_t dst_origin[3] = { dst_x, dst_y, dst_z }; 1521 const size_t dst_origin[3] = { dst_x, dst_y, dst_z };
1410 const size_t region[3] = { width, height, depth }; 1522 const size_t region[3] = { width, height, depth };
1411 EVENT_LIST (16, items - 16); 1523 EVENT_LIST (16);
1412 1524
1413 NEED_SUCCESS (EnqueueCopyBufferRect, (self, src, dst, src_origin, dst_origin, region, src_row_pitch, src_slice_pitch, dst_row_pitch, dst_slice_pitch, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1525 NEED_SUCCESS (EnqueueCopyBufferRect, (self, src, dst, src_origin, dst_origin, region, src_row_pitch, src_slice_pitch, dst_row_pitch, dst_slice_pitch, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1414 1526
1415 if (ev) 1527 if (ev)
1416 XPUSH_CLOBJ ("OpenCL::Event", ev); 1528 XPUSH_CLOBJ (stash_event, ev);
1417 1529
1418void 1530void
1419enqueue_read_image (OpenCL::Queue self, OpenCL::Image src, cl_bool blocking, size_t src_x, size_t src_y, size_t src_z, size_t width, size_t height, size_t depth, size_t row_pitch, size_t slice_pitch, SV *data, ...) 1531read_image (OpenCL::Queue self, OpenCL::Image src, cl_bool blocking, size_t src_x, size_t src_y, size_t src_z, size_t width, size_t height, size_t depth, size_t row_pitch, size_t slice_pitch, SV *data, ...)
1532 ALIAS:
1533 enqueue_read_image = 0
1420 PPCODE: 1534 PPCODE:
1421 cl_event ev = 0; 1535 cl_event ev = 0;
1422 const size_t src_origin[3] = { src_x, src_y, src_z }; 1536 const size_t src_origin[3] = { src_x, src_y, src_z };
1423 const size_t region[3] = { width, height, depth }; 1537 const size_t region[3] = { width, height, depth };
1424 EVENT_LIST (12, items - 12); 1538 EVENT_LIST (12);
1425 1539
1426 if (!row_pitch) 1540 if (!row_pitch)
1427 row_pitch = img_row_pitch (src); 1541 row_pitch = img_row_pitch (src);
1428 1542
1429 if (depth > 1 && !slice_pitch) 1543 if (depth > 1 && !slice_pitch)
1436 SvPOK_only (data); 1550 SvPOK_only (data);
1437 SvCUR_set (data, len); 1551 SvCUR_set (data, len);
1438 NEED_SUCCESS (EnqueueReadImage, (self, src, blocking, src_origin, region, row_pitch, slice_pitch, SvPVX (data), event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1552 NEED_SUCCESS (EnqueueReadImage, (self, src, blocking, src_origin, region, row_pitch, slice_pitch, SvPVX (data), event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1439 1553
1440 if (ev) 1554 if (ev)
1441 XPUSH_CLOBJ ("OpenCL::Event", ev); 1555 XPUSH_CLOBJ (stash_event, ev);
1442 1556
1443void 1557void
1444enqueue_write_image (OpenCL::Queue self, OpenCL::Image dst, cl_bool blocking, size_t dst_x, size_t dst_y, size_t dst_z, size_t width, size_t height, size_t depth, size_t row_pitch, size_t slice_pitch, SV *data, ...) 1558write_image (OpenCL::Queue self, OpenCL::Image dst, cl_bool blocking, size_t dst_x, size_t dst_y, size_t dst_z, size_t width, size_t height, size_t depth, size_t row_pitch, size_t slice_pitch, SV *data, ...)
1559 ALIAS:
1560 enqueue_write_image = 0
1445 PPCODE: 1561 PPCODE:
1446 cl_event ev = 0; 1562 cl_event ev = 0;
1447 const size_t dst_origin[3] = { dst_x, dst_y, dst_z }; 1563 const size_t dst_origin[3] = { dst_x, dst_y, dst_z };
1448 const size_t region[3] = { width, height, depth }; 1564 const size_t region[3] = { width, height, depth };
1449 STRLEN len; 1565 STRLEN len;
1450 char *ptr = SvPVbyte (data, len); 1566 char *ptr = SvPVbyte (data, len);
1451 EVENT_LIST (12, items - 12); 1567 EVENT_LIST (12);
1452 1568
1453 if (!row_pitch) 1569 if (!row_pitch)
1454 row_pitch = img_row_pitch (dst); 1570 row_pitch = img_row_pitch (dst);
1455 1571
1456 if (depth > 1 && !slice_pitch) 1572 if (depth > 1 && !slice_pitch)
1462 croak ("clEnqueueWriteImage: data string is shorter than what would be transferred"); 1578 croak ("clEnqueueWriteImage: data string is shorter than what would be transferred");
1463 1579
1464 NEED_SUCCESS (EnqueueWriteImage, (self, dst, blocking, dst_origin, region, row_pitch, slice_pitch, ptr, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1580 NEED_SUCCESS (EnqueueWriteImage, (self, dst, blocking, dst_origin, region, row_pitch, slice_pitch, ptr, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1465 1581
1466 if (ev) 1582 if (ev)
1467 XPUSH_CLOBJ ("OpenCL::Event", ev); 1583 XPUSH_CLOBJ (stash_event, ev);
1468 1584
1469void 1585void
1470enqueue_copy_image (OpenCL::Queue self, OpenCL::Image src, OpenCL::Image dst, size_t src_x, size_t src_y, size_t src_z, size_t dst_x, size_t dst_y, size_t dst_z, size_t width, size_t height, size_t depth, ...) 1586copy_image (OpenCL::Queue self, OpenCL::Image src, OpenCL::Image dst, size_t src_x, size_t src_y, size_t src_z, size_t dst_x, size_t dst_y, size_t dst_z, size_t width, size_t height, size_t depth, ...)
1587 ALIAS:
1588 enqueue_copy_image = 0
1471 PPCODE: 1589 PPCODE:
1472 cl_event ev = 0; 1590 cl_event ev = 0;
1473 const size_t src_origin[3] = { src_x, src_y, src_z }; 1591 const size_t src_origin[3] = { src_x, src_y, src_z };
1474 const size_t dst_origin[3] = { dst_x, dst_y, dst_z }; 1592 const size_t dst_origin[3] = { dst_x, dst_y, dst_z };
1475 const size_t region[3] = { width, height, depth }; 1593 const size_t region[3] = { width, height, depth };
1476 EVENT_LIST (12, items - 12); 1594 EVENT_LIST (12);
1477 1595
1478 NEED_SUCCESS (EnqueueCopyImage, (self, src, dst, src_origin, dst_origin, region, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1596 NEED_SUCCESS (EnqueueCopyImage, (self, src, dst, src_origin, dst_origin, region, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1479 1597
1480 if (ev) 1598 if (ev)
1481 XPUSH_CLOBJ ("OpenCL::Event", ev); 1599 XPUSH_CLOBJ (stash_event, ev);
1482 1600
1483void 1601void
1484enqueue_copy_image_to_buffer (OpenCL::Queue self, OpenCL::Image src, OpenCL::Buffer dst, size_t src_x, size_t src_y, size_t src_z, size_t width, size_t height, size_t depth, size_t dst_offset, ...) 1602copy_image_to_buffer (OpenCL::Queue self, OpenCL::Image src, OpenCL::Buffer dst, size_t src_x, size_t src_y, size_t src_z, size_t width, size_t height, size_t depth, size_t dst_offset, ...)
1603 ALIAS:
1604 enqueue_copy_image_to_buffer = 0
1485 PPCODE: 1605 PPCODE:
1486 cl_event ev = 0; 1606 cl_event ev = 0;
1487 const size_t src_origin[3] = { src_x, src_y, src_z }; 1607 const size_t src_origin[3] = { src_x, src_y, src_z };
1608 const size_t region [3] = { width, height, depth };
1609 EVENT_LIST (10);
1610
1611 NEED_SUCCESS (EnqueueCopyImageToBuffer, (self, src, dst, src_origin, region, dst_offset, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1612
1613 if (ev)
1614 XPUSH_CLOBJ (stash_event, ev);
1615
1616void
1617copy_buffer_to_image (OpenCL::Queue self, OpenCL::Buffer src, OpenCL::Image dst, size_t src_offset, size_t dst_x, size_t dst_y, size_t dst_z, size_t width, size_t height, size_t depth, ...)
1618 ALIAS:
1619 enqueue_copy_buffer_to_image = 0
1620 PPCODE:
1621 cl_event ev = 0;
1622 const size_t dst_origin[3] = { dst_x, dst_y, dst_z };
1623 const size_t region [3] = { width, height, depth };
1624 EVENT_LIST (10);
1625
1626 NEED_SUCCESS (EnqueueCopyBufferToImage, (self, src, dst, src_offset, dst_origin, region, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1627
1628 if (ev)
1629 XPUSH_CLOBJ (stash_event, ev);
1630
1631void
1632map_buffer (OpenCL::Queue self, OpenCL::Buffer buf, cl_bool blocking = 1, cl_map_flags map_flags = CL_MAP_READ | CL_MAP_WRITE, size_t offset = 0, size_t cb = 0, ...)
1633 ALIAS:
1634 enqueue_map_buffer = 0
1635 PPCODE:
1636 cl_event ev;
1637 EVENT_LIST (6);
1638
1639 NEED_SUCCESS_ARG (void *ptr, EnqueueMapBuffer, (self, buf, blocking, map_flags, offset, cb, event_list_count, event_list_ptr, &ev, &res));
1640 XPUSHs (mapped_new (stash_mappedbuffer, self, buf, map_flags, ptr, cb, ev, 0, 0));
1641
1642void
1643map_image (OpenCL::Queue self, OpenCL::Image img, cl_bool blocking = 1, cl_map_flags map_flags = CL_MAP_READ | CL_MAP_WRITE, size_t x = 0, size_t y = 0, size_t z = 0, size_t width = 0, size_t height = 0, size_t depth = 0, ...)
1644 ALIAS:
1645 enqueue_map_image = 0
1646 PPCODE:
1647 cl_event ev;
1648 const size_t origin[3] = { x, y, z };
1488 const size_t region[3] = { width, height, depth }; 1649 const size_t region[3] = { width, height, depth };
1650 size_t row_pitch, slice_pitch;
1489 EVENT_LIST (10, items - 10); 1651 EVENT_LIST (10);
1490 1652
1653 NEED_SUCCESS_ARG (void *ptr, EnqueueMapImage, (self, img, blocking, map_flags, origin, region, &row_pitch, &slice_pitch, event_list_count, event_list_ptr, &ev, &res));
1654
1655 size_t cb = slice_pitch ? slice_pitch * region [2]
1656 : row_pitch ? row_pitch * region [1]
1657 : region [0];
1658
1659 XPUSHs (mapped_new (stash_mappedimage, self, img, map_flags, ptr, cb, ev, row_pitch, slice_pitch));
1660
1661void
1662unmap (OpenCL::Queue self, OpenCL::Mapped mapped, ...)
1663 PPCODE:
1664 mapped_unmap (ST (1), mapped, self, &ST (2), items - 2);
1665 if (GIMME_V != G_VOID)
1666 {
1667 clRetainEvent (mapped->event);
1668 XPUSH_CLOBJ (stash_event, mapped->event);
1669 }
1670
1671void
1672task (OpenCL::Queue self, OpenCL::Kernel kernel, ...)
1673 ALIAS:
1674 enqueue_task = 0
1675 PPCODE:
1676 cl_event ev = 0;
1677 EVENT_LIST (2);
1678
1491 NEED_SUCCESS (EnqueueCopyImageToBuffer, (self, src, dst, src_origin, region, dst_offset, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1679 NEED_SUCCESS (EnqueueTask, (self, kernel, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1492 1680
1493 if (ev) 1681 if (ev)
1494 XPUSH_CLOBJ ("OpenCL::Event", ev); 1682 XPUSH_CLOBJ (stash_event, ev);
1495 1683
1496void 1684void
1497enqueue_copy_buffer_to_image (OpenCL::Queue self, OpenCL::Buffer src, OpenCL::Image dst, size_t src_offset, size_t dst_x, size_t dst_y, size_t dst_z, size_t width, size_t height, size_t depth, ...)
1498 PPCODE:
1499 cl_event ev = 0;
1500 const size_t dst_origin[3] = { dst_x, dst_y, dst_z };
1501 const size_t region[3] = { width, height, depth };
1502 EVENT_LIST (10, items - 10);
1503
1504 NEED_SUCCESS (EnqueueCopyBufferToImage, (self, src, dst, src_offset, dst_origin, region, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1505
1506 if (ev)
1507 XPUSH_CLOBJ ("OpenCL::Event", ev);
1508
1509void
1510enqueue_task (OpenCL::Queue self, OpenCL::Kernel kernel, ...)
1511 PPCODE:
1512 cl_event ev = 0;
1513 EVENT_LIST (2, items - 2);
1514
1515 NEED_SUCCESS (EnqueueTask, (self, kernel, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1516
1517 if (ev)
1518 XPUSH_CLOBJ ("OpenCL::Event", ev);
1519
1520void
1521enqueue_nd_range_kernel (OpenCL::Queue self, OpenCL::Kernel kernel, SV *global_work_offset, SV *global_work_size, SV *local_work_size = &PL_sv_undef, ...) 1685nd_range_kernel (OpenCL::Queue self, OpenCL::Kernel kernel, SV *global_work_offset, SV *global_work_size, SV *local_work_size = &PL_sv_undef, ...)
1686 ALIAS:
1687 enqueue_nd_range_kernel = 0
1522 PPCODE: 1688 PPCODE:
1523 cl_event ev = 0; 1689 cl_event ev = 0;
1524 size_t *gwo = 0, *gws, *lws = 0; 1690 size_t *gwo = 0, *gws, *lws = 0;
1525 int gws_len; 1691 int gws_len;
1526 size_t *lists; 1692 size_t *lists;
1527 int i; 1693 int i;
1528 EVENT_LIST (5, items - 5); 1694 EVENT_LIST (5);
1529 1695
1530 if (!SvROK (global_work_size) || SvTYPE (SvRV (global_work_size)) != SVt_PVAV) 1696 if (!SvROK (global_work_size) || SvTYPE (SvRV (global_work_size)) != SVt_PVAV)
1531 croak ("clEnqueueNDRangeKernel: global_work_size must be an array reference"); 1697 croak ("clEnqueueNDRangeKernel: global_work_size must be an array reference");
1532 1698
1533 gws_len = AvFILLp (SvRV (global_work_size)) + 1; 1699 gws_len = AvFILLp (SvRV (global_work_size)) + 1;
1534 1700
1535 lists = tmpbuf (sizeof (size_t) * 3 * gws_len); 1701 lists = tmpbuf (sizeof (size_t) * 3 * gws_len);
1536 1702
1537 gws = lists + gws_len * 0; 1703 gws = lists + gws_len * 0;
1538 for (i = 0; i < gws_len; ++i) 1704 for (i = 0; i < gws_len; ++i)
1705 {
1539 gws [i] = SvIV (AvARRAY (SvRV (global_work_size))[i]); 1706 gws [i] = SvIV (AvARRAY (SvRV (global_work_size))[i]);
1707 // at least nvidia crashes for 0-sized work group sizes, work around
1708 if (!gws [i])
1709 croak ("clEnqueueNDRangeKernel: global_work_size[%d] is zero, must be non-zero", i);
1710 }
1540 1711
1541 if (SvOK (global_work_offset)) 1712 if (SvOK (global_work_offset))
1542 { 1713 {
1543 if (!SvROK (global_work_offset) || SvTYPE (SvRV (global_work_offset)) != SVt_PVAV) 1714 if (!SvROK (global_work_offset) || SvTYPE (SvRV (global_work_offset)) != SVt_PVAV)
1544 croak ("clEnqueueNDRangeKernel: global_work_offset must be undef or an array reference"); 1715 croak ("clEnqueueNDRangeKernel: global_work_offset must be undef or an array reference");
1552 } 1723 }
1553 1724
1554 if (SvOK (local_work_size)) 1725 if (SvOK (local_work_size))
1555 { 1726 {
1556 if ((SvOK (local_work_size) && !SvROK (local_work_size)) || SvTYPE (SvRV (local_work_size)) != SVt_PVAV) 1727 if ((SvOK (local_work_size) && !SvROK (local_work_size)) || SvTYPE (SvRV (local_work_size)) != SVt_PVAV)
1557 croak ("clEnqueueNDRangeKernel: global_work_size must be undef or an array reference"); 1728 croak ("clEnqueueNDRangeKernel: local_work_size must be undef or an array reference");
1558 1729
1559 if (AvFILLp (SvRV (local_work_size)) + 1 != gws_len) 1730 if (AvFILLp (SvRV (local_work_size)) + 1 != gws_len)
1560 croak ("clEnqueueNDRangeKernel: local_work_local must be undef or an array of same size as global_work_size"); 1731 croak ("clEnqueueNDRangeKernel: local_work_local must be undef or an array of same size as global_work_size");
1561 1732
1562 lws = lists + gws_len * 2; 1733 lws = lists + gws_len * 2;
1563 for (i = 0; i < gws_len; ++i) 1734 for (i = 0; i < gws_len; ++i)
1735 {
1564 lws [i] = SvIV (AvARRAY (SvRV (local_work_size))[i]); 1736 lws [i] = SvIV (AvARRAY (SvRV (local_work_size))[i]);
1737 // at least nvidia crashes for 0-sized work group sizes, work around
1738 if (!lws [i])
1739 croak ("clEnqueueNDRangeKernel: local_work_size[%d] is zero, must be non-zero", i);
1740 }
1565 } 1741 }
1566 1742
1567 NEED_SUCCESS (EnqueueNDRangeKernel, (self, kernel, gws_len, gwo, gws, lws, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1743 NEED_SUCCESS (EnqueueNDRangeKernel, (self, kernel, gws_len, gwo, gws, lws, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1568 1744
1569 if (ev) 1745 if (ev)
1570 XPUSH_CLOBJ ("OpenCL::Event", ev); 1746 XPUSH_CLOBJ (stash_event, ev);
1571 1747
1572#if cl_apple_gl_sharing || cl_khr_gl_sharing 1748#if cl_apple_gl_sharing || cl_khr_gl_sharing
1573 1749
1574void 1750void
1575enqueue_acquire_gl_objects (OpenCL::Queue self, SV *objects, ...) 1751acquire_gl_objects (OpenCL::Queue self, SV *objects, ...)
1576 ALIAS: 1752 ALIAS:
1753 release_gl_objects = 1
1754 enqueue_acquire_gl_objects = 0
1577 enqueue_release_gl_objects = 1 1755 enqueue_release_gl_objects = 1
1578 PPCODE: 1756 PPCODE:
1579 if (!SvROK (objects) || SvTYPE (SvRV (objects)) != SVt_PVAV) 1757 if (!SvROK (objects) || SvTYPE (SvRV (objects)) != SVt_PVAV)
1580 croak ("OpenCL::Queue::enqueue_acquire/release_gl_objects argument 'objects' must be an arrayref with memory objects, in call"); 1758 croak ("OpenCL::Queue::enqueue_acquire/release_gl_objects argument 'objects' must be an arrayref with memory objects, in call");
1581 1759
1582 cl_event ev = 0; 1760 cl_event ev = 0;
1583 EVENT_LIST (2, items - 2); 1761 EVENT_LIST (2);
1584 AV *av = (AV *)SvRV (objects); 1762 AV *av = (AV *)SvRV (objects);
1585 cl_uint num_objects = av_len (av) + 1; 1763 cl_uint num_objects = av_len (av) + 1;
1586 cl_mem *object_list = tmpbuf (sizeof (cl_mem) * num_objects); 1764 cl_mem *object_list = tmpbuf (sizeof (cl_mem) * num_objects);
1587 int i; 1765 int i;
1588 1766
1593 NEED_SUCCESS (EnqueueReleaseGLObjects, (self, num_objects, object_list, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1771 NEED_SUCCESS (EnqueueReleaseGLObjects, (self, num_objects, object_list, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1594 else 1772 else
1595 NEED_SUCCESS (EnqueueAcquireGLObjects, (self, num_objects, object_list, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1773 NEED_SUCCESS (EnqueueAcquireGLObjects, (self, num_objects, object_list, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1596 1774
1597 if (ev) 1775 if (ev)
1598 XPUSH_CLOBJ ("OpenCL::Event", ev); 1776 XPUSH_CLOBJ (stash_event, ev);
1599 1777
1600#endif 1778#endif
1601 1779
1602void 1780void
1603enqueue_wait_for_events (OpenCL::Queue self, ...) 1781wait_for_events (OpenCL::Queue self, ...)
1782 ALIAS:
1783 enqueue_wait_for_events = 0
1604 CODE: 1784 CODE:
1605 EVENT_LIST (1, items - 1); 1785 EVENT_LIST (1);
1606#if PREFER_1_1 1786#if PREFER_1_1
1607 NEED_SUCCESS (EnqueueWaitForEvents, (self, event_list_count, event_list_ptr)); 1787 NEED_SUCCESS (EnqueueWaitForEvents, (self, event_list_count, event_list_ptr));
1608#else 1788#else
1609 NEED_SUCCESS (EnqueueBarrierWithWaitList, (self, event_list_count, event_list_ptr, 0)); 1789 NEED_SUCCESS (EnqueueBarrierWithWaitList, (self, event_list_count, event_list_ptr, 0));
1610#endif 1790#endif
1611 1791
1612void 1792void
1613enqueue_marker (OpenCL::Queue self, ...) 1793marker (OpenCL::Queue self, ...)
1794 ALIAS:
1795 enqueue_marker = 0
1614 PPCODE: 1796 PPCODE:
1615 cl_event ev = 0; 1797 cl_event ev = 0;
1616 EVENT_LIST (1, items - 1); 1798 EVENT_LIST (1);
1617#if PREFER_1_1 1799#if PREFER_1_1
1618 if (!event_list_count) 1800 if (!event_list_count)
1619 NEED_SUCCESS (EnqueueMarker, (self, GIMME_V != G_VOID ? &ev : 0)); 1801 NEED_SUCCESS (EnqueueMarker, (self, GIMME_V != G_VOID ? &ev : 0));
1620 else 1802 else
1621#if CL_VERSION_1_2 1803#if CL_VERSION_1_2
1628#endif 1810#endif
1629#else 1811#else
1630 NEED_SUCCESS (EnqueueMarkerWithWaitList, (self, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1812 NEED_SUCCESS (EnqueueMarkerWithWaitList, (self, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1631#endif 1813#endif
1632 if (ev) 1814 if (ev)
1633 XPUSH_CLOBJ ("OpenCL::Event", ev); 1815 XPUSH_CLOBJ (stash_event, ev);
1634 1816
1635void 1817void
1636enqueue_barrier (OpenCL::Queue self, ...) 1818barrier (OpenCL::Queue self, ...)
1819 ALIAS:
1820 enqueue_barrier = 0
1637 PPCODE: 1821 PPCODE:
1638 cl_event ev = 0; 1822 cl_event ev = 0;
1639 EVENT_LIST (1, items - 1); 1823 EVENT_LIST (1);
1640#if PREFER_1_1 1824#if PREFER_1_1
1641 if (!event_list_count && GIMME_V == G_VOID) 1825 if (!event_list_count && GIMME_V == G_VOID)
1642 NEED_SUCCESS (EnqueueBarrier, (self)); 1826 NEED_SUCCESS (EnqueueBarrier, (self));
1643 else 1827 else
1644#if CL_VERSION_1_2 1828#if CL_VERSION_1_2
1654#endif 1838#endif
1655#else 1839#else
1656 NEED_SUCCESS (EnqueueBarrierWithWaitList, (self, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1840 NEED_SUCCESS (EnqueueBarrierWithWaitList, (self, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1657#endif 1841#endif
1658 if (ev) 1842 if (ev)
1659 XPUSH_CLOBJ ("OpenCL::Event", ev); 1843 XPUSH_CLOBJ (stash_event, ev);
1660 1844
1661void 1845void
1662flush (OpenCL::Queue self) 1846flush (OpenCL::Queue self)
1663 CODE: 1847 CODE:
1664 NEED_SUCCESS (Flush, (self)); 1848 NEED_SUCCESS (Flush, (self));
1682 NEED_SUCCESS (GetCommandQueueInfo, (self, CL_QUEUE_CONTEXT, sizeof (value), value, 0)); 1866 NEED_SUCCESS (GetCommandQueueInfo, (self, CL_QUEUE_CONTEXT, sizeof (value), value, 0));
1683 EXTEND (SP, 1); 1867 EXTEND (SP, 1);
1684 const int i = 0; 1868 const int i = 0;
1685 { 1869 {
1686 NEED_SUCCESS (RetainContext, (value [i])); 1870 NEED_SUCCESS (RetainContext, (value [i]));
1687 PUSH_CLOBJ ("OpenCL::Context", value [i]); 1871 PUSH_CLOBJ (stash_context, value [i]);
1688 } 1872 }
1689 1873
1690void 1874void
1691device (OpenCL::Queue self) 1875device (OpenCL::Queue self)
1692 PPCODE: 1876 PPCODE:
1693 cl_device_id value [1]; 1877 cl_device_id value [1];
1694 NEED_SUCCESS (GetCommandQueueInfo, (self, CL_QUEUE_DEVICE, sizeof (value), value, 0)); 1878 NEED_SUCCESS (GetCommandQueueInfo, (self, CL_QUEUE_DEVICE, sizeof (value), value, 0));
1695 EXTEND (SP, 1); 1879 EXTEND (SP, 1);
1696 const int i = 0; 1880 const int i = 0;
1697 { 1881 {
1698 PUSH_CLOBJ ("OpenCL::Device", value [i]); 1882 PUSH_CLOBJ (stash_device, value [i]);
1699 } 1883 }
1700 1884
1701void 1885void
1702reference_count (OpenCL::Queue self) 1886reference_count (OpenCL::Queue self)
1703 PPCODE: 1887 PPCODE:
1712 PPCODE: 1896 PPCODE:
1713 cl_command_queue_properties value [1]; 1897 cl_command_queue_properties value [1];
1714 NEED_SUCCESS (GetCommandQueueInfo, (self, CL_QUEUE_PROPERTIES, sizeof (value), value, 0)); 1898 NEED_SUCCESS (GetCommandQueueInfo, (self, CL_QUEUE_PROPERTIES, sizeof (value), value, 0));
1715 EXTEND (SP, 1); 1899 EXTEND (SP, 1);
1716 const int i = 0; 1900 const int i = 0;
1717 PUSHs (sv_2mortal (newSViv (value [i]))); 1901 PUSHs (sv_2mortal (newSVuv (value [i])));
1718 1902
1719#END:command_queue 1903#END:command_queue
1720 1904
1721MODULE = OpenCL PACKAGE = OpenCL::Memory 1905MODULE = OpenCL PACKAGE = OpenCL::Memory
1722 1906
1737 PPCODE: 1921 PPCODE:
1738 cl_mem_object_type value [1]; 1922 cl_mem_object_type value [1];
1739 NEED_SUCCESS (GetMemObjectInfo, (self, CL_MEM_TYPE, sizeof (value), value, 0)); 1923 NEED_SUCCESS (GetMemObjectInfo, (self, CL_MEM_TYPE, sizeof (value), value, 0));
1740 EXTEND (SP, 1); 1924 EXTEND (SP, 1);
1741 const int i = 0; 1925 const int i = 0;
1742 PUSHs (sv_2mortal (newSViv (value [i]))); 1926 PUSHs (sv_2mortal (newSVuv (value [i])));
1743 1927
1744void 1928void
1745flags (OpenCL::Memory self) 1929flags (OpenCL::Memory self)
1746 PPCODE: 1930 PPCODE:
1747 cl_mem_flags value [1]; 1931 cl_mem_flags value [1];
1748 NEED_SUCCESS (GetMemObjectInfo, (self, CL_MEM_FLAGS, sizeof (value), value, 0)); 1932 NEED_SUCCESS (GetMemObjectInfo, (self, CL_MEM_FLAGS, sizeof (value), value, 0));
1749 EXTEND (SP, 1); 1933 EXTEND (SP, 1);
1750 const int i = 0; 1934 const int i = 0;
1751 PUSHs (sv_2mortal (newSViv (value [i]))); 1935 PUSHs (sv_2mortal (newSVuv (value [i])));
1752 1936
1753void 1937void
1754size (OpenCL::Memory self) 1938size (OpenCL::Memory self)
1755 ALIAS: 1939 ALIAS:
1756 size = CL_MEM_SIZE 1940 size = CL_MEM_SIZE
1790 NEED_SUCCESS (GetMemObjectInfo, (self, CL_MEM_CONTEXT, sizeof (value), value, 0)); 1974 NEED_SUCCESS (GetMemObjectInfo, (self, CL_MEM_CONTEXT, sizeof (value), value, 0));
1791 EXTEND (SP, 1); 1975 EXTEND (SP, 1);
1792 const int i = 0; 1976 const int i = 0;
1793 { 1977 {
1794 NEED_SUCCESS (RetainContext, (value [i])); 1978 NEED_SUCCESS (RetainContext, (value [i]));
1795 PUSH_CLOBJ ("OpenCL::Context", value [i]); 1979 PUSH_CLOBJ (stash_context, value [i]);
1796 } 1980 }
1797 1981
1798void 1982void
1799associated_memobject (OpenCL::Memory self) 1983associated_memobject (OpenCL::Memory self)
1800 PPCODE: 1984 PPCODE:
1802 NEED_SUCCESS (GetMemObjectInfo, (self, CL_MEM_ASSOCIATED_MEMOBJECT, sizeof (value), value, 0)); 1986 NEED_SUCCESS (GetMemObjectInfo, (self, CL_MEM_ASSOCIATED_MEMOBJECT, sizeof (value), value, 0));
1803 EXTEND (SP, 1); 1987 EXTEND (SP, 1);
1804 const int i = 0; 1988 const int i = 0;
1805 { 1989 {
1806 NEED_SUCCESS (RetainMemObject, (value [i])); 1990 NEED_SUCCESS (RetainMemObject, (value [i]));
1807 PUSH_CLOBJ ("OpenCL::Memory", value [i]); 1991 PUSH_CLOBJ (stash_memory, value [i]);
1808 } 1992 }
1809 1993
1810#END:mem 1994#END:mem
1811 1995
1812#if cl_apple_gl_sharing || cl_khr_gl_sharing 1996#if cl_apple_gl_sharing || cl_khr_gl_sharing
1832 croak ("clCreateSubBuffer: cannot use/copy/alloc host ptr, doesn't make sense, check your flags!"); 2016 croak ("clCreateSubBuffer: cannot use/copy/alloc host ptr, doesn't make sense, check your flags!");
1833 2017
1834 cl_buffer_region crdata = { origin, size }; 2018 cl_buffer_region crdata = { origin, size };
1835 2019
1836 NEED_SUCCESS_ARG (cl_mem mem, CreateSubBuffer, (self, flags, CL_BUFFER_CREATE_TYPE_REGION, &crdata, &res)); 2020 NEED_SUCCESS_ARG (cl_mem mem, CreateSubBuffer, (self, flags, CL_BUFFER_CREATE_TYPE_REGION, &crdata, &res));
1837 XPUSH_CLOBJ ("OpenCL::Buffer", mem); 2021 XPUSH_CLOBJ (stash_buffer, mem);
1838 2022
1839MODULE = OpenCL PACKAGE = OpenCL::Image 2023MODULE = OpenCL PACKAGE = OpenCL::Image
1840 2024
1841void 2025void
1842image_info (OpenCL::Image self, cl_image_info name) 2026image_info (OpenCL::Image self, cl_image_info name)
1928 NEED_SUCCESS (GetSamplerInfo, (self, CL_SAMPLER_CONTEXT, sizeof (value), value, 0)); 2112 NEED_SUCCESS (GetSamplerInfo, (self, CL_SAMPLER_CONTEXT, sizeof (value), value, 0));
1929 EXTEND (SP, 1); 2113 EXTEND (SP, 1);
1930 const int i = 0; 2114 const int i = 0;
1931 { 2115 {
1932 NEED_SUCCESS (RetainContext, (value [i])); 2116 NEED_SUCCESS (RetainContext, (value [i]));
1933 PUSH_CLOBJ ("OpenCL::Context", value [i]); 2117 PUSH_CLOBJ (stash_context, value [i]);
1934 } 2118 }
1935 2119
1936void 2120void
1937normalized_coords (OpenCL::Sampler self) 2121normalized_coords (OpenCL::Sampler self)
1938 PPCODE: 2122 PPCODE:
1939 cl_addressing_mode value [1]; 2123 cl_addressing_mode value [1];
1940 NEED_SUCCESS (GetSamplerInfo, (self, CL_SAMPLER_NORMALIZED_COORDS, sizeof (value), value, 0)); 2124 NEED_SUCCESS (GetSamplerInfo, (self, CL_SAMPLER_NORMALIZED_COORDS, sizeof (value), value, 0));
1941 EXTEND (SP, 1); 2125 EXTEND (SP, 1);
1942 const int i = 0; 2126 const int i = 0;
1943 PUSHs (sv_2mortal (newSViv (value [i]))); 2127 PUSHs (sv_2mortal (newSVuv (value [i])));
1944 2128
1945void 2129void
1946addressing_mode (OpenCL::Sampler self) 2130addressing_mode (OpenCL::Sampler self)
1947 PPCODE: 2131 PPCODE:
1948 cl_filter_mode value [1]; 2132 cl_filter_mode value [1];
1949 NEED_SUCCESS (GetSamplerInfo, (self, CL_SAMPLER_ADDRESSING_MODE, sizeof (value), value, 0)); 2133 NEED_SUCCESS (GetSamplerInfo, (self, CL_SAMPLER_ADDRESSING_MODE, sizeof (value), value, 0));
1950 EXTEND (SP, 1); 2134 EXTEND (SP, 1);
1951 const int i = 0; 2135 const int i = 0;
1952 PUSHs (sv_2mortal (newSViv (value [i]))); 2136 PUSHs (sv_2mortal (newSVuv (value [i])));
1953 2137
1954void 2138void
1955filter_mode (OpenCL::Sampler self) 2139filter_mode (OpenCL::Sampler self)
1956 PPCODE: 2140 PPCODE:
1957 cl_bool value [1]; 2141 cl_bool value [1];
2049 2233
2050void 2234void
2051kernel (OpenCL::Program program, SV *function) 2235kernel (OpenCL::Program program, SV *function)
2052 PPCODE: 2236 PPCODE:
2053 NEED_SUCCESS_ARG (cl_kernel kernel, CreateKernel, (program, SvPVbyte_nolen (function), &res)); 2237 NEED_SUCCESS_ARG (cl_kernel kernel, CreateKernel, (program, SvPVbyte_nolen (function), &res));
2054 XPUSH_CLOBJ ("OpenCL::Kernel", kernel); 2238 XPUSH_CLOBJ (stash_kernel, kernel);
2055 2239
2056void 2240void
2057kernels_in_program (OpenCL::Program program) 2241kernels_in_program (OpenCL::Program program)
2058 PPCODE: 2242 PPCODE:
2059 cl_uint num_kernels; 2243 cl_uint num_kernels;
2062 NEED_SUCCESS (CreateKernelsInProgram, (program, num_kernels, kernels, 0)); 2246 NEED_SUCCESS (CreateKernelsInProgram, (program, num_kernels, kernels, 0));
2063 2247
2064 int i; 2248 int i;
2065 EXTEND (SP, num_kernels); 2249 EXTEND (SP, num_kernels);
2066 for (i = 0; i < num_kernels; ++i) 2250 for (i = 0; i < num_kernels; ++i)
2067 PUSH_CLOBJ ("OpenCL::Kernel", kernels [i]); 2251 PUSH_CLOBJ (stash_kernel, kernels [i]);
2068 2252
2069void 2253void
2070info (OpenCL::Program self, cl_program_info name) 2254info (OpenCL::Program self, cl_program_info name)
2071 PPCODE: 2255 PPCODE:
2072 INFO (Program) 2256 INFO (Program)
2120 NEED_SUCCESS (GetProgramInfo, (self, CL_PROGRAM_CONTEXT, sizeof (value), value, 0)); 2304 NEED_SUCCESS (GetProgramInfo, (self, CL_PROGRAM_CONTEXT, sizeof (value), value, 0));
2121 EXTEND (SP, 1); 2305 EXTEND (SP, 1);
2122 const int i = 0; 2306 const int i = 0;
2123 { 2307 {
2124 NEED_SUCCESS (RetainContext, (value [i])); 2308 NEED_SUCCESS (RetainContext, (value [i]));
2125 PUSH_CLOBJ ("OpenCL::Context", value [i]); 2309 PUSH_CLOBJ (stash_context, value [i]);
2126 } 2310 }
2127 2311
2128void 2312void
2129devices (OpenCL::Program self) 2313devices (OpenCL::Program self)
2130 PPCODE: 2314 PPCODE:
2134 NEED_SUCCESS (GetProgramInfo, (self, CL_PROGRAM_DEVICES, size, value, 0)); 2318 NEED_SUCCESS (GetProgramInfo, (self, CL_PROGRAM_DEVICES, size, value, 0));
2135 int i, n = size / sizeof (*value); 2319 int i, n = size / sizeof (*value);
2136 EXTEND (SP, n); 2320 EXTEND (SP, n);
2137 for (i = 0; i < n; ++i) 2321 for (i = 0; i < n; ++i)
2138 { 2322 {
2139 PUSH_CLOBJ ("OpenCL::Device", value [i]); 2323 PUSH_CLOBJ (stash_device, value [i]);
2140 } 2324 }
2141 2325
2142void 2326void
2143source (OpenCL::Program self) 2327source (OpenCL::Program self)
2144 PPCODE: 2328 PPCODE:
2170DESTROY (OpenCL::Kernel self) 2354DESTROY (OpenCL::Kernel self)
2171 CODE: 2355 CODE:
2172 clReleaseKernel (self); 2356 clReleaseKernel (self);
2173 2357
2174void 2358void
2359setf (OpenCL::Kernel self, const char *format, ...)
2360 CODE:
2361 int i;
2362 for (i = 2; ; ++i)
2363 {
2364 while (*format == ' ')
2365 ++format;
2366
2367 char type = *format++;
2368
2369 if (!type)
2370 break;
2371
2372 if (i >= items)
2373 croak ("OpenCL::Kernel::setf format string too long (not enough arguments)");
2374
2375 SV *sv = ST (i);
2376
2377 union
2378 {
2379 cl_char cc; cl_uchar cC; cl_short cs; cl_ushort cS;
2380 cl_int ci; cl_uint cI; cl_long cl; cl_ulong cL;
2381 cl_half ch; cl_float cf; cl_double cd;
2382 cl_mem cm;
2383 cl_sampler ca;
2384 size_t cz;
2385 cl_event ce;
2386 } arg;
2387 size_t size;
2388 int nullarg = 0;
2389
2390 switch (type)
2391 {
2392 case 'c': arg.cc = SvIV (sv); size = sizeof (arg.cc); break;
2393 case 'C': arg.cC = SvUV (sv); size = sizeof (arg.cC); break;
2394 case 's': arg.cs = SvIV (sv); size = sizeof (arg.cs); break;
2395 case 'S': arg.cS = SvUV (sv); size = sizeof (arg.cS); break;
2396 case 'i': arg.ci = SvIV (sv); size = sizeof (arg.ci); break;
2397 case 'I': arg.cI = SvUV (sv); size = sizeof (arg.cI); break;
2398 case 'l': arg.cl = SvIV (sv); size = sizeof (arg.cl); break;
2399 case 'L': arg.cL = SvUV (sv); size = sizeof (arg.cL); break;
2400
2401 case 'h': arg.ch = SvUV (sv); size = sizeof (arg.ch); break;
2402 case 'f': arg.cf = SvNV (sv); size = sizeof (arg.cf); break;
2403 case 'd': arg.cd = SvNV (sv); size = sizeof (arg.cd); break;
2404
2405 case 'z': nullarg = 1; size = SvIV (sv); break;
2406
2407 case 'm': nullarg = !SvOK (sv); arg.cm = SvCLOBJ ("OpenCL::Kernel::setf", "m", sv, "OpenCL::Memory" ); size = sizeof (arg.cm); break;
2408 case 'a': nullarg = !SvOK (sv); arg.ca = SvCLOBJ ("OpenCL::Kernel::setf", "a", sv, "OpenCL::Sampler"); size = sizeof (arg.ca); break;
2409 case 'e': nullarg = !SvOK (sv); arg.ca = SvCLOBJ ("OpenCL::Kernel::setf", "e", sv, "OpenCL::Event" ); size = sizeof (arg.ce); break;
2410
2411 default:
2412 croak ("OpenCL::Kernel::setf format character '%c' not supported", type);
2413 }
2414
2415 res = clSetKernelArg (self, i - 2, size, nullarg ? 0 : &arg);
2416 if (res)
2417 croak ("OpenCL::Kernel::setf kernel parameter '%c' (#%d): %s", type, i - 2, err2str (res));
2418 }
2419
2420 if (i != items)
2421 croak ("OpenCL::Kernel::setf format string too short (too many arguments)");
2422
2423void
2175set_char (OpenCL::Kernel self, cl_uint idx, cl_char value) 2424set_char (OpenCL::Kernel self, cl_uint idx, cl_char value)
2176 CODE: 2425 CODE:
2177 clSetKernelArg (self, idx, sizeof (value), &value); 2426 clSetKernelArg (self, idx, sizeof (value), &value);
2178 2427
2179void 2428void
2227 clSetKernelArg (self, idx, sizeof (value), &value); 2476 clSetKernelArg (self, idx, sizeof (value), &value);
2228 2477
2229void 2478void
2230set_memory (OpenCL::Kernel self, cl_uint idx, OpenCL::Memory_ornull value) 2479set_memory (OpenCL::Kernel self, cl_uint idx, OpenCL::Memory_ornull value)
2231 CODE: 2480 CODE:
2232 clSetKernelArg (self, idx, sizeof (value), &value); 2481 clSetKernelArg (self, idx, sizeof (value), value ? &value : 0);
2233 2482
2234void 2483void
2235set_buffer (OpenCL::Kernel self, cl_uint idx, OpenCL::Buffer_ornull value) 2484set_buffer (OpenCL::Kernel self, cl_uint idx, OpenCL::Buffer_ornull value)
2236 CODE: 2485 CODE:
2237 clSetKernelArg (self, idx, sizeof (value), &value); 2486 clSetKernelArg (self, idx, sizeof (value), value ? &value : 0);
2238 2487
2239void 2488void
2240set_image2d (OpenCL::Kernel self, cl_uint idx, OpenCL::Image2D_ornull value) 2489set_image (OpenCL::Kernel self, cl_uint idx, OpenCL::Image_ornull value)
2490 ALIAS:
2491 set_image2d = 0
2492 set_image3d = 0
2241 CODE: 2493 CODE:
2242 clSetKernelArg (self, idx, sizeof (value), &value); 2494 clSetKernelArg (self, idx, sizeof (value), value ? &value : 0);
2243
2244void
2245set_image3d (OpenCL::Kernel self, cl_uint idx, OpenCL::Image3D_ornull value)
2246 CODE:
2247 clSetKernelArg (self, idx, sizeof (value), &value);
2248 2495
2249void 2496void
2250set_sampler (OpenCL::Kernel self, cl_uint idx, OpenCL::Sampler value) 2497set_sampler (OpenCL::Kernel self, cl_uint idx, OpenCL::Sampler value)
2251 CODE: 2498 CODE:
2252 clSetKernelArg (self, idx, sizeof (value), &value); 2499 clSetKernelArg (self, idx, sizeof (value), &value);
2298 NEED_SUCCESS (GetKernelInfo, (self, CL_KERNEL_CONTEXT, sizeof (value), value, 0)); 2545 NEED_SUCCESS (GetKernelInfo, (self, CL_KERNEL_CONTEXT, sizeof (value), value, 0));
2299 EXTEND (SP, 1); 2546 EXTEND (SP, 1);
2300 const int i = 0; 2547 const int i = 0;
2301 { 2548 {
2302 NEED_SUCCESS (RetainContext, (value [i])); 2549 NEED_SUCCESS (RetainContext, (value [i]));
2303 PUSH_CLOBJ ("OpenCL::Context", value [i]); 2550 PUSH_CLOBJ (stash_context, value [i]);
2304 } 2551 }
2305 2552
2306void 2553void
2307program (OpenCL::Kernel self) 2554program (OpenCL::Kernel self)
2308 PPCODE: 2555 PPCODE:
2310 NEED_SUCCESS (GetKernelInfo, (self, CL_KERNEL_PROGRAM, sizeof (value), value, 0)); 2557 NEED_SUCCESS (GetKernelInfo, (self, CL_KERNEL_PROGRAM, sizeof (value), value, 0));
2311 EXTEND (SP, 1); 2558 EXTEND (SP, 1);
2312 const int i = 0; 2559 const int i = 0;
2313 { 2560 {
2314 NEED_SUCCESS (RetainProgram, (value [i])); 2561 NEED_SUCCESS (RetainProgram, (value [i]));
2315 PUSH_CLOBJ ("OpenCL::Program", value [i]); 2562 PUSH_CLOBJ (stash_program, value [i]);
2316 } 2563 }
2317 2564
2318#END:kernel 2565#END:kernel
2319 2566
2320void 2567void
2400 NEED_SUCCESS (GetEventInfo, (self, CL_EVENT_COMMAND_QUEUE, sizeof (value), value, 0)); 2647 NEED_SUCCESS (GetEventInfo, (self, CL_EVENT_COMMAND_QUEUE, sizeof (value), value, 0));
2401 EXTEND (SP, 1); 2648 EXTEND (SP, 1);
2402 const int i = 0; 2649 const int i = 0;
2403 { 2650 {
2404 NEED_SUCCESS (RetainCommandQueue, (value [i])); 2651 NEED_SUCCESS (RetainCommandQueue, (value [i]));
2405 PUSH_CLOBJ ("OpenCL::Queue", value [i]); 2652 PUSH_CLOBJ (stash_queue, value [i]);
2406 } 2653 }
2407 2654
2408void 2655void
2409command_type (OpenCL::Event self) 2656command_type (OpenCL::Event self)
2410 PPCODE: 2657 PPCODE:
2433 NEED_SUCCESS (GetEventInfo, (self, CL_EVENT_CONTEXT, sizeof (value), value, 0)); 2680 NEED_SUCCESS (GetEventInfo, (self, CL_EVENT_CONTEXT, sizeof (value), value, 0));
2434 EXTEND (SP, 1); 2681 EXTEND (SP, 1);
2435 const int i = 0; 2682 const int i = 0;
2436 { 2683 {
2437 NEED_SUCCESS (RetainContext, (value [i])); 2684 NEED_SUCCESS (RetainContext, (value [i]));
2438 PUSH_CLOBJ ("OpenCL::Context", value [i]); 2685 PUSH_CLOBJ (stash_context, value [i]);
2439 } 2686 }
2440 2687
2441#END:event 2688#END:event
2442 2689
2443void 2690void
2468void 2715void
2469set_status (OpenCL::UserEvent self, cl_int execution_status) 2716set_status (OpenCL::UserEvent self, cl_int execution_status)
2470 CODE: 2717 CODE:
2471 clSetUserEventStatus (self, execution_status); 2718 clSetUserEventStatus (self, execution_status);
2472 2719
2720MODULE = OpenCL PACKAGE = OpenCL::Mapped
2721
2722void
2723DESTROY (SV *self)
2724 CODE:
2725 OpenCL__Mapped mapped = SvMAPPED (self);
2726
2727 clEnqueueUnmapMemObject (mapped->queue, mapped->memobj, mapped->ptr, 1, &mapped->event, 0);
2728 mapped_detach (self, mapped);
2729
2730 clReleaseCommandQueue (mapped->queue);
2731 clReleaseEvent (mapped->event);
2732 Safefree (mapped);
2733
2734void
2735unmap (OpenCL::Mapped self, ...)
2736 CODE:
2737 mapped_unmap (ST (0), self, self->queue, &ST (1), items - 1);
2738
2739bool
2740mapped (OpenCL::Mapped self)
2741 CODE:
2742 RETVAL = !!self->ptr;
2743 OUTPUT:
2744 RETVAL
2745
2746void
2747wait (OpenCL::Mapped self)
2748 PPCODE:
2749 if (self->event)
2750 NEED_SUCCESS (WaitForEvents, (1, &self->event));
2751
2752void
2753event (OpenCL::Mapped self)
2754 PPCODE:
2755 if (!self->event)
2756 XSRETURN_UNDEF;
2757
2758 clRetainEvent (self->event);
2759 XPUSH_CLOBJ (stash_event, self->event);
2760
2761size_t
2762size (OpenCL::Mapped self)
2763 CODE:
2764 RETVAL = self->cb;
2765 OUTPUT:
2766 RETVAL
2767
2768IV
2769ptr (OpenCL::Mapped self)
2770 CODE:
2771 RETVAL = PTR2IV (self->ptr);
2772 OUTPUT:
2773 RETVAL
2774
2775MODULE = OpenCL PACKAGE = OpenCL::MappedBuffer
2776
2777MODULE = OpenCL PACKAGE = OpenCL::MappedImage
2778
2779IV
2780row_pitch (OpenCL::Mapped self)
2781 ALIAS:
2782 slice_pitch = 1
2783 CODE:
2784 RETVAL = ix ? self->slice_pitch : self->row_pitch;
2785 OUTPUT:
2786 RETVAL
2787
2788
2789

Diff Legend

Removed lines
+ Added lines
< Changed lines
> Changed lines