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

Comparing OpenCL/OpenCL.xs (file contents):
Revision 1.43 by root, Sat Apr 21 22:16:09 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"
10
11#define X_STACKSIZE sizeof (void *) * 512 * 1024 // 2-4mb should be enough, really
12#include "xthread.h"
13#include "schmorp.h"
4 14
5#ifdef I_DLFCN 15#ifdef I_DLFCN
6 #include <dlfcn.h> 16 #include <dlfcn.h>
7#endif 17#endif
8 18
… …
11// this. This breaks 100% of the opencl 1.1 apps, for what reason? 21// this. This breaks 100% of the opencl 1.1 apps, for what reason?
12// after all, the functions are deprecated, not removed. 22// after all, the functions are deprecated, not removed.
13// in addition, you cannot test for this in any future-proof way. 23// in addition, you cannot test for this in any future-proof way.
14// each time a new opencl version comes out, you need to make a new 24// each time a new opencl version comes out, you need to make a new
15// release. 25// release.
16#define CL_USE_DEPRECATED_OPENCL_1_1_APIS
17#define CL_USE_DEPRECATED_OPENCL_1_2_APIS /* just guessing, you stupid idiots */ 26#define CL_USE_DEPRECATED_OPENCL_1_2_APIS /* just guessing, you stupid idiots */
18 27
28#ifndef PREFER_1_1
29 #define PREFER_1_1 1
30#endif
31
32#if PREFER_1_1
33 #define CL_USE_DEPRECATED_OPENCL_1_1_APIS
34#endif
35
19#ifdef __APPLE__ 36#ifdef __APPLE__
20 #include <OpenCL/opencl.h> 37 #define CLHDR(name) <OpenCL/name>
21#else 38#else
22 #include <CL/opencl.h> 39 #define CLHDR(name) <CL/name>
23#endif 40#endif
24 41
25#if 0 // testing 42#include CLHDR(opencl.h)
26#undef CL_USE_DEPRECATED_OPENCL_1_1_APIS 43
44#if 0
27#undef CL_VERSION_1_2 45#ifndef CL_VERSION_1_2
46 #include CLHDR(cl_d3d9.h)
28#endif 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)
56#endif
57
58#ifndef CL_VERSION_1_2
59 #undef PREFER_1_1
60 #define PREFER_1_1 1
61#endif
62
63// make sure all constants we might use are actually defined
64#include "default.h"
29 65
30typedef cl_platform_id OpenCL__Platform; 66typedef cl_platform_id OpenCL__Platform;
31typedef cl_device_id OpenCL__Device; 67typedef cl_device_id OpenCL__Device;
68typedef cl_device_id OpenCL__SubDevice;
32typedef cl_context OpenCL__Context; 69typedef cl_context OpenCL__Context;
33typedef cl_command_queue OpenCL__Queue; 70typedef cl_command_queue OpenCL__Queue;
34typedef cl_mem OpenCL__Memory; 71typedef cl_mem OpenCL__Memory;
35typedef cl_mem OpenCL__Buffer; 72typedef cl_mem OpenCL__Buffer;
36typedef cl_mem OpenCL__BufferObj; 73typedef cl_mem OpenCL__BufferObj;
37typedef cl_mem OpenCL__Image; 74typedef cl_mem OpenCL__Image;
38typedef cl_mem OpenCL__Image2D;
39typedef cl_mem OpenCL__Image3D;
40typedef cl_mem OpenCL__Memory_ornull; 75typedef cl_mem OpenCL__Memory_ornull;
41typedef cl_mem OpenCL__Buffer_ornull; 76typedef cl_mem OpenCL__Buffer_ornull;
42typedef cl_mem OpenCL__Image_ornull; 77typedef cl_mem OpenCL__Image_ornull;
43typedef cl_mem OpenCL__Image2D_ornull;
44typedef cl_mem OpenCL__Image3D_ornull;
45typedef cl_sampler OpenCL__Sampler; 78typedef cl_sampler OpenCL__Sampler;
46typedef cl_program OpenCL__Program; 79typedef cl_program OpenCL__Program;
47typedef cl_kernel OpenCL__Kernel; 80typedef cl_kernel OpenCL__Kernel;
48typedef cl_event OpenCL__Event; 81typedef cl_event OpenCL__Event;
49typedef cl_event OpenCL__UserEvent; 82typedef cl_event OpenCL__UserEvent;
50 83
51typedef 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;
52 110
53/*****************************************************************************/ 111/*****************************************************************************/
54 112
55// name must include a leading underscore 113// name must include a leading underscore
56// 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
… …
96} 154}
97 155
98/*****************************************************************************/ 156/*****************************************************************************/
99 157
100/* up to two temporary buffers */ 158/* up to two temporary buffers */
101static void * 159static void * ecb_noinline
102tmpbuf (size_t size) 160tmpbuf (size_t size)
103{ 161{
104 enum { buffers = 3 }; 162 enum { buffers = 4 };
105 static int idx; 163 static int idx;
106 static void *buf [buffers]; 164 static void *buf [buffers];
107 static size_t len [buffers]; 165 static size_t len [buffers];
108 166
109 idx = (idx + 1) % buffers; 167 idx = (idx + 1) % buffers;
… …
116 } 174 }
117 175
118 return buf [idx]; 176 return buf [idx];
119} 177}
120 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
121/*****************************************************************************/ 194/*****************************************************************************/
122 195
123typedef struct 196typedef struct
124{ 197{
125 IV iv; 198 IV iv;
126 const char *name; 199 const char *name;
127 #define const_iv(name) { (IV)CL_ ## name, # name }, 200 #define const_iv(name) { (IV)CL_ ## name, # name },
128} 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;
129 209
130static const char * 210static const char *
131iv2str (IV value, const ivstr *base, int count, const char *fallback) 211iv2str (IV value, const ivstr *base, int count, const char *fallback)
132{ 212{
133 int i; 213 int i;
… …
182 if (res) \ 262 if (res) \
183 FAIL (name); 263 FAIL (name);
184 264
185/*****************************************************************************/ 265/*****************************************************************************/
186 266
187#define NEW_MORTAL_OBJ(class,ptr) sv_setref_pv (sv_newmortal (), class, ptr)
188#define XPUSH_NEW_OBJ(class,ptr) XPUSHs (NEW_MORTAL_OBJ (class, ptr))
189
190static void * 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
191SvPTROBJ (const char *func, const char *svname, SV *sv, const char *pkg) 279SvCLOBJ (CV *cv, const char *svname, SV *sv, const char *pkg)
192{ 280{
281 // sv_derived_from is quite slow :(
193 if (SvROK (sv) && sv_derived_from (sv, pkg)) 282 if (SvROK (sv) && sv_derived_from (sv, pkg))
194 return (void *)SvIV (SvRV (sv)); 283 return (void *)SvIV (SvRV (sv));
195 284
196 croak ("%s: %s is not of type %s", func, svname, pkg); 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));
197} 296}
198 297
199/*****************************************************************************/ 298/*****************************************************************************/
200 299
201static cl_context_properties * 300static cl_context_properties * ecb_noinline
202SvCONTEXTPROPERTIES (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)
203{ 302{
204 if (!sv || !SvOK (sv)) 303 if (!sv || !SvOK (sv))
205 if (extra) 304 if (extra)
206 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
207 else 306 else
… …
213 int i, len = av_len (av) + 1; 312 int i, len = av_len (av) + 1;
214 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));
215 cl_context_properties *l = p; 314 cl_context_properties *l = p;
216 315
217 if (len & 1) 316 if (len & 1)
218 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);
219 318
220 while (extracount--) 319 while (extracount--)
221 *l++ = *extra++; 320 *l++ = *extra++;
222 321
223 for (i = 0; i < len; i += 2) 322 for (i = 0; i < len; i += 2)
… …
226 SV *p_sv = *av_fetch (av, i + 1, 0); 325 SV *p_sv = *av_fetch (av, i + 1, 0);
227 cl_context_properties v = SvIV (p_sv); // code below can override 326 cl_context_properties v = SvIV (p_sv); // code below can override
228 327
229 switch (t) 328 switch (t)
230 { 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
231 case CL_GLX_DISPLAY_KHR: 335 case CL_GLX_DISPLAY_KHR:
232 if (!SvOK (p_sv)) 336 if (!SvOK (p_sv))
233 { 337 {
234 void *func = glsym ("_glXGetCurrentDisplay"); 338 void *func = glsym ("_glXGetCurrentDisplay");
235 if (func) 339 if (func)
… …
258 *l = 0; 362 *l = 0;
259 363
260 return p; 364 return p;
261 } 365 }
262 366
263 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);
368}
369
370// parse an array of CLOBJ into a void ** array in C - works only for CLOBJs whose representation
371// is a pointer (and only on well-behaved systems).
372static void * ecb_noinline
373object_list (CV *cv, int or_undef, const char *argname, SV *arg, const char *klass, cl_uint *rcount)
374{
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)
384 {
385 list = tmpbuf (sizeof (*list) * count);
386 int i;
387 for (i = 0; i < count; ++i)
388 list [i] = SvCLOBJ (cv, argname, *av_fetch (av, i, 1), klass);
389 }
390
391 if (!count && !or_undef)
392 croak ("%s: '%s' must contain at least one %s object",
393 cv_get_name (cv), argname, klass);
394
395 *rcount = count;
396 return (void *)list;
264} 397}
265 398
266/*****************************************************************************/ 399/*****************************************************************************/
400/* callback stuff */
401
402/* default context callback, log to stderr */
403static void CL_CALLBACK
404context_default_notify (const char *msg, const void *info, size_t cb, void *data)
405{
406 fprintf (stderr, "OpenCL Context Notify: %s\n", msg);
407}
408
409typedef struct
410{
411 int free_cb;
412 void (*push)(void *data1, void *data2, void *data3);
413} eq_vtbl;
414
415typedef struct eq_item
416{
417 struct eq_item *next;
418 eq_vtbl *vtbl;
419 SV *cb;
420 void *data1, *data2, *data3;
421} eq_item;
422
423static void (*eq_signal_func)(void *signal_arg, int value);
424static void *eq_signal_arg;
425static xmutex_t eq_lock = X_MUTEX_INIT;
426static eq_item *eq_head, *eq_tail;
427
428static void ecb_noinline
429eq_enq (eq_vtbl *vtbl, SV *cb, void *data1, void *data2, void *data3)
430{
431 eq_item *item = malloc (sizeof (eq_item));
432
433 item->next = 0;
434 item->vtbl = vtbl;
435 item->cb = cb;
436 item->data1 = data1;
437 item->data2 = data2;
438 item->data3 = data3;
439
440 X_LOCK (eq_lock);
441
442 *(eq_head ? &eq_tail->next : &eq_head) = item;
443 eq_tail = item;
444
445 X_UNLOCK (eq_lock);
446
447 eq_signal_func (eq_signal_arg, 0);
448}
449
450static eq_item *
451eq_dec (void)
452{
453 eq_item *res;
454
455 X_LOCK (eq_lock);
456
457 res = eq_head;
458 if (res)
459 eq_head = res->next;
460
461 X_UNLOCK (eq_lock);
462
463 return res;
464}
465
466static void
467eq_poll (void)
468{
469 eq_item *item;
470
471 while ((item = eq_dec ()))
472 {
473 ENTER;
474 SAVETMPS;
475
476 dSP;
477 PUSHMARK (SP);
478 EXTEND (SP, 2);
479
480 if (item->vtbl->free_cb)
481 sv_2mortal (item->cb);
482
483 PUTBACK;
484 item->vtbl->push (item->data1, item->data2, item->data3);
485
486 SV *cb = item->cb;
487 free (item);
488
489 call_sv (cb, G_DISCARD | G_VOID);
490
491 FREETMPS;
492 LEAVE;
493 }
494}
495
496static void
497eq_poll_interrupt (pTHX_ void *c_arg, int value)
498{
499 eq_poll ();
500}
501
502/*****************************************************************************/
503/* context notify */
504
505static void ecb_noinline
506eq_context_push (void *data1, void *data2, void *data3)
507{
508 dSP;
509 PUSHs (sv_2mortal (newSVpv (data1, 0)));
510 PUSHs (sv_2mortal (newSVpvn (data2, (STRLEN)data3)));
511 PUTBACK;
512
513 free (data1);
514 free (data2);
515}
516
517static eq_vtbl eq_context_vtbl = { 0, eq_context_push };
518
519static void CL_CALLBACK
520eq_context_notify (const char *msg, const void *pvt, size_t cb, void *user_data)
521{
522 void *pvt_copy = malloc (cb);
523 memcpy (pvt_copy, pvt, cb);
524 eq_enq (&eq_context_vtbl, user_data, strdup (msg), pvt_copy, (void *)cb);
525}
526
527#define CONTEXT_NOTIFY_CALLBACK \
528 void (CL_CALLBACK *pfn_notify)(const char *, const void *, size_t, void *) = context_default_notify; \
529 void *user_data = 0; \
530 \
531 if (SvOK (notify)) \
532 { \
533 pfn_notify = eq_context_notify; \
534 user_data = s_get_cv (notify); \
535 }
536
537static SV * ecb_noinline
538new_clobj_context (cl_context ctx, void *user_data)
539{
540 SV *sv = new_clobj (stash_context, (IV)ctx);
541
542 if (user_data)
543 sv_magicext (SvRV (sv), user_data, PERL_MAGIC_ext, 0, 0, 0);
544
545 return sv;
546}
547
548#define XPUSH_CLOBJ_CONTEXT XPUSHs (new_clobj_context (ctx, user_data));
549
550/*****************************************************************************/
551/* build/compile/link notify */
552
553static void
554eq_program_push (void *data1, void *data2, void *data3)
555{
556 dSP;
557 PUSH_CLOBJ (stash_program, data1);
558 PUTBACK;
559}
560
561static eq_vtbl eq_program_vtbl = { 1, eq_program_push };
562
563static void CL_CALLBACK
564eq_program_notify (cl_program program, void *user_data)
565{
566 clRetainProgram (program);
567
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 }
586}
587
588struct build_args
589{
590 cl_program program;
591 char *options;
592 void *user_data;
593 cl_uint num_devices;
594};
595
596X_THREAD_PROC (build_program_thread)
597{
598 struct build_args *arg = thr_arg;
599
600 clBuildProgram (arg->program, arg->num_devices, arg->num_devices ? (void *)(arg + 1) : 0, arg->options, 0, 0);
601
602 if (arg->user_data)
603 eq_program_notify (arg->program, arg->user_data);
604 else
605 clReleaseProgram (arg->program);
606
607 free (arg->options);
608 free (arg);
609
610 return 0;
611}
612
613static void
614build_program_async (cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options, void *user_data)
615{
616 struct build_args *arg = malloc (sizeof (struct build_args) + sizeof (*device_list) * num_devices);
617
618 arg->program = program;
619 arg->options = strdup (options);
620 arg->user_data = user_data;
621 arg->num_devices = num_devices;
622 memcpy (arg + 1, device_list, sizeof (*device_list) * num_devices);
623
624 xthread_t id;
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);
642}
643
644/*****************************************************************************/
645/* event objects */
646
647static void
648eq_event_push (void *data1, void *data2, void *data3)
649{
650 dSP;
651 PUSH_CLOBJ (stash_event, data1);
652 PUSHs (sv_2mortal (newSViv ((IV)data2)));
653 PUTBACK;
654}
655
656static eq_vtbl eq_event_vtbl = { 1, eq_event_push };
657
658static void CL_CALLBACK
659eq_event_notify (cl_event event, cl_int event_command_exec_status, void *user_data)
660{
661 clRetainEvent (event);
662 eq_enq (&eq_event_vtbl, user_data, (void *)event, (void *)(IV)event_command_exec_status, 0);
663}
664
665/*****************************************************************************/
666/* utilities for XS code */
267 667
268static size_t 668static size_t
269img_row_pitch (cl_mem img) 669img_row_pitch (cl_mem img)
270{ 670{
271 size_t res; 671 size_t res;
272 clGetImageInfo (img, CL_IMAGE_ROW_PITCH, sizeof (res), &res, 0); 672 clGetImageInfo (img, CL_IMAGE_ROW_PITCH, sizeof (res), &res, 0);
273 return res; 673 return res;
274} 674}
275 675
276static cl_event * 676static cl_event * ecb_noinline
277event_list (SV **items, cl_uint *rcount) 677event_list (CV *cv, SV **items, cl_uint *rcount, cl_event extra)
278{ 678{
279 cl_uint count = *rcount; 679 cl_uint count = *rcount;
280 680
281 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)
282 return 0; 685 return 0;
283 686
284 cl_event *list = tmpbuf (sizeof (cl_event) * count); 687 cl_event *list = tmpbuf (sizeof (cl_event) * (count + 1));
285 int i = 0; 688 int i = 0;
286 689
287 do 690 while (count--)
288 {
289 --count;
290 if (SvOK (items [count])) 691 if (SvOK (items [count]))
291 list [i++] = SvPTROBJ ("clEnqueue", "wait_events", items [count], "OpenCL::Event"); 692 list [i++] = SvCLOBJ (cv, "wait_events", items [count], "OpenCL::Event");
292 } 693
293 while (count); 694 if (extra)
695 list [i++] = extra;
294 696
295 *rcount = i; 697 *rcount = i;
296 698
297 return i ? list : 0; 699 return i ? list : 0;
298} 700}
299 701
300#define EVENT_LIST(items,count) \ 702#define EVENT_LIST(skip) \
301 cl_uint event_list_count = (count); \ 703 cl_uint event_list_count = items - (skip); \
302 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)
303 705
304#define INFO(class) \ 706#define INFO(class) \
305{ \ 707{ \
306 size_t size; \ 708 size_t size; \
307 NEED_SUCCESS (Get ## class ## Info, (self, name, 0, 0, &size)); \ 709 NEED_SUCCESS (Get ## class ## Info, (self, name, 0, 0, &size)); \
… …
311 SvCUR_set (sv, size); \ 713 SvCUR_set (sv, size); \
312 NEED_SUCCESS (Get ## class ## Info, (self, name, size, SvPVX (sv), 0)); \ 714 NEED_SUCCESS (Get ## class ## Info, (self, name, size, SvPVX (sv), 0)); \
313 XPUSHs (sv); \ 715 XPUSHs (sv); \
314} 716}
315 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
316MODULE = OpenCL PACKAGE = OpenCL 833MODULE = OpenCL PACKAGE = OpenCL
317 834
318PROTOTYPES: ENABLE 835PROTOTYPES: ENABLE
319 836
837void
838poll ()
839 CODE:
840 eq_poll ();
841
842void
843_eq_initialise (IV func, IV arg)
844 CODE:
845 eq_signal_func = (void (*)(void *, int))func;
846 eq_signal_arg = (void*)arg;
847
320BOOT: 848BOOT:
321{ 849{
322 HV *stash = gv_stashpv ("OpenCL", 1); 850 HV *stash = gv_stashpv ("OpenCL", 1);
851
323 static const ivstr *civ, const_iv[] = { 852 static const ivstr *civ, const_iv[] = {
324 { sizeof (cl_char ), "SIZEOF_CHAR" }, 853 { sizeof (cl_char ), "SIZEOF_CHAR" },
325 { sizeof (cl_uchar ), "SIZEOF_UCHAR" }, 854 { sizeof (cl_uchar ), "SIZEOF_UCHAR" },
326 { sizeof (cl_short ), "SIZEOF_SHORT" }, 855 { sizeof (cl_short ), "SIZEOF_SHORT" },
327 { sizeof (cl_ushort), "SIZEOF_USHORT" }, 856 { sizeof (cl_ushort), "SIZEOF_USHORT" },
… …
330 { sizeof (cl_long ), "SIZEOF_LONG" }, 859 { sizeof (cl_long ), "SIZEOF_LONG" },
331 { sizeof (cl_ulong ), "SIZEOF_ULONG" }, 860 { sizeof (cl_ulong ), "SIZEOF_ULONG" },
332 { sizeof (cl_half ), "SIZEOF_HALF" }, 861 { sizeof (cl_half ), "SIZEOF_HALF" },
333 { sizeof (cl_float ), "SIZEOF_FLOAT" }, 862 { sizeof (cl_float ), "SIZEOF_FLOAT" },
334 { sizeof (cl_double), "SIZEOF_DOUBLE" }, 863 { sizeof (cl_double), "SIZEOF_DOUBLE" },
864 { PREFER_1_1 , "PREFER_1_1" },
335#include "constiv.h" 865#include "constiv.h"
336 }; 866 };
867
337 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--)
338 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);
903
904 sv_setiv (perl_get_sv ("OpenCL::POLL_FUNC", TRUE), (IV)eq_poll_interrupt);
339} 905}
340 906
341cl_int 907cl_int
342errno () 908errno ()
343 CODE: 909 CODE:
344 RETVAL = res; 910 RETVAL = res;
345 OUTPUT: 911 OUTPUT:
346 RETVAL 912 RETVAL
347 913
348const char * 914const char *
349err2str (cl_int err) 915err2str (cl_int err = res)
350 916
351const char * 917const char *
352enum2str (cl_uint value) 918enum2str (cl_uint value)
353 919
354void 920void
… …
362 list = tmpbuf (sizeof (*list) * count); 928 list = tmpbuf (sizeof (*list) * count);
363 NEED_SUCCESS (GetPlatformIDs, (count, list, 0)); 929 NEED_SUCCESS (GetPlatformIDs, (count, list, 0));
364 930
365 EXTEND (SP, count); 931 EXTEND (SP, count);
366 for (i = 0; i < count; ++i) 932 for (i = 0; i < count; ++i)
367 PUSHs (NEW_MORTAL_OBJ ("OpenCL::Platform", list [i])); 933 PUSH_CLOBJ (stash_platform, list [i]);
368 934
369void 935void
370context_from_type (cl_context_properties *properties = 0, cl_device_type type = CL_DEVICE_TYPE_DEFAULT, FUTURE notify = 0) 936context_from_type (cl_context_properties *properties = 0, cl_device_type type = CL_DEVICE_TYPE_DEFAULT, SV *notify = &PL_sv_undef)
371 PPCODE: 937 PPCODE:
938 CONTEXT_NOTIFY_CALLBACK;
372 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));
373 XPUSH_NEW_OBJ ("OpenCL::Context", ctx); 940 XPUSH_CLOBJ_CONTEXT;
374 941
375#if 0
376
377void 942void
378context (cl_context_properties *properties = 0, FUTURE devices, FUTURE notify = 0) 943context (cl_context_properties *properties, SV *devices, SV *notify = &PL_sv_undef)
944 PPCODE:
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;
951
952void
953wait_for_events (...)
379 PPCODE: 954 CODE:
380 /* der Gipfel der Kunst */ 955 EVENT_LIST (0);
956 NEED_SUCCESS (WaitForEvents, (event_list_count, event_list_ptr));
381 957
958PROTOTYPES: DISABLE
959
960MODULE = OpenCL PACKAGE = OpenCL::Platform
961
962void
963info (OpenCL::Platform self, cl_platform_info name)
964 PPCODE:
965 INFO (Platform)
966
967void
968unload_compiler (OpenCL::Platform self)
969 CODE:
970#if CL_VERSION_1_2
971 clUnloadPlatformCompiler (self);
382#endif 972#endif
383
384void
385wait_for_events (...)
386 CODE:
387 EVENT_LIST (0, items);
388 NEED_SUCCESS (WaitForEvents, (event_list_count, event_list_ptr));
389
390PROTOTYPES: DISABLE
391
392MODULE = OpenCL PACKAGE = OpenCL::Platform
393
394void
395info (OpenCL::Platform self, cl_platform_info name)
396 PPCODE:
397 INFO (Platform)
398 973
399#BEGIN:platform 974#BEGIN:platform
400 975
401void 976void
402profile (OpenCL::Platform self) 977profile (OpenCL::Platform self)
… …
428 list = tmpbuf (sizeof (*list) * count); 1003 list = tmpbuf (sizeof (*list) * count);
429 NEED_SUCCESS (GetDeviceIDs, (self, type, count, list, 0)); 1004 NEED_SUCCESS (GetDeviceIDs, (self, type, count, list, 0));
430 1005
431 EXTEND (SP, count); 1006 EXTEND (SP, count);
432 for (i = 0; i < count; ++i) 1007 for (i = 0; i < count; ++i)
433 PUSHs (sv_setref_pv (sv_newmortal (), "OpenCL::Device", list [i])); 1008 PUSH_CLOBJ (stash_device, list [i]);
434 1009
435void 1010void
436context (OpenCL::Platform self, cl_context_properties *properties = 0, SV *devices, FUTURE notify = 0) 1011context (OpenCL::Platform self, SV *properties, SV *devices, SV *notify = &PL_sv_undef)
437 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:
438 if (!SvROK (devices) || SvTYPE (SvRV (devices)) != SVt_PVAV) 1045 if (!SvROK (properties) || SvTYPE (SvRV (properties)) != SVt_PVAV)
439 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");
440 1047
441 AV *av = (AV *)SvRV (devices); 1048 properties = SvRV (properties);
442 cl_uint num_devices = av_len (av) + 1; 1049
443 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);
1052
444 int i; 1053 int i;
1054 for (i = 0; i < count; ++i)
1055 props [i] = (cl_device_partition_property)SvIV (*av_fetch ((AV *)properties, i, 0));
445 1056
446 for (i = num_devices; i--; ) 1057 props [count] = 0;
447 device_list [i] = SvPTROBJ ("clCreateContext", "devices", *av_fetch (av, i, 0), "OpenCL::Device");
448 1058
449 NEED_SUCCESS_ARG (cl_context ctx, CreateContext, (properties, num_devices, device_list, 0, 0, &res)); 1059 cl_uint num_devices;
450 XPUSH_NEW_OBJ ("OpenCL::Context", ctx); 1060 NEED_SUCCESS (CreateSubDevices, (self, props, 0, 0, &num_devices));
1061 cl_device_id *list = tmpbuf (sizeof (*list) * num_devices);
1062 NEED_SUCCESS (CreateSubDevices, (self, props, num_devices, list, 0));
451 1063
452void 1064 EXTEND (SP, num_devices);
453context_from_type (OpenCL::Platform self, SV *properties = 0, cl_device_type type = CL_DEVICE_TYPE_DEFAULT, FUTURE notify = 0) 1065 for (i = 0; i < count; ++i)
454 PPCODE: 1066 PUSH_CLOBJ (stash_subdevice, list [i]);
455 cl_context_properties extra[] = { CL_CONTEXT_PLATFORM, (cl_context_properties)self };
456 cl_context_properties *props = SvCONTEXTPROPERTIES ("OpenCL::Platform::context_from_type", "properties", properties, extra, 2);
457 NEED_SUCCESS_ARG (cl_context ctx, CreateContextFromType, (props, type, 0, 0, &res));
458 XPUSH_NEW_OBJ ("OpenCL::Context", ctx);
459 1067
460MODULE = OpenCL PACKAGE = OpenCL::Device 1068#endif
461
462void
463info (OpenCL::Device self, cl_device_info name)
464 PPCODE:
465 INFO (Device)
466 1069
467#BEGIN:device 1070#BEGIN:device
468 1071
469void 1072void
470type (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
471 PPCODE: 1087 PPCODE:
472 cl_device_type value [1]; 1088 cl_ulong value [1];
473 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_TYPE, sizeof (value), value, 0)); 1089 NEED_SUCCESS (GetDeviceInfo, (self, ix, sizeof (value), value, 0));
474 EXTEND (SP, 1); 1090 EXTEND (SP, 1);
475 const int i = 0; 1091 const int i = 0;
476 PUSHs (sv_2mortal (newSViv (value [i]))); 1092 PUSHs (sv_2mortal (newSVuv (value [i])));
477 1093
478void 1094void
479vendor_id (OpenCL::Device self) 1095vendor_id (OpenCL::Device self)
480 ALIAS: 1096 ALIAS:
481 vendor_id = CL_DEVICE_VENDOR_ID 1097 vendor_id = CL_DEVICE_VENDOR_ID
… …
492 max_write_image_args = CL_DEVICE_MAX_WRITE_IMAGE_ARGS 1108 max_write_image_args = CL_DEVICE_MAX_WRITE_IMAGE_ARGS
493 image_support = CL_DEVICE_IMAGE_SUPPORT 1109 image_support = CL_DEVICE_IMAGE_SUPPORT
494 max_samplers = CL_DEVICE_MAX_SAMPLERS 1110 max_samplers = CL_DEVICE_MAX_SAMPLERS
495 mem_base_addr_align = CL_DEVICE_MEM_BASE_ADDR_ALIGN 1111 mem_base_addr_align = CL_DEVICE_MEM_BASE_ADDR_ALIGN
496 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
497 global_mem_cacheline_size = CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE 1114 global_mem_cacheline_size = CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE
498 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
499 preferred_vector_width_half = CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF 1117 preferred_vector_width_half = CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF
500 native_vector_width_char = CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR 1118 native_vector_width_char = CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR
501 native_vector_width_short = CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT 1119 native_vector_width_short = CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT
502 native_vector_width_int = CL_DEVICE_NATIVE_VECTOR_WIDTH_INT 1120 native_vector_width_int = CL_DEVICE_NATIVE_VECTOR_WIDTH_INT
503 native_vector_width_long = CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG 1121 native_vector_width_long = CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG
… …
541 EXTEND (SP, n); 1159 EXTEND (SP, n);
542 for (i = 0; i < n; ++i) 1160 for (i = 0; i < n; ++i)
543 PUSHs (sv_2mortal (newSVuv (value [i]))); 1161 PUSHs (sv_2mortal (newSVuv (value [i])));
544 1162
545void 1163void
546address_bits (OpenCL::Device self)
547 PPCODE:
548 cl_bitfield value [1];
549 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_ADDRESS_BITS, sizeof (value), value, 0));
550 EXTEND (SP, 1);
551 const int i = 0;
552 PUSHs (sv_2mortal (newSVuv (value [i])));
553
554void
555max_mem_alloc_size (OpenCL::Device self)
556 ALIAS:
557 max_mem_alloc_size = CL_DEVICE_MAX_MEM_ALLOC_SIZE
558 global_mem_cache_size = CL_DEVICE_GLOBAL_MEM_CACHE_SIZE
559 global_mem_size = CL_DEVICE_GLOBAL_MEM_SIZE
560 max_constant_buffer_size = CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE
561 local_mem_size = CL_DEVICE_LOCAL_MEM_SIZE
562 PPCODE:
563 cl_ulong value [1];
564 NEED_SUCCESS (GetDeviceInfo, (self, ix, sizeof (value), value, 0));
565 EXTEND (SP, 1);
566 const int i = 0;
567 PUSHs (sv_2mortal (newSVuv (value [i])));
568
569void
570single_fp_config (OpenCL::Device self)
571 ALIAS:
572 single_fp_config = CL_DEVICE_SINGLE_FP_CONFIG
573 double_fp_config = CL_DEVICE_DOUBLE_FP_CONFIG
574 half_fp_config = CL_DEVICE_HALF_FP_CONFIG
575 PPCODE:
576 cl_device_fp_config value [1];
577 NEED_SUCCESS (GetDeviceInfo, (self, ix, sizeof (value), value, 0));
578 EXTEND (SP, 1);
579 const int i = 0;
580 PUSHs (sv_2mortal (newSVuv (value [i])));
581
582void
583global_mem_cache_type (OpenCL::Device self)
584 PPCODE:
585 cl_device_mem_cache_type value [1];
586 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_GLOBAL_MEM_CACHE_TYPE, sizeof (value), value, 0));
587 EXTEND (SP, 1);
588 const int i = 0;
589 PUSHs (sv_2mortal (newSVuv (value [i])));
590
591void
592local_mem_type (OpenCL::Device self)
593 PPCODE:
594 cl_device_local_mem_type value [1];
595 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_LOCAL_MEM_TYPE, sizeof (value), value, 0));
596 EXTEND (SP, 1);
597 const int i = 0;
598 PUSHs (sv_2mortal (newSVuv (value [i])));
599
600void
601error_correction_support (OpenCL::Device self) 1164error_correction_support (OpenCL::Device self)
602 ALIAS: 1165 ALIAS:
603 error_correction_support = CL_DEVICE_ERROR_CORRECTION_SUPPORT 1166 error_correction_support = CL_DEVICE_ERROR_CORRECTION_SUPPORT
604 endian_little = CL_DEVICE_ENDIAN_LITTLE 1167 endian_little = CL_DEVICE_ENDIAN_LITTLE
605 available = CL_DEVICE_AVAILABLE 1168 available = CL_DEVICE_AVAILABLE
606 compiler_available = CL_DEVICE_COMPILER_AVAILABLE 1169 compiler_available = CL_DEVICE_COMPILER_AVAILABLE
607 host_unified_memory = CL_DEVICE_HOST_UNIFIED_MEMORY 1170 host_unified_memory = CL_DEVICE_HOST_UNIFIED_MEMORY
608 PPCODE: 1171 PPCODE:
609 cl_bool value [1]; 1172 cl_uint value [1];
610 NEED_SUCCESS (GetDeviceInfo, (self, ix, sizeof (value), value, 0)); 1173 NEED_SUCCESS (GetDeviceInfo, (self, ix, sizeof (value), value, 0));
611 EXTEND (SP, 1); 1174 EXTEND (SP, 1);
612 const int i = 0; 1175 const int i = 0;
613 PUSHs (sv_2mortal (value [i] ? &PL_sv_yes : &PL_sv_no)); 1176 PUSHs (sv_2mortal (value [i] ? &PL_sv_yes : &PL_sv_no));
614
615void
616execution_capabilities (OpenCL::Device self)
617 PPCODE:
618 cl_device_exec_capabilities value [1];
619 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_EXECUTION_CAPABILITIES, sizeof (value), value, 0));
620 EXTEND (SP, 1);
621 const int i = 0;
622 PUSHs (sv_2mortal (newSVuv (value [i])));
623
624void
625properties (OpenCL::Device self)
626 PPCODE:
627 cl_command_queue_properties value [1];
628 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_QUEUE_PROPERTIES, sizeof (value), value, 0));
629 EXTEND (SP, 1);
630 const int i = 0;
631 PUSHs (sv_2mortal (newSViv (value [i])));
632 1177
633void 1178void
634platform (OpenCL::Device self) 1179platform (OpenCL::Device self)
635 PPCODE: 1180 PPCODE:
636 cl_platform_id value [1]; 1181 cl_platform_id value [1];
637 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_PLATFORM, sizeof (value), value, 0)); 1182 NEED_SUCCESS (GetDeviceInfo, (self, CL_DEVICE_PLATFORM, sizeof (value), value, 0));
638 EXTEND (SP, 1); 1183 EXTEND (SP, 1);
639 const int i = 0; 1184 const int i = 0;
640 { 1185 PUSH_CLOBJ (stash_platform, value [i]);
641 PUSHs (NEW_MORTAL_OBJ ("OpenCL::Platform", value [i]));
642 }
643 1186
644void 1187void
645name (OpenCL::Device self) 1188name (OpenCL::Device self)
646 ALIAS: 1189 ALIAS:
647 name = CL_DEVICE_NAME 1190 name = CL_DEVICE_NAME
… …
664 PPCODE: 1207 PPCODE:
665 cl_device_id value [1]; 1208 cl_device_id value [1];
666 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));
667 EXTEND (SP, 1); 1210 EXTEND (SP, 1);
668 const int i = 0; 1211 const int i = 0;
669 { 1212 PUSH_CLOBJ (stash_device, value [i]);
670 PUSHs (NEW_MORTAL_OBJ ("OpenCL::Device", value [i]));
671 }
672 1213
673void 1214void
674partition_types_ext (OpenCL::Device self) 1215partition_types_ext (OpenCL::Device self)
675 ALIAS: 1216 ALIAS:
676 partition_types_ext = CL_DEVICE_PARTITION_TYPES_EXT 1217 partition_types_ext = CL_DEVICE_PARTITION_TYPES_EXT
… …
686 for (i = 0; i < n; ++i) 1227 for (i = 0; i < n; ++i)
687 PUSHs (sv_2mortal (newSVuv (value [i]))); 1228 PUSHs (sv_2mortal (newSVuv (value [i])));
688 1229
689#END:device 1230#END:device
690 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
691MODULE = OpenCL PACKAGE = OpenCL::Context 1243MODULE = OpenCL PACKAGE = OpenCL::Context
692 1244
693void 1245void
694DESTROY (OpenCL::Context context) 1246DESTROY (OpenCL::Context self)
695 CODE: 1247 CODE:
696 clReleaseContext (context); 1248 clReleaseContext (self);
697 1249
698void 1250void
699info (OpenCL::Context self, cl_context_info name) 1251info (OpenCL::Context self, cl_context_info name)
700 PPCODE: 1252 PPCODE:
701 INFO (Context) 1253 INFO (Context)
702 1254
703void 1255void
704queue (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)
705 PPCODE: 1257 PPCODE:
706 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));
707 XPUSH_NEW_OBJ ("OpenCL::Queue", queue); 1259 XPUSH_CLOBJ (stash_queue, queue);
708 1260
709void 1261void
710user_event (OpenCL::Context self) 1262user_event (OpenCL::Context self)
711 PPCODE: 1263 PPCODE:
712 NEED_SUCCESS_ARG (cl_event ev, CreateUserEvent, (self, &res)); 1264 NEED_SUCCESS_ARG (cl_event ev, CreateUserEvent, (self, &res));
713 XPUSH_NEW_OBJ ("OpenCL::UserEvent", ev); 1265 XPUSH_CLOBJ (stash_userevent, ev);
714 1266
715void 1267void
716buffer (OpenCL::Context self, cl_mem_flags flags, size_t len) 1268buffer (OpenCL::Context self, cl_mem_flags flags, size_t len)
717 PPCODE: 1269 PPCODE:
718 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))
719 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?");
720 1272
721 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));
722 XPUSH_NEW_OBJ ("OpenCL::BufferObj", mem); 1274 XPUSH_CLOBJ (stash_bufferobj, mem);
723 1275
724void 1276void
725buffer_sv (OpenCL::Context self, cl_mem_flags flags, SV *data) 1277buffer_sv (OpenCL::Context self, cl_mem_flags flags, SV *data)
726 PPCODE: 1278 PPCODE:
727 STRLEN len; 1279 STRLEN len;
728 char *ptr = SvOK (data) ? SvPVbyte (data, len) : 0; 1280 char *ptr = SvOK (data) ? SvPVbyte (data, len) : 0;
729 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)))
730 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?");
731 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));
732 XPUSH_NEW_OBJ ("OpenCL::BufferObj", mem); 1284 XPUSH_CLOBJ (stash_bufferobj, mem);
733 1285
734#if CL_VERSION_1_2 1286#if CL_VERSION_1_2
735 1287
736void 1288void
737image (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)
738 PPCODE: 1290 PPCODE:
739 STRLEN len; 1291 STRLEN len;
740 char *ptr = SvOK (data) ? SvPVbyte (data, len) : 0; 1292 char *ptr = SvOK (data) ? SvPVbyte (data, len) : 0;
741 const cl_image_format format = { channel_order, channel_type }; 1293 const cl_image_format format = { channel_order, channel_type };
742 const cl_image_desc desc = { 1294 const cl_image_desc desc = {
743 type, 1295 type,
744 width, height, depth, 1296 width, height, depth,
745 array_size, row_pitch, slice_pitch, 1297 array_size, row_pitch, slice_pitch,
746 num_mip_level, num_samples, 1298 num_mip_level, num_samples,
747 type == CL_MEM_OBJECT_IMAGE1D_BUFFER ? (cl_mem)SvPTROBJ ("OpenCL::Context::Image", "data", data, "OpenCL::Buffer") : 0 1299 type == CL_MEM_OBJECT_IMAGE1D_BUFFER ? (cl_mem)SvCLOBJ (cv, "data", data, "OpenCL::Buffer") : 0
748 }; 1300 };
749 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));
750 char *klass = "OpenCL::Image"; 1302 HV *stash = stash_image;
751 switch (type) 1303 switch (type)
752 { 1304 {
753 case CL_MEM_OBJECT_IMAGE1D_BUFFER: klass = "OpenCL::Image1DBuffer"; break; 1305 case CL_MEM_OBJECT_IMAGE1D_BUFFER: stash = stash_image1dbuffer; break;
754 case CL_MEM_OBJECT_IMAGE1D: klass = "OpenCL::Image1D"; break; 1306 case CL_MEM_OBJECT_IMAGE1D: stash = stash_image1d; break;
755 case CL_MEM_OBJECT_IMAGE1D_ARRAY: klass = "OpenCL::Image2DArray"; break; 1307 case CL_MEM_OBJECT_IMAGE1D_ARRAY: stash = stash_image2darray; break;
756 case CL_MEM_OBJECT_IMAGE2D: klass = "OpenCL::Image2D"; break; 1308 case CL_MEM_OBJECT_IMAGE2D: stash = stash_image2d; break;
757 case CL_MEM_OBJECT_IMAGE2D_ARRAY: klass = "OpenCL::Image2DArray"; break; 1309 case CL_MEM_OBJECT_IMAGE2D_ARRAY: stash = stash_image2darray; break;
758 case CL_MEM_OBJECT_IMAGE3D: klass = "OpenCL::Image3D"; break; 1310 case CL_MEM_OBJECT_IMAGE3D: stash = stash_image3d; break;
759 } 1311 }
760 XPUSH_NEW_OBJ (klass, mem); 1312 XPUSH_CLOBJ (stash, mem);
761 1313
762#endif 1314#endif
763 1315
764void 1316void
765image2d (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)
766 PPCODE: 1318 PPCODE:
767 STRLEN len; 1319 STRLEN len;
768 char *ptr = SvOK (data) ? SvPVbyte (data, len) : 0; 1320 char *ptr = SvOK (data) ? SvPVbyte (data, len) : 0;
769 const cl_image_format format = { channel_order, channel_type }; 1321 const cl_image_format format = { channel_order, channel_type };
770#if CL_VERSION_1_2 1322#if PREFER_1_1
1323 NEED_SUCCESS_ARG (cl_mem mem, CreateImage2D, (self, flags, &format, width, height, row_pitch, ptr, &res));
1324#else
771 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 };
772 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));
773#else
774 NEED_SUCCESS_ARG (cl_mem mem, CreateImage2D, (self, flags, &format, width, height, row_pitch, ptr, &res));
775#endif 1327#endif
776 XPUSH_NEW_OBJ ("OpenCL::Image2D", mem); 1328 XPUSH_CLOBJ (stash_image2d, mem);
777 1329
778void 1330void
779image3d (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)
780 PPCODE: 1332 PPCODE:
781 STRLEN len; 1333 STRLEN len;
782 char *ptr = SvOK (data) ? SvPVbyte (data, len) : 0; 1334 char *ptr = SvOK (data) ? SvPVbyte (data, len) : 0;
783 const cl_image_format format = { channel_order, channel_type }; 1335 const cl_image_format format = { channel_order, channel_type };
784#if CL_VERSION_1_2 1336#if PREFER_1_1
1337 NEED_SUCCESS_ARG (cl_mem mem, CreateImage3D, (self, flags, &format, width, height, depth, row_pitch, slice_pitch, ptr, &res));
1338#else
785 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 };
786 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));
787#else
788 NEED_SUCCESS_ARG (cl_mem mem, CreateImage3D, (self, flags, &format, width, height, depth, row_pitch, slice_pitch, ptr, &res));
789#endif 1341#endif
790 XPUSH_NEW_OBJ ("OpenCL::Image3D", mem); 1342 XPUSH_CLOBJ (stash_image3d, mem);
791 1343
792#if cl_apple_gl_sharing || cl_khr_gl_sharing 1344#if cl_apple_gl_sharing || cl_khr_gl_sharing
793 1345
794void 1346void
795gl_buffer (OpenCL::Context self, cl_mem_flags flags, cl_GLuint bufobj) 1347gl_buffer (OpenCL::Context self, cl_mem_flags flags, cl_GLuint bufobj)
796 PPCODE: 1348 PPCODE:
797 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLBuffer, (self, flags, bufobj, &res)); 1349 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLBuffer, (self, flags, bufobj, &res));
798 XPUSH_NEW_OBJ ("OpenCL::BufferObj", mem); 1350 XPUSH_CLOBJ (stash_bufferobj, mem);
799 1351
800void 1352void
801gl_renderbuffer (OpenCL::Context self, cl_mem_flags flags, cl_GLuint renderbuffer) 1353gl_renderbuffer (OpenCL::Context self, cl_mem_flags flags, cl_GLuint renderbuffer)
802 PPCODE: 1354 PPCODE:
803 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLRenderbuffer, (self, flags, renderbuffer, &res)); 1355 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLRenderbuffer, (self, flags, renderbuffer, &res));
804 XPUSH_NEW_OBJ ("OpenCL::Image2D", mem); 1356 XPUSH_CLOBJ (stash_image2d, mem);
805 1357
806#if CL_VERSION_1_2 1358#if CL_VERSION_1_2
807 1359
808void 1360void
809gl_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)
810 ALIAS: 1362 ALIAS:
811 gl_texture2d = 0
812 gl_texture3d = 0
813 PPCODE: 1363 PPCODE:
814 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));
815 cl_gl_object_type type; 1365 cl_gl_object_type type;
816 NEED_SUCCESS (GetGLObjectInfo, (mem, &type, 0)); // TODO: use target instead? 1366 NEED_SUCCESS (GetGLObjectInfo, (mem, &type, 0)); // TODO: use target instead?
817 char *klass = "OpenCL::Memory"; 1367 HV *stash = stash_memory;
818 switch (type) 1368 switch (type)
819 { 1369 {
820 case CL_GL_OBJECT_TEXTURE_BUFFER: klass = "OpenCL::Image1DBuffer"; break; 1370 case CL_GL_OBJECT_TEXTURE_BUFFER: stash = stash_image1dbuffer; break;
821 case CL_GL_OBJECT_TEXTURE1D: klass = "OpenCL::Image1D"; break; 1371 case CL_GL_OBJECT_TEXTURE1D: stash = stash_image1d; break;
822 case CL_GL_OBJECT_TEXTURE1D_ARRAY: klass = "OpenCL::Image2DArray"; break; 1372 case CL_GL_OBJECT_TEXTURE1D_ARRAY: stash = stash_image2darray; break;
823 case CL_GL_OBJECT_TEXTURE2D: klass = "OpenCL::Image2D"; break; 1373 case CL_GL_OBJECT_TEXTURE2D: stash = stash_image2d; break;
824 case CL_GL_OBJECT_TEXTURE2D_ARRAY: klass = "OpenCL::Image2DArray"; break; 1374 case CL_GL_OBJECT_TEXTURE2D_ARRAY: stash = stash_image2darray; break;
825 case CL_GL_OBJECT_TEXTURE3D: klass = "OpenCL::Image3D"; break; 1375 case CL_GL_OBJECT_TEXTURE3D: stash = stash_image3d; break;
826 } 1376 }
827 XPUSH_NEW_OBJ (klass, mem); 1377 XPUSH_CLOBJ (stash, mem);
828 1378
1379#endif
1380
1381void
1382gl_texture2d (OpenCL::Context self, cl_mem_flags flags, cl_GLenum target, cl_GLint miplevel, cl_GLuint texture)
1383 PPCODE:
1384#if PREFER_1_1
1385 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture2D, (self, flags, target, miplevel, texture, &res));
829#else 1386#else
830
831void
832gl_texture2d (OpenCL::Context self, cl_mem_flags flags, cl_GLenum target, cl_GLint miplevel, cl_GLuint texture)
833 PPCODE:
834 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture2D, (self, flags, target, miplevel, texture, &res)); 1387 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture , (self, flags, target, miplevel, texture, &res));
835 XPUSH_NEW_OBJ ("OpenCL::Image2D", mem); 1388#endif
1389 XPUSH_CLOBJ (stash_image2d, mem);
836 1390
837void 1391void
838gl_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)
839 PPCODE: 1393 PPCODE:
1394#if PREFER_1_1
840 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));
841 XPUSH_NEW_OBJ ("OpenCL::Image3D", mem); 1396#else
842 1397 NEED_SUCCESS_ARG (cl_mem mem, CreateFromGLTexture , (self, flags, target, miplevel, texture, &res));
843#endif 1398#endif
1399 XPUSH_CLOBJ (stash_image3d, mem);
844 1400
845#endif 1401#endif
846 1402
847void 1403void
848supported_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)
… …
868 1424
869void 1425void
870sampler (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)
871 PPCODE: 1427 PPCODE:
872 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));
873 XPUSH_NEW_OBJ ("OpenCL::Sampler", sampler); 1429 XPUSH_CLOBJ (stash_sampler, sampler);
874 1430
875void 1431void
876program_with_source (OpenCL::Context self, SV *program) 1432program_with_source (OpenCL::Context self, SV *program)
877 PPCODE: 1433 PPCODE:
878 STRLEN len; 1434 STRLEN len;
879 size_t len2; 1435 size_t len2;
880 const char *ptr = SvPVbyte (program, len); 1436 const char *ptr = SvPVbyte (program, len);
881 1437
882 len2 = len; 1438 len2 = len;
883 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));
884 XPUSH_NEW_OBJ ("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
885 1517
886#BEGIN:context 1518#BEGIN:context
887 1519
888void 1520void
889reference_count (OpenCL::Context self) 1521reference_count (OpenCL::Context self)
… …
905 cl_device_id *value = tmpbuf (size); 1537 cl_device_id *value = tmpbuf (size);
906 NEED_SUCCESS (GetContextInfo, (self, CL_CONTEXT_DEVICES, size, value, 0)); 1538 NEED_SUCCESS (GetContextInfo, (self, CL_CONTEXT_DEVICES, size, value, 0));
907 int i, n = size / sizeof (*value); 1539 int i, n = size / sizeof (*value);
908 EXTEND (SP, n); 1540 EXTEND (SP, n);
909 for (i = 0; i < n; ++i) 1541 for (i = 0; i < n; ++i)
910 { 1542 PUSH_CLOBJ (stash_device, value [i]);
911 PUSHs (NEW_MORTAL_OBJ ("OpenCL::Device", value [i]));
912 }
913 1543
914void 1544void
915properties (OpenCL::Context self) 1545properties (OpenCL::Context self)
916 PPCODE: 1546 PPCODE:
917 size_t size; 1547 size_t size;
… …
931DESTROY (OpenCL::Queue self) 1561DESTROY (OpenCL::Queue self)
932 CODE: 1562 CODE:
933 clReleaseCommandQueue (self); 1563 clReleaseCommandQueue (self);
934 1564
935void 1565void
936enqueue_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
937 PPCODE: 1569 PPCODE:
938 cl_event ev = 0;
939 EVENT_LIST (6, items - 6); 1570 EVENT_LIST (6);
940 1571
941 SvUPGRADE (data, SVt_PV); 1572 SvUPGRADE (data, SVt_PV);
942 SvGROW (data, len); 1573 SvGROW (data, len);
943 SvPOK_only (data); 1574 SvPOK_only (data);
944 SvCUR_set (data, len); 1575 SvCUR_set (data, len);
1576
1577 cl_event ev = 0;
945 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));
946 1579
947 if (ev) 1580 if (ev)
948 XPUSH_NEW_OBJ ("OpenCL::Event", ev); 1581 XPUSH_CLOBJ (stash_event, ev);
949 1582
950void 1583void
951enqueue_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
952 PPCODE: 1587 PPCODE:
953 cl_event ev = 0; 1588 EVENT_LIST (5);
1589
954 STRLEN len; 1590 STRLEN len;
955 char *ptr = SvPVbyte (data, len); 1591 char *ptr = SvPVbyte (data, len);
956 EVENT_LIST (5, items - 5);
957 1592
1593 cl_event ev = 0;
958 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));
959 1595
960 if (ev) 1596 if (ev)
961 XPUSH_NEW_OBJ ("OpenCL::Event", ev); 1597 XPUSH_CLOBJ (stash_event, ev);
962 1598
1599#if CL_VERSION_1_2
1600
963void 1601void
964enqueue_copy_buffer (OpenCL::Queue self, OpenCL::Buffer src, OpenCL::Buffer dst, size_t src_offset, size_t dst_offset, size_t len, ...) 1602fill_buffer (OpenCL::Queue self, OpenCL::Buffer mem, SV *data, size_t offset, size_t size, ...)
1603 ALIAS:
1604 enqueue_fill_buffer = 0
965 PPCODE: 1605 PPCODE:
1606 EVENT_LIST (5);
1607
1608 STRLEN len;
1609 char *ptr = SvPVbyte (data, len);
1610
966 cl_event ev = 0; 1611 cl_event ev = 0;
1612 NEED_SUCCESS (EnqueueFillBuffer, (self, mem, ptr, len, offset, size, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1613
1614 if (ev)
1615 XPUSH_CLOBJ (stash_event, ev);
1616
1617void
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
1621 PPCODE:
1622 EVENT_LIST (12);
1623
1624 const size_t origin [3] = { x, y, z };
1625 const size_t region [3] = { width, height, depth };
1626
1627 const cl_float c_f [4] = { r, g, b, a };
1628 const cl_uint c_u [4] = { r, g, b, a };
1629 const cl_int c_s [4] = { r, g, b, a };
1630 const void *c_fus [3] = { &c_f, &c_u, &c_s };
1631 static const char fus [] = { 0, 0, 0, 0, 0, 0, 0, 2, 2, 2, 1, 1, 1, 0, 0 };
1632 cl_image_format format;
1633 NEED_SUCCESS (GetImageInfo, (img, CL_IMAGE_FORMAT, sizeof (format), &format, 0));
1634 assert (sizeof (fus) == CL_FLOAT + 1 - CL_SNORM_INT8);
1635 if (format.image_channel_data_type < CL_SNORM_INT8 || CL_FLOAT < format.image_channel_data_type)
1636 croak ("enqueue_fill_image: image has unsupported channel type, only opencl 1.2 channel types supported.");
1637
1638 cl_event ev = 0;
1639 NEED_SUCCESS (EnqueueFillImage, (self, img, c_fus [fus [format.image_channel_data_type - CL_SNORM_INT8]],
1640 origin, region, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1641
1642 if (ev)
1643 XPUSH_CLOBJ (stash_event, ev);
1644
1645#endif
1646
1647void
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
1651 PPCODE:
967 EVENT_LIST (6, items - 6); 1652 EVENT_LIST (6);
968 1653
1654 cl_event ev = 0;
969 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));
970 1656
971 if (ev) 1657 if (ev)
972 XPUSH_NEW_OBJ ("OpenCL::Event", ev); 1658 XPUSH_CLOBJ (stash_event, ev);
973 1659
974void 1660void
975enqueue_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
976 PPCODE: 1664 PPCODE:
977 cl_event ev = 0; 1665 EVENT_LIST (17);
1666
978 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 };
979 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 };
980 const size_t region[3] = { width, height, depth }; 1669 const size_t region[3] = { width, height, depth };
981 EVENT_LIST (17, items - 17);
982 1670
983 if (!buf_row_pitch) 1671 if (!buf_row_pitch)
984 buf_row_pitch = region [0]; 1672 buf_row_pitch = region [0];
985 1673
986 if (!buf_slice_pitch) 1674 if (!buf_slice_pitch)
… …
996 1684
997 SvUPGRADE (data, SVt_PV); 1685 SvUPGRADE (data, SVt_PV);
998 SvGROW (data, len); 1686 SvGROW (data, len);
999 SvPOK_only (data); 1687 SvPOK_only (data);
1000 SvCUR_set (data, len); 1688 SvCUR_set (data, len);
1689
1690 cl_event ev = 0;
1001 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));
1002 1692
1003 if (ev) 1693 if (ev)
1004 XPUSH_NEW_OBJ ("OpenCL::Event", ev); 1694 XPUSH_CLOBJ (stash_event, ev);
1005 1695
1006void 1696void
1007enqueue_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
1008 PPCODE: 1700 PPCODE:
1009 cl_event ev = 0; 1701 EVENT_LIST (17);
1702
1010 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 };
1011 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 };
1012 const size_t region[3] = { width, height, depth }; 1705 const size_t region[3] = { width, height, depth };
1013 STRLEN len; 1706 STRLEN len;
1014 char *ptr = SvPVbyte (data, len); 1707 char *ptr = SvPVbyte (data, len);
1015 EVENT_LIST (17, items - 17);
1016 1708
1017 if (!buf_row_pitch) 1709 if (!buf_row_pitch)
1018 buf_row_pitch = region [0]; 1710 buf_row_pitch = region [0];
1019 1711
1020 if (!buf_slice_pitch) 1712 if (!buf_slice_pitch)
… …
1029 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];
1030 1722
1031 if (len < min_len) 1723 if (len < min_len)
1032 croak ("clEnqueueWriteImage: data string is shorter than what would be transferred"); 1724 croak ("clEnqueueWriteImage: data string is shorter than what would be transferred");
1033 1725
1726 cl_event ev = 0;
1034 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));
1035 1728
1036 if (ev) 1729 if (ev)
1037 XPUSH_NEW_OBJ ("OpenCL::Event", ev); 1730 XPUSH_CLOBJ (stash_event, ev);
1038 1731
1039void 1732void
1040enqueue_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
1041 PPCODE: 1736 PPCODE:
1042 cl_event ev = 0; 1737 EVENT_LIST (16);
1738
1043 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 };
1044 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 };
1045 const size_t region[3] = { width, height, depth }; 1741 const size_t region[3] = { width, height, depth };
1046 EVENT_LIST (16, items - 16);
1047 1742
1743 cl_event ev = 0;
1048 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));
1049 1745
1050 if (ev) 1746 if (ev)
1051 XPUSH_NEW_OBJ ("OpenCL::Event", ev); 1747 XPUSH_CLOBJ (stash_event, ev);
1052 1748
1053void 1749void
1054enqueue_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
1055 PPCODE: 1753 PPCODE:
1056 cl_event ev = 0; 1754 EVENT_LIST (12);
1755
1057 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 };
1058 const size_t region[3] = { width, height, depth }; 1757 const size_t region[3] = { width, height, depth };
1059 EVENT_LIST (12, items - 12);
1060 1758
1061 if (!row_pitch) 1759 if (!row_pitch)
1062 row_pitch = img_row_pitch (src); 1760 row_pitch = img_row_pitch (src);
1063 1761
1064 if (depth > 1 && !slice_pitch) 1762 if (depth > 1 && !slice_pitch)
… …
1068 1766
1069 SvUPGRADE (data, SVt_PV); 1767 SvUPGRADE (data, SVt_PV);
1070 SvGROW (data, len); 1768 SvGROW (data, len);
1071 SvPOK_only (data); 1769 SvPOK_only (data);
1072 SvCUR_set (data, len); 1770 SvCUR_set (data, len);
1771
1772 cl_event ev = 0;
1073 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));
1074 1774
1075 if (ev) 1775 if (ev)
1076 XPUSH_NEW_OBJ ("OpenCL::Event", ev); 1776 XPUSH_CLOBJ (stash_event, ev);
1077 1777
1078void 1778void
1079enqueue_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
1080 PPCODE: 1782 PPCODE:
1081 cl_event ev = 0; 1783 EVENT_LIST (12);
1784
1082 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 };
1083 const size_t region[3] = { width, height, depth }; 1786 const size_t region[3] = { width, height, depth };
1084 STRLEN len; 1787 STRLEN len;
1085 char *ptr = SvPVbyte (data, len); 1788 char *ptr = SvPVbyte (data, len);
1086 EVENT_LIST (12, items - 12);
1087 1789
1088 if (!row_pitch) 1790 if (!row_pitch)
1089 row_pitch = img_row_pitch (dst); 1791 row_pitch = img_row_pitch (dst);
1090 1792
1091 if (depth > 1 && !slice_pitch) 1793 if (depth > 1 && !slice_pitch)
… …
1094 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;
1095 1797
1096 if (len < min_len) 1798 if (len < min_len)
1097 croak ("clEnqueueWriteImage: data string is shorter than what would be transferred"); 1799 croak ("clEnqueueWriteImage: data string is shorter than what would be transferred");
1098 1800
1801 cl_event ev = 0;
1099 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));
1100 1803
1101 if (ev) 1804 if (ev)
1102 XPUSH_NEW_OBJ ("OpenCL::Event", ev); 1805 XPUSH_CLOBJ (stash_event, ev);
1103 1806
1104void 1807void
1105enqueue_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
1106 PPCODE: 1811 PPCODE:
1107 cl_event ev = 0; 1812 EVENT_LIST (12);
1813
1108 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 };
1109 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 };
1110 const size_t region[3] = { width, height, depth }; 1816 const size_t region[3] = { width, height, depth };
1111 EVENT_LIST (12, items - 12);
1112 1817
1818 cl_event ev = 0;
1113 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));
1114 1820
1115 if (ev) 1821 if (ev)
1116 XPUSH_NEW_OBJ ("OpenCL::Event", ev); 1822 XPUSH_CLOBJ (stash_event, ev);
1117 1823
1118void 1824void
1119enqueue_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
1120 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
1121 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
1122 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 };
1123 const size_t region[3] = { width, height, depth }; 1910 const size_t region[3] = { width, height, depth };
1911 size_t row_pitch, slice_pitch;
1124 EVENT_LIST (10, items - 10); 1912 EVENT_LIST (10);
1125 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;
1126 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));
1127 1942
1128 if (ev) 1943 if (ev)
1129 XPUSH_NEW_OBJ ("OpenCL::Event", ev); 1944 XPUSH_CLOBJ (stash_event, ev);
1130 1945
1131void 1946void
1132enqueue_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, ...)
1133 PPCODE:
1134 cl_event ev = 0;
1135 const size_t dst_origin[3] = { dst_x, dst_y, dst_z };
1136 const size_t region[3] = { width, height, depth };
1137 EVENT_LIST (10, items - 10);
1138
1139 NEED_SUCCESS (EnqueueCopyBufferToImage, (self, src, dst, src_offset, dst_origin, region, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1140
1141 if (ev)
1142 XPUSH_NEW_OBJ ("OpenCL::Event", ev);
1143
1144void
1145enqueue_task (OpenCL::Queue self, OpenCL::Kernel kernel, ...)
1146 PPCODE:
1147 cl_event ev = 0;
1148 EVENT_LIST (2, items - 2);
1149
1150 NEED_SUCCESS (EnqueueTask, (self, kernel, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1151
1152 if (ev)
1153 XPUSH_NEW_OBJ ("OpenCL::Event", ev);
1154
1155void
1156enqueue_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
1157 PPCODE: 1950 PPCODE:
1158 cl_event ev = 0; 1951 EVENT_LIST (5);
1952
1159 size_t *gwo = 0, *gws, *lws = 0; 1953 size_t *gwo = 0, *gws, *lws = 0;
1160 int gws_len; 1954 int gws_len;
1161 size_t *lists; 1955 size_t *lists;
1162 int i; 1956 int i;
1163 EVENT_LIST (5, items - 5);
1164 1957
1165 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)
1166 croak ("clEnqueueNDRangeKernel: global_work_size must be an array reference"); 1959 croak ("clEnqueueNDRangeKernel: global_work_size must be an array reference");
1167 1960
1168 gws_len = AvFILLp (SvRV (global_work_size)) + 1; 1961 gws_len = AvFILLp (SvRV (global_work_size)) + 1;
1169 1962
1170 lists = tmpbuf (sizeof (size_t) * 3 * gws_len); 1963 lists = tmpbuf (sizeof (size_t) * 3 * gws_len);
1171 1964
1172 gws = lists + gws_len * 0; 1965 gws = lists + gws_len * 0;
1173 for (i = 0; i < gws_len; ++i) 1966 for (i = 0; i < gws_len; ++i)
1967 {
1174 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 }
1175 1973
1176 if (SvOK (global_work_offset)) 1974 if (SvOK (global_work_offset))
1177 { 1975 {
1178 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)
1179 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");
… …
1187 } 1985 }
1188 1986
1189 if (SvOK (local_work_size)) 1987 if (SvOK (local_work_size))
1190 { 1988 {
1191 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)
1192 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");
1193 1991
1194 if (AvFILLp (SvRV (local_work_size)) + 1 != gws_len) 1992 if (AvFILLp (SvRV (local_work_size)) + 1 != gws_len)
1195 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");
1196 1994
1197 lws = lists + gws_len * 2; 1995 lws = lists + gws_len * 2;
1198 for (i = 0; i < gws_len; ++i) 1996 for (i = 0; i < gws_len; ++i)
1997 {
1199 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 }
1200 } 2003 }
1201 2004
2005 cl_event ev = 0;
1202 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));
1203 2007
1204 if (ev) 2008 if (ev)
1205 XPUSH_NEW_OBJ ("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
1206 2030
1207#if cl_apple_gl_sharing || cl_khr_gl_sharing 2031#if cl_apple_gl_sharing || cl_khr_gl_sharing
1208 2032
1209void 2033void
1210enqueue_acquire_gl_objects (OpenCL::Queue self, SV *objects, ...) 2034acquire_gl_objects (OpenCL::Queue self, SV *objects, ...)
1211 ALIAS: 2035 ALIAS:
2036 release_gl_objects = 1
2037 enqueue_acquire_gl_objects = 0
1212 enqueue_release_gl_objects = 1 2038 enqueue_release_gl_objects = 1
1213 PPCODE: 2039 PPCODE:
1214 if (!SvROK (objects) || SvTYPE (SvRV (objects)) != SVt_PVAV) 2040 EVENT_LIST (2);
1215 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);
1216 2044
1217 cl_event ev = 0; 2045 cl_event ev = 0;
1218 EVENT_LIST (2, items - 2);
1219 AV *av = (AV *)SvRV (objects);
1220 cl_uint num_objects = av_len (av) + 1;
1221 cl_mem *object_list = tmpbuf (sizeof (cl_mem) * num_objects);
1222 int i;
1223
1224 for (i = num_objects; i--; )
1225 object_list [i] = SvPTROBJ ("OpenCL::Queue::enqueue_acquire/release_gl_objects", "objects", *av_fetch (av, i, 0), "OpenCL::Memory");
1226 2046
1227 if (ix) 2047 if (ix)
1228 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));
1229 else 2049 else
1230 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));
1231 2051
1232 if (ev) 2052 if (ev)
1233 XPUSH_NEW_OBJ ("OpenCL::Event", ev); 2053 XPUSH_CLOBJ (stash_event, ev);
1234 2054
1235#endif 2055#endif
1236 2056
1237#if !CL_VERSION_1_2 || defined CL_USE_DEPRECATED_OPENCL_1_1_APIS
1238
1239void 2057void
1240enqueue_wait_for_events (OpenCL::Queue self, ...) 2058wait_for_events (OpenCL::Queue self, ...)
2059 ALIAS:
2060 enqueue_wait_for_events = 0
1241 CODE: 2061 CODE:
1242 EVENT_LIST (1, items - 1); 2062 EVENT_LIST (1);
2063#if PREFER_1_1
1243 NEED_SUCCESS (EnqueueWaitForEvents, (self, event_list_count, event_list_ptr)); 2064 NEED_SUCCESS (EnqueueWaitForEvents, (self, event_list_count, event_list_ptr));
1244 2065#else
2066 NEED_SUCCESS (EnqueueBarrierWithWaitList, (self, event_list_count, event_list_ptr, 0));
1245#endif 2067#endif
1246 2068
1247void 2069void
1248enqueue_marker (OpenCL::Queue self, ...) 2070marker (OpenCL::Queue self, ...)
2071 ALIAS:
2072 enqueue_marker = 0
1249 PPCODE: 2073 PPCODE:
2074 EVENT_LIST (1);
1250 cl_event ev = 0; 2075 cl_event ev = 0;
1251 EVENT_LIST (1, items - 1); 2076#if PREFER_1_1
2077 if (!event_list_count)
2078 NEED_SUCCESS (EnqueueMarker, (self, GIMME_V != G_VOID ? &ev : 0));
2079 else
1252#if CL_VERSION_1_2 2080#if CL_VERSION_1_2
2081 NEED_SUCCESS (EnqueueMarkerWithWaitList, (self, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
2082#else
2083 {
2084 NEED_SUCCESS (EnqueueWaitForEvents, (self, event_list_count, event_list_ptr)); // also a barrier
2085 NEED_SUCCESS (EnqueueMarker, (self, GIMME_V != G_VOID ? &ev : 0));
2086 }
2087#endif
2088#else
1253 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));
1254#else
1255 if (event_list_count)
1256 croak ("OpenCL::Queue->enqueue_marker does not support a wait list in OpenCL 1.1 - upgrade to 1.2");
1257 NEED_SUCCESS (EnqueueMarker, (self, GIMME_V != G_VOID ? &ev : 0));
1258#endif 2090#endif
1259 if (ev) 2091 if (ev)
1260 XPUSH_NEW_OBJ ("OpenCL::Event", ev); 2092 XPUSH_CLOBJ (stash_event, ev);
1261 2093
1262void 2094void
1263enqueue_barrier (OpenCL::Queue self, ...) 2095barrier (OpenCL::Queue self, ...)
2096 ALIAS:
2097 enqueue_barrier = 0
1264 PPCODE: 2098 PPCODE:
2099 EVENT_LIST (1);
1265 cl_event ev = 0; 2100 cl_event ev = 0;
1266 EVENT_LIST (1, items - 1); 2101#if PREFER_1_1
2102 if (!event_list_count && GIMME_V == G_VOID)
2103 NEED_SUCCESS (EnqueueBarrier, (self));
2104 else
1267#if CL_VERSION_1_2 2105#if CL_VERSION_1_2
1268 NEED_SUCCESS (EnqueueBarrierWithWaitList, (self, event_list_count, event_list_ptr, &ev)); 2106 NEED_SUCCESS (EnqueueBarrierWithWaitList, (self, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1269#else 2107#else
2108 {
1270 if (event_list_count) 2109 if (event_list_count)
1271 croak ("OpenCL::Queue->enqueue_barrier does not support a wait list in OpenCL 1.1 - upgrade to 1.2"); 2110 NEED_SUCCESS (EnqueueWaitForEvents, (self, event_list_count, event_list_ptr));
2111
1272 if (GIMME_V != G_VOID) 2112 if (GIMME_V != G_VOID)
1273 croak ("OpenCL::Queue->enqueue_barrier does not return an event object in OpenCL 1.1 - upgrade to 1.2");
1274 NEED_SUCCESS (EnqueueBarrier, (self)); 2113 NEED_SUCCESS (EnqueueMarker, (self, &ev));
2114 }
2115#endif
2116#else
2117 NEED_SUCCESS (EnqueueBarrierWithWaitList, (self, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
1275#endif 2118#endif
1276 if (ev) 2119 if (ev)
1277 XPUSH_NEW_OBJ ("OpenCL::Event", ev); 2120 XPUSH_CLOBJ (stash_event, ev);
1278 2121
1279void 2122void
1280flush (OpenCL::Queue self) 2123flush (OpenCL::Queue self)
1281 CODE: 2124 CODE:
1282 NEED_SUCCESS (Flush, (self)); 2125 NEED_SUCCESS (Flush, (self));
… …
1298 PPCODE: 2141 PPCODE:
1299 cl_context value [1]; 2142 cl_context value [1];
1300 NEED_SUCCESS (GetCommandQueueInfo, (self, CL_QUEUE_CONTEXT, sizeof (value), value, 0)); 2143 NEED_SUCCESS (GetCommandQueueInfo, (self, CL_QUEUE_CONTEXT, sizeof (value), value, 0));
1301 EXTEND (SP, 1); 2144 EXTEND (SP, 1);
1302 const int i = 0; 2145 const int i = 0;
1303 {
1304 NEED_SUCCESS (RetainContext, (value [i])); 2146 NEED_SUCCESS (RetainContext, (value [i]));
1305 PUSHs (NEW_MORTAL_OBJ ("OpenCL::Context", value [i])); 2147 PUSH_CLOBJ (stash_context, value [i]);
1306 }
1307 2148
1308void 2149void
1309device (OpenCL::Queue self) 2150device (OpenCL::Queue self)
1310 PPCODE: 2151 PPCODE:
1311 cl_device_id value [1]; 2152 cl_device_id value [1];
1312 NEED_SUCCESS (GetCommandQueueInfo, (self, CL_QUEUE_DEVICE, sizeof (value), value, 0)); 2153 NEED_SUCCESS (GetCommandQueueInfo, (self, CL_QUEUE_DEVICE, sizeof (value), value, 0));
1313 EXTEND (SP, 1); 2154 EXTEND (SP, 1);
1314 const int i = 0; 2155 const int i = 0;
1315 { 2156 PUSH_CLOBJ (stash_device, value [i]);
1316 PUSHs (NEW_MORTAL_OBJ ("OpenCL::Device", value [i]));
1317 }
1318 2157
1319void 2158void
1320reference_count (OpenCL::Queue self) 2159reference_count (OpenCL::Queue self)
1321 PPCODE: 2160 PPCODE:
1322 cl_uint value [1]; 2161 cl_uint value [1];
… …
1326 PUSHs (sv_2mortal (newSVuv (value [i]))); 2165 PUSHs (sv_2mortal (newSVuv (value [i])));
1327 2166
1328void 2167void
1329properties (OpenCL::Queue self) 2168properties (OpenCL::Queue self)
1330 PPCODE: 2169 PPCODE:
1331 cl_command_queue_properties value [1]; 2170 cl_ulong value [1];
1332 NEED_SUCCESS (GetCommandQueueInfo, (self, CL_QUEUE_PROPERTIES, sizeof (value), value, 0)); 2171 NEED_SUCCESS (GetCommandQueueInfo, (self, CL_QUEUE_PROPERTIES, sizeof (value), value, 0));
1333 EXTEND (SP, 1); 2172 EXTEND (SP, 1);
1334 const int i = 0; 2173 const int i = 0;
1335 PUSHs (sv_2mortal (newSViv (value [i]))); 2174 PUSHs (sv_2mortal (newSVuv (value [i])));
1336 2175
1337#END:command_queue 2176#END:command_queue
1338 2177
1339MODULE = OpenCL PACKAGE = OpenCL::Memory 2178MODULE = OpenCL PACKAGE = OpenCL::Memory
1340 2179
… …
1346void 2185void
1347info (OpenCL::Memory self, cl_mem_info name) 2186info (OpenCL::Memory self, cl_mem_info name)
1348 PPCODE: 2187 PPCODE:
1349 INFO (MemObject) 2188 INFO (MemObject)
1350 2189
2190void
2191destructor_callback (OpenCL::Memory self, SV *notify)
2192 PPCODE:
2193 clSetMemObjectDestructorCallback (self, eq_destructor_notify, SvREFCNT_inc (s_get_cv (notify)));
2194
1351#BEGIN:mem 2195#BEGIN:mem
1352 2196
1353void 2197void
1354type (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
1355 PPCODE: 2203 PPCODE:
1356 cl_mem_object_type value [1]; 2204 cl_uint value [1];
1357 NEED_SUCCESS (GetMemObjectInfo, (self, CL_MEM_TYPE, sizeof (value), value, 0)); 2205 NEED_SUCCESS (GetMemObjectInfo, (self, ix, sizeof (value), value, 0));
1358 EXTEND (SP, 1); 2206 EXTEND (SP, 1);
1359 const int i = 0; 2207 const int i = 0;
1360 PUSHs (sv_2mortal (newSViv (value [i]))); 2208 PUSHs (sv_2mortal (newSVuv (value [i])));
1361 2209
1362void 2210void
1363flags (OpenCL::Memory self) 2211flags (OpenCL::Memory self)
1364 PPCODE: 2212 PPCODE:
1365 cl_mem_flags value [1]; 2213 cl_ulong value [1];
1366 NEED_SUCCESS (GetMemObjectInfo, (self, CL_MEM_FLAGS, sizeof (value), value, 0)); 2214 NEED_SUCCESS (GetMemObjectInfo, (self, CL_MEM_FLAGS, sizeof (value), value, 0));
1367 EXTEND (SP, 1); 2215 EXTEND (SP, 1);
1368 const int i = 0; 2216 const int i = 0;
1369 PUSHs (sv_2mortal (newSViv (value [i]))); 2217 PUSHs (sv_2mortal (newSVuv (value [i])));
1370 2218
1371void 2219void
1372size (OpenCL::Memory self) 2220size (OpenCL::Memory self)
1373 ALIAS: 2221 ALIAS:
1374 size = CL_MEM_SIZE 2222 size = CL_MEM_SIZE
… …
1388 EXTEND (SP, 1); 2236 EXTEND (SP, 1);
1389 const int i = 0; 2237 const int i = 0;
1390 PUSHs (sv_2mortal (newSVuv ((IV)(intptr_t)value [i]))); 2238 PUSHs (sv_2mortal (newSVuv ((IV)(intptr_t)value [i])));
1391 2239
1392void 2240void
1393map_count (OpenCL::Memory self)
1394 ALIAS:
1395 map_count = CL_MEM_MAP_COUNT
1396 reference_count = CL_MEM_REFERENCE_COUNT
1397 PPCODE:
1398 cl_uint value [1];
1399 NEED_SUCCESS (GetMemObjectInfo, (self, ix, sizeof (value), value, 0));
1400 EXTEND (SP, 1);
1401 const int i = 0;
1402 PUSHs (sv_2mortal (newSVuv (value [i])));
1403
1404void
1405context (OpenCL::Memory self) 2241context (OpenCL::Memory self)
1406 PPCODE: 2242 PPCODE:
1407 cl_context value [1]; 2243 cl_context value [1];
1408 NEED_SUCCESS (GetMemObjectInfo, (self, CL_MEM_CONTEXT, sizeof (value), value, 0)); 2244 NEED_SUCCESS (GetMemObjectInfo, (self, CL_MEM_CONTEXT, sizeof (value), value, 0));
1409 EXTEND (SP, 1); 2245 EXTEND (SP, 1);
1410 const int i = 0; 2246 const int i = 0;
1411 {
1412 NEED_SUCCESS (RetainContext, (value [i])); 2247 NEED_SUCCESS (RetainContext, (value [i]));
1413 PUSHs (NEW_MORTAL_OBJ ("OpenCL::Context", value [i])); 2248 PUSH_CLOBJ (stash_context, value [i]);
1414 }
1415 2249
1416void 2250void
1417associated_memobject (OpenCL::Memory self) 2251associated_memobject (OpenCL::Memory self)
1418 PPCODE: 2252 PPCODE:
1419 cl_mem value [1]; 2253 cl_mem value [1];
1420 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));
1421 EXTEND (SP, 1); 2255 EXTEND (SP, 1);
1422 const int i = 0; 2256 const int i = 0;
1423 {
1424 NEED_SUCCESS (RetainMemObject, (value [i])); 2257 NEED_SUCCESS (RetainMemObject, (value [i]));
1425 PUSHs (NEW_MORTAL_OBJ ("OpenCL::Memory", value [i])); 2258 PUSH_CLOBJ (stash_memory, value [i]);
1426 }
1427 2259
1428#END:mem 2260#END:mem
1429 2261
1430#if cl_apple_gl_sharing || cl_khr_gl_sharing 2262#if cl_apple_gl_sharing || cl_khr_gl_sharing
1431 2263
… …
1450 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!");
1451 2283
1452 cl_buffer_region crdata = { origin, size }; 2284 cl_buffer_region crdata = { origin, size };
1453 2285
1454 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));
1455 XPUSH_NEW_OBJ ("OpenCL::Buffer", mem); 2287 XPUSH_CLOBJ (stash_buffer, mem);
1456 2288
1457MODULE = OpenCL PACKAGE = OpenCL::Image 2289MODULE = OpenCL PACKAGE = OpenCL::Image
1458 2290
1459void 2291void
1460image_info (OpenCL::Image self, cl_image_info name) 2292image_info (OpenCL::Image self, cl_image_info name)
1461 PPCODE: 2293 PPCODE:
1462 INFO (Image) 2294 INFO (Image)
2295
2296void
2297format (OpenCL::Image self)
2298 PPCODE:
2299 cl_image_format format;
2300 NEED_SUCCESS (GetImageInfo, (self, CL_IMAGE_FORMAT, sizeof (format), &format, 0));
2301 EXTEND (SP, 2);
2302 PUSHs (sv_2mortal (newSVuv (format.image_channel_order)));
2303 PUSHs (sv_2mortal (newSVuv (format.image_channel_data_type)));
1463 2304
1464#BEGIN:image 2305#BEGIN:image
1465 2306
1466void 2307void
1467element_size (OpenCL::Image self) 2308element_size (OpenCL::Image self)
… …
1521 2362
1522#BEGIN:sampler 2363#BEGIN:sampler
1523 2364
1524void 2365void
1525reference_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
1526 PPCODE: 2371 PPCODE:
1527 cl_uint value [1]; 2372 cl_uint value [1];
1528 NEED_SUCCESS (GetSamplerInfo, (self, CL_SAMPLER_REFERENCE_COUNT, sizeof (value), value, 0)); 2373 NEED_SUCCESS (GetSamplerInfo, (self, ix, sizeof (value), value, 0));
1529 EXTEND (SP, 1); 2374 EXTEND (SP, 1);
1530 const int i = 0; 2375 const int i = 0;
1531 PUSHs (sv_2mortal (newSVuv (value [i]))); 2376 PUSHs (sv_2mortal (newSVuv (value [i])));
1532 2377
1533void 2378void
… …
1535 PPCODE: 2380 PPCODE:
1536 cl_context value [1]; 2381 cl_context value [1];
1537 NEED_SUCCESS (GetSamplerInfo, (self, CL_SAMPLER_CONTEXT, sizeof (value), value, 0)); 2382 NEED_SUCCESS (GetSamplerInfo, (self, CL_SAMPLER_CONTEXT, sizeof (value), value, 0));
1538 EXTEND (SP, 1); 2383 EXTEND (SP, 1);
1539 const int i = 0; 2384 const int i = 0;
1540 {
1541 NEED_SUCCESS (RetainContext, (value [i])); 2385 NEED_SUCCESS (RetainContext, (value [i]));
1542 PUSHs (NEW_MORTAL_OBJ ("OpenCL::Context", value [i])); 2386 PUSH_CLOBJ (stash_context, value [i]);
1543 }
1544
1545void
1546normalized_coords (OpenCL::Sampler self)
1547 PPCODE:
1548 cl_addressing_mode value [1];
1549 NEED_SUCCESS (GetSamplerInfo, (self, CL_SAMPLER_NORMALIZED_COORDS, sizeof (value), value, 0));
1550 EXTEND (SP, 1);
1551 const int i = 0;
1552 PUSHs (sv_2mortal (newSViv (value [i])));
1553
1554void
1555addressing_mode (OpenCL::Sampler self)
1556 PPCODE:
1557 cl_filter_mode value [1];
1558 NEED_SUCCESS (GetSamplerInfo, (self, CL_SAMPLER_ADDRESSING_MODE, sizeof (value), value, 0));
1559 EXTEND (SP, 1);
1560 const int i = 0;
1561 PUSHs (sv_2mortal (newSViv (value [i])));
1562 2387
1563void 2388void
1564filter_mode (OpenCL::Sampler self) 2389filter_mode (OpenCL::Sampler self)
1565 PPCODE: 2390 PPCODE:
1566 cl_bool value [1]; 2391 cl_uint value [1];
1567 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));
1568 EXTEND (SP, 1); 2393 EXTEND (SP, 1);
1569 const int i = 0; 2394 const int i = 0;
1570 PUSHs (sv_2mortal (value [i] ? &PL_sv_yes : &PL_sv_no)); 2395 PUSHs (sv_2mortal (value [i] ? &PL_sv_yes : &PL_sv_no));
1571 2396
… …
1577DESTROY (OpenCL::Program self) 2402DESTROY (OpenCL::Program self)
1578 CODE: 2403 CODE:
1579 clReleaseProgram (self); 2404 clReleaseProgram (self);
1580 2405
1581void 2406void
1582build (OpenCL::Program self, OpenCL::Device device, SV *options = &PL_sv_undef) 2407build (OpenCL::Program self, SV *devices = &PL_sv_undef, SV *options = &PL_sv_undef, SV *notify = &PL_sv_undef)
2408 ALIAS:
2409 build_async = 1
1583 CODE: 2410 CODE:
2411 cl_uint device_count = 0;
2412 cl_device_id *device_list = 0;
2413
2414 if (SvOK (devices))
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 {
2442 if (!SvROK (devices) || SvTYPE (SvRV (devices)) != SVt_PVHV)
2443 croak ("clCompileProgram: headers must be undef or a hashref of name => OpenCL::Program pairs");
2444
2445 HV *hv = (HV *)SvRV (devices);
2446
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))
2454 {
2455 header_name [i] = SvPVbyte_nolen (HeSVKEY_force (he));
2456 header_list [i] = SvCLOBJ (cv, "headers", HeVAL (he), "OpenCL::Program");
2457 ++i;
2458 }
2459 }
2460
2461 void *user_data;
2462 program_callback pfn_notify = make_program_callback (notify, &user_data);
2463
1584 NEED_SUCCESS (BuildProgram, (self, 1, &device, SvPVbyte_nolen (options), 0, 0)); 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
1585 2468
1586void 2469void
1587build_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)
1588 PPCODE: 2471 PPCODE:
1589 size_t size; 2472 size_t size;
… …
1598#BEGIN:program_build 2481#BEGIN:program_build
1599 2482
1600void 2483void
1601build_status (OpenCL::Program self, OpenCL::Device device) 2484build_status (OpenCL::Program self, OpenCL::Device device)
1602 PPCODE: 2485 PPCODE:
1603 cl_build_status value [1]; 2486 cl_int value [1];
1604 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));
1605 EXTEND (SP, 1); 2488 EXTEND (SP, 1);
1606 const int i = 0; 2489 const int i = 0;
1607 PUSHs (sv_2mortal (newSViv (value [i]))); 2490 PUSHs (sv_2mortal (newSViv (value [i])));
1608 2491
… …
1618 NEED_SUCCESS (GetProgramBuildInfo, (self, device, ix, size, value, 0)); 2501 NEED_SUCCESS (GetProgramBuildInfo, (self, device, ix, size, value, 0));
1619 EXTEND (SP, 1); 2502 EXTEND (SP, 1);
1620 const int i = 0; 2503 const int i = 0;
1621 PUSHs (sv_2mortal (newSVpv (value, 0))); 2504 PUSHs (sv_2mortal (newSVpv (value, 0)));
1622 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
1623#END:program_build 2515#END:program_build
1624 2516
1625void 2517void
1626kernel (OpenCL::Program program, SV *function) 2518kernel (OpenCL::Program program, SV *function)
1627 PPCODE: 2519 PPCODE:
1628 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));
1629 XPUSH_NEW_OBJ ("OpenCL::Kernel", kernel); 2521 XPUSH_CLOBJ (stash_kernel, kernel);
2522
2523void
2524kernels_in_program (OpenCL::Program program)
2525 PPCODE:
2526 cl_uint num_kernels;
2527 NEED_SUCCESS (CreateKernelsInProgram, (program, 0, 0, &num_kernels));
2528 cl_kernel *kernels = tmpbuf (sizeof (cl_kernel) * num_kernels);
2529 NEED_SUCCESS (CreateKernelsInProgram, (program, num_kernels, kernels, 0));
2530
2531 int i;
2532 EXTEND (SP, num_kernels);
2533 for (i = 0; i < num_kernels; ++i)
2534 PUSH_CLOBJ (stash_kernel, kernels [i]);
1630 2535
1631void 2536void
1632info (OpenCL::Program self, cl_program_info name) 2537info (OpenCL::Program self, cl_program_info name)
1633 PPCODE: 2538 PPCODE:
1634 INFO (Program) 2539 INFO (Program)
… …
1680 PPCODE: 2585 PPCODE:
1681 cl_context value [1]; 2586 cl_context value [1];
1682 NEED_SUCCESS (GetProgramInfo, (self, CL_PROGRAM_CONTEXT, sizeof (value), value, 0)); 2587 NEED_SUCCESS (GetProgramInfo, (self, CL_PROGRAM_CONTEXT, sizeof (value), value, 0));
1683 EXTEND (SP, 1); 2588 EXTEND (SP, 1);
1684 const int i = 0; 2589 const int i = 0;
1685 {
1686 NEED_SUCCESS (RetainContext, (value [i])); 2590 NEED_SUCCESS (RetainContext, (value [i]));
1687 PUSHs (NEW_MORTAL_OBJ ("OpenCL::Context", value [i])); 2591 PUSH_CLOBJ (stash_context, value [i]);
1688 }
1689 2592
1690void 2593void
1691devices (OpenCL::Program self) 2594devices (OpenCL::Program self)
1692 PPCODE: 2595 PPCODE:
1693 size_t size; 2596 size_t size;
… …
1695 cl_device_id *value = tmpbuf (size); 2598 cl_device_id *value = tmpbuf (size);
1696 NEED_SUCCESS (GetProgramInfo, (self, CL_PROGRAM_DEVICES, size, value, 0)); 2599 NEED_SUCCESS (GetProgramInfo, (self, CL_PROGRAM_DEVICES, size, value, 0));
1697 int i, n = size / sizeof (*value); 2600 int i, n = size / sizeof (*value);
1698 EXTEND (SP, n); 2601 EXTEND (SP, n);
1699 for (i = 0; i < n; ++i) 2602 for (i = 0; i < n; ++i)
1700 { 2603 PUSH_CLOBJ (stash_device, value [i]);
1701 PUSHs (NEW_MORTAL_OBJ ("OpenCL::Device", value [i]));
1702 }
1703 2604
1704void 2605void
1705source (OpenCL::Program self) 2606source (OpenCL::Program self)
1706 PPCODE: 2607 PPCODE:
1707 size_t size; 2608 size_t size;
… …
1732DESTROY (OpenCL::Kernel self) 2633DESTROY (OpenCL::Kernel self)
1733 CODE: 2634 CODE:
1734 clReleaseKernel (self); 2635 clReleaseKernel (self);
1735 2636
1736void 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
1737set_char (OpenCL::Kernel self, cl_uint idx, cl_char value) 2703set_char (OpenCL::Kernel self, cl_uint idx, cl_char value)
1738 CODE: 2704 CODE:
1739 clSetKernelArg (self, idx, sizeof (value), &value); 2705 clSetKernelArg (self, idx, sizeof (value), &value);
1740 2706
1741void 2707void
… …
1789 clSetKernelArg (self, idx, sizeof (value), &value); 2755 clSetKernelArg (self, idx, sizeof (value), &value);
1790 2756
1791void 2757void
1792set_memory (OpenCL::Kernel self, cl_uint idx, OpenCL::Memory_ornull value) 2758set_memory (OpenCL::Kernel self, cl_uint idx, OpenCL::Memory_ornull value)
1793 CODE: 2759 CODE:
1794 clSetKernelArg (self, idx, sizeof (value), &value); 2760 clSetKernelArg (self, idx, sizeof (value), value ? &value : 0);
1795 2761
1796void 2762void
1797set_buffer (OpenCL::Kernel self, cl_uint idx, OpenCL::Buffer_ornull value) 2763set_buffer (OpenCL::Kernel self, cl_uint idx, OpenCL::Buffer_ornull value)
1798 CODE: 2764 CODE:
1799 clSetKernelArg (self, idx, sizeof (value), &value); 2765 clSetKernelArg (self, idx, sizeof (value), value ? &value : 0);
1800 2766
1801void 2767void
1802set_image2d (OpenCL::Kernel self, cl_uint idx, OpenCL::Image2D_ornull value) 2768set_image (OpenCL::Kernel self, cl_uint idx, OpenCL::Image_ornull value)
1803 CODE: 2769 CODE:
1804 clSetKernelArg (self, idx, sizeof (value), &value); 2770 clSetKernelArg (self, idx, sizeof (value), value ? &value : 0);
1805
1806void
1807set_image3d (OpenCL::Kernel self, cl_uint idx, OpenCL::Image3D_ornull value)
1808 CODE:
1809 clSetKernelArg (self, idx, sizeof (value), &value);
1810 2771
1811void 2772void
1812set_sampler (OpenCL::Kernel self, cl_uint idx, OpenCL::Sampler value) 2773set_sampler (OpenCL::Kernel self, cl_uint idx, OpenCL::Sampler value)
1813 CODE: 2774 CODE:
1814 clSetKernelArg (self, idx, sizeof (value), &value); 2775 clSetKernelArg (self, idx, sizeof (value), &value);
… …
1858 PPCODE: 2819 PPCODE:
1859 cl_context value [1]; 2820 cl_context value [1];
1860 NEED_SUCCESS (GetKernelInfo, (self, CL_KERNEL_CONTEXT, sizeof (value), value, 0)); 2821 NEED_SUCCESS (GetKernelInfo, (self, CL_KERNEL_CONTEXT, sizeof (value), value, 0));
1861 EXTEND (SP, 1); 2822 EXTEND (SP, 1);
1862 const int i = 0; 2823 const int i = 0;
1863 {
1864 NEED_SUCCESS (RetainContext, (value [i])); 2824 NEED_SUCCESS (RetainContext, (value [i]));
1865 PUSHs (NEW_MORTAL_OBJ ("OpenCL::Context", value [i])); 2825 PUSH_CLOBJ (stash_context, value [i]);
1866 }
1867 2826
1868void 2827void
1869program (OpenCL::Kernel self) 2828program (OpenCL::Kernel self)
1870 PPCODE: 2829 PPCODE:
1871 cl_program value [1]; 2830 cl_program value [1];
1872 NEED_SUCCESS (GetKernelInfo, (self, CL_KERNEL_PROGRAM, sizeof (value), value, 0)); 2831 NEED_SUCCESS (GetKernelInfo, (self, CL_KERNEL_PROGRAM, sizeof (value), value, 0));
1873 EXTEND (SP, 1); 2832 EXTEND (SP, 1);
1874 const int i = 0; 2833 const int i = 0;
1875 {
1876 NEED_SUCCESS (RetainProgram, (value [i])); 2834 NEED_SUCCESS (RetainProgram, (value [i]));
1877 PUSHs (NEW_MORTAL_OBJ ("OpenCL::Program", value [i])); 2835 PUSH_CLOBJ (stash_program, value [i]);
1878 }
1879 2836
1880#END:kernel 2837#END:kernel
1881 2838
1882void 2839void
1883work_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)
… …
1929 const int i = 0; 2886 const int i = 0;
1930 PUSHs (sv_2mortal (newSVuv (value [i]))); 2887 PUSHs (sv_2mortal (newSVuv (value [i])));
1931 2888
1932#END:kernel_work_group 2889#END:kernel_work_group
1933 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
1934MODULE = OpenCL PACKAGE = OpenCL::Event 2946MODULE = OpenCL PACKAGE = OpenCL::Event
1935 2947
1936void 2948void
1937DESTROY (OpenCL::Event self) 2949DESTROY (OpenCL::Event self)
1938 CODE: 2950 CODE:
… …
1940 2952
1941void 2953void
1942wait (OpenCL::Event self) 2954wait (OpenCL::Event self)
1943 CODE: 2955 CODE:
1944 clWaitForEvents (1, &self); 2956 clWaitForEvents (1, &self);
2957
2958void
2959cb (OpenCL::Event self, cl_int command_exec_callback_type, SV *cb)
2960 CODE:
2961 clSetEventCallback (self, command_exec_callback_type, eq_event_notify, SvREFCNT_inc (s_get_cv (cb)));
1945 2962
1946void 2963void
1947info (OpenCL::Event self, cl_event_info name) 2964info (OpenCL::Event self, cl_event_info name)
1948 PPCODE: 2965 PPCODE:
1949 INFO (Event) 2966 INFO (Event)
… …
1955 PPCODE: 2972 PPCODE:
1956 cl_command_queue value [1]; 2973 cl_command_queue value [1];
1957 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));
1958 EXTEND (SP, 1); 2975 EXTEND (SP, 1);
1959 const int i = 0; 2976 const int i = 0;
1960 {
1961 NEED_SUCCESS (RetainCommandQueue, (value [i])); 2977 NEED_SUCCESS (RetainCommandQueue, (value [i]));
1962 PUSHs (NEW_MORTAL_OBJ ("OpenCL::Queue", value [i])); 2978 PUSH_CLOBJ (stash_queue, value [i]);
1963 }
1964 2979
1965void 2980void
1966command_type (OpenCL::Event self) 2981command_type (OpenCL::Event self)
1967 PPCODE:
1968 cl_command_type value [1];
1969 NEED_SUCCESS (GetEventInfo, (self, CL_EVENT_COMMAND_TYPE, sizeof (value), value, 0));
1970 EXTEND (SP, 1);
1971 const int i = 0;
1972 PUSHs (sv_2mortal (newSVuv (value [i])));
1973
1974void
1975reference_count (OpenCL::Event self)
1976 ALIAS: 2982 ALIAS:
2983 command_type = CL_EVENT_COMMAND_TYPE
1977 reference_count = CL_EVENT_REFERENCE_COUNT 2984 reference_count = CL_EVENT_REFERENCE_COUNT
1978 command_execution_status = CL_EVENT_COMMAND_EXECUTION_STATUS 2985 command_execution_status = CL_EVENT_COMMAND_EXECUTION_STATUS
1979 PPCODE: 2986 PPCODE:
1980 cl_uint value [1]; 2987 cl_uint value [1];
1981 NEED_SUCCESS (GetEventInfo, (self, ix, sizeof (value), value, 0)); 2988 NEED_SUCCESS (GetEventInfo, (self, ix, sizeof (value), value, 0));
… …
1988 PPCODE: 2995 PPCODE:
1989 cl_context value [1]; 2996 cl_context value [1];
1990 NEED_SUCCESS (GetEventInfo, (self, CL_EVENT_CONTEXT, sizeof (value), value, 0)); 2997 NEED_SUCCESS (GetEventInfo, (self, CL_EVENT_CONTEXT, sizeof (value), value, 0));
1991 EXTEND (SP, 1); 2998 EXTEND (SP, 1);
1992 const int i = 0; 2999 const int i = 0;
1993 {
1994 NEED_SUCCESS (RetainContext, (value [i])); 3000 NEED_SUCCESS (RetainContext, (value [i]));
1995 PUSHs (NEW_MORTAL_OBJ ("OpenCL::Context", value [i])); 3001 PUSH_CLOBJ (stash_context, value [i]);
1996 }
1997 3002
1998#END:event 3003#END:event
1999 3004
2000void 3005void
2001profiling_info (OpenCL::Event self, cl_profiling_info name) 3006profiling_info (OpenCL::Event self, cl_profiling_info name)
… …
2025void 3030void
2026set_status (OpenCL::UserEvent self, cl_int execution_status) 3031set_status (OpenCL::UserEvent self, cl_int execution_status)
2027 CODE: 3032 CODE:
2028 clSetUserEventStatus (self, execution_status); 3033 clSetUserEventStatus (self, execution_status);
2029 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