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.83 by root, Mon Apr 7 22:36:54 2014 UTC

1#include <stddef.h>
2
1#include "EXTERN.h" 3#include "EXTERN.h"
2#include "perl.h" 4#include "perl.h"
3#include "XSUB.h" 5#include "XSUB.h"
6
7#define ECB_NO_THREADS 1
8#define ECB_NO_LIBM 1
9#include "ecb.h"
4 10
5#define X_STACKSIZE sizeof (void *) * 512 * 1024 // 2-4mb should be enough, really 11#define X_STACKSIZE sizeof (void *) * 512 * 1024 // 2-4mb should be enough, really
6#include "xthread.h" 12#include "xthread.h"
7#include "schmorp.h" 13#include "schmorp.h"
8 14
… …
26#if PREFER_1_1 32#if PREFER_1_1
27 #define CL_USE_DEPRECATED_OPENCL_1_1_APIS 33 #define CL_USE_DEPRECATED_OPENCL_1_1_APIS
28#endif 34#endif
29 35
30#ifdef __APPLE__ 36#ifdef __APPLE__
31 #include <OpenCL/opencl.h> 37 #define CLHDR(name) <OpenCL/name>
32#else 38#else
39 #define CLHDR(name) <CL/name>
40#endif
41
33 #include <CL/opencl.h> 42#include CLHDR(opencl.h)
43
44#if 0
45#ifndef CL_VERSION_1_2
46 #include CLHDR(cl_d3d9.h)
47#endif
48#endif
49
50#if _WIN32
51 #include CLHDR(cl_d3d10.h)
52 #if CL_VERSION_1_2
53 #include CLHDR(cl_d3d11.h)
54 #endif
55 #include CLHDR(cl_dx9_media_sharing.h)
34#endif 56#endif
35 57
36#ifndef CL_VERSION_1_2 58#ifndef CL_VERSION_1_2
37 #undef PREFER_1_1 59 #undef PREFER_1_1
38 #define PREFER_1_1 1 60 #define PREFER_1_1 1
39#endif 61#endif
40 62
63// make sure all constants we might use are actually defined
64#include "default.h"
65
41typedef cl_platform_id OpenCL__Platform; 66typedef cl_platform_id OpenCL__Platform;
42typedef cl_device_id OpenCL__Device; 67typedef cl_device_id OpenCL__Device;
68typedef cl_device_id OpenCL__SubDevice;
43typedef cl_context OpenCL__Context; 69typedef cl_context OpenCL__Context;
44typedef cl_command_queue OpenCL__Queue; 70typedef cl_command_queue OpenCL__Queue;
45typedef cl_mem OpenCL__Memory; 71typedef cl_mem OpenCL__Memory;
46typedef cl_mem OpenCL__Buffer; 72typedef cl_mem OpenCL__Buffer;
47typedef cl_mem OpenCL__BufferObj; 73typedef cl_mem OpenCL__BufferObj;
48typedef cl_mem OpenCL__Image; 74typedef cl_mem OpenCL__Image;
49typedef cl_mem OpenCL__Image2D;
50typedef cl_mem OpenCL__Image3D;
51typedef cl_mem OpenCL__Memory_ornull; 75typedef cl_mem OpenCL__Memory_ornull;
52typedef cl_mem OpenCL__Buffer_ornull; 76typedef cl_mem OpenCL__Buffer_ornull;
53typedef cl_mem OpenCL__Image_ornull; 77typedef cl_mem OpenCL__Image_ornull;
54typedef cl_mem OpenCL__Image2D_ornull;
55typedef cl_mem OpenCL__Image3D_ornull;
56typedef cl_sampler OpenCL__Sampler; 78typedef cl_sampler OpenCL__Sampler;
57typedef cl_program OpenCL__Program; 79typedef cl_program OpenCL__Program;
58typedef cl_kernel OpenCL__Kernel; 80typedef cl_kernel OpenCL__Kernel;
59typedef cl_event OpenCL__Event; 81typedef cl_event OpenCL__Event;
60typedef cl_event OpenCL__UserEvent; 82typedef cl_event OpenCL__UserEvent;
61 83
62typedef SV *FUTURE; 84typedef struct mapped * OpenCL__Mapped;
85
86static HV
87 *stash_platform,
88 *stash_device,
89 *stash_subdevice,
90 *stash_context,
91 *stash_queue,
92 *stash_program,
93 *stash_kernel,
94 *stash_sampler,
95 *stash_event,
96 *stash_userevent,
97 *stash_memory,
98 *stash_buffer,
99 *stash_bufferobj,
100 *stash_image,
101 *stash_image1d,
102 *stash_image1darray,
103 *stash_image1dbuffer,
104 *stash_image2d,
105 *stash_image2darray,
106 *stash_image3d,
107 *stash_mapped,
108 *stash_mappedbuffer,
109 *stash_mappedimage;
63 110
64/*****************************************************************************/ 111/*****************************************************************************/
65 112
66// name must include a leading underscore 113// name must include a leading underscore
67// all of this horrors would be unneceesary if somebody wrote a proper OpenGL module 114// all of this horrors would be unneceesary if somebody wrote a proper OpenGL module
… …
107} 154}
108 155
109/*****************************************************************************/ 156/*****************************************************************************/
110 157
111/* up to two temporary buffers */ 158/* up to two temporary buffers */
112static void * 159static void * ecb_noinline
113tmpbuf (size_t size) 160tmpbuf (size_t size)
114{ 161{
115 enum { buffers = 3 }; 162 enum { buffers = 4 };
116 static int idx; 163 static int idx;
117 static void *buf [buffers]; 164 static void *buf [buffers];
118 static size_t len [buffers]; 165 static size_t len [buffers];
119 166
120 idx = (idx + 1) % buffers; 167 idx = (idx + 1) % buffers;
… …
127 } 174 }
128 175
129 return buf [idx]; 176 return buf [idx];
130} 177}
131 178
179static const char * ecb_noinline
180cv_get_name (CV *cv)
181{
182 static char fullname [256];
183
184 GV *gv = CvGV (cv); // gv better be nonzero
185
186 HV *stash = GvSTASH (gv);
187 const char *hvname = HvNAME_get (stash); // stash also better be nonzero
188 const char *gvname = GvNAME (gv);
189
190 snprintf (fullname, sizeof (fullname), "%s::%s", hvname, gvname);
191 return fullname;
192}
193
132/*****************************************************************************/ 194/*****************************************************************************/
133 195
134typedef struct 196typedef struct
135{ 197{
136 IV iv; 198 IV iv;
137 const char *name; 199 const char *name;
138 #define const_iv(name) { (IV)CL_ ## name, # name }, 200 #define const_iv(name) { (IV)CL_ ## name, # name },
139} ivstr; 201} ivstr;
202
203typedef struct
204{
205 NV nv;
206 const char *name;
207 #define const_nv(name) { (NV)CL_ ## name, # name },
208} nvstr;
140 209
141static const char * 210static const char *
142iv2str (IV value, const ivstr *base, int count, const char *fallback) 211iv2str (IV value, const ivstr *base, int count, const char *fallback)
143{ 212{
144 int i; 213 int i;
… …
193 if (res) \ 262 if (res) \
194 FAIL (name); 263 FAIL (name);
195 264
196/*****************************************************************************/ 265/*****************************************************************************/
197 266
267static SV *
268new_clobj (HV *stash, IV id)
269{
270 return sv_2mortal (sv_bless (newRV_noinc (newSViv (id)), stash));
271}
272
273#define PUSH_CLOBJ(stash,id) PUSHs (new_clobj ((stash), (IV)(id)))
274#define XPUSH_CLOBJ(stash,id) XPUSHs (new_clobj ((stash), (IV)(id)))
275
276/* cl objects are either \$iv, or [$iv, ...] */
277/* they can be upgraded at runtime to the array form */
278static void * ecb_noinline
279SvCLOBJ (CV *cv, const char *svname, SV *sv, const char *pkg)
280{
281 // sv_derived_from is quite slow :(
282 if (SvROK (sv) && sv_derived_from (sv, pkg))
283 return (void *)SvIV (SvRV (sv));
284
285 croak ("%s: %s is not of type %s", cv_get_name (cv), svname, pkg);
286}
287
288// the "no-inherit" version of the above
289static void * ecb_noinline
290SvCLOBJ_ni (CV *cv, const char *svname, SV *sv, HV *stash)
291{
292 if (SvROK (sv) && SvSTASH (SvRV (sv)) == stash)
293 return (void *)SvIV (SvRV (sv));
294
295 croak ("%s: %s is not of type %s", cv_get_name (cv), svname, HvNAME (stash));
296}
297
298/*****************************************************************************/
299
198static cl_context_properties * 300static cl_context_properties * ecb_noinline
199SvCONTEXTPROPERTIES (const char *func, const char *svname, SV *sv, cl_context_properties *extra, int extracount) 301SvCONTEXTPROPERTIES (CV *cv, const char *svname, SV *sv, cl_context_properties *extra, int extracount)
200{ 302{
201 if (!sv || !SvOK (sv)) 303 if (!sv || !SvOK (sv))
202 if (extra) 304 if (extra)
203 sv = sv_2mortal (newRV_noinc ((SV *)newAV ())); // slow, but rarely used hopefully 305 sv = sv_2mortal (newRV_noinc ((SV *)newAV ())); // slow, but rarely used hopefully
204 else 306 else
… …
210 int i, len = av_len (av) + 1; 312 int i, len = av_len (av) + 1;
211 cl_context_properties *p = tmpbuf (sizeof (cl_context_properties) * (len + extracount + 1)); 313 cl_context_properties *p = tmpbuf (sizeof (cl_context_properties) * (len + extracount + 1));
212 cl_context_properties *l = p; 314 cl_context_properties *l = p;
213 315
214 if (len & 1) 316 if (len & 1)
215 croak ("%s: %s is not a property list (must be even number of elements)", func, svname); 317 croak ("%s: %s is not a property list (must contain an even number of elements)", cv_get_name (cv), svname);
216 318
217 while (extracount--) 319 while (extracount--)
218 *l++ = *extra++; 320 *l++ = *extra++;
219 321
220 for (i = 0; i < len; i += 2) 322 for (i = 0; i < len; i += 2)
… …
223 SV *p_sv = *av_fetch (av, i + 1, 0); 325 SV *p_sv = *av_fetch (av, i + 1, 0);
224 cl_context_properties v = SvIV (p_sv); // code below can override 326 cl_context_properties v = SvIV (p_sv); // code below can override
225 327
226 switch (t) 328 switch (t)
227 { 329 {
330 case CL_CONTEXT_PLATFORM:
331 if (SvROK (p_sv))
332 v = (cl_context_properties)SvCLOBJ (cv, svname, p_sv, "OpenCL::Platform");
333 break;
334
228 case CL_GLX_DISPLAY_KHR: 335 case CL_GLX_DISPLAY_KHR:
229 if (!SvOK (p_sv)) 336 if (!SvOK (p_sv))
230 { 337 {
231 void *func = glsym ("_glXGetCurrentDisplay"); 338 void *func = glsym ("_glXGetCurrentDisplay");
232 if (func) 339 if (func)
… …
255 *l = 0; 362 *l = 0;
256 363
257 return p; 364 return p;
258 } 365 }
259 366
260 croak ("%s: %s is not a property list (either undef or [type => value, ...])", func, svname); 367 croak ("%s: %s is not a property list (either undef or [type => value, ...])", cv_get_name (cv), svname);
261} 368}
262 369
263/*****************************************************************************/ 370// parse an array of CLOBJ into a void ** array in C - works only for CLOBJs whose representation
264 371// is a pointer (and only on well-behaved systems).
265#define NEW_CLOBJ(class,ptr) sv_setref_pv (sv_newmortal (), class, ptr) 372static void * ecb_noinline
266#define PUSH_CLOBJ(class,ptr) PUSHs (NEW_CLOBJ (class, ptr)) 373object_list (CV *cv, int or_undef, const char *argname, SV *arg, const char *klass, cl_uint *rcount)
267#define XPUSH_CLOBJ(class,ptr) XPUSHs (NEW_CLOBJ (class, ptr))
268
269/* cl objects are either \$iv, or [$iv, ...] */
270/* they can be upgraded at runtime to the array form */
271static void *
272SvCLOBJ (const char *func, const char *svname, SV *sv, const char *pkg)
273{ 374{
274 if (SvROK (sv) && sv_derived_from (sv, pkg)) 375 if (!SvROK (arg) || SvTYPE (SvRV (arg)) != SVt_PVAV)
376 croak ("%s: '%s' parameter must be %sa reference to an array of %s objects",
377 cv_get_name (cv), argname, or_undef ? "undef or " : "", klass);
378
379 AV *av = (AV *)SvRV (arg);
380 void **list = 0;
381 cl_uint count = av_len (av) + 1;
382
383 if (count)
275 { 384 {
276 SV *rv = SvRV (sv); 385 list = tmpbuf (sizeof (*list) * count);
277 386 int i;
278#if CLOBJ_PUSH 387 for (i = 0; i < count; ++i)
279 if (SvTYPE (rv) == SVt_PVAV) 388 list [i] = SvCLOBJ (cv, argname, *av_fetch (av, i, 1), klass);
280 rv = AvARRAY (rv)[0];
281#endif
282
283 return (void *)SvIV (rv);
284 } 389 }
285 390
286 croak ("%s: %s is not of type %s", func, svname, pkg); 391 if (!count && !or_undef)
287} 392 croak ("%s: '%s' must contain at least one %s object",
393 cv_get_name (cv), argname, klass);
288 394
289#if CLOBJ_PUSH 395 *rcount = count;
290 396 return (void *)list;
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} 397}
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 398
325/*****************************************************************************/ 399/*****************************************************************************/
326/* callback stuff */ 400/* callback stuff */
327 401
328/* default context callback, log to stderr */ 402/* default context callback, log to stderr */
… …
349static void (*eq_signal_func)(void *signal_arg, int value); 423static void (*eq_signal_func)(void *signal_arg, int value);
350static void *eq_signal_arg; 424static void *eq_signal_arg;
351static xmutex_t eq_lock = X_MUTEX_INIT; 425static xmutex_t eq_lock = X_MUTEX_INIT;
352static eq_item *eq_head, *eq_tail; 426static eq_item *eq_head, *eq_tail;
353 427
354static void 428static void ecb_noinline
355eq_enq (eq_vtbl *vtbl, SV *cb, void *data1, void *data2, void *data3) 429eq_enq (eq_vtbl *vtbl, SV *cb, void *data1, void *data2, void *data3)
356{ 430{
357 eq_item *item = malloc (sizeof (eq_item)); 431 eq_item *item = malloc (sizeof (eq_item));
358 432
359 item->next = 0; 433 item->next = 0;
… …
387 X_UNLOCK (eq_lock); 461 X_UNLOCK (eq_lock);
388 462
389 return res; 463 return res;
390} 464}
391 465
392#if 0
393static void
394mem_free (pTHX_ void *p)
395{
396 free (p);
397}
398//SAVEDESTRUCTOR_X (mem_free, item);
399#endif
400
401static void 466static void
402eq_poll (void) 467eq_poll (void)
403{ 468{
404 eq_item *item; 469 eq_item *item;
405 470
… …
435} 500}
436 501
437/*****************************************************************************/ 502/*****************************************************************************/
438/* context notify */ 503/* context notify */
439 504
440static void 505static void ecb_noinline
441eq_context_push (void *data1, void *data2, void *data3) 506eq_context_push (void *data1, void *data2, void *data3)
442{ 507{
443 dSP; 508 dSP;
444 PUSHs (sv_2mortal (newSVpv (data1, 0))); 509 PUSHs (sv_2mortal (newSVpv (data1, 0)));
445 PUSHs (sv_2mortal (newSVpvn (data2, (STRLEN)data3))); 510 PUSHs (sv_2mortal (newSVpvn (data2, (STRLEN)data3)));
… …
467 { \ 532 { \
468 pfn_notify = eq_context_notify; \ 533 pfn_notify = eq_context_notify; \
469 user_data = s_get_cv (notify); \ 534 user_data = s_get_cv (notify); \
470 } 535 }
471 536
472static SV * 537static SV * ecb_noinline
473new_clobj_context (cl_context ctx, void *user_data) 538new_clobj_context (cl_context ctx, void *user_data)
474{ 539{
475 SV *sv = NEW_CLOBJ ("OpenCL::Context", ctx); 540 SV *sv = new_clobj (stash_context, (IV)ctx);
476 541
477 if (user_data) 542 if (user_data)
478 sv_magicext (SvRV (sv), user_data, PERL_MAGIC_ext, 0, 0, 0); 543 sv_magicext (SvRV (sv), user_data, PERL_MAGIC_ext, 0, 0, 0);
479 544
480 return sv; 545 return sv;
… …
487 552
488static void 553static void
489eq_program_push (void *data1, void *data2, void *data3) 554eq_program_push (void *data1, void *data2, void *data3)
490{ 555{
491 dSP; 556 dSP;
492 PUSH_CLOBJ ("OpenCL::Program", data1); 557 PUSH_CLOBJ (stash_program, data1);
493 PUTBACK; 558 PUTBACK;
494} 559}
495 560
496static eq_vtbl eq_program_vtbl = { 1, eq_program_push }; 561static eq_vtbl eq_program_vtbl = { 1, eq_program_push };
497 562
498static void CL_CALLBACK 563static void CL_CALLBACK
499eq_program_notify (cl_program program, void *user_data) 564eq_program_notify (cl_program program, void *user_data)
500{ 565{
566 clRetainProgram (program);
567
501 eq_enq (&eq_program_vtbl, user_data, (void *)program, 0, 0); 568 eq_enq (&eq_program_vtbl, user_data, (void *)program, 0, 0);
569}
570
571typedef void (CL_CALLBACK *program_callback)(cl_program program, void *user_data);
572
573static program_callback ecb_noinline
574make_program_callback (SV *notify, void **ruser_data)
575{
576 if (SvOK (notify))
577 {
578 *ruser_data = SvREFCNT_inc (s_get_cv (notify));
579 return eq_program_notify;
580 }
581 else
582 {
583 *ruser_data = 0;
584 return 0;
585 }
502} 586}
503 587
504struct build_args 588struct build_args
505{ 589{
506 cl_program program; 590 cl_program program;
… …
520 else 604 else
521 clReleaseProgram (arg->program); 605 clReleaseProgram (arg->program);
522 606
523 free (arg->options); 607 free (arg->options);
524 free (arg); 608 free (arg);
609
610 return 0;
525} 611}
526 612
527static void 613static void
528build_program_async (cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options, void *user_data) 614build_program_async (cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options, void *user_data)
529{ 615{
… …
534 arg->user_data = user_data; 620 arg->user_data = user_data;
535 arg->num_devices = num_devices; 621 arg->num_devices = num_devices;
536 memcpy (arg + 1, device_list, sizeof (*device_list) * num_devices); 622 memcpy (arg + 1, device_list, sizeof (*device_list) * num_devices);
537 623
538 xthread_t id; 624 xthread_t id;
539 thread_create (&id, build_program_thread, arg); 625 xthread_create (&id, build_program_thread, arg);
626}
627
628/*****************************************************************************/
629/* mem object destructor notify */
630
631static void ecb_noinline
632eq_destructor_push (void *data1, void *data2, void *data3)
633{
634}
635
636static eq_vtbl eq_destructor_vtbl = { 0, eq_destructor_push };
637
638static void CL_CALLBACK
639eq_destructor_notify (cl_mem memobj, void *user_data)
640{
641 eq_enq (&eq_destructor_vtbl, (SV *)user_data, (void *)memobj, 0, 0);
540} 642}
541 643
542/*****************************************************************************/ 644/*****************************************************************************/
543/* event objects */ 645/* event objects */
544 646
545static void 647static void
546eq_event_push (void *data1, void *data2, void *data3) 648eq_event_push (void *data1, void *data2, void *data3)
547{ 649{
548 dSP; 650 dSP;
549 PUSH_CLOBJ ("OpenCL::Event", data1); 651 PUSH_CLOBJ (stash_event, data1);
550 PUSHs (sv_2mortal (newSViv ((IV)data2))); 652 PUSHs (sv_2mortal (newSViv ((IV)data2)));
551 PUTBACK; 653 PUTBACK;
552} 654}
553 655
554static eq_vtbl eq_event_vtbl = { 1, eq_event_push }; 656static eq_vtbl eq_event_vtbl = { 1, eq_event_push };
… …
559 clRetainEvent (event); 661 clRetainEvent (event);
560 eq_enq (&eq_event_vtbl, user_data, (void *)event, (void *)(IV)event_command_exec_status, 0); 662 eq_enq (&eq_event_vtbl, user_data, (void *)event, (void *)(IV)event_command_exec_status, 0);
561} 663}
562 664
563/*****************************************************************************/ 665/*****************************************************************************/
666/* utilities for XS code */
564 667
565static size_t 668static size_t
566img_row_pitch (cl_mem img) 669img_row_pitch (cl_mem img)
567{ 670{
568 size_t res; 671 size_t res;
569 clGetImageInfo (img, CL_IMAGE_ROW_PITCH, sizeof (res), &res, 0); 672 clGetImageInfo (img, CL_IMAGE_ROW_PITCH, sizeof (res), &res, 0);
570 return res; 673 return res;
571} 674}
572 675
573static cl_event * 676static cl_event * ecb_noinline
574event_list (SV **items, cl_uint *rcount) 677event_list (CV *cv, SV **items, cl_uint *rcount, cl_event extra)
575{ 678{
576 cl_uint count = *rcount; 679 cl_uint count = *rcount;
577 680
578 if (!count) 681 if (count > 0x7fffffffU) // yeah, it's a hack - the caller might have underflowed
682 *rcount = count = 0;
683
684 if (!count && !extra)
579 return 0; 685 return 0;
580 686
581 cl_event *list = tmpbuf (sizeof (cl_event) * count); 687 cl_event *list = tmpbuf (sizeof (cl_event) * (count + 1));
582 int i = 0; 688 int i = 0;
583 689
584 do 690 while (count--)
585 {
586 --count;
587 if (SvOK (items [count])) 691 if (SvOK (items [count]))
588 list [i++] = SvCLOBJ ("clEnqueue", "wait_events", items [count], "OpenCL::Event"); 692 list [i++] = SvCLOBJ (cv, "wait_events", items [count], "OpenCL::Event");
589 } 693
590 while (count); 694 if (extra)
695 list [i++] = extra;
591 696
592 *rcount = i; 697 *rcount = i;
593 698
594 return i ? list : 0; 699 return i ? list : 0;
595} 700}
596 701
597#define EVENT_LIST(items,count) \ 702#define EVENT_LIST(skip) \
598 cl_uint event_list_count = (count); \ 703 cl_uint event_list_count = items - (skip); \
599 cl_event *event_list_ptr = event_list (&ST (items), &event_list_count) 704 cl_event *event_list_ptr = event_list (cv, &ST (skip), &event_list_count, 0)
600 705
601#define INFO(class) \ 706#define INFO(class) \
602{ \ 707{ \
603 size_t size; \ 708 size_t size; \
604 NEED_SUCCESS (Get ## class ## Info, (self, name, 0, 0, &size)); \ 709 NEED_SUCCESS (Get ## class ## Info, (self, name, 0, 0, &size)); \
… …
608 SvCUR_set (sv, size); \ 713 SvCUR_set (sv, size); \
609 NEED_SUCCESS (Get ## class ## Info, (self, name, size, SvPVX (sv), 0)); \ 714 NEED_SUCCESS (Get ## class ## Info, (self, name, size, SvPVX (sv), 0)); \
610 XPUSHs (sv); \ 715 XPUSHs (sv); \
611} 716}
612 717
718/*****************************************************************************/
719/* mapped_xxx */
720
721static OpenCL__Mapped
722SvMAPPED (SV *sv)
723{
724 // no typechecking atm., keep your fingers crossed
725 return (OpenCL__Mapped)SvMAGIC (SvRV (sv))->mg_ptr;
726}
727
728struct mapped
729{
730 cl_command_queue queue;
731 cl_mem memobj;
732 void *ptr;
733 size_t cb;
734 cl_event event;
735 size_t row_pitch;
736 size_t slice_pitch;
737
738 size_t element_size;
739 size_t width, height, depth;
740};
741
742static SV *
743mapped_new (
744 HV *stash, cl_command_queue queue, cl_mem memobj, cl_map_flags flags,
745 void *ptr, size_t cb, cl_event ev,
746 size_t row_pitch, size_t slice_pitch, size_t element_size,
747 size_t width, size_t height, size_t depth
748)
749{
750 SV *data = newSV (0);
751 SvUPGRADE (data, SVt_PVMG);
752
753 OpenCL__Mapped mapped;
754 New (0, mapped, 1, struct mapped);
755
756 clRetainCommandQueue (queue);
757
758 mapped->queue = queue;
759 mapped->memobj = memobj;
760 mapped->ptr = ptr;
761 mapped->cb = cb;
762 mapped->event = ev;
763 mapped->row_pitch = row_pitch;
764 mapped->slice_pitch = slice_pitch;
765
766 mapped->element_size = element_size;
767 mapped->width = width;
768 mapped->height = height;
769 mapped->depth = depth;
770
771 sv_magicext (data, 0, PERL_MAGIC_ext, 0, (char *)mapped, 0);
772
773 if (SvLEN (data))
774 Safefree (data);
775
776 SvPVX (data) = (char *)ptr;
777 SvCUR_set (data, cb);
778 SvLEN_set (data, 0);
779 SvPOK_only (data);
780
781 SV *obj = sv_2mortal (sv_bless (newRV_noinc (data), stash));
782
783 if (!(flags & CL_MAP_WRITE))
784 SvREADONLY_on (data);
785
786 return obj;
787}
788
789static void
790mapped_detach (SV *sv, OpenCL__Mapped mapped)
791{
792 SV *data = SvRV (sv);
793
794 // the next check checks both before AND after detach, where SvPVX should STILL be 0
795 if (SvPVX (data) != (char *)mapped->ptr)
796 warn ("FATAL: OpenCL memory mapped scalar changed location, detected");
797 else
798 {
799 SvREADONLY_off (data);
800 SvCUR_set (data, 0);
801 SvPVX (data) = 0;
802 SvOK_off (data);
803 }
804
805 mapped->ptr = 0;
806}
807
808static void
809mapped_unmap (CV *cv, SV *self, OpenCL__Mapped mapped, cl_command_queue queue, SV **wait_list, cl_uint event_list_count)
810{
811 cl_event *event_list_ptr = event_list (cv, wait_list, &event_list_count, mapped->event);
812 cl_event ev;
813
814 NEED_SUCCESS (EnqueueUnmapMemObject, (queue, mapped->memobj, mapped->ptr, event_list_count, event_list_ptr, &ev));
815
816 clReleaseEvent (mapped->event);
817 mapped->event = ev;
818
819 mapped_detach (self, mapped);
820}
821
822static size_t
823mapped_element_size (OpenCL__Mapped self)
824{
825 if (!self->element_size)
826 clGetImageInfo (self->memobj, CL_IMAGE_ELEMENT_SIZE, sizeof (self->element_size), &self->element_size, 0);
827
828 return self->element_size;
829}
830
831/*****************************************************************************/
832
613MODULE = OpenCL PACKAGE = OpenCL 833MODULE = OpenCL PACKAGE = OpenCL
614 834
615PROTOTYPES: ENABLE 835PROTOTYPES: ENABLE
616 836
617void 837void
… …
626 eq_signal_arg = (void*)arg; 846 eq_signal_arg = (void*)arg;
627 847
628BOOT: 848BOOT:
629{ 849{
630 HV *stash = gv_stashpv ("OpenCL", 1); 850 HV *stash = gv_stashpv ("OpenCL", 1);
851
631 static const ivstr *civ, const_iv[] = { 852 static const ivstr *civ, const_iv[] = {
632 { sizeof (cl_char ), "SIZEOF_CHAR" }, 853 { sizeof (cl_char ), "SIZEOF_CHAR" },
633 { sizeof (cl_uchar ), "SIZEOF_UCHAR" }, 854 { sizeof (cl_uchar ), "SIZEOF_UCHAR" },
634 { sizeof (cl_short ), "SIZEOF_SHORT" }, 855 { sizeof (cl_short ), "SIZEOF_SHORT" },
635 { sizeof (cl_ushort), "SIZEOF_USHORT" }, 856 { sizeof (cl_ushort), "SIZEOF_USHORT" },
… …
638 { sizeof (cl_long ), "SIZEOF_LONG" }, 859 { sizeof (cl_long ), "SIZEOF_LONG" },
639 { sizeof (cl_ulong ), "SIZEOF_ULONG" }, 860 { sizeof (cl_ulong ), "SIZEOF_ULONG" },
640 { sizeof (cl_half ), "SIZEOF_HALF" }, 861 { sizeof (cl_half ), "SIZEOF_HALF" },
641 { sizeof (cl_float ), "SIZEOF_FLOAT" }, 862 { sizeof (cl_float ), "SIZEOF_FLOAT" },
642 { sizeof (cl_double), "SIZEOF_DOUBLE" }, 863 { sizeof (cl_double), "SIZEOF_DOUBLE" },
864 { PREFER_1_1 , "PREFER_1_1" },
643#include "constiv.h" 865#include "constiv.h"
644 }; 866 };
645 867
646 for (civ = const_iv + sizeof (const_iv) / sizeof (const_iv [0]); civ > const_iv; civ--) 868 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)); 869 newCONSTSUB (stash, (char *)civ[-1].name, newSViv (civ[-1].iv));
870
871 static const nvstr *cnv, const_nv[] = {
872#include "constnv.h"
873 };
874
875 for (cnv = const_nv + sizeof (const_nv) / sizeof (const_nv [0]); cnv > const_nv; cnv--)
876 newCONSTSUB (stash, (char *)cnv[-1].name, newSVnv (cnv[-1].nv));
877
878 newCONSTSUB (stash, "NAN", newSVnv (CL_NAN)); // CL_NAN might be a function call
879
880 stash_platform = gv_stashpv ("OpenCL::Platform", GV_ADD);
881 stash_device = gv_stashpv ("OpenCL::Device", GV_ADD);
882 stash_subdevice = gv_stashpv ("OpenCL::SubDevice", GV_ADD);
883 stash_context = gv_stashpv ("OpenCL::Context", GV_ADD);
884 stash_queue = gv_stashpv ("OpenCL::Queue", GV_ADD);
885 stash_program = gv_stashpv ("OpenCL::Program", GV_ADD);
886 stash_kernel = gv_stashpv ("OpenCL::Kernel", GV_ADD);
887 stash_sampler = gv_stashpv ("OpenCL::Sampler", GV_ADD);
888 stash_event = gv_stashpv ("OpenCL::Event", GV_ADD);
889 stash_userevent = gv_stashpv ("OpenCL::UserEvent", GV_ADD);
890 stash_memory = gv_stashpv ("OpenCL::Memory", GV_ADD);
891 stash_buffer = gv_stashpv ("OpenCL::Buffer", GV_ADD);
892 stash_bufferobj = gv_stashpv ("OpenCL::BufferObj", GV_ADD);
893 stash_image = gv_stashpv ("OpenCL::Image", GV_ADD);
894 stash_image1d = gv_stashpv ("OpenCL::Image1D", GV_ADD);
895 stash_image1darray = gv_stashpv ("OpenCL::Image1DArray", GV_ADD);
896 stash_image1dbuffer = gv_stashpv ("OpenCL::Image1DBuffer", GV_ADD);
897 stash_image2d = gv_stashpv ("OpenCL::Image2D", GV_ADD);
898 stash_image2darray = gv_stashpv ("OpenCL::Image2DArray", GV_ADD);
899 stash_image3d = gv_stashpv ("OpenCL::Image3D", GV_ADD);
900 stash_mapped = gv_stashpv ("OpenCL::Mapped", GV_ADD);
901 stash_mappedbuffer = gv_stashpv ("OpenCL::MappedBuffer", GV_ADD);
902 stash_mappedimage = gv_stashpv ("OpenCL::MappedImage", GV_ADD);
648 903
649 sv_setiv (perl_get_sv ("OpenCL::POLL_FUNC", TRUE), (IV)eq_poll_interrupt); 904 sv_setiv (perl_get_sv ("OpenCL::POLL_FUNC", TRUE), (IV)eq_poll_interrupt);
650} 905}
651 906
652cl_int 907cl_int
… …
655 RETVAL = res; 910 RETVAL = res;
656 OUTPUT: 911 OUTPUT:
657 RETVAL 912 RETVAL
658 913
659const char * 914const char *
660err2str (cl_int err) 915err2str (cl_int err = res)
661 916
662const char * 917const char *
663enum2str (cl_uint value) 918enum2str (cl_uint value)
664 919
665void 920void
… …
673 list = tmpbuf (sizeof (*list) * count); 928 list = tmpbuf (sizeof (*list) * count);
674 NEED_SUCCESS (GetPlatformIDs, (count, list, 0)); 929 NEED_SUCCESS (GetPlatformIDs, (count, list, 0));
675 930
676 EXTEND (SP, count); 931 EXTEND (SP, count);
677 for (i = 0; i < count; ++i) 932 for (i = 0; i < count; ++i)
678 PUSH_CLOBJ ("OpenCL::Platform", list [i]); 933 PUSH_CLOBJ (stash_platform, list [i]);
679 934
680void 935void
681context_from_type (cl_context_properties *properties = 0, cl_device_type type = CL_DEVICE_TYPE_DEFAULT, SV *notify = &PL_sv_undef) 936context_from_type (cl_context_properties *properties = 0, cl_device_type type = CL_DEVICE_TYPE_DEFAULT, SV *notify = &PL_sv_undef)
682 PPCODE: 937 PPCODE:
683 CONTEXT_NOTIFY_CALLBACK; 938 CONTEXT_NOTIFY_CALLBACK;
684 NEED_SUCCESS_ARG (cl_context ctx, CreateContextFromType, (properties, type, 0, 0, &res)); 939 NEED_SUCCESS_ARG (cl_context ctx, CreateContextFromType, (properties, type, pfn_notify, user_data, &res));
685 XPUSH_CLOBJ_CONTEXT; 940 XPUSH_CLOBJ_CONTEXT;
686 941
687void 942void
688context (FUTURE properties, FUTURE devices, FUTURE notify) 943context (cl_context_properties *properties, SV *devices, SV *notify = &PL_sv_undef)
689 PPCODE: 944 PPCODE:
690 /* der Gipfel der Kunst */ 945 cl_uint device_count;
946 cl_device_id *device_list = object_list (cv, 0, "devices", devices, "OpenCL::Device", &device_count);
947
948 CONTEXT_NOTIFY_CALLBACK;
949 NEED_SUCCESS_ARG (cl_context ctx, CreateContext, (properties, device_count, device_list, pfn_notify, user_data, &res));
950 XPUSH_CLOBJ_CONTEXT;
691 951
692void 952void
693wait_for_events (...) 953wait_for_events (...)
694 CODE: 954 CODE:
695 EVENT_LIST (0, items); 955 EVENT_LIST (0);
696 NEED_SUCCESS (WaitForEvents, (event_list_count, event_list_ptr)); 956 NEED_SUCCESS (WaitForEvents, (event_list_count, event_list_ptr));
697 957
698PROTOTYPES: DISABLE 958PROTOTYPES: DISABLE
699 959
700MODULE = OpenCL PACKAGE = OpenCL::Platform 960MODULE = OpenCL PACKAGE = OpenCL::Platform
… …
743 list = tmpbuf (sizeof (*list) * count); 1003 list = tmpbuf (sizeof (*list) * count);
744 NEED_SUCCESS (GetDeviceIDs, (self, type, count, list, 0)); 1004 NEED_SUCCESS (GetDeviceIDs, (self, type, count, list, 0));
745 1005
746 EXTEND (SP, count); 1006 EXTEND (SP, count);
747 for (i = 0; i < count; ++i) 1007 for (i = 0; i < count; ++i)
748 PUSHs (sv_setref_pv (sv_newmortal (), "OpenCL::Device", list [i])); 1008 PUSH_CLOBJ (stash_device, list [i]);
749 1009
750void 1010void
751context (OpenCL::Platform self, cl_context_properties *properties, SV *devices, SV *notify = &PL_sv_undef) 1011context (OpenCL::Platform self, SV *properties, SV *devices, SV *notify = &PL_sv_undef)
752 PPCODE: 1012 PPCODE:
1013 cl_context_properties extra[] = { CL_CONTEXT_PLATFORM, (cl_context_properties)self };
1014 cl_context_properties *props = SvCONTEXTPROPERTIES (cv, "properties", properties, extra, 2);
1015
1016 cl_uint device_count;
1017 cl_device_id *device_list = object_list (cv, 0, "devices", devices, "OpenCL::Device", &device_count);
1018
1019 CONTEXT_NOTIFY_CALLBACK;
1020 NEED_SUCCESS_ARG (cl_context ctx, CreateContext, (props, device_count, device_list, pfn_notify, user_data, &res));
1021 XPUSH_CLOBJ_CONTEXT;
1022
1023void
1024context_from_type (OpenCL::Platform self, SV *properties = 0, cl_device_type type = CL_DEVICE_TYPE_DEFAULT, SV *notify = &PL_sv_undef)
1025 PPCODE:
1026 cl_context_properties extra[] = { CL_CONTEXT_PLATFORM, (cl_context_properties)self };
1027 cl_context_properties *props = SvCONTEXTPROPERTIES (cv, "properties", properties, extra, 2);
1028
1029 CONTEXT_NOTIFY_CALLBACK;
1030 NEED_SUCCESS_ARG (cl_context ctx, CreateContextFromType, (props, type, pfn_notify, user_data, &res));
1031 XPUSH_CLOBJ_CONTEXT;
1032
1033MODULE = OpenCL PACKAGE = OpenCL::Device
1034
1035void
1036info (OpenCL::Device self, cl_device_info name)
1037 PPCODE:
1038 INFO (Device)
1039
1040#if CL_VERSION_1_2
1041
1042void
1043sub_devices (OpenCL::Device self, SV *properties)
1044 PPCODE:
753 if (!SvROK (devices) || SvTYPE (SvRV (devices)) != SVt_PVAV) 1045 if (!SvROK (properties) || SvTYPE (SvRV (properties)) != SVt_PVAV)
754 croak ("OpenCL::Platform::context argument 'device' must be an arrayref with device objects, in call"); 1046 croak ("OpenCL::Device::sub_devices: properties must be specified as reference to an array of property-value pairs");
755 1047
756 AV *av = (AV *)SvRV (devices); 1048 properties = SvRV (properties);
757 cl_uint num_devices = av_len (av) + 1; 1049
758 cl_device_id *device_list = tmpbuf (sizeof (cl_device_id) * num_devices); 1050 cl_uint count = av_len ((AV *)properties) + 1;
1051 cl_device_partition_property *props = tmpbuf (sizeof (*props) * count + 1);
759 1052
760 int i; 1053 int i;
761 for (i = num_devices; i--; ) 1054 for (i = 0; i < count; ++i)
762 device_list [i] = SvCLOBJ ("clCreateContext", "devices", *av_fetch (av, i, 0), "OpenCL::Device"); 1055 props [i] = (cl_device_partition_property)SvIV (*av_fetch ((AV *)properties, i, 0));
763 1056
764 CONTEXT_NOTIFY_CALLBACK; 1057 props [count] = 0;
765 NEED_SUCCESS_ARG (cl_context ctx, CreateContext, (properties, num_devices, device_list, pfn_notify, user_data, &res));
766 XPUSH_CLOBJ_CONTEXT;
767 1058
768void 1059 cl_uint num_devices;
769context_from_type (OpenCL::Platform self, SV *properties = 0, cl_device_type type = CL_DEVICE_TYPE_DEFAULT, SV *notify = &PL_sv_undef) 1060 NEED_SUCCESS (CreateSubDevices, (self, props, 0, 0, &num_devices));
770 PPCODE: 1061 cl_device_id *list = tmpbuf (sizeof (*list) * num_devices);
771 cl_context_properties extra[] = { CL_CONTEXT_PLATFORM, (cl_context_properties)self }; 1062 NEED_SUCCESS (CreateSubDevices, (self, props, num_devices, list, 0));
772 cl_context_properties *props = SvCONTEXTPROPERTIES ("OpenCL::Platform::context_from_type", "properties", properties, extra, 2);
773 1063
774 CONTEXT_NOTIFY_CALLBACK; 1064 EXTEND (SP, num_devices);
775 NEED_SUCCESS_ARG (cl_context ctx, CreateContextFromType, (props, type, 0, 0, &res)); 1065 for (i = 0; i < count; ++i)
776 XPUSH_CLOBJ_CONTEXT; 1066 PUSH_CLOBJ (stash_subdevice, list [i]);
777 1067
778MODULE = OpenCL PACKAGE = OpenCL::Device 1068#endif
779
780void
781info (OpenCL::Device self, cl_device_info name)
782 PPCODE:
783 INFO (Device)
784 1069
785#BEGIN:device 1070#BEGIN:device
786 1071
787void 1072void
788type (OpenCL::Device self) 1073type (OpenCL::Device self)
1074 ALIAS:
1075 type = CL_DEVICE_TYPE
1076 address_bits = CL_DEVICE_ADDRESS_BITS
1077 max_mem_alloc_size = CL_DEVICE_MAX_MEM_ALLOC_SIZE
1078 single_fp_config = CL_DEVICE_SINGLE_FP_CONFIG
1079 global_mem_cache_size = CL_DEVICE_GLOBAL_MEM_CACHE_SIZE
1080 global_mem_size = CL_DEVICE_GLOBAL_MEM_SIZE
1081 max_constant_buffer_size = CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE
1082 local_mem_size = CL_DEVICE_LOCAL_MEM_SIZE
1083 execution_capabilities = CL_DEVICE_EXECUTION_CAPABILITIES
1084 properties = CL_DEVICE_QUEUE_PROPERTIES
1085 double_fp_config = CL_DEVICE_DOUBLE_FP_CONFIG
1086 half_fp_config = CL_DEVICE_HALF_FP_CONFIG
789 PPCODE: 1087 PPCODE:
790 cl_device_type value [1]; 1088 cl_ulong value [1];
791 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_TYPE, sizeof (value), value, 0)); 1089 NEED_SUCCESS (GetDeviceInfo, (self, ix, sizeof (value), value, 0));
792 EXTEND (SP, 1); 1090 EXTEND (SP, 1);
793 const int i = 0; 1091 const int i = 0;
794 PUSHs (sv_2mortal (newSViv (value [i]))); 1092 PUSHs (sv_2mortal (newSVuv (value [i])));
795 1093
796void 1094void
797vendor_id (OpenCL::Device self) 1095vendor_id (OpenCL::Device self)
798 ALIAS: 1096 ALIAS:
799 vendor_id = CL_DEVICE_VENDOR_ID 1097 vendor_id = CL_DEVICE_VENDOR_ID
… …
810 max_write_image_args = CL_DEVICE_MAX_WRITE_IMAGE_ARGS 1108 max_write_image_args = CL_DEVICE_MAX_WRITE_IMAGE_ARGS
811 image_support = CL_DEVICE_IMAGE_SUPPORT 1109 image_support = CL_DEVICE_IMAGE_SUPPORT
812 max_samplers = CL_DEVICE_MAX_SAMPLERS 1110 max_samplers = CL_DEVICE_MAX_SAMPLERS
813 mem_base_addr_align = CL_DEVICE_MEM_BASE_ADDR_ALIGN 1111 mem_base_addr_align = CL_DEVICE_MEM_BASE_ADDR_ALIGN
814 min_data_type_align_size = CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE 1112 min_data_type_align_size = CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE
1113 global_mem_cache_type = CL_DEVICE_GLOBAL_MEM_CACHE_TYPE
815 global_mem_cacheline_size = CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE 1114 global_mem_cacheline_size = CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE
816 max_constant_args = CL_DEVICE_MAX_CONSTANT_ARGS 1115 max_constant_args = CL_DEVICE_MAX_CONSTANT_ARGS
1116 local_mem_type = CL_DEVICE_LOCAL_MEM_TYPE
817 preferred_vector_width_half = CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF 1117 preferred_vector_width_half = CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF
818 native_vector_width_char = CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR 1118 native_vector_width_char = CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR
819 native_vector_width_short = CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT 1119 native_vector_width_short = CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT
820 native_vector_width_int = CL_DEVICE_NATIVE_VECTOR_WIDTH_INT 1120 native_vector_width_int = CL_DEVICE_NATIVE_VECTOR_WIDTH_INT
821 native_vector_width_long = CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG 1121 native_vector_width_long = CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG
… …
859 EXTEND (SP, n); 1159 EXTEND (SP, n);
860 for (i = 0; i < n; ++i) 1160 for (i = 0; i < n; ++i)
861 PUSHs (sv_2mortal (newSVuv (value [i]))); 1161 PUSHs (sv_2mortal (newSVuv (value [i])));
862 1162
863void 1163void
864address_bits (OpenCL::Device self)
865 PPCODE:
866 cl_bitfield value [1];
867 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_ADDRESS_BITS, sizeof (value), value, 0));
868 EXTEND (SP, 1);
869 const int i = 0;
870 PUSHs (sv_2mortal (newSVuv (value [i])));
871
872void
873max_mem_alloc_size (OpenCL::Device self)
874 ALIAS:
875 max_mem_alloc_size = CL_DEVICE_MAX_MEM_ALLOC_SIZE
876 global_mem_cache_size = CL_DEVICE_GLOBAL_MEM_CACHE_SIZE
877 global_mem_size = CL_DEVICE_GLOBAL_MEM_SIZE
878 max_constant_buffer_size = CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE
879 local_mem_size = CL_DEVICE_LOCAL_MEM_SIZE
880 PPCODE:
881 cl_ulong value [1];
882 NEED_SUCCESS (GetDeviceInfo, (self, ix, sizeof (value), value, 0));
883 EXTEND (SP, 1);
884 const int i = 0;
885 PUSHs (sv_2mortal (newSVuv (value [i])));
886
887void
888single_fp_config (OpenCL::Device self)
889 ALIAS:
890 single_fp_config = CL_DEVICE_SINGLE_FP_CONFIG
891 double_fp_config = CL_DEVICE_DOUBLE_FP_CONFIG
892 half_fp_config = CL_DEVICE_HALF_FP_CONFIG
893 PPCODE:
894 cl_device_fp_config value [1];
895 NEED_SUCCESS (GetDeviceInfo, (self, ix, sizeof (value), value, 0));
896 EXTEND (SP, 1);
897 const int i = 0;
898 PUSHs (sv_2mortal (newSVuv (value [i])));
899
900void
901global_mem_cache_type (OpenCL::Device self)
902 PPCODE:
903 cl_device_mem_cache_type value [1];
904 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_GLOBAL_MEM_CACHE_TYPE, sizeof (value), value, 0));
905 EXTEND (SP, 1);
906 const int i = 0;
907 PUSHs (sv_2mortal (newSVuv (value [i])));
908
909void
910local_mem_type (OpenCL::Device self)
911 PPCODE:
912 cl_device_local_mem_type value [1];
913 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_LOCAL_MEM_TYPE, sizeof (value), value, 0));
914 EXTEND (SP, 1);
915 const int i = 0;
916 PUSHs (sv_2mortal (newSVuv (value [i])));
917
918void
919error_correction_support (OpenCL::Device self) 1164error_correction_support (OpenCL::Device self)
920 ALIAS: 1165 ALIAS:
921 error_correction_support = CL_DEVICE_ERROR_CORRECTION_SUPPORT 1166 error_correction_support = CL_DEVICE_ERROR_CORRECTION_SUPPORT
922 endian_little = CL_DEVICE_ENDIAN_LITTLE 1167 endian_little = CL_DEVICE_ENDIAN_LITTLE
923 available = CL_DEVICE_AVAILABLE 1168 available = CL_DEVICE_AVAILABLE
924 compiler_available = CL_DEVICE_COMPILER_AVAILABLE 1169 compiler_available = CL_DEVICE_COMPILER_AVAILABLE
925 host_unified_memory = CL_DEVICE_HOST_UNIFIED_MEMORY 1170 host_unified_memory = CL_DEVICE_HOST_UNIFIED_MEMORY
926 PPCODE: 1171 PPCODE:
927 cl_bool value [1]; 1172 cl_uint value [1];
928 NEED_SUCCESS (GetDeviceInfo, (self, ix, sizeof (value), value, 0)); 1173 NEED_SUCCESS (GetDeviceInfo, (self, ix, sizeof (value), value, 0));
929 EXTEND (SP, 1); 1174 EXTEND (SP, 1);
930 const int i = 0; 1175 const int i = 0;
931 PUSHs (sv_2mortal (value [i] ? &PL_sv_yes : &PL_sv_no)); 1176 PUSHs (sv_2mortal (value [i] ? &PL_sv_yes : &PL_sv_no));
932
933void
934execution_capabilities (OpenCL::Device self)
935 PPCODE:
936 cl_device_exec_capabilities value [1];
937 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_EXECUTION_CAPABILITIES, sizeof (value), value, 0));
938 EXTEND (SP, 1);
939 const int i = 0;
940 PUSHs (sv_2mortal (newSVuv (value [i])));
941
942void
943properties (OpenCL::Device self)
944 PPCODE:
945 cl_command_queue_properties value [1];
946 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_QUEUE_PROPERTIES, sizeof (value), value, 0));
947 EXTEND (SP, 1);
948 const int i = 0;
949 PUSHs (sv_2mortal (newSViv (value [i])));
950 1177
951void 1178void
952platform (OpenCL::Device self) 1179platform (OpenCL::Device self)
953 PPCODE: 1180 PPCODE:
954 cl_platform_id value [1]; 1181 cl_platform_id value [1];
955 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_PLATFORM, sizeof (value), value, 0)); 1182 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_PLATFORM, sizeof (value), value, 0));
956 EXTEND (SP, 1); 1183 EXTEND (SP, 1);
957 const int i = 0; 1184 const int i = 0;
958 {
959 PUSH_CLOBJ ("OpenCL::Platform", value [i]); 1185 PUSH_CLOBJ (stash_platform, value [i]);
960 }
961 1186
962void 1187void
963name (OpenCL::Device self) 1188name (OpenCL::Device self)
964 ALIAS: 1189 ALIAS:
965 name = CL_DEVICE_NAME 1190 name = CL_DEVICE_NAME
… …
982 PPCODE: 1207 PPCODE:
983 cl_device_id value [1]; 1208 cl_device_id value [1];
984 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_PARENT_DEVICE_EXT, sizeof (value), value, 0)); 1209 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_PARENT_DEVICE_EXT, sizeof (value), value, 0));
985 EXTEND (SP, 1); 1210 EXTEND (SP, 1);
986 const int i = 0; 1211 const int i = 0;
987 {
988 PUSH_CLOBJ ("OpenCL::Device", value [i]); 1212 PUSH_CLOBJ (stash_device, value [i]);
989 }
990 1213
991void 1214void
992partition_types_ext (OpenCL::Device self) 1215partition_types_ext (OpenCL::Device self)
993 ALIAS: 1216 ALIAS:
994 partition_types_ext = CL_DEVICE_PARTITION_TYPES_EXT 1217 partition_types_ext = CL_DEVICE_PARTITION_TYPES_EXT
… …
1004 for (i = 0; i < n; ++i) 1227 for (i = 0; i < n; ++i)
1005 PUSHs (sv_2mortal (newSVuv (value [i]))); 1228 PUSHs (sv_2mortal (newSVuv (value [i])));
1006 1229
1007#END:device 1230#END:device
1008 1231
1232MODULE = OpenCL PACKAGE = OpenCL::SubDevice
1233
1234#if CL_VERSION_1_2
1235
1236void
1237DESTROY (OpenCL::SubDevice self)
1238 CODE:
1239 clReleaseDevice (self);
1240
1241#endif
1242
1009MODULE = OpenCL PACKAGE = OpenCL::Context 1243MODULE = OpenCL PACKAGE = OpenCL::Context
1010 1244
1011void 1245void
1012DESTROY (OpenCL::Context context) 1246DESTROY (OpenCL::Context self)
1013 CODE: 1247 CODE:
1014 clReleaseContext (context); 1248 clReleaseContext (self);
1015 1249
1016void 1250void
1017info (OpenCL::Context self, cl_context_info name) 1251info (OpenCL::Context self, cl_context_info name)
1018 PPCODE: 1252 PPCODE:
1019 INFO (Context) 1253 INFO (Context)
1020 1254
1021void 1255void
1022queue (OpenCL::Context self, OpenCL::Device device, cl_command_queue_properties properties = 0) 1256queue (OpenCL::Context self, OpenCL::Device device, cl_command_queue_properties properties = 0)
1023 PPCODE: 1257 PPCODE:
1024 NEED_SUCCESS_ARG (cl_command_queue queue, CreateCommandQueue, (self, device, properties, &res)); 1258 NEED_SUCCESS_ARG (cl_command_queue queue, CreateCommandQueue, (self, device, properties, &res));
1025 XPUSH_CLOBJ ("OpenCL::Queue", queue); 1259 XPUSH_CLOBJ (stash_queue, queue);
1026 1260
1027void 1261void
1028user_event (OpenCL::Context self) 1262user_event (OpenCL::Context self)
1029 PPCODE: 1263 PPCODE:
1030 NEED_SUCCESS_ARG (cl_event ev, CreateUserEvent, (self, &res)); 1264 NEED_SUCCESS_ARG (cl_event ev, CreateUserEvent, (self, &res));
1031 XPUSH_CLOBJ ("OpenCL::UserEvent", ev); 1265 XPUSH_CLOBJ (stash_userevent, ev);
1032 1266
1033void 1267void
1034buffer (OpenCL::Context self, cl_mem_flags flags, size_t len) 1268buffer (OpenCL::Context self, cl_mem_flags flags, size_t len)
1035 PPCODE: 1269 PPCODE:
1036 if (flags & (CL_MEM_USE_HOST_PTR | CL_MEM_COPY_HOST_PTR)) 1270 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?"); 1271 croak ("OpenCL::Context::buffer: cannot use/copy host ptr when no data is given, use $context->buffer_sv instead?");
1038 1272
1039 NEED_SUCCESS_ARG (cl_mem mem, CreateBuffer, (self, flags, len, 0, &res)); 1273 NEED_SUCCESS_ARG (cl_mem mem, CreateBuffer, (self, flags, len, 0, &res));
1040 XPUSH_CLOBJ ("OpenCL::BufferObj", mem); 1274 XPUSH_CLOBJ (stash_bufferobj, mem);
1041 1275
1042void 1276void
1043buffer_sv (OpenCL::Context self, cl_mem_flags flags, SV *data) 1277buffer_sv (OpenCL::Context self, cl_mem_flags flags, SV *data)
1044 PPCODE: 1278 PPCODE:
1045 STRLEN len; 1279 STRLEN len;
1046 char *ptr = SvOK (data) ? SvPVbyte (data, len) : 0; 1280 char *ptr = SvOK (data) ? SvPVbyte (data, len) : 0;
1047 if (!(flags & (CL_MEM_USE_HOST_PTR | CL_MEM_COPY_HOST_PTR))) 1281 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?"); 1282 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)); 1283 NEED_SUCCESS_ARG (cl_mem mem, CreateBuffer, (self, flags, len, ptr, &res));
1050 XPUSH_CLOBJ ("OpenCL::BufferObj", mem); 1284 XPUSH_CLOBJ (stash_bufferobj, mem);
1051 1285
1052#if CL_VERSION_1_2 1286#if CL_VERSION_1_2
1053 1287
1054void 1288void
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) 1289image (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: 1290 PPCODE:
1057 STRLEN len; 1291 STRLEN len;
1058 char *ptr = SvOK (data) ? SvPVbyte (data, len) : 0; 1292 char *ptr = SvOK (data) ? SvPVbyte (data, len) : 0;
1059 const cl_image_format format = { channel_order, channel_type }; 1293 const cl_image_format format = { channel_order, channel_type };
1060 const cl_image_desc desc = { 1294 const cl_image_desc desc = {
1061 type, 1295 type,
1062 width, height, depth, 1296 width, height, depth,
1063 array_size, row_pitch, slice_pitch, 1297 array_size, row_pitch, slice_pitch,
1064 num_mip_level, num_samples, 1298 num_mip_level, num_samples,
1065 type == CL_MEM_OBJECT_IMAGE1D_BUFFER ? (cl_mem)SvCLOBJ ("OpenCL::Context::Image", "data", data, "OpenCL::Buffer") : 0 1299 type == CL_MEM_OBJECT_IMAGE1D_BUFFER ? (cl_mem)SvCLOBJ (cv, "data", data, "OpenCL::Buffer") : 0
1066 }; 1300 };
1067 NEED_SUCCESS_ARG (cl_mem mem, CreateImage, (self, flags, &format, &desc, ptr, &res)); 1301 NEED_SUCCESS_ARG (cl_mem mem, CreateImage, (self, flags, &format, &desc, ptr, &res));
1068 char *klass = "OpenCL::Image"; 1302 HV *stash = stash_image;
1069 switch (type) 1303 switch (type)
1070 { 1304 {
1071 case CL_MEM_OBJECT_IMAGE1D_BUFFER: klass = "OpenCL::Image1DBuffer"; break; 1305 case CL_MEM_OBJECT_IMAGE1D_BUFFER: stash = stash_image1dbuffer; break;
1072 case CL_MEM_OBJECT_IMAGE1D: klass = "OpenCL::Image1D"; break; 1306 case CL_MEM_OBJECT_IMAGE1D: stash = stash_image1d; break;
1073 case CL_MEM_OBJECT_IMAGE1D_ARRAY: klass = "OpenCL::Image2DArray"; break; 1307 case CL_MEM_OBJECT_IMAGE1D_ARRAY: stash = stash_image2darray; break;
1074 case CL_MEM_OBJECT_IMAGE2D: klass = "OpenCL::Image2D"; break; 1308 case CL_MEM_OBJECT_IMAGE2D: stash = stash_image2d; break;
1075 case CL_MEM_OBJECT_IMAGE2D_ARRAY: klass = "OpenCL::Image2DArray"; break; 1309 case CL_MEM_OBJECT_IMAGE2D_ARRAY: stash = stash_image2darray; break;
1076 case CL_MEM_OBJECT_IMAGE3D: klass = "OpenCL::Image3D"; break; 1310 case CL_MEM_OBJECT_IMAGE3D: stash = stash_image3d; break;
1077 } 1311 }
1078 XPUSH_CLOBJ (klass, mem); 1312 XPUSH_CLOBJ (stash, mem);
1079 1313
1080#endif 1314#endif
1081 1315
1082void 1316void
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) 1317image2d (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)); 1323 NEED_SUCCESS_ARG (cl_mem mem, CreateImage2D, (self, flags, &format, width, height, row_pitch, ptr, &res));
1090#else 1324#else
1091 const cl_image_desc desc = { CL_MEM_OBJECT_IMAGE2D, width, height, 0, 0, row_pitch, 0, 0, 0, 0 }; 1325 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)); 1326 NEED_SUCCESS_ARG (cl_mem mem, CreateImage, (self, flags, &format, &desc, ptr, &res));
1093#endif 1327#endif
1094 XPUSH_CLOBJ ("OpenCL::Image2D", mem); 1328 XPUSH_CLOBJ (stash_image2d, mem);
1095 1329
1096void 1330void
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) 1331image3d (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: 1332 PPCODE:
1099 STRLEN len; 1333 STRLEN len;
… …
1103 NEED_SUCCESS_ARG (cl_mem mem, CreateImage3D, (self, flags, &format, width, height, depth, row_pitch, slice_pitch, ptr, &res)); 1337 NEED_SUCCESS_ARG (cl_mem mem, CreateImage3D, (self, flags, &format, width, height, depth, row_pitch, slice_pitch, ptr, &res));
1104#else 1338#else
1105 const cl_image_desc desc = { CL_MEM_OBJECT_IMAGE3D, width, height, depth, 0, row_pitch, slice_pitch, 0, 0, 0 }; 1339 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)); 1340 NEED_SUCCESS_ARG (cl_mem mem, CreateImage, (self, flags, &format, &desc, ptr, &res));
1107#endif 1341#endif
1108 XPUSH_CLOBJ ("OpenCL::Image3D", mem); 1342 XPUSH_CLOBJ (stash_image3d, mem);
1109 1343
1110#if cl_apple_gl_sharing || cl_khr_gl_sharing 1344#if cl_apple_gl_sharing || cl_khr_gl_sharing
1111 1345
1112void 1346void
1113gl_buffer (OpenCL::Context self, cl_mem_flags flags, cl_GLuint bufobj) 1347gl_buffer (OpenCL::Context self, cl_mem_flags flags, cl_GLuint bufobj)
1114 PPCODE: 1348 PPCODE:
1115 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLBuffer, (self, flags, bufobj, &res)); 1349 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLBuffer, (self, flags, bufobj, &res));
1116 XPUSH_CLOBJ ("OpenCL::BufferObj", mem); 1350 XPUSH_CLOBJ (stash_bufferobj, mem);
1117 1351
1118void 1352void
1119gl_renderbuffer (OpenCL::Context self, cl_mem_flags flags, cl_GLuint renderbuffer) 1353gl_renderbuffer (OpenCL::Context self, cl_mem_flags flags, cl_GLuint renderbuffer)
1120 PPCODE: 1354 PPCODE:
1121 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLRenderbuffer, (self, flags, renderbuffer, &res)); 1355 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLRenderbuffer, (self, flags, renderbuffer, &res));
1122 XPUSH_CLOBJ ("OpenCL::Image2D", mem); 1356 XPUSH_CLOBJ (stash_image2d, mem);
1123 1357
1124#if CL_VERSION_1_2 1358#if CL_VERSION_1_2
1125 1359
1126void 1360void
1127gl_texture (OpenCL::Context self, cl_mem_flags flags, cl_GLenum target, cl_GLint miplevel, cl_GLuint texture) 1361gl_texture (OpenCL::Context self, cl_mem_flags flags, cl_GLenum target, cl_GLint miplevel, cl_GLuint texture)
1128 ALIAS: 1362 ALIAS:
1129 PPCODE: 1363 PPCODE:
1130 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture, (self, flags, target, miplevel, texture, &res)); 1364 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture, (self, flags, target, miplevel, texture, &res));
1131 cl_gl_object_type type; 1365 cl_gl_object_type type;
1132 NEED_SUCCESS (GetGLObjectInfo, (mem, &type, 0)); // TODO: use target instead? 1366 NEED_SUCCESS (GetGLObjectInfo, (mem, &type, 0)); // TODO: use target instead?
1133 char *klass = "OpenCL::Memory"; 1367 HV *stash = stash_memory;
1134 switch (type) 1368 switch (type)
1135 { 1369 {
1136 case CL_GL_OBJECT_TEXTURE_BUFFER: klass = "OpenCL::Image1DBuffer"; break; 1370 case CL_GL_OBJECT_TEXTURE_BUFFER: stash = stash_image1dbuffer; break;
1137 case CL_GL_OBJECT_TEXTURE1D: klass = "OpenCL::Image1D"; break; 1371 case CL_GL_OBJECT_TEXTURE1D: stash = stash_image1d; break;
1138 case CL_GL_OBJECT_TEXTURE1D_ARRAY: klass = "OpenCL::Image2DArray"; break; 1372 case CL_GL_OBJECT_TEXTURE1D_ARRAY: stash = stash_image2darray; break;
1139 case CL_GL_OBJECT_TEXTURE2D: klass = "OpenCL::Image2D"; break; 1373 case CL_GL_OBJECT_TEXTURE2D: stash = stash_image2d; break;
1140 case CL_GL_OBJECT_TEXTURE2D_ARRAY: klass = "OpenCL::Image2DArray"; break; 1374 case CL_GL_OBJECT_TEXTURE2D_ARRAY: stash = stash_image2darray; break;
1141 case CL_GL_OBJECT_TEXTURE3D: klass = "OpenCL::Image3D"; break; 1375 case CL_GL_OBJECT_TEXTURE3D: stash = stash_image3d; break;
1142 } 1376 }
1143 XPUSH_CLOBJ (klass, mem); 1377 XPUSH_CLOBJ (stash, mem);
1144 1378
1145#endif 1379#endif
1146 1380
1147void 1381void
1148gl_texture2d (OpenCL::Context self, cl_mem_flags flags, cl_GLenum target, cl_GLint miplevel, cl_GLuint texture) 1382gl_texture2d (OpenCL::Context self, cl_mem_flags flags, cl_GLenum target, cl_GLint miplevel, cl_GLuint texture)
… …
1150#if PREFER_1_1 1384#if PREFER_1_1
1151 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture2D, (self, flags, target, miplevel, texture, &res)); 1385 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture2D, (self, flags, target, miplevel, texture, &res));
1152#else 1386#else
1153 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture , (self, flags, target, miplevel, texture, &res)); 1387 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture , (self, flags, target, miplevel, texture, &res));
1154#endif 1388#endif
1155 XPUSH_CLOBJ ("OpenCL::Image2D", mem); 1389 XPUSH_CLOBJ (stash_image2d, mem);
1156 1390
1157void 1391void
1158gl_texture3d (OpenCL::Context self, cl_mem_flags flags, cl_GLenum target, cl_GLint miplevel, cl_GLuint texture) 1392gl_texture3d (OpenCL::Context self, cl_mem_flags flags, cl_GLenum target, cl_GLint miplevel, cl_GLuint texture)
1159 PPCODE: 1393 PPCODE:
1160#if PREFER_1_1 1394#if PREFER_1_1
1161 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture3D, (self, flags, target, miplevel, texture, &res)); 1395 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture3D, (self, flags, target, miplevel, texture, &res));
1162#else 1396#else
1163 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture , (self, flags, target, miplevel, texture, &res)); 1397 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture , (self, flags, target, miplevel, texture, &res));
1164#endif 1398#endif
1165 XPUSH_CLOBJ ("OpenCL::Image3D", mem); 1399 XPUSH_CLOBJ (stash_image3d, mem);
1166 1400
1167#endif 1401#endif
1168 1402
1169void 1403void
1170supported_image_formats (OpenCL::Context self, cl_mem_flags flags, cl_mem_object_type image_type) 1404supported_image_formats (OpenCL::Context self, cl_mem_flags flags, cl_mem_object_type image_type)
… …
1190 1424
1191void 1425void
1192sampler (OpenCL::Context self, cl_bool normalized_coords, cl_addressing_mode addressing_mode, cl_filter_mode filter_mode) 1426sampler (OpenCL::Context self, cl_bool normalized_coords, cl_addressing_mode addressing_mode, cl_filter_mode filter_mode)
1193 PPCODE: 1427 PPCODE:
1194 NEED_SUCCESS_ARG (cl_sampler sampler, CreateSampler, (self, normalized_coords, addressing_mode, filter_mode, &res)); 1428 NEED_SUCCESS_ARG (cl_sampler sampler, CreateSampler, (self, normalized_coords, addressing_mode, filter_mode, &res));
1195 XPUSH_CLOBJ ("OpenCL::Sampler", sampler); 1429 XPUSH_CLOBJ (stash_sampler, sampler);
1196 1430
1197void 1431void
1198program_with_source (OpenCL::Context self, SV *program) 1432program_with_source (OpenCL::Context self, SV *program)
1199 PPCODE: 1433 PPCODE:
1200 STRLEN len; 1434 STRLEN len;
1201 size_t len2; 1435 size_t len2;
1202 const char *ptr = SvPVbyte (program, len); 1436 const char *ptr = SvPVbyte (program, len);
1203 1437
1204 len2 = len; 1438 len2 = len;
1205 NEED_SUCCESS_ARG (cl_program prog, CreateProgramWithSource, (self, 1, &ptr, &len2, &res)); 1439 NEED_SUCCESS_ARG (cl_program prog, CreateProgramWithSource, (self, 1, &ptr, &len2, &res));
1206 XPUSH_CLOBJ ("OpenCL::Program", prog); 1440 XPUSH_CLOBJ (stash_program, prog);
1441
1442void
1443program_with_binary (OpenCL::Context self, SV *devices, SV *binaries)
1444 PPCODE:
1445 cl_uint device_count;
1446 cl_device_id *device_list = object_list (cv, 0, "devices", devices, "OpenCL::Device", &device_count);
1447
1448 if (!SvROK (binaries) || SvTYPE (SvRV (binaries)) != SVt_PVAV)
1449 croak ("OpenCL::Context::program_with_binary: binaries must be specified as reference to an array of strings");
1450
1451 binaries = SvRV (binaries);
1452
1453 if (device_count != av_len ((AV *)binaries) + 1)
1454 croak ("OpenCL::Context::program_with_binary: differing numbers of devices and binaries are not allowed");
1455
1456 size_t *length_list = tmpbuf (sizeof (*length_list) * device_count);
1457 const unsigned char **binary_list = tmpbuf (sizeof (*binary_list) * device_count);
1458 cl_int *status_list = tmpbuf (sizeof (*status_list) * device_count);
1459
1460 int i;
1461 for (i = 0; i < device_count; ++i)
1462 {
1463 STRLEN len;
1464 binary_list [i] = (const unsigned char *)SvPVbyte (*av_fetch ((AV *)binaries, i, 0), len);
1465 length_list [i] = len;
1466 }
1467
1468 NEED_SUCCESS_ARG (cl_program prog, CreateProgramWithBinary, (self, device_count, device_list,
1469 length_list, binary_list,
1470 GIMME_V == G_ARRAY ? status_list : 0, &res));
1471
1472 EXTEND (SP, 2);
1473 PUSH_CLOBJ (stash_program, prog);
1474
1475 if (GIMME_V == G_ARRAY)
1476 {
1477 AV *av = newAV ();
1478 PUSHs (sv_2mortal (newRV_noinc ((SV *)av)));
1479
1480 for (i = device_count; i--; )
1481 av_store (av, i, newSViv (status_list [i]));
1482 }
1483
1484#if CL_VERSION_1_2
1485
1486void
1487program_with_built_in_kernels (OpenCL::Context self, SV *devices, SV *kernel_names)
1488 PPCODE:
1489 cl_uint device_count;
1490 cl_device_id *device_list = object_list (cv, 0, "devices", devices, "OpenCL::Device", &device_count);
1491
1492 NEED_SUCCESS_ARG (cl_program prog, CreateProgramWithBuiltInKernels, (self, device_count, device_list, SvPVbyte_nolen (kernel_names), &res));
1493
1494 XPUSH_CLOBJ (stash_program, prog);
1495
1496void
1497link_program (OpenCL::Context self, SV *devices, SV *options, SV *programs, SV *notify = &PL_sv_undef)
1498 CODE:
1499 cl_uint device_count = 0;
1500 cl_device_id *device_list = 0;
1501
1502 if (SvOK (devices))
1503 device_list = object_list (cv, 1, "devices", devices, "OpenCL::Device", &device_count);
1504
1505 cl_uint program_count;
1506 cl_program *program_list = object_list (cv, 0, "programs", programs, "OpenCL::Program", &program_count);
1507
1508 void *user_data;
1509 program_callback pfn_notify = make_program_callback (notify, &user_data);
1510
1511 NEED_SUCCESS_ARG (cl_program prog, LinkProgram, (self, device_count, device_list, SvPVbyte_nolen (options),
1512 program_count, program_list, pfn_notify, user_data, &res));
1513
1514 XPUSH_CLOBJ (stash_program, prog);
1515
1516#endif
1207 1517
1208#BEGIN:context 1518#BEGIN:context
1209 1519
1210void 1520void
1211reference_count (OpenCL::Context self) 1521reference_count (OpenCL::Context self)
… …
1227 cl_device_id *value = tmpbuf (size); 1537 cl_device_id *value = tmpbuf (size);
1228 NEED_SUCCESS (GetContextInfo, (self, CL_CONTEXT_DEVICES, size, value, 0)); 1538 NEED_SUCCESS (GetContextInfo, (self, CL_CONTEXT_DEVICES, size, value, 0));
1229 int i, n = size / sizeof (*value); 1539 int i, n = size / sizeof (*value);
1230 EXTEND (SP, n); 1540 EXTEND (SP, n);
1231 for (i = 0; i < n; ++i) 1541 for (i = 0; i < n; ++i)
1232 {
1233 PUSH_CLOBJ ("OpenCL::Device", value [i]); 1542 PUSH_CLOBJ (stash_device, value [i]);
1234 }
1235 1543
1236void 1544void
1237properties (OpenCL::Context self) 1545properties (OpenCL::Context self)
1238 PPCODE: 1546 PPCODE:
1239 size_t size; 1547 size_t size;
… …
1253DESTROY (OpenCL::Queue self) 1561DESTROY (OpenCL::Queue self)
1254 CODE: 1562 CODE:
1255 clReleaseCommandQueue (self); 1563 clReleaseCommandQueue (self);
1256 1564
1257void 1565void
1258enqueue_read_buffer (OpenCL::Queue self, OpenCL::Buffer mem, cl_bool blocking, size_t offset, size_t len, SV *data, ...) 1566read_buffer (OpenCL::Queue self, OpenCL::Buffer mem, cl_bool blocking, size_t offset, size_t len, SV *data, ...)
1567 ALIAS:
1568 enqueue_read_buffer = 0
1259 PPCODE: 1569 PPCODE:
1260 cl_event ev = 0;
1261 EVENT_LIST (6, items - 6); 1570 EVENT_LIST (6);
1262 1571
1263 SvUPGRADE (data, SVt_PV); 1572 SvUPGRADE (data, SVt_PV);
1264 SvGROW (data, len); 1573 SvGROW (data, len);
1265 SvPOK_only (data); 1574 SvPOK_only (data);
1266 SvCUR_set (data, len); 1575 SvCUR_set (data, len);
1576
1577 cl_event ev = 0;
1267 NEED_SUCCESS (EnqueueReadBuffer, (self, mem, blocking, offset, len, SvPVX (data), event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1578 NEED_SUCCESS (EnqueueReadBuffer, (self, mem, blocking, offset, len, SvPVX (data), event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1268 1579
1269 if (ev) 1580 if (ev)
1270 XPUSH_CLOBJ ("OpenCL::Event", ev); 1581 XPUSH_CLOBJ (stash_event, ev);
1271 1582
1272void 1583void
1273enqueue_write_buffer (OpenCL::Queue self, OpenCL::Buffer mem, cl_bool blocking, size_t offset, SV *data, ...) 1584write_buffer (OpenCL::Queue self, OpenCL::Buffer mem, cl_bool blocking, size_t offset, SV *data, ...)
1585 ALIAS:
1586 enqueue_write_buffer = 0
1274 PPCODE: 1587 PPCODE:
1275 cl_event ev = 0; 1588 EVENT_LIST (5);
1589
1276 STRLEN len; 1590 STRLEN len;
1277 char *ptr = SvPVbyte (data, len); 1591 char *ptr = SvPVbyte (data, len);
1278 EVENT_LIST (5, items - 5);
1279 1592
1593 cl_event ev = 0;
1280 NEED_SUCCESS (EnqueueWriteBuffer, (self, mem, blocking, offset, len, ptr, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1594 NEED_SUCCESS (EnqueueWriteBuffer, (self, mem, blocking, offset, len, ptr, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1281 1595
1282 if (ev) 1596 if (ev)
1283 XPUSH_CLOBJ ("OpenCL::Event", ev); 1597 XPUSH_CLOBJ (stash_event, ev);
1284 1598
1285#if CL_VERSION_1_2 1599#if CL_VERSION_1_2
1286 1600
1287void 1601void
1288enqueue_fill_buffer (OpenCL::Queue self, OpenCL::Buffer mem, SV *data, size_t offset, size_t size, ...) 1602fill_buffer (OpenCL::Queue self, OpenCL::Buffer mem, SV *data, size_t offset, size_t size, ...)
1603 ALIAS:
1604 enqueue_fill_buffer = 0
1289 PPCODE: 1605 PPCODE:
1290 cl_event ev = 0; 1606 EVENT_LIST (5);
1607
1291 STRLEN len; 1608 STRLEN len;
1292 char *ptr = SvPVbyte (data, len); 1609 char *ptr = SvPVbyte (data, len);
1293 EVENT_LIST (5, items - 5);
1294 1610
1611 cl_event ev = 0;
1295 NEED_SUCCESS (EnqueueFillBuffer, (self, mem, ptr, len, offset, size, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1612 NEED_SUCCESS (EnqueueFillBuffer, (self, mem, ptr, len, offset, size, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1296 1613
1297 if (ev) 1614 if (ev)
1298 XPUSH_CLOBJ ("OpenCL::Event", ev); 1615 XPUSH_CLOBJ (stash_event, ev);
1299 1616
1300void 1617void
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, ...) 1618fill_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, ...)
1619 ALIAS:
1620 enqueue_fill_image = 0
1302 PPCODE: 1621 PPCODE:
1303 cl_event ev = 0; 1622 EVENT_LIST (12);
1304 STRLEN len; 1623
1305 const size_t origin [3] = { x, y, z }; 1624 const size_t origin [3] = { x, y, z };
1306 const size_t region [3] = { width, height, depth }; 1625 const size_t region [3] = { width, height, depth };
1307 EVENT_LIST (12, items - 12);
1308 1626
1309 const cl_float c_f [4] = { r, g, b, a }; 1627 const cl_float c_f [4] = { r, g, b, a };
1310 const cl_uint c_u [4] = { r, g, b, a }; 1628 const cl_uint c_u [4] = { r, g, b, a };
1311 const cl_int c_s [4] = { r, g, b, a }; 1629 const cl_int c_s [4] = { r, g, b, a };
1312 const void *c_fus [3] = { &c_f, &c_u, &c_s }; 1630 const void *c_fus [3] = { &c_f, &c_u, &c_s };
… …
1315 NEED_SUCCESS (GetImageInfo, (img, CL_IMAGE_FORMAT, sizeof (format), &format, 0)); 1633 NEED_SUCCESS (GetImageInfo, (img, CL_IMAGE_FORMAT, sizeof (format), &format, 0));
1316 assert (sizeof (fus) == CL_FLOAT + 1 - CL_SNORM_INT8); 1634 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) 1635 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."); 1636 croak ("enqueue_fill_image: image has unsupported channel type, only opencl 1.2 channel types supported.");
1319 1637
1638 cl_event ev = 0;
1320 NEED_SUCCESS (EnqueueFillImage, (self, img, c_fus [fus [format.image_channel_data_type]], 1639 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)); 1640 origin, region, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1322 1641
1323 if (ev) 1642 if (ev)
1324 XPUSH_CLOBJ ("OpenCL::Event", ev); 1643 XPUSH_CLOBJ (stash_event, ev);
1325 1644
1326#endif 1645#endif
1327 1646
1328void 1647void
1329enqueue_copy_buffer (OpenCL::Queue self, OpenCL::Buffer src, OpenCL::Buffer dst, size_t src_offset, size_t dst_offset, size_t len, ...) 1648copy_buffer (OpenCL::Queue self, OpenCL::Buffer src, OpenCL::Buffer dst, size_t src_offset, size_t dst_offset, size_t len, ...)
1649 ALIAS:
1650 enqueue_copy_buffer = 0
1330 PPCODE: 1651 PPCODE:
1652 EVENT_LIST (6);
1653
1331 cl_event ev = 0; 1654 cl_event ev = 0;
1332 EVENT_LIST (6, items - 6);
1333
1334 NEED_SUCCESS (EnqueueCopyBuffer, (self, src, dst, src_offset, dst_offset, len, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1655 NEED_SUCCESS (EnqueueCopyBuffer, (self, src, dst, src_offset, dst_offset, len, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1335 1656
1336 if (ev) 1657 if (ev)
1337 XPUSH_CLOBJ ("OpenCL::Event", ev); 1658 XPUSH_CLOBJ (stash_event, ev);
1338 1659
1339void 1660void
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, ...) 1661read_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, ...)
1662 ALIAS:
1663 enqueue_read_buffer_rect = 0
1341 PPCODE: 1664 PPCODE:
1342 cl_event ev = 0; 1665 EVENT_LIST (17);
1666
1343 const size_t buf_origin [3] = { buf_x , buf_y , buf_z }; 1667 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 }; 1668 const size_t host_origin[3] = { host_x, host_y, host_z };
1345 const size_t region[3] = { width, height, depth }; 1669 const size_t region[3] = { width, height, depth };
1346 EVENT_LIST (17, items - 17);
1347 1670
1348 if (!buf_row_pitch) 1671 if (!buf_row_pitch)
1349 buf_row_pitch = region [0]; 1672 buf_row_pitch = region [0];
1350 1673
1351 if (!buf_slice_pitch) 1674 if (!buf_slice_pitch)
… …
1361 1684
1362 SvUPGRADE (data, SVt_PV); 1685 SvUPGRADE (data, SVt_PV);
1363 SvGROW (data, len); 1686 SvGROW (data, len);
1364 SvPOK_only (data); 1687 SvPOK_only (data);
1365 SvCUR_set (data, len); 1688 SvCUR_set (data, len);
1689
1690 cl_event ev = 0;
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)); 1691 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 1692
1368 if (ev) 1693 if (ev)
1369 XPUSH_CLOBJ ("OpenCL::Event", ev); 1694 XPUSH_CLOBJ (stash_event, ev);
1370 1695
1371void 1696void
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, ...) 1697write_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, ...)
1698 ALIAS:
1699 enqueue_write_buffer_rect = 0
1373 PPCODE: 1700 PPCODE:
1374 cl_event ev = 0; 1701 EVENT_LIST (17);
1702
1375 const size_t buf_origin [3] = { buf_x , buf_y , buf_z }; 1703 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 }; 1704 const size_t host_origin[3] = { host_x, host_y, host_z };
1377 const size_t region[3] = { width, height, depth }; 1705 const size_t region[3] = { width, height, depth };
1378 STRLEN len; 1706 STRLEN len;
1379 char *ptr = SvPVbyte (data, len); 1707 char *ptr = SvPVbyte (data, len);
1380 EVENT_LIST (17, items - 17);
1381 1708
1382 if (!buf_row_pitch) 1709 if (!buf_row_pitch)
1383 buf_row_pitch = region [0]; 1710 buf_row_pitch = region [0];
1384 1711
1385 if (!buf_slice_pitch) 1712 if (!buf_slice_pitch)
… …
1394 size_t min_len = host_row_pitch * host_slice_pitch * region [2]; 1721 size_t min_len = host_row_pitch * host_slice_pitch * region [2];
1395 1722
1396 if (len < min_len) 1723 if (len < min_len)
1397 croak ("clEnqueueWriteImage: data string is shorter than what would be transferred"); 1724 croak ("clEnqueueWriteImage: data string is shorter than what would be transferred");
1398 1725
1726 cl_event ev = 0;
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)); 1727 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 1728
1401 if (ev) 1729 if (ev)
1402 XPUSH_CLOBJ ("OpenCL::Event", ev); 1730 XPUSH_CLOBJ (stash_event, ev);
1403 1731
1404void 1732void
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, ...) 1733copy_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, ...)
1734 ALIAS:
1735 enqueue_copy_buffer_rect = 0
1406 PPCODE: 1736 PPCODE:
1407 cl_event ev = 0; 1737 EVENT_LIST (16);
1738
1408 const size_t src_origin[3] = { src_x, src_y, src_z }; 1739 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 }; 1740 const size_t dst_origin[3] = { dst_x, dst_y, dst_z };
1410 const size_t region[3] = { width, height, depth }; 1741 const size_t region[3] = { width, height, depth };
1411 EVENT_LIST (16, items - 16);
1412 1742
1743 cl_event ev = 0;
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)); 1744 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 1745
1415 if (ev) 1746 if (ev)
1416 XPUSH_CLOBJ ("OpenCL::Event", ev); 1747 XPUSH_CLOBJ (stash_event, ev);
1417 1748
1418void 1749void
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, ...) 1750read_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, ...)
1751 ALIAS:
1752 enqueue_read_image = 0
1420 PPCODE: 1753 PPCODE:
1421 cl_event ev = 0; 1754 EVENT_LIST (12);
1755
1422 const size_t src_origin[3] = { src_x, src_y, src_z }; 1756 const size_t src_origin[3] = { src_x, src_y, src_z };
1423 const size_t region[3] = { width, height, depth }; 1757 const size_t region[3] = { width, height, depth };
1424 EVENT_LIST (12, items - 12);
1425 1758
1426 if (!row_pitch) 1759 if (!row_pitch)
1427 row_pitch = img_row_pitch (src); 1760 row_pitch = img_row_pitch (src);
1428 1761
1429 if (depth > 1 && !slice_pitch) 1762 if (depth > 1 && !slice_pitch)
… …
1433 1766
1434 SvUPGRADE (data, SVt_PV); 1767 SvUPGRADE (data, SVt_PV);
1435 SvGROW (data, len); 1768 SvGROW (data, len);
1436 SvPOK_only (data); 1769 SvPOK_only (data);
1437 SvCUR_set (data, len); 1770 SvCUR_set (data, len);
1771
1772 cl_event ev = 0;
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)); 1773 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 1774
1440 if (ev) 1775 if (ev)
1441 XPUSH_CLOBJ ("OpenCL::Event", ev); 1776 XPUSH_CLOBJ (stash_event, ev);
1442 1777
1443void 1778void
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, ...) 1779write_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, ...)
1780 ALIAS:
1781 enqueue_write_image = 0
1445 PPCODE: 1782 PPCODE:
1446 cl_event ev = 0; 1783 EVENT_LIST (12);
1784
1447 const size_t dst_origin[3] = { dst_x, dst_y, dst_z }; 1785 const size_t dst_origin[3] = { dst_x, dst_y, dst_z };
1448 const size_t region[3] = { width, height, depth }; 1786 const size_t region[3] = { width, height, depth };
1449 STRLEN len; 1787 STRLEN len;
1450 char *ptr = SvPVbyte (data, len); 1788 char *ptr = SvPVbyte (data, len);
1451 EVENT_LIST (12, items - 12);
1452 1789
1453 if (!row_pitch) 1790 if (!row_pitch)
1454 row_pitch = img_row_pitch (dst); 1791 row_pitch = img_row_pitch (dst);
1455 1792
1456 if (depth > 1 && !slice_pitch) 1793 if (depth > 1 && !slice_pitch)
… …
1459 size_t min_len = slice_pitch ? slice_pitch * depth : row_pitch * height; 1796 size_t min_len = slice_pitch ? slice_pitch * depth : row_pitch * height;
1460 1797
1461 if (len < min_len) 1798 if (len < min_len)
1462 croak ("clEnqueueWriteImage: data string is shorter than what would be transferred"); 1799 croak ("clEnqueueWriteImage: data string is shorter than what would be transferred");
1463 1800
1801 cl_event ev = 0;
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)); 1802 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 1803
1466 if (ev) 1804 if (ev)
1467 XPUSH_CLOBJ ("OpenCL::Event", ev); 1805 XPUSH_CLOBJ (stash_event, ev);
1468 1806
1469void 1807void
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, ...) 1808copy_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, ...)
1809 ALIAS:
1810 enqueue_copy_image = 0
1471 PPCODE: 1811 PPCODE:
1472 cl_event ev = 0; 1812 EVENT_LIST (12);
1813
1473 const size_t src_origin[3] = { src_x, src_y, src_z }; 1814 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 }; 1815 const size_t dst_origin[3] = { dst_x, dst_y, dst_z };
1475 const size_t region[3] = { width, height, depth }; 1816 const size_t region[3] = { width, height, depth };
1476 EVENT_LIST (12, items - 12);
1477 1817
1818 cl_event ev = 0;
1478 NEED_SUCCESS (EnqueueCopyImage, (self, src, dst, src_origin, dst_origin, region, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1819 NEED_SUCCESS (EnqueueCopyImage, (self, src, dst, src_origin, dst_origin, region, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1479 1820
1480 if (ev) 1821 if (ev)
1481 XPUSH_CLOBJ ("OpenCL::Event", ev); 1822 XPUSH_CLOBJ (stash_event, ev);
1482 1823
1483void 1824void
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, ...) 1825copy_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, ...)
1826 ALIAS:
1827 enqueue_copy_image_to_buffer = 0
1485 PPCODE: 1828 PPCODE:
1829 EVENT_LIST (10);
1830
1831 const size_t src_origin[3] = { src_x, src_y, src_z };
1832 const size_t region [3] = { width, height, depth };
1833
1486 cl_event ev = 0; 1834 cl_event ev = 0;
1835 NEED_SUCCESS (EnqueueCopyImageToBuffer, (self, src, dst, src_origin, region, dst_offset, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1836
1837 if (ev)
1838 XPUSH_CLOBJ (stash_event, ev);
1839
1840void
1841copy_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, ...)
1842 ALIAS:
1843 enqueue_copy_buffer_to_image = 0
1844 PPCODE:
1845 EVENT_LIST (10);
1846
1487 const size_t src_origin[3] = { src_x, src_y, src_z }; 1847 const size_t dst_origin[3] = { dst_x, dst_y, dst_z };
1848 const size_t region [3] = { width, height, depth };
1849
1850 cl_event ev = 0;
1851 NEED_SUCCESS (EnqueueCopyBufferToImage, (self, src, dst, src_offset, dst_origin, region, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1852
1853 if (ev)
1854 XPUSH_CLOBJ (stash_event, ev);
1855
1856void
1857map_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, SV *cb_ = &PL_sv_undef, ...)
1858 ALIAS:
1859 enqueue_map_buffer = 0
1860 PPCODE:
1861 EVENT_LIST (6);
1862
1863 size_t cb = SvIV (cb_);
1864
1865 if (!SvOK (cb_))
1866 {
1867 NEED_SUCCESS (GetMemObjectInfo, (buf, CL_MEM_SIZE, sizeof (cb), &cb, 0));
1868 cb -= offset;
1869 }
1870
1871 cl_event ev;
1872 NEED_SUCCESS_ARG (void *ptr, EnqueueMapBuffer, (self, buf, blocking, map_flags, offset, cb, event_list_count, event_list_ptr, &ev, &res));
1873 XPUSHs (mapped_new (stash_mappedbuffer, self, buf, map_flags, ptr, cb, ev, 0, 0, 1, cb, 1, 1));
1874
1875void
1876map_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, SV *width_ = &PL_sv_undef, SV *height_ = &PL_sv_undef, SV *depth_ = &PL_sv_undef, ...)
1877 ALIAS:
1878 enqueue_map_image = 0
1879 PPCODE:
1880 size_t width = SvIV (width_);
1881 if (!SvOK (width_))
1882 {
1883 NEED_SUCCESS (GetImageInfo, (img, CL_IMAGE_WIDTH, sizeof (width), &width, 0));
1884 width -= x;
1885 }
1886
1887 size_t height = SvIV (width_);
1888 if (!SvOK (height_))
1889 {
1890 NEED_SUCCESS (GetImageInfo, (img, CL_IMAGE_HEIGHT, sizeof (height), &height, 0));
1891 height -= y;
1892
1893 // stupid opencl returns 0 for depth, but requires 1 for 2d images
1894 if (!height)
1895 height = 1;
1896 }
1897
1898 size_t depth = SvIV (width_);
1899 if (!SvOK (depth_))
1900 {
1901 NEED_SUCCESS (GetImageInfo, (img, CL_IMAGE_DEPTH, sizeof (depth), &depth, 0));
1902 depth -= z;
1903
1904 // stupid opencl returns 0 for depth, but requires 1 for 2d images
1905 if (!depth)
1906 depth = 1;
1907 }
1908
1909 const size_t origin[3] = { x, y, z };
1488 const size_t region[3] = { width, height, depth }; 1910 const size_t region[3] = { width, height, depth };
1911 size_t row_pitch, slice_pitch;
1489 EVENT_LIST (10, items - 10); 1912 EVENT_LIST (10);
1490 1913
1914 cl_event ev;
1915 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));
1916
1917 size_t cb = slice_pitch ? slice_pitch * region [2]
1918 : row_pitch ? row_pitch * region [1]
1919 : region [0];
1920
1921 XPUSHs (mapped_new (stash_mappedimage, self, img, map_flags, ptr, cb, ev, row_pitch, slice_pitch, 0, width, height, depth));
1922
1923void
1924unmap (OpenCL::Queue self, OpenCL::Mapped mapped, ...)
1925 PPCODE:
1926 mapped_unmap (cv, ST (1), mapped, self, &ST (2), items - 2);
1927 if (GIMME_V != G_VOID)
1928 {
1929 clRetainEvent (mapped->event);
1930 XPUSH_CLOBJ (stash_event, mapped->event);
1931 }
1932
1933void
1934task (OpenCL::Queue self, OpenCL::Kernel kernel, ...)
1935 ALIAS:
1936 enqueue_task = 0
1937 PPCODE:
1938 EVENT_LIST (2);
1939
1940 cl_event ev = 0;
1491 NEED_SUCCESS (EnqueueCopyImageToBuffer, (self, src, dst, src_origin, region, dst_offset, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 1941 NEED_SUCCESS (EnqueueTask, (self, kernel, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1492 1942
1493 if (ev) 1943 if (ev)
1494 XPUSH_CLOBJ ("OpenCL::Event", ev); 1944 XPUSH_CLOBJ (stash_event, ev);
1495 1945
1496void 1946void
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, ...) 1947nd_range_kernel (OpenCL::Queue self, OpenCL::Kernel kernel, SV *global_work_offset, SV *global_work_size, SV *local_work_size = &PL_sv_undef, ...)
1948 ALIAS:
1949 enqueue_nd_range_kernel = 0
1522 PPCODE: 1950 PPCODE:
1523 cl_event ev = 0; 1951 EVENT_LIST (5);
1952
1524 size_t *gwo = 0, *gws, *lws = 0; 1953 size_t *gwo = 0, *gws, *lws = 0;
1525 int gws_len; 1954 int gws_len;
1526 size_t *lists; 1955 size_t *lists;
1527 int i; 1956 int i;
1528 EVENT_LIST (5, items - 5);
1529 1957
1530 if (!SvROK (global_work_size) || SvTYPE (SvRV (global_work_size)) != SVt_PVAV) 1958 if (!SvROK (global_work_size) || SvTYPE (SvRV (global_work_size)) != SVt_PVAV)
1531 croak ("clEnqueueNDRangeKernel: global_work_size must be an array reference"); 1959 croak ("clEnqueueNDRangeKernel: global_work_size must be an array reference");
1532 1960
1533 gws_len = AvFILLp (SvRV (global_work_size)) + 1; 1961 gws_len = AvFILLp (SvRV (global_work_size)) + 1;
1534 1962
1535 lists = tmpbuf (sizeof (size_t) * 3 * gws_len); 1963 lists = tmpbuf (sizeof (size_t) * 3 * gws_len);
1536 1964
1537 gws = lists + gws_len * 0; 1965 gws = lists + gws_len * 0;
1538 for (i = 0; i < gws_len; ++i) 1966 for (i = 0; i < gws_len; ++i)
1967 {
1539 gws [i] = SvIV (AvARRAY (SvRV (global_work_size))[i]); 1968 gws [i] = SvIV (AvARRAY (SvRV (global_work_size))[i]);
1969 // at least nvidia crashes for 0-sized work group sizes, work around
1970 if (!gws [i])
1971 croak ("clEnqueueNDRangeKernel: global_work_size[%d] is zero, must be non-zero", i);
1972 }
1540 1973
1541 if (SvOK (global_work_offset)) 1974 if (SvOK (global_work_offset))
1542 { 1975 {
1543 if (!SvROK (global_work_offset) || SvTYPE (SvRV (global_work_offset)) != SVt_PVAV) 1976 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"); 1977 croak ("clEnqueueNDRangeKernel: global_work_offset must be undef or an array reference");
… …
1552 } 1985 }
1553 1986
1554 if (SvOK (local_work_size)) 1987 if (SvOK (local_work_size))
1555 { 1988 {
1556 if ((SvOK (local_work_size) && !SvROK (local_work_size)) || SvTYPE (SvRV (local_work_size)) != SVt_PVAV) 1989 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"); 1990 croak ("clEnqueueNDRangeKernel: local_work_size must be undef or an array reference");
1558 1991
1559 if (AvFILLp (SvRV (local_work_size)) + 1 != gws_len) 1992 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"); 1993 croak ("clEnqueueNDRangeKernel: local_work_local must be undef or an array of same size as global_work_size");
1561 1994
1562 lws = lists + gws_len * 2; 1995 lws = lists + gws_len * 2;
1563 for (i = 0; i < gws_len; ++i) 1996 for (i = 0; i < gws_len; ++i)
1997 {
1564 lws [i] = SvIV (AvARRAY (SvRV (local_work_size))[i]); 1998 lws [i] = SvIV (AvARRAY (SvRV (local_work_size))[i]);
1999 // at least nvidia crashes for 0-sized work group sizes, work around
2000 if (!lws [i])
2001 croak ("clEnqueueNDRangeKernel: local_work_size[%d] is zero, must be non-zero", i);
2002 }
1565 } 2003 }
1566 2004
2005 cl_event ev = 0;
1567 NEED_SUCCESS (EnqueueNDRangeKernel, (self, kernel, gws_len, gwo, gws, lws, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 2006 NEED_SUCCESS (EnqueueNDRangeKernel, (self, kernel, gws_len, gwo, gws, lws, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1568 2007
1569 if (ev) 2008 if (ev)
1570 XPUSH_CLOBJ ("OpenCL::Event", ev); 2009 XPUSH_CLOBJ (stash_event, ev);
2010
2011#if CL_VERSION_1_2
2012
2013void
2014migrate_mem_objects (OpenCL::Queue self, SV *objects, cl_mem_migration_flags flags, ...)
2015 ALIAS:
2016 enqueue_migrate_mem_objects = 0
2017 PPCODE:
2018 EVENT_LIST (3);
2019
2020 cl_uint obj_count;
2021 cl_mem *obj_list = object_list (cv, 0, "objects", objects, "OpenCL::Memory", &obj_count);
2022
2023 cl_event ev = 0;
2024 NEED_SUCCESS (EnqueueMigrateMemObjects, (self, obj_count, obj_list, flags, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
2025
2026 if (ev)
2027 XPUSH_CLOBJ (stash_event, ev);
2028
2029#endif
1571 2030
1572#if cl_apple_gl_sharing || cl_khr_gl_sharing 2031#if cl_apple_gl_sharing || cl_khr_gl_sharing
1573 2032
1574void 2033void
1575enqueue_acquire_gl_objects (OpenCL::Queue self, SV *objects, ...) 2034acquire_gl_objects (OpenCL::Queue self, SV *objects, ...)
1576 ALIAS: 2035 ALIAS:
2036 release_gl_objects = 1
2037 enqueue_acquire_gl_objects = 0
1577 enqueue_release_gl_objects = 1 2038 enqueue_release_gl_objects = 1
1578 PPCODE: 2039 PPCODE:
1579 if (!SvROK (objects) || SvTYPE (SvRV (objects)) != SVt_PVAV) 2040 EVENT_LIST (2);
1580 croak ("OpenCL::Queue::enqueue_acquire/release_gl_objects argument 'objects' must be an arrayref with memory objects, in call"); 2041
2042 cl_uint obj_count;
2043 cl_mem *obj_list = object_list (cv, 0, "objects", objects, "OpenCL::Memory", &obj_count);
1581 2044
1582 cl_event ev = 0; 2045 cl_event ev = 0;
1583 EVENT_LIST (2, items - 2);
1584 AV *av = (AV *)SvRV (objects);
1585 cl_uint num_objects = av_len (av) + 1;
1586 cl_mem *object_list = tmpbuf (sizeof (cl_mem) * num_objects);
1587 int i;
1588
1589 for (i = num_objects; i--; )
1590 object_list [i] = SvCLOBJ ("OpenCL::Queue::enqueue_acquire/release_gl_objects", "objects", *av_fetch (av, i, 0), "OpenCL::Memory");
1591 2046
1592 if (ix) 2047 if (ix)
1593 NEED_SUCCESS (EnqueueReleaseGLObjects, (self, num_objects, object_list, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 2048 NEED_SUCCESS (EnqueueReleaseGLObjects, (self, obj_count, obj_list, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1594 else 2049 else
1595 NEED_SUCCESS (EnqueueAcquireGLObjects, (self, num_objects, object_list, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 2050 NEED_SUCCESS (EnqueueAcquireGLObjects, (self, obj_count, obj_list, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1596 2051
1597 if (ev) 2052 if (ev)
1598 XPUSH_CLOBJ ("OpenCL::Event", ev); 2053 XPUSH_CLOBJ (stash_event, ev);
1599 2054
1600#endif 2055#endif
1601 2056
1602void 2057void
1603enqueue_wait_for_events (OpenCL::Queue self, ...) 2058wait_for_events (OpenCL::Queue self, ...)
2059 ALIAS:
2060 enqueue_wait_for_events = 0
1604 CODE: 2061 CODE:
1605 EVENT_LIST (1, items - 1); 2062 EVENT_LIST (1);
1606#if PREFER_1_1 2063#if PREFER_1_1
1607 NEED_SUCCESS (EnqueueWaitForEvents, (self, event_list_count, event_list_ptr)); 2064 NEED_SUCCESS (EnqueueWaitForEvents, (self, event_list_count, event_list_ptr));
1608#else 2065#else
1609 NEED_SUCCESS (EnqueueBarrierWithWaitList, (self, event_list_count, event_list_ptr, 0)); 2066 NEED_SUCCESS (EnqueueBarrierWithWaitList, (self, event_list_count, event_list_ptr, 0));
1610#endif 2067#endif
1611 2068
1612void 2069void
1613enqueue_marker (OpenCL::Queue self, ...) 2070marker (OpenCL::Queue self, ...)
2071 ALIAS:
2072 enqueue_marker = 0
1614 PPCODE: 2073 PPCODE:
2074 EVENT_LIST (1);
1615 cl_event ev = 0; 2075 cl_event ev = 0;
1616 EVENT_LIST (1, items - 1);
1617#if PREFER_1_1 2076#if PREFER_1_1
1618 if (!event_list_count) 2077 if (!event_list_count)
1619 NEED_SUCCESS (EnqueueMarker, (self, GIMME_V != G_VOID ? &ev : 0)); 2078 NEED_SUCCESS (EnqueueMarker, (self, GIMME_V != G_VOID ? &ev : 0));
1620 else 2079 else
1621#if CL_VERSION_1_2 2080#if CL_VERSION_1_2
… …
1628#endif 2087#endif
1629#else 2088#else
1630 NEED_SUCCESS (EnqueueMarkerWithWaitList, (self, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 2089 NEED_SUCCESS (EnqueueMarkerWithWaitList, (self, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1631#endif 2090#endif
1632 if (ev) 2091 if (ev)
1633 XPUSH_CLOBJ ("OpenCL::Event", ev); 2092 XPUSH_CLOBJ (stash_event, ev);
1634 2093
1635void 2094void
1636enqueue_barrier (OpenCL::Queue self, ...) 2095barrier (OpenCL::Queue self, ...)
2096 ALIAS:
2097 enqueue_barrier = 0
1637 PPCODE: 2098 PPCODE:
2099 EVENT_LIST (1);
1638 cl_event ev = 0; 2100 cl_event ev = 0;
1639 EVENT_LIST (1, items - 1);
1640#if PREFER_1_1 2101#if PREFER_1_1
1641 if (!event_list_count && GIMME_V == G_VOID) 2102 if (!event_list_count && GIMME_V == G_VOID)
1642 NEED_SUCCESS (EnqueueBarrier, (self)); 2103 NEED_SUCCESS (EnqueueBarrier, (self));
1643 else 2104 else
1644#if CL_VERSION_1_2 2105#if CL_VERSION_1_2
… …
1654#endif 2115#endif
1655#else 2116#else
1656 NEED_SUCCESS (EnqueueBarrierWithWaitList, (self, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0)); 2117 NEED_SUCCESS (EnqueueBarrierWithWaitList, (self, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1657#endif 2118#endif
1658 if (ev) 2119 if (ev)
1659 XPUSH_CLOBJ ("OpenCL::Event", ev); 2120 XPUSH_CLOBJ (stash_event, ev);
1660 2121
1661void 2122void
1662flush (OpenCL::Queue self) 2123flush (OpenCL::Queue self)
1663 CODE: 2124 CODE:
1664 NEED_SUCCESS (Flush, (self)); 2125 NEED_SUCCESS (Flush, (self));
… …
1680 PPCODE: 2141 PPCODE:
1681 cl_context value [1]; 2142 cl_context value [1];
1682 NEED_SUCCESS (GetCommandQueueInfo, (self, CL_QUEUE_CONTEXT, sizeof (value), value, 0)); 2143 NEED_SUCCESS (GetCommandQueueInfo, (self, CL_QUEUE_CONTEXT, sizeof (value), value, 0));
1683 EXTEND (SP, 1); 2144 EXTEND (SP, 1);
1684 const int i = 0; 2145 const int i = 0;
1685 {
1686 NEED_SUCCESS (RetainContext, (value [i])); 2146 NEED_SUCCESS (RetainContext, (value [i]));
1687 PUSH_CLOBJ ("OpenCL::Context", value [i]); 2147 PUSH_CLOBJ (stash_context, value [i]);
1688 }
1689 2148
1690void 2149void
1691device (OpenCL::Queue self) 2150device (OpenCL::Queue self)
1692 PPCODE: 2151 PPCODE:
1693 cl_device_id value [1]; 2152 cl_device_id value [1];
1694 NEED_SUCCESS (GetCommandQueueInfo, (self, CL_QUEUE_DEVICE, sizeof (value), value, 0)); 2153 NEED_SUCCESS (GetCommandQueueInfo, (self, CL_QUEUE_DEVICE, sizeof (value), value, 0));
1695 EXTEND (SP, 1); 2154 EXTEND (SP, 1);
1696 const int i = 0; 2155 const int i = 0;
1697 {
1698 PUSH_CLOBJ ("OpenCL::Device", value [i]); 2156 PUSH_CLOBJ (stash_device, value [i]);
1699 }
1700 2157
1701void 2158void
1702reference_count (OpenCL::Queue self) 2159reference_count (OpenCL::Queue self)
1703 PPCODE: 2160 PPCODE:
1704 cl_uint value [1]; 2161 cl_uint value [1];
… …
1708 PUSHs (sv_2mortal (newSVuv (value [i]))); 2165 PUSHs (sv_2mortal (newSVuv (value [i])));
1709 2166
1710void 2167void
1711properties (OpenCL::Queue self) 2168properties (OpenCL::Queue self)
1712 PPCODE: 2169 PPCODE:
1713 cl_command_queue_properties value [1]; 2170 cl_ulong value [1];
1714 NEED_SUCCESS (GetCommandQueueInfo, (self, CL_QUEUE_PROPERTIES, sizeof (value), value, 0)); 2171 NEED_SUCCESS (GetCommandQueueInfo, (self, CL_QUEUE_PROPERTIES, sizeof (value), value, 0));
1715 EXTEND (SP, 1); 2172 EXTEND (SP, 1);
1716 const int i = 0; 2173 const int i = 0;
1717 PUSHs (sv_2mortal (newSViv (value [i]))); 2174 PUSHs (sv_2mortal (newSVuv (value [i])));
1718 2175
1719#END:command_queue 2176#END:command_queue
1720 2177
1721MODULE = OpenCL PACKAGE = OpenCL::Memory 2178MODULE = OpenCL PACKAGE = OpenCL::Memory
1722 2179
… …
1728void 2185void
1729info (OpenCL::Memory self, cl_mem_info name) 2186info (OpenCL::Memory self, cl_mem_info name)
1730 PPCODE: 2187 PPCODE:
1731 INFO (MemObject) 2188 INFO (MemObject)
1732 2189
2190void
2191destructor_callback (OpenCL::Memory self, SV *notify)
2192 PPCODE:
2193 clSetMemObjectDestructorCallback (self, eq_destructor_notify, SvREFCNT_inc (s_get_cv (notify)));
2194
1733#BEGIN:mem 2195#BEGIN:mem
1734 2196
1735void 2197void
1736type (OpenCL::Memory self) 2198type (OpenCL::Memory self)
2199 ALIAS:
2200 type = CL_MEM_TYPE
2201 map_count = CL_MEM_MAP_COUNT
2202 reference_count = CL_MEM_REFERENCE_COUNT
1737 PPCODE: 2203 PPCODE:
1738 cl_mem_object_type value [1]; 2204 cl_uint value [1];
1739 NEED_SUCCESS (GetMemObjectInfo, (self, CL_MEM_TYPE, sizeof (value), value, 0)); 2205 NEED_SUCCESS (GetMemObjectInfo, (self, ix, sizeof (value), value, 0));
1740 EXTEND (SP, 1); 2206 EXTEND (SP, 1);
1741 const int i = 0; 2207 const int i = 0;
1742 PUSHs (sv_2mortal (newSViv (value [i]))); 2208 PUSHs (sv_2mortal (newSVuv (value [i])));
1743 2209
1744void 2210void
1745flags (OpenCL::Memory self) 2211flags (OpenCL::Memory self)
1746 PPCODE: 2212 PPCODE:
1747 cl_mem_flags value [1]; 2213 cl_ulong value [1];
1748 NEED_SUCCESS (GetMemObjectInfo, (self, CL_MEM_FLAGS, sizeof (value), value, 0)); 2214 NEED_SUCCESS (GetMemObjectInfo, (self, CL_MEM_FLAGS, sizeof (value), value, 0));
1749 EXTEND (SP, 1); 2215 EXTEND (SP, 1);
1750 const int i = 0; 2216 const int i = 0;
1751 PUSHs (sv_2mortal (newSViv (value [i]))); 2217 PUSHs (sv_2mortal (newSVuv (value [i])));
1752 2218
1753void 2219void
1754size (OpenCL::Memory self) 2220size (OpenCL::Memory self)
1755 ALIAS: 2221 ALIAS:
1756 size = CL_MEM_SIZE 2222 size = CL_MEM_SIZE
… …
1770 EXTEND (SP, 1); 2236 EXTEND (SP, 1);
1771 const int i = 0; 2237 const int i = 0;
1772 PUSHs (sv_2mortal (newSVuv ((IV)(intptr_t)value [i]))); 2238 PUSHs (sv_2mortal (newSVuv ((IV)(intptr_t)value [i])));
1773 2239
1774void 2240void
1775map_count (OpenCL::Memory self)
1776 ALIAS:
1777 map_count = CL_MEM_MAP_COUNT
1778 reference_count = CL_MEM_REFERENCE_COUNT
1779 PPCODE:
1780 cl_uint value [1];
1781 NEED_SUCCESS (GetMemObjectInfo, (self, ix, sizeof (value), value, 0));
1782 EXTEND (SP, 1);
1783 const int i = 0;
1784 PUSHs (sv_2mortal (newSVuv (value [i])));
1785
1786void
1787context (OpenCL::Memory self) 2241context (OpenCL::Memory self)
1788 PPCODE: 2242 PPCODE:
1789 cl_context value [1]; 2243 cl_context value [1];
1790 NEED_SUCCESS (GetMemObjectInfo, (self, CL_MEM_CONTEXT, sizeof (value), value, 0)); 2244 NEED_SUCCESS (GetMemObjectInfo, (self, CL_MEM_CONTEXT, sizeof (value), value, 0));
1791 EXTEND (SP, 1); 2245 EXTEND (SP, 1);
1792 const int i = 0; 2246 const int i = 0;
1793 {
1794 NEED_SUCCESS (RetainContext, (value [i])); 2247 NEED_SUCCESS (RetainContext, (value [i]));
1795 PUSH_CLOBJ ("OpenCL::Context", value [i]); 2248 PUSH_CLOBJ (stash_context, value [i]);
1796 }
1797 2249
1798void 2250void
1799associated_memobject (OpenCL::Memory self) 2251associated_memobject (OpenCL::Memory self)
1800 PPCODE: 2252 PPCODE:
1801 cl_mem value [1]; 2253 cl_mem value [1];
1802 NEED_SUCCESS (GetMemObjectInfo, (self, CL_MEM_ASSOCIATED_MEMOBJECT, sizeof (value), value, 0)); 2254 NEED_SUCCESS (GetMemObjectInfo, (self, CL_MEM_ASSOCIATED_MEMOBJECT, sizeof (value), value, 0));
1803 EXTEND (SP, 1); 2255 EXTEND (SP, 1);
1804 const int i = 0; 2256 const int i = 0;
1805 {
1806 NEED_SUCCESS (RetainMemObject, (value [i])); 2257 NEED_SUCCESS (RetainMemObject, (value [i]));
1807 PUSH_CLOBJ ("OpenCL::Memory", value [i]); 2258 PUSH_CLOBJ (stash_memory, value [i]);
1808 }
1809 2259
1810#END:mem 2260#END:mem
1811 2261
1812#if cl_apple_gl_sharing || cl_khr_gl_sharing 2262#if cl_apple_gl_sharing || cl_khr_gl_sharing
1813 2263
… …
1832 croak ("clCreateSubBuffer: cannot use/copy/alloc host ptr, doesn't make sense, check your flags!"); 2282 croak ("clCreateSubBuffer: cannot use/copy/alloc host ptr, doesn't make sense, check your flags!");
1833 2283
1834 cl_buffer_region crdata = { origin, size }; 2284 cl_buffer_region crdata = { origin, size };
1835 2285
1836 NEED_SUCCESS_ARG (cl_mem mem, CreateSubBuffer, (self, flags, CL_BUFFER_CREATE_TYPE_REGION, &crdata, &res)); 2286 NEED_SUCCESS_ARG (cl_mem mem, CreateSubBuffer, (self, flags, CL_BUFFER_CREATE_TYPE_REGION, &crdata, &res));
1837 XPUSH_CLOBJ ("OpenCL::Buffer", mem); 2287 XPUSH_CLOBJ (stash_buffer, mem);
1838 2288
1839MODULE = OpenCL PACKAGE = OpenCL::Image 2289MODULE = OpenCL PACKAGE = OpenCL::Image
1840 2290
1841void 2291void
1842image_info (OpenCL::Image self, cl_image_info name) 2292image_info (OpenCL::Image self, cl_image_info name)
… …
1912 2362
1913#BEGIN:sampler 2363#BEGIN:sampler
1914 2364
1915void 2365void
1916reference_count (OpenCL::Sampler self) 2366reference_count (OpenCL::Sampler self)
2367 ALIAS:
2368 reference_count = CL_SAMPLER_REFERENCE_COUNT
2369 normalized_coords = CL_SAMPLER_NORMALIZED_COORDS
2370 addressing_mode = CL_SAMPLER_ADDRESSING_MODE
1917 PPCODE: 2371 PPCODE:
1918 cl_uint value [1]; 2372 cl_uint value [1];
1919 NEED_SUCCESS (GetSamplerInfo, (self, CL_SAMPLER_REFERENCE_COUNT, sizeof (value), value, 0)); 2373 NEED_SUCCESS (GetSamplerInfo, (self, ix, sizeof (value), value, 0));
1920 EXTEND (SP, 1); 2374 EXTEND (SP, 1);
1921 const int i = 0; 2375 const int i = 0;
1922 PUSHs (sv_2mortal (newSVuv (value [i]))); 2376 PUSHs (sv_2mortal (newSVuv (value [i])));
1923 2377
1924void 2378void
… …
1926 PPCODE: 2380 PPCODE:
1927 cl_context value [1]; 2381 cl_context value [1];
1928 NEED_SUCCESS (GetSamplerInfo, (self, CL_SAMPLER_CONTEXT, sizeof (value), value, 0)); 2382 NEED_SUCCESS (GetSamplerInfo, (self, CL_SAMPLER_CONTEXT, sizeof (value), value, 0));
1929 EXTEND (SP, 1); 2383 EXTEND (SP, 1);
1930 const int i = 0; 2384 const int i = 0;
1931 {
1932 NEED_SUCCESS (RetainContext, (value [i])); 2385 NEED_SUCCESS (RetainContext, (value [i]));
1933 PUSH_CLOBJ ("OpenCL::Context", value [i]); 2386 PUSH_CLOBJ (stash_context, value [i]);
1934 }
1935
1936void
1937normalized_coords (OpenCL::Sampler self)
1938 PPCODE:
1939 cl_addressing_mode value [1];
1940 NEED_SUCCESS (GetSamplerInfo, (self, CL_SAMPLER_NORMALIZED_COORDS, sizeof (value), value, 0));
1941 EXTEND (SP, 1);
1942 const int i = 0;
1943 PUSHs (sv_2mortal (newSViv (value [i])));
1944
1945void
1946addressing_mode (OpenCL::Sampler self)
1947 PPCODE:
1948 cl_filter_mode value [1];
1949 NEED_SUCCESS (GetSamplerInfo, (self, CL_SAMPLER_ADDRESSING_MODE, sizeof (value), value, 0));
1950 EXTEND (SP, 1);
1951 const int i = 0;
1952 PUSHs (sv_2mortal (newSViv (value [i])));
1953 2387
1954void 2388void
1955filter_mode (OpenCL::Sampler self) 2389filter_mode (OpenCL::Sampler self)
1956 PPCODE: 2390 PPCODE:
1957 cl_bool value [1]; 2391 cl_uint value [1];
1958 NEED_SUCCESS (GetSamplerInfo, (self, CL_SAMPLER_FILTER_MODE, sizeof (value), value, 0)); 2392 NEED_SUCCESS (GetSamplerInfo, (self, CL_SAMPLER_FILTER_MODE, sizeof (value), value, 0));
1959 EXTEND (SP, 1); 2393 EXTEND (SP, 1);
1960 const int i = 0; 2394 const int i = 0;
1961 PUSHs (sv_2mortal (value [i] ? &PL_sv_yes : &PL_sv_no)); 2395 PUSHs (sv_2mortal (value [i] ? &PL_sv_yes : &PL_sv_no));
1962 2396
… …
1972void 2406void
1973build (OpenCL::Program self, SV *devices = &PL_sv_undef, SV *options = &PL_sv_undef, SV *notify = &PL_sv_undef) 2407build (OpenCL::Program self, SV *devices = &PL_sv_undef, SV *options = &PL_sv_undef, SV *notify = &PL_sv_undef)
1974 ALIAS: 2408 ALIAS:
1975 build_async = 1 2409 build_async = 1
1976 CODE: 2410 CODE:
1977 void (CL_CALLBACK *pfn_notify)(cl_program program, void *user_data) = 0;
1978 void *user_data = 0;
1979 cl_uint num_devices = 0; 2411 cl_uint device_count = 0;
1980 cl_device_id *device_list = 0; 2412 cl_device_id *device_list = 0;
1981 2413
1982 if (SvOK (devices)) 2414 if (SvOK (devices))
1983 { 2415 device_list = object_list (cv, 1, "devices", devices, "OpenCL::Device", &device_count);
2416
2417 void *user_data;
2418 program_callback pfn_notify = make_program_callback (notify, &user_data);
2419
2420 if (ix)
2421 build_program_async (self, device_count, device_list, SvPVbyte_nolen (options), user_data);
2422 else
2423 NEED_SUCCESS (BuildProgram, (self, device_count, device_list, SvPVbyte_nolen (options), pfn_notify, user_data));
2424
2425#if CL_VERSION_1_2
2426
2427void
2428compile (OpenCL::Program self, SV *devices, SV *options = &PL_sv_undef, SV *headers = &PL_sv_undef, SV *notify = &PL_sv_undef)
2429 CODE:
2430 cl_uint device_count = 0;
2431 cl_device_id *device_list = 0;
2432
2433 if (SvOK (devices))
2434 device_list = object_list (cv, 1, "devices", devices, "OpenCL::Device", &device_count);
2435
2436 cl_uint header_count = 0;
2437 cl_program *header_list = 0;
2438 const char **header_name = 0;
2439
2440 if (SvOK (headers))
2441 {
1984 if (!SvROK (devices) || SvTYPE (SvRV (devices)) != SVt_PVAV) 2442 if (!SvROK (devices) || SvTYPE (SvRV (devices)) != SVt_PVHV)
1985 croak ("clProgramBuild: devices must be undef or an array of OpenCL::Device objects."); 2443 croak ("clCompileProgram: headers must be undef or a hashref of name => OpenCL::Program pairs");
1986 2444
1987 AV *av = (AV *)SvRV (devices); 2445 HV *hv = (HV *)SvRV (devices);
1988 num_devices = av_len (av) + 1;
1989 2446
1990 if (num_devices) 2447 header_count = hv_iterinit (hv);
2448 header_list = tmpbuf (sizeof (*header_list) * header_count);
2449 header_name = tmpbuf (sizeof (*header_name) * header_count);
2450
2451 HE *he;
2452 int i = 0;
2453 while (he = hv_iternext (hv))
1991 { 2454 {
1992 device_list = tmpbuf (sizeof (*device_list) * num_devices); 2455 header_name [i] = SvPVbyte_nolen (HeSVKEY_force (he));
2456 header_list [i] = SvCLOBJ (cv, "headers", HeVAL (he), "OpenCL::Program");
1993 int count; 2457 ++i;
1994 for (count = 0; count < num_devices; ++count)
1995 device_list [count] = SvCLOBJ ("clBuildProgram", "devices", *av_fetch (av, count, 1), "OpenCL::Device");
1996 } 2458 }
1997 } 2459 }
1998 2460
1999 if (SvOK (notify)) 2461 void *user_data;
2000 { 2462 program_callback pfn_notify = make_program_callback (notify, &user_data);
2001 NEED_SUCCESS (RetainProgram, (self));
2002 pfn_notify = eq_program_notify;
2003 user_data = SvREFCNT_inc (s_get_cv (notify));
2004 }
2005 2463
2006 if (ix)
2007 build_program_async (self, num_devices, device_list, SvPVbyte_nolen (options), user_data);
2008 else
2009 NEED_SUCCESS (BuildProgram, (self, num_devices, device_list, SvPVbyte_nolen (options), pfn_notify, user_data)); 2464 NEED_SUCCESS (CompileProgram, (self, device_count, device_list, SvPVbyte_nolen (options),
2465 header_count, header_list, header_name, pfn_notify, user_data));
2466
2467#endif
2010 2468
2011void 2469void
2012build_info (OpenCL::Program self, OpenCL::Device device, cl_program_build_info name) 2470build_info (OpenCL::Program self, OpenCL::Device device, cl_program_build_info name)
2013 PPCODE: 2471 PPCODE:
2014 size_t size; 2472 size_t size;
… …
2023#BEGIN:program_build 2481#BEGIN:program_build
2024 2482
2025void 2483void
2026build_status (OpenCL::Program self, OpenCL::Device device) 2484build_status (OpenCL::Program self, OpenCL::Device device)
2027 PPCODE: 2485 PPCODE:
2028 cl_build_status value [1]; 2486 cl_int value [1];
2029 NEED_SUCCESS (GetProgramBuildInfo, (self, device, CL_PROGRAM_BUILD_STATUS, sizeof (value), value, 0)); 2487 NEED_SUCCESS (GetProgramBuildInfo, (self, device, CL_PROGRAM_BUILD_STATUS, sizeof (value), value, 0));
2030 EXTEND (SP, 1); 2488 EXTEND (SP, 1);
2031 const int i = 0; 2489 const int i = 0;
2032 PUSHs (sv_2mortal (newSViv (value [i]))); 2490 PUSHs (sv_2mortal (newSViv (value [i])));
2033 2491
… …
2043 NEED_SUCCESS (GetProgramBuildInfo, (self, device, ix, size, value, 0)); 2501 NEED_SUCCESS (GetProgramBuildInfo, (self, device, ix, size, value, 0));
2044 EXTEND (SP, 1); 2502 EXTEND (SP, 1);
2045 const int i = 0; 2503 const int i = 0;
2046 PUSHs (sv_2mortal (newSVpv (value, 0))); 2504 PUSHs (sv_2mortal (newSVpv (value, 0)));
2047 2505
2506void
2507binary_type (OpenCL::Program self, OpenCL::Device device)
2508 PPCODE:
2509 cl_uint value [1];
2510 NEED_SUCCESS (GetProgramBuildInfo, (self, device, CL_PROGRAM_BINARY_TYPE, sizeof (value), value, 0));
2511 EXTEND (SP, 1);
2512 const int i = 0;
2513 PUSHs (sv_2mortal (newSVuv ((UV)value [i])));
2514
2048#END:program_build 2515#END:program_build
2049 2516
2050void 2517void
2051kernel (OpenCL::Program program, SV *function) 2518kernel (OpenCL::Program program, SV *function)
2052 PPCODE: 2519 PPCODE:
2053 NEED_SUCCESS_ARG (cl_kernel kernel, CreateKernel, (program, SvPVbyte_nolen (function), &res)); 2520 NEED_SUCCESS_ARG (cl_kernel kernel, CreateKernel, (program, SvPVbyte_nolen (function), &res));
2054 XPUSH_CLOBJ ("OpenCL::Kernel", kernel); 2521 XPUSH_CLOBJ (stash_kernel, kernel);
2055 2522
2056void 2523void
2057kernels_in_program (OpenCL::Program program) 2524kernels_in_program (OpenCL::Program program)
2058 PPCODE: 2525 PPCODE:
2059 cl_uint num_kernels; 2526 cl_uint num_kernels;
… …
2062 NEED_SUCCESS (CreateKernelsInProgram, (program, num_kernels, kernels, 0)); 2529 NEED_SUCCESS (CreateKernelsInProgram, (program, num_kernels, kernels, 0));
2063 2530
2064 int i; 2531 int i;
2065 EXTEND (SP, num_kernels); 2532 EXTEND (SP, num_kernels);
2066 for (i = 0; i < num_kernels; ++i) 2533 for (i = 0; i < num_kernels; ++i)
2067 PUSH_CLOBJ ("OpenCL::Kernel", kernels [i]); 2534 PUSH_CLOBJ (stash_kernel, kernels [i]);
2068 2535
2069void 2536void
2070info (OpenCL::Program self, cl_program_info name) 2537info (OpenCL::Program self, cl_program_info name)
2071 PPCODE: 2538 PPCODE:
2072 INFO (Program) 2539 INFO (Program)
… …
2118 PPCODE: 2585 PPCODE:
2119 cl_context value [1]; 2586 cl_context value [1];
2120 NEED_SUCCESS (GetProgramInfo, (self, CL_PROGRAM_CONTEXT, sizeof (value), value, 0)); 2587 NEED_SUCCESS (GetProgramInfo, (self, CL_PROGRAM_CONTEXT, sizeof (value), value, 0));
2121 EXTEND (SP, 1); 2588 EXTEND (SP, 1);
2122 const int i = 0; 2589 const int i = 0;
2123 {
2124 NEED_SUCCESS (RetainContext, (value [i])); 2590 NEED_SUCCESS (RetainContext, (value [i]));
2125 PUSH_CLOBJ ("OpenCL::Context", value [i]); 2591 PUSH_CLOBJ (stash_context, value [i]);
2126 }
2127 2592
2128void 2593void
2129devices (OpenCL::Program self) 2594devices (OpenCL::Program self)
2130 PPCODE: 2595 PPCODE:
2131 size_t size; 2596 size_t size;
… …
2133 cl_device_id *value = tmpbuf (size); 2598 cl_device_id *value = tmpbuf (size);
2134 NEED_SUCCESS (GetProgramInfo, (self, CL_PROGRAM_DEVICES, size, value, 0)); 2599 NEED_SUCCESS (GetProgramInfo, (self, CL_PROGRAM_DEVICES, size, value, 0));
2135 int i, n = size / sizeof (*value); 2600 int i, n = size / sizeof (*value);
2136 EXTEND (SP, n); 2601 EXTEND (SP, n);
2137 for (i = 0; i < n; ++i) 2602 for (i = 0; i < n; ++i)
2138 {
2139 PUSH_CLOBJ ("OpenCL::Device", value [i]); 2603 PUSH_CLOBJ (stash_device, value [i]);
2140 }
2141 2604
2142void 2605void
2143source (OpenCL::Program self) 2606source (OpenCL::Program self)
2144 PPCODE: 2607 PPCODE:
2145 size_t size; 2608 size_t size;
… …
2170DESTROY (OpenCL::Kernel self) 2633DESTROY (OpenCL::Kernel self)
2171 CODE: 2634 CODE:
2172 clReleaseKernel (self); 2635 clReleaseKernel (self);
2173 2636
2174void 2637void
2638setf (OpenCL::Kernel self, const char *format, ...)
2639 CODE:
2640 int i;
2641 for (i = 2; ; ++i)
2642 {
2643 while (*format == ' ')
2644 ++format;
2645
2646 char type = *format++;
2647
2648 if (!type)
2649 break;
2650
2651 if (i >= items)
2652 croak ("OpenCL::Kernel::setf format string too long (not enough arguments)");
2653
2654 SV *sv = ST (i);
2655
2656 union
2657 {
2658 cl_char cc; cl_uchar cC; cl_short cs; cl_ushort cS;
2659 cl_int ci; cl_uint cI; cl_long cl; cl_ulong cL;
2660 cl_half ch; cl_float cf; cl_double cd;
2661 cl_mem cm;
2662 cl_sampler ca;
2663 size_t cz;
2664 cl_event ce;
2665 } arg;
2666 size_t size;
2667 int nullarg = 0;
2668
2669 switch (type)
2670 {
2671 case 'c': arg.cc = SvIV (sv); size = sizeof (arg.cc); break;
2672 case 'C': arg.cC = SvUV (sv); size = sizeof (arg.cC); break;
2673 case 's': arg.cs = SvIV (sv); size = sizeof (arg.cs); break;
2674 case 'S': arg.cS = SvUV (sv); size = sizeof (arg.cS); break;
2675 case 'i': arg.ci = SvIV (sv); size = sizeof (arg.ci); break;
2676 case 'I': arg.cI = SvUV (sv); size = sizeof (arg.cI); break;
2677 case 'l': arg.cl = SvIV (sv); size = sizeof (arg.cl); break;
2678 case 'L': arg.cL = SvUV (sv); size = sizeof (arg.cL); break;
2679
2680 case 'h': arg.ch = SvUV (sv); size = sizeof (arg.ch); break;
2681 case 'f': arg.cf = SvNV (sv); size = sizeof (arg.cf); break;
2682 case 'd': arg.cd = SvNV (sv); size = sizeof (arg.cd); break;
2683
2684 case 'z': nullarg = 1; size = SvIV (sv); break;
2685
2686 case 'm': nullarg = !SvOK (sv); arg.cm = SvCLOBJ (cv, "m", sv, "OpenCL::Memory" ); size = sizeof (arg.cm); break;
2687 case 'a': nullarg = !SvOK (sv); arg.ca = SvCLOBJ (cv, "a", sv, "OpenCL::Sampler"); size = sizeof (arg.ca); break;
2688 case 'e': nullarg = !SvOK (sv); arg.ca = SvCLOBJ (cv, "e", sv, "OpenCL::Event" ); size = sizeof (arg.ce); break;
2689
2690 default:
2691 croak ("OpenCL::Kernel::setf format character '%c' not supported", type);
2692 }
2693
2694 res = clSetKernelArg (self, i - 2, size, nullarg ? 0 : &arg);
2695 if (res)
2696 croak ("OpenCL::Kernel::setf kernel parameter '%c' (#%d): %s", type, i - 2, err2str (res));
2697 }
2698
2699 if (i != items)
2700 croak ("OpenCL::Kernel::setf format string too short (too many arguments)");
2701
2702void
2175set_char (OpenCL::Kernel self, cl_uint idx, cl_char value) 2703set_char (OpenCL::Kernel self, cl_uint idx, cl_char value)
2176 CODE: 2704 CODE:
2177 clSetKernelArg (self, idx, sizeof (value), &value); 2705 clSetKernelArg (self, idx, sizeof (value), &value);
2178 2706
2179void 2707void
… …
2227 clSetKernelArg (self, idx, sizeof (value), &value); 2755 clSetKernelArg (self, idx, sizeof (value), &value);
2228 2756
2229void 2757void
2230set_memory (OpenCL::Kernel self, cl_uint idx, OpenCL::Memory_ornull value) 2758set_memory (OpenCL::Kernel self, cl_uint idx, OpenCL::Memory_ornull value)
2231 CODE: 2759 CODE:
2232 clSetKernelArg (self, idx, sizeof (value), &value); 2760 clSetKernelArg (self, idx, sizeof (value), value ? &value : 0);
2233 2761
2234void 2762void
2235set_buffer (OpenCL::Kernel self, cl_uint idx, OpenCL::Buffer_ornull value) 2763set_buffer (OpenCL::Kernel self, cl_uint idx, OpenCL::Buffer_ornull value)
2236 CODE: 2764 CODE:
2237 clSetKernelArg (self, idx, sizeof (value), &value); 2765 clSetKernelArg (self, idx, sizeof (value), value ? &value : 0);
2238 2766
2239void 2767void
2240set_image2d (OpenCL::Kernel self, cl_uint idx, OpenCL::Image2D_ornull value) 2768set_image (OpenCL::Kernel self, cl_uint idx, OpenCL::Image_ornull value)
2241 CODE: 2769 CODE:
2242 clSetKernelArg (self, idx, sizeof (value), &value); 2770 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 2771
2249void 2772void
2250set_sampler (OpenCL::Kernel self, cl_uint idx, OpenCL::Sampler value) 2773set_sampler (OpenCL::Kernel self, cl_uint idx, OpenCL::Sampler value)
2251 CODE: 2774 CODE:
2252 clSetKernelArg (self, idx, sizeof (value), &value); 2775 clSetKernelArg (self, idx, sizeof (value), &value);
… …
2296 PPCODE: 2819 PPCODE:
2297 cl_context value [1]; 2820 cl_context value [1];
2298 NEED_SUCCESS (GetKernelInfo, (self, CL_KERNEL_CONTEXT, sizeof (value), value, 0)); 2821 NEED_SUCCESS (GetKernelInfo, (self, CL_KERNEL_CONTEXT, sizeof (value), value, 0));
2299 EXTEND (SP, 1); 2822 EXTEND (SP, 1);
2300 const int i = 0; 2823 const int i = 0;
2301 {
2302 NEED_SUCCESS (RetainContext, (value [i])); 2824 NEED_SUCCESS (RetainContext, (value [i]));
2303 PUSH_CLOBJ ("OpenCL::Context", value [i]); 2825 PUSH_CLOBJ (stash_context, value [i]);
2304 }
2305 2826
2306void 2827void
2307program (OpenCL::Kernel self) 2828program (OpenCL::Kernel self)
2308 PPCODE: 2829 PPCODE:
2309 cl_program value [1]; 2830 cl_program value [1];
2310 NEED_SUCCESS (GetKernelInfo, (self, CL_KERNEL_PROGRAM, sizeof (value), value, 0)); 2831 NEED_SUCCESS (GetKernelInfo, (self, CL_KERNEL_PROGRAM, sizeof (value), value, 0));
2311 EXTEND (SP, 1); 2832 EXTEND (SP, 1);
2312 const int i = 0; 2833 const int i = 0;
2313 {
2314 NEED_SUCCESS (RetainProgram, (value [i])); 2834 NEED_SUCCESS (RetainProgram, (value [i]));
2315 PUSH_CLOBJ ("OpenCL::Program", value [i]); 2835 PUSH_CLOBJ (stash_program, value [i]);
2316 }
2317 2836
2318#END:kernel 2837#END:kernel
2319 2838
2320void 2839void
2321work_group_info (OpenCL::Kernel self, OpenCL::Device device, cl_kernel_work_group_info name) 2840work_group_info (OpenCL::Kernel self, OpenCL::Device device, cl_kernel_work_group_info name)
… …
2367 const int i = 0; 2886 const int i = 0;
2368 PUSHs (sv_2mortal (newSVuv (value [i]))); 2887 PUSHs (sv_2mortal (newSVuv (value [i])));
2369 2888
2370#END:kernel_work_group 2889#END:kernel_work_group
2371 2890
2891#if CL_VERSION_1_2
2892
2893void
2894arg_info (OpenCL::Kernel self, cl_uint idx, cl_kernel_arg_info name)
2895 PPCODE:
2896 size_t size;
2897 NEED_SUCCESS (GetKernelArgInfo, (self, idx, name, 0, 0, &size));
2898 SV *sv = sv_2mortal (newSV (size));
2899 SvUPGRADE (sv, SVt_PV);
2900 SvPOK_only (sv);
2901 SvCUR_set (sv, size);
2902 NEED_SUCCESS (GetKernelArgInfo, (self, idx, name, size, SvPVX (sv), 0));
2903 XPUSHs (sv);
2904
2905#BEGIN:kernel_arg
2906
2907void
2908arg_address_qualifier (OpenCL::Kernel self, cl_uint idx)
2909 ALIAS:
2910 arg_address_qualifier = CL_KERNEL_ARG_ADDRESS_QUALIFIER
2911 arg_access_qualifier = CL_KERNEL_ARG_ACCESS_QUALIFIER
2912 PPCODE:
2913 cl_uint value [1];
2914 NEED_SUCCESS (GetKernelArgInfo, (self, idx, ix, sizeof (value), value, 0));
2915 EXTEND (SP, 1);
2916 const int i = 0;
2917 PUSHs (sv_2mortal (newSVuv (value [i])));
2918
2919void
2920arg_type_name (OpenCL::Kernel self, cl_uint idx)
2921 ALIAS:
2922 arg_type_name = CL_KERNEL_ARG_TYPE_NAME
2923 arg_name = CL_KERNEL_ARG_NAME
2924 PPCODE:
2925 size_t size;
2926 NEED_SUCCESS (GetKernelArgInfo, (self, idx, ix, 0, 0, &size));
2927 char *value = tmpbuf (size);
2928 NEED_SUCCESS (GetKernelArgInfo, (self, idx, ix, size, value, 0));
2929 EXTEND (SP, 1);
2930 const int i = 0;
2931 PUSHs (sv_2mortal (newSVpv (value, 0)));
2932
2933void
2934arg_type_qualifier (OpenCL::Kernel self, cl_uint idx)
2935 PPCODE:
2936 cl_ulong value [1];
2937 NEED_SUCCESS (GetKernelArgInfo, (self, idx, CL_KERNEL_ARG_TYPE_QUALIFIER, sizeof (value), value, 0));
2938 EXTEND (SP, 1);
2939 const int i = 0;
2940 PUSHs (sv_2mortal (newSVuv (value [i])));
2941
2942#END:kernel_arg
2943
2944#endif
2945
2372MODULE = OpenCL PACKAGE = OpenCL::Event 2946MODULE = OpenCL PACKAGE = OpenCL::Event
2373 2947
2374void 2948void
2375DESTROY (OpenCL::Event self) 2949DESTROY (OpenCL::Event self)
2376 CODE: 2950 CODE:
… …
2398 PPCODE: 2972 PPCODE:
2399 cl_command_queue value [1]; 2973 cl_command_queue value [1];
2400 NEED_SUCCESS (GetEventInfo, (self, CL_EVENT_COMMAND_QUEUE, sizeof (value), value, 0)); 2974 NEED_SUCCESS (GetEventInfo, (self, CL_EVENT_COMMAND_QUEUE, sizeof (value), value, 0));
2401 EXTEND (SP, 1); 2975 EXTEND (SP, 1);
2402 const int i = 0; 2976 const int i = 0;
2403 {
2404 NEED_SUCCESS (RetainCommandQueue, (value [i])); 2977 NEED_SUCCESS (RetainCommandQueue, (value [i]));
2405 PUSH_CLOBJ ("OpenCL::Queue", value [i]); 2978 PUSH_CLOBJ (stash_queue, value [i]);
2406 }
2407 2979
2408void 2980void
2409command_type (OpenCL::Event self) 2981command_type (OpenCL::Event self)
2410 PPCODE:
2411 cl_command_type value [1];
2412 NEED_SUCCESS (GetEventInfo, (self, CL_EVENT_COMMAND_TYPE, sizeof (value), value, 0));
2413 EXTEND (SP, 1);
2414 const int i = 0;
2415 PUSHs (sv_2mortal (newSVuv (value [i])));
2416
2417void
2418reference_count (OpenCL::Event self)
2419 ALIAS: 2982 ALIAS:
2983 command_type = CL_EVENT_COMMAND_TYPE
2420 reference_count = CL_EVENT_REFERENCE_COUNT 2984 reference_count = CL_EVENT_REFERENCE_COUNT
2421 command_execution_status = CL_EVENT_COMMAND_EXECUTION_STATUS 2985 command_execution_status = CL_EVENT_COMMAND_EXECUTION_STATUS
2422 PPCODE: 2986 PPCODE:
2423 cl_uint value [1]; 2987 cl_uint value [1];
2424 NEED_SUCCESS (GetEventInfo, (self, ix, sizeof (value), value, 0)); 2988 NEED_SUCCESS (GetEventInfo, (self, ix, sizeof (value), value, 0));
… …
2431 PPCODE: 2995 PPCODE:
2432 cl_context value [1]; 2996 cl_context value [1];
2433 NEED_SUCCESS (GetEventInfo, (self, CL_EVENT_CONTEXT, sizeof (value), value, 0)); 2997 NEED_SUCCESS (GetEventInfo, (self, CL_EVENT_CONTEXT, sizeof (value), value, 0));
2434 EXTEND (SP, 1); 2998 EXTEND (SP, 1);
2435 const int i = 0; 2999 const int i = 0;
2436 {
2437 NEED_SUCCESS (RetainContext, (value [i])); 3000 NEED_SUCCESS (RetainContext, (value [i]));
2438 PUSH_CLOBJ ("OpenCL::Context", value [i]); 3001 PUSH_CLOBJ (stash_context, value [i]);
2439 }
2440 3002
2441#END:event 3003#END:event
2442 3004
2443void 3005void
2444profiling_info (OpenCL::Event self, cl_profiling_info name) 3006profiling_info (OpenCL::Event self, cl_profiling_info name)
… …
2468void 3030void
2469set_status (OpenCL::UserEvent self, cl_int execution_status) 3031set_status (OpenCL::UserEvent self, cl_int execution_status)
2470 CODE: 3032 CODE:
2471 clSetUserEventStatus (self, execution_status); 3033 clSetUserEventStatus (self, execution_status);
2472 3034
3035MODULE = OpenCL PACKAGE = OpenCL::Mapped
3036
3037void
3038DESTROY (SV *self)
3039 CODE:
3040 OpenCL__Mapped mapped = SvMAPPED (self);
3041
3042 clEnqueueUnmapMemObject (mapped->queue, mapped->memobj, mapped->ptr, 1, &mapped->event, 0);
3043 mapped_detach (self, mapped);
3044
3045 clReleaseCommandQueue (mapped->queue);
3046 clReleaseEvent (mapped->event);
3047 Safefree (mapped);
3048
3049void
3050unmap (OpenCL::Mapped self, ...)
3051 CODE:
3052 mapped_unmap (cv, ST (0), self, self->queue, &ST (1), items - 1);
3053
3054bool
3055mapped (OpenCL::Mapped self)
3056 CODE:
3057 RETVAL = !!self->ptr;
3058 OUTPUT:
3059 RETVAL
3060
3061void
3062wait (OpenCL::Mapped self)
3063 PPCODE:
3064 if (self->event)
3065 NEED_SUCCESS (WaitForEvents, (1, &self->event));
3066
3067void
3068event (OpenCL::Mapped self)
3069 PPCODE:
3070 if (!self->event)
3071 XSRETURN_UNDEF;
3072
3073 clRetainEvent (self->event);
3074 XPUSH_CLOBJ (stash_event, self->event);
3075
3076#define MAPPED_OFFSET_CB offsetof (struct mapped, cb)
3077#define MAPPED_OFFSET_ROW_PITCH offsetof (struct mapped, row_pitch)
3078#define MAPPED_OFFSET_SLICE_PITCH offsetof (struct mapped, slice_pitch)
3079#define MAPPED_OFFSET_WIDTH offsetof (struct mapped, width)
3080#define MAPPED_OFFSET_HEIGHT offsetof (struct mapped, height)
3081#define MAPPED_OFFSET_DEPTH offsetof (struct mapped, depth)
3082
3083IV
3084size (OpenCL::Mapped self)
3085 ALIAS:
3086 size = MAPPED_OFFSET_CB
3087 row_pitch = MAPPED_OFFSET_ROW_PITCH
3088 slice_pitch = MAPPED_OFFSET_SLICE_PITCH
3089 width = MAPPED_OFFSET_WIDTH
3090 height = MAPPED_OFFSET_HEIGHT
3091 depth = MAPPED_OFFSET_DEPTH
3092 CODE:
3093 RETVAL = *(size_t *)((char *)self + ix);
3094 OUTPUT:
3095 RETVAL
3096
3097IV
3098ptr (OpenCL::Mapped self)
3099 CODE:
3100 RETVAL = PTR2IV (self->ptr);
3101 OUTPUT:
3102 RETVAL
3103
3104void
3105set (OpenCL::Mapped self, size_t offset, SV *data)
3106 CODE:
3107 STRLEN len;
3108 const char *ptr = SvPVbyte (data, len);
3109
3110 if (offset + len > self->cb)
3111 croak ("OpenCL::Mapped::set out of bound condition detected");
3112
3113 memcpy (offset + (char *)self->ptr, ptr, len);
3114
3115void
3116get_row (OpenCL::Mapped self, size_t count, size_t x = 0, size_t y = 0, size_t z = 0)
3117 PPCODE:
3118 if (!SvOK (ST (1)))
3119 count = self->width - x;
3120
3121 if (x + count > self->width)
3122 croak ("OpenCL::Mapped::get: x + count crosses a row boundary");
3123
3124 if (y >= self->height)
3125 croak ("OpenCL::Mapped::get: y coordinate out of bounds");
3126
3127 if (z >= self->depth)
3128 croak ("OpenCL::Mapped::get: z coordinate out of bounds");
3129
3130 size_t element = mapped_element_size (self);
3131
3132 count *= element;
3133 x *= element;
3134
3135 char *ptr = (char *)self->ptr + x + y * self->row_pitch + z * self->slice_pitch;
3136 XPUSHs (sv_2mortal (newSVpvn (ptr, count)));
3137
3138void
3139set_row (OpenCL::Mapped self, SV *data, size_t x = 0, size_t y = 0, size_t z = 0)
3140 PPCODE:
3141 STRLEN count;
3142 char *dataptr = SvPVbyte (data, count);
3143 size_t element = mapped_element_size (self);
3144
3145 x *= element;
3146
3147 if (x + count > self->width * element)
3148 croak ("OpenCL::Mapped::set: x + data size crosses a row boundary");
3149
3150 if (y >= self->height)
3151 croak ("OpenCL::Mapped::set: y coordinate out of bounds");
3152
3153 if (z >= self->depth)
3154 croak ("OpenCL::Mapped::set: z coordinate out of bounds");
3155
3156 char *ptr = (char *)self->ptr + x + y * self->row_pitch + z * self->slice_pitch;
3157 memcpy (ptr, dataptr, count);
3158
3159MODULE = OpenCL PACKAGE = OpenCL::MappedBuffer
3160
3161MODULE = OpenCL PACKAGE = OpenCL::MappedImage
3162
3163IV
3164element_size (OpenCL::Mapped self)
3165 CODE:
3166 RETVAL = mapped_element_size (self);
3167 OUTPUT:
3168 RETVAL
3169

Diff Legend

– Removed lines
+ Added lines
< Changed lines
> Changed lines