1 ! Copyright (C) 2010 Erik Charlebois.
2 ! See http://factorcode.org/license.txt for BSD license.
3 USING: accessors alien alien.c-types alien.data arrays
4 byte-arrays combinators combinators.smart destructors
5 io.encodings.ascii io.encodings.string kernel libc locals make
6 math namespaces opencl.ffi sequences shuffle specialized-arrays
9 SPECIALIZED-ARRAYS: void* char size_t ;
14 : cl-success ( err -- )
15 dup CL_SUCCESS = [ drop ] [ cl-error ] if ; inline
17 : cl-not-null ( err -- )
18 dup f = [ cl-error ] [ drop ] if ; inline
20 : info-data-size ( handle name info-quot -- size_t )
21 [ 0 f 0 size_t <ref> ] dip [ call cl-success ] 2keep drop size_t deref ; inline
23 : info-data-bytes ( handle name info-quot size -- bytes )
24 swap [ dup <byte-array> f ] dip [ call cl-success ] 3keep 2drop ; inline
26 : info ( handle name info-quot lift-quot -- value )
27 [ 3dup info-data-size info-data-bytes ] dip call ; inline
29 : 2info-data-size ( handle1 handle2 name info-quot -- size_t )
30 [ 0 f 0 size_t <ref> ] dip [ call cl-success ] 2keep drop size_t deref ; inline
32 : 2info-data-bytes ( handle1 handle2 name info-quot size -- bytes )
33 swap [ dup <byte-array> f ] dip [ call cl-success ] 3keep 2drop ; inline
35 : 2info ( handle1 handle2 name info_quot lift_quot -- value )
36 [ 4dup 2info-data-size 2info-data-bytes ] dip call ; inline
38 : info-bool ( handle name quot -- ? )
39 [ uint deref CL_TRUE = ] info ; inline
41 : info-ulong ( handle name quot -- ulong )
42 [ ulonglong deref ] info ; inline
44 : info-int ( handle name quot -- int )
45 [ int deref ] info ; inline
47 : info-uint ( handle name quot -- uint )
48 [ uint deref ] info ; inline
50 : info-size_t ( handle name quot -- size_t )
51 [ size_t deref ] info ; inline
53 : 2info-size_t ( handle1 handle2 name quot -- size_t )
54 [ size_t deref ] 2info ; inline
56 : info-string ( handle name quot -- string )
57 [ ascii decode but-last ] info ; inline
59 : 2info-string ( handle name quot -- string )
60 [ ascii decode but-last ] 2info ; inline
62 : info-size_t-array ( handle name quot -- size_t-array )
63 [ [ length size_t heap-size / ] keep swap size_t <c-direct-array> ] info ; inline
65 TUPLE: cl-handle < disposable handle ;
69 VARIANT: cl-device-type
70 cl-device-default cl-device-cpu cl-device-gpu cl-device-accelerator ;
72 : size_t>cl-device-type ( size_t -- cl-device-type )
74 { CL_DEVICE_TYPE_DEFAULT [ cl-device-default ] }
75 { CL_DEVICE_TYPE_CPU [ cl-device-cpu ] }
76 { CL_DEVICE_TYPE_GPU [ cl-device-gpu ] }
77 { CL_DEVICE_TYPE_ACCELERATOR [ cl-device-accelerator ] }
80 VARIANT: cl-fp-feature
81 cl-denorm cl-inf-and-nan cl-round-to-nearest cl-round-to-zero cl-round-to-inf cl-fma ;
83 VARIANT: cl-cache-type
84 cl-no-cache cl-read-only-cache cl-read-write-cache ;
86 VARIANT: cl-buffer-access-mode
87 cl-read-access cl-write-access cl-read-write-access ;
89 VARIANT: cl-image-channel-order
90 cl-channel-order-r cl-channel-order-a cl-channel-order-rg cl-channel-order-ra
91 cl-channel-order-rga cl-channel-order-rgba cl-channel-order-bgra cl-channel-order-argb
92 cl-channel-order-intensity cl-channel-order-luminance ;
94 VARIANT: cl-image-channel-type
95 cl-channel-type-snorm-int8 cl-channel-type-snorm-int16 cl-channel-type-unorm-int8
96 cl-channel-type-unorm-int16 cl-channel-type-unorm-short-565
97 cl-channel-type-unorm-short-555 cl-channel-type-unorm-int-101010
98 cl-channel-type-signed-int8 cl-channel-type-signed-int16 cl-channel-type-signed-int32
99 cl-channel-type-unsigned-int8 cl-channel-type-unsigned-int16
100 cl-channel-type-unsigned-int32 cl-channel-type-half-float cl-channel-type-float ;
102 VARIANT: cl-addressing-mode
103 cl-repeat-addressing cl-clamp-to-edge-addressing cl-clamp-addressing cl-no-addressing ;
105 VARIANT: cl-filter-mode
106 cl-filter-nearest cl-filter-linear ;
108 VARIANT: cl-command-type
109 cl-ndrange-kernel-command cl-task-command cl-native-kernel-command cl-read-buffer-command
110 cl-write-buffer-command cl-copy-buffer-command cl-read-image-command cl-write-image-command
111 cl-copy-image-command cl-copy-buffer-to-image-command cl-copy-image-to-buffer-command
112 cl-map-buffer-command cl-map-image-command cl-unmap-mem-object-command
113 cl-marker-command cl-acquire-gl-objects-command cl-release-gl-objects-command ;
115 VARIANT: cl-execution-status
116 cl-queued cl-submitted cl-running cl-complete cl-failure ;
119 id profile version name vendor extensions devices ;
122 id type vendor-id max-compute-units max-work-item-dimensions
123 max-work-item-sizes max-work-group-size preferred-vector-width-char
124 preferred-vector-width-short preferred-vector-width-int
125 preferred-vector-width-long preferred-vector-width-float
126 preferred-vector-width-double max-clock-frequency address-bits
127 max-mem-alloc-size image-support max-read-image-args max-write-image-args
128 image2d-max-width image2d-max-height image3d-max-width image3d-max-height
129 image3d-max-depth max-samplers max-parameter-size mem-base-addr-align
130 min-data-type-align-size single-fp-config global-mem-cache-type
131 global-mem-cacheline-size global-mem-cache-size global-mem-size
132 max-constant-buffer-size max-constant-args local-mem? local-mem-size
133 error-correction-support profiling-timer-resolution endian-little
134 available compiler-available execute-kernels? execute-native-kernels?
135 out-of-order-exec-available? profiling-available?
136 name vendor driver-version profile version extensions ;
138 TUPLE: cl-context < cl-handle ;
139 TUPLE: cl-queue < cl-handle ;
140 TUPLE: cl-buffer < cl-handle ;
141 TUPLE: cl-sampler < cl-handle ;
142 TUPLE: cl-program < cl-handle ;
143 TUPLE: cl-kernel < cl-handle ;
144 TUPLE: cl-event < cl-handle ;
146 M: cl-context dispose* handle>> clReleaseContext cl-success ;
147 M: cl-queue dispose* handle>> clReleaseCommandQueue cl-success ;
148 M: cl-buffer dispose* handle>> clReleaseMemObject cl-success ;
149 M: cl-sampler dispose* handle>> clReleaseSampler cl-success ;
150 M: cl-program dispose* handle>> clReleaseProgram cl-success ;
151 M: cl-kernel dispose* handle>> clReleaseKernel cl-success ;
152 M: cl-event dispose* handle>> clReleaseEvent cl-success ;
155 { buffer cl-buffer read-only }
156 { offset integer read-only } ;
157 C: <cl-buffer-ptr> cl-buffer-ptr
159 TUPLE: cl-buffer-range
160 { buffer cl-buffer read-only }
161 { offset integer read-only }
162 { size integer read-only } ;
163 C: <cl-buffer-range> cl-buffer-range
165 SYMBOLS: cl-current-context cl-current-queue cl-current-device ;
169 : (current-cl-context) ( -- cl-context )
170 cl-current-context get ; inline
172 : (current-cl-queue) ( -- cl-queue )
173 cl-current-queue get ; inline
175 : (current-cl-device) ( -- cl-device )
176 cl-current-device get ; inline
178 GENERIC: buffer-access-constant ( buffer-access-mode -- n )
179 M: cl-read-write-access buffer-access-constant drop CL_MEM_READ_WRITE ;
180 M: cl-read-access buffer-access-constant drop CL_MEM_READ_ONLY ;
181 M: cl-write-access buffer-access-constant drop CL_MEM_WRITE_ONLY ;
183 GENERIC: buffer-map-flags ( buffer-access-mode -- n )
184 M: cl-read-write-access buffer-map-flags drop CL_MAP_READ CL_MAP_WRITE bitor ;
185 M: cl-read-access buffer-map-flags drop CL_MAP_READ ;
186 M: cl-write-access buffer-map-flags drop CL_MAP_WRITE ;
188 GENERIC: addressing-mode-constant ( addressing-mode -- n )
189 M: cl-repeat-addressing addressing-mode-constant drop CL_ADDRESS_REPEAT ;
190 M: cl-clamp-to-edge-addressing addressing-mode-constant drop CL_ADDRESS_CLAMP_TO_EDGE ;
191 M: cl-clamp-addressing addressing-mode-constant drop CL_ADDRESS_CLAMP ;
192 M: cl-no-addressing addressing-mode-constant drop CL_ADDRESS_NONE ;
194 GENERIC: filter-mode-constant ( filter-mode -- n )
195 M: cl-filter-nearest filter-mode-constant drop CL_FILTER_NEAREST ;
196 M: cl-filter-linear filter-mode-constant drop CL_FILTER_LINEAR ;
198 : cl_addressing_mode>addressing-mode ( cl_addressing_mode -- addressing-mode )
200 { CL_ADDRESS_REPEAT [ cl-repeat-addressing ] }
201 { CL_ADDRESS_CLAMP_TO_EDGE [ cl-clamp-to-edge-addressing ] }
202 { CL_ADDRESS_CLAMP [ cl-clamp-addressing ] }
203 { CL_ADDRESS_NONE [ cl-no-addressing ] }
206 : cl_filter_mode>filter-mode ( cl_filter_mode -- filter-mode )
208 { CL_FILTER_LINEAR [ cl-filter-linear ] }
209 { CL_FILTER_NEAREST [ cl-filter-nearest ] }
212 : platform-info-string ( handle name -- string )
213 [ clGetPlatformInfo ] info-string ;
215 : platform-info ( id -- profile version name vendor extensions )
217 [ CL_PLATFORM_PROFILE platform-info-string ]
218 [ CL_PLATFORM_VERSION platform-info-string ]
219 [ CL_PLATFORM_NAME platform-info-string ]
220 [ CL_PLATFORM_VENDOR platform-info-string ]
221 [ CL_PLATFORM_EXTENSIONS platform-info-string ]
224 : cl_device_fp_config>flags ( ulong -- sequence )
226 [ CL_FP_DENORM bitand 0 = [ f ] [ cl-denorm ] if ]
227 [ CL_FP_INF_NAN bitand 0 = [ f ] [ cl-inf-and-nan ] if ]
228 [ CL_FP_ROUND_TO_NEAREST bitand 0 = [ f ] [ cl-round-to-nearest ] if ]
229 [ CL_FP_ROUND_TO_ZERO bitand 0 = [ f ] [ cl-round-to-zero ] if ]
230 [ CL_FP_ROUND_TO_INF bitand 0 = [ f ] [ cl-round-to-inf ] if ]
231 [ CL_FP_FMA bitand 0 = [ f ] [ cl-fma ] if ]
232 } cleave ] { } output>sequence sift ;
234 : cl_device_mem_cache_type>cache-type ( uint -- cache-type )
236 { CL_NONE [ cl-no-cache ] }
237 { CL_READ_ONLY_CACHE [ cl-read-only-cache ] }
238 { CL_READ_WRITE_CACHE [ cl-read-write-cache ] }
241 : device-info-bool ( handle name -- ? )
242 [ clGetDeviceInfo ] info-bool ;
244 : device-info-ulong ( handle name -- ulong )
245 [ clGetDeviceInfo ] info-ulong ;
247 : device-info-uint ( handle name -- uint )
248 [ clGetDeviceInfo ] info-uint ;
250 : device-info-string ( handle name -- string )
251 [ clGetDeviceInfo ] info-string ;
253 : device-info-size_t ( handle name -- size_t )
254 [ clGetDeviceInfo ] info-size_t ;
256 : device-info-size_t-array ( handle name -- size_t-array )
257 [ clGetDeviceInfo ] info-size_t-array ;
259 : device-info ( device-id -- device )
261 [ CL_DEVICE_TYPE device-info-size_t size_t>cl-device-type ]
262 [ CL_DEVICE_VENDOR_ID device-info-uint ]
263 [ CL_DEVICE_MAX_COMPUTE_UNITS device-info-uint ]
264 [ CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS device-info-uint ]
265 [ CL_DEVICE_MAX_WORK_ITEM_SIZES device-info-size_t-array ]
266 [ CL_DEVICE_MAX_WORK_GROUP_SIZE device-info-size_t ]
267 [ CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR device-info-uint ]
268 [ CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT device-info-uint ]
269 [ CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT device-info-uint ]
270 [ CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG device-info-uint ]
271 [ CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT device-info-uint ]
272 [ CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE device-info-uint ]
273 [ CL_DEVICE_MAX_CLOCK_FREQUENCY device-info-uint ]
274 [ CL_DEVICE_ADDRESS_BITS device-info-uint ]
275 [ CL_DEVICE_MAX_MEM_ALLOC_SIZE device-info-ulong ]
276 [ CL_DEVICE_IMAGE_SUPPORT device-info-bool ]
277 [ CL_DEVICE_MAX_READ_IMAGE_ARGS device-info-uint ]
278 [ CL_DEVICE_MAX_WRITE_IMAGE_ARGS device-info-uint ]
279 [ CL_DEVICE_IMAGE2D_MAX_WIDTH device-info-size_t ]
280 [ CL_DEVICE_IMAGE2D_MAX_HEIGHT device-info-size_t ]
281 [ CL_DEVICE_IMAGE3D_MAX_WIDTH device-info-size_t ]
282 [ CL_DEVICE_IMAGE3D_MAX_HEIGHT device-info-size_t ]
283 [ CL_DEVICE_IMAGE3D_MAX_DEPTH device-info-size_t ]
284 [ CL_DEVICE_MAX_SAMPLERS device-info-uint ]
285 [ CL_DEVICE_MAX_PARAMETER_SIZE device-info-size_t ]
286 [ CL_DEVICE_MEM_BASE_ADDR_ALIGN device-info-uint ]
287 [ CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE device-info-uint ]
288 [ CL_DEVICE_SINGLE_FP_CONFIG device-info-ulong cl_device_fp_config>flags ]
289 [ CL_DEVICE_GLOBAL_MEM_CACHE_TYPE device-info-uint cl_device_mem_cache_type>cache-type ]
290 [ CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE device-info-uint ]
291 [ CL_DEVICE_GLOBAL_MEM_CACHE_SIZE device-info-ulong ]
292 [ CL_DEVICE_GLOBAL_MEM_SIZE device-info-ulong ]
293 [ CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE device-info-ulong ]
294 [ CL_DEVICE_MAX_CONSTANT_ARGS device-info-uint ]
295 [ CL_DEVICE_LOCAL_MEM_TYPE device-info-uint CL_LOCAL = ]
296 [ CL_DEVICE_LOCAL_MEM_SIZE device-info-ulong ]
297 [ CL_DEVICE_ERROR_CORRECTION_SUPPORT device-info-bool ]
298 [ CL_DEVICE_PROFILING_TIMER_RESOLUTION device-info-size_t ]
299 [ CL_DEVICE_ENDIAN_LITTLE device-info-bool ]
300 [ CL_DEVICE_AVAILABLE device-info-bool ]
301 [ CL_DEVICE_COMPILER_AVAILABLE device-info-bool ]
302 [ CL_DEVICE_EXECUTION_CAPABILITIES device-info-ulong CL_EXEC_KERNEL bitand 0 = not ]
303 [ CL_DEVICE_EXECUTION_CAPABILITIES device-info-ulong CL_EXEC_NATIVE_KERNEL bitand 0 = not ]
304 [ CL_DEVICE_QUEUE_PROPERTIES device-info-ulong CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE bitand 0 = not ]
305 [ CL_DEVICE_QUEUE_PROPERTIES device-info-ulong CL_QUEUE_PROFILING_ENABLE bitand 0 = not ]
306 [ CL_DEVICE_NAME device-info-string ]
307 [ CL_DEVICE_VENDOR device-info-string ]
308 [ CL_DRIVER_VERSION device-info-string ]
309 [ CL_DEVICE_PROFILE device-info-string ]
310 [ CL_DEVICE_VERSION device-info-string ]
311 [ CL_DEVICE_EXTENSIONS device-info-string ]
312 } cleave cl-device boa ;
314 : platform-devices ( platform-id -- devices )
316 0 f 0 uint <ref> [ clGetDeviceIDs cl-success ] keep uint deref
318 rot dup void* <c-array> [ f clGetDeviceIDs cl-success ] keep
321 : command-queue-info-ulong ( handle name -- ulong )
322 [ clGetCommandQueueInfo ] info-ulong ;
324 : sampler-info-bool ( handle name -- ? )
325 [ clGetSamplerInfo ] info-bool ;
327 : sampler-info-uint ( handle name -- uint )
328 [ clGetSamplerInfo ] info-uint ;
330 : program-build-info-string ( program-handle device-handle name -- string )
331 [ clGetProgramBuildInfo ] 2info-string ;
333 : program-build-log ( program-handle device-handle -- string )
334 CL_PROGRAM_BUILD_LOG program-build-info-string ;
336 : strings>char*-array ( strings -- char*-array )
337 [ ascii encode dup length dup malloc [ cl-not-null ]
338 keep &free [ -rot memcpy ] keep ] void*-array{ } map-as ;
340 : (program) ( cl-context sources -- program-handle )
343 [ strings>char*-array ]
344 [ [ length ] size_t-array{ } map-as ] tri
345 0 int <ref> [ clCreateProgramWithSource ] keep int deref cl-success
348 :: (build-program) ( program-handle device options -- program )
349 program-handle 1 device 1array [ id>> ] void*-array{ } map-as
350 options ascii encode 0 suffix f f clBuildProgram
352 { CL_BUILD_PROGRAM_FAILURE [
353 program-handle device id>> program-build-log program-handle
354 clReleaseProgram cl-success cl-error f ] }
355 { CL_SUCCESS [ cl-program new-disposable program-handle >>handle ] }
356 [ program-handle clReleaseProgram cl-success cl-success f ]
359 : kernel-info-string ( handle name -- string )
360 [ clGetKernelInfo ] info-string ;
362 : kernel-info-uint ( handle name -- uint )
363 [ clGetKernelInfo ] info-uint ;
365 : kernel-work-group-info-size_t ( handle1 handle2 name -- size_t )
366 [ clGetKernelWorkGroupInfo ] 2info-size_t ;
368 : event-info-uint ( handle name -- uint )
369 [ clGetEventInfo ] info-uint ;
371 : event-info-int ( handle name -- int )
372 [ clGetEventInfo ] info-int ;
374 : cl_command_type>command-type ( cl_command-type -- command-type )
376 { CL_COMMAND_NDRANGE_KERNEL [ cl-ndrange-kernel-command ] }
377 { CL_COMMAND_TASK [ cl-task-command ] }
378 { CL_COMMAND_NATIVE_KERNEL [ cl-native-kernel-command ] }
379 { CL_COMMAND_READ_BUFFER [ cl-read-buffer-command ] }
380 { CL_COMMAND_WRITE_BUFFER [ cl-write-buffer-command ] }
381 { CL_COMMAND_COPY_BUFFER [ cl-copy-buffer-command ] }
382 { CL_COMMAND_READ_IMAGE [ cl-read-image-command ] }
383 { CL_COMMAND_WRITE_IMAGE [ cl-write-image-command ] }
384 { CL_COMMAND_COPY_IMAGE [ cl-copy-image-command ] }
385 { CL_COMMAND_COPY_BUFFER_TO_IMAGE [ cl-copy-buffer-to-image-command ] }
386 { CL_COMMAND_COPY_IMAGE_TO_BUFFER [ cl-copy-image-to-buffer-command ] }
387 { CL_COMMAND_MAP_BUFFER [ cl-map-buffer-command ] }
388 { CL_COMMAND_MAP_IMAGE [ cl-map-image-command ] }
389 { CL_COMMAND_UNMAP_MEM_OBJECT [ cl-unmap-mem-object-command ] }
390 { CL_COMMAND_MARKER [ cl-marker-command ] }
391 { CL_COMMAND_ACQUIRE_GL_OBJECTS [ cl-acquire-gl-objects-command ] }
392 { CL_COMMAND_RELEASE_GL_OBJECTS [ cl-release-gl-objects-command ] }
395 : cl_int>execution-status ( clint -- execution-status )
397 { CL_QUEUED [ cl-queued ] }
398 { CL_SUBMITTED [ cl-submitted ] }
399 { CL_RUNNING [ cl-running ] }
400 { CL_COMPLETE [ cl-complete ] }
404 : profiling-info-ulong ( handle name -- ulong )
405 [ clGetEventProfilingInfo ] info-ulong ;
407 : bind-kernel-arg-buffer ( kernel index buffer -- )
408 [ handle>> ] [ cl_mem heap-size ] [ handle>> void* deref ] tri*
409 clSetKernelArg cl-success ; inline
411 : bind-kernel-arg-data ( kernel index byte-array -- )
413 [ byte-length ] keep clSetKernelArg cl-success ; inline
415 GENERIC: bind-kernel-arg ( kernel index data -- )
416 M: cl-buffer bind-kernel-arg bind-kernel-arg-buffer ;
417 M: byte-array bind-kernel-arg bind-kernel-arg-data ;
421 : with-cl-state ( context/f device/f queue/f quot -- )
424 [ cl-current-queue ,, ] when*
425 [ cl-current-device ,, ] when*
426 [ cl-current-context ,, ] when*
428 ] dip with-variables ; inline
430 : cl-platforms ( -- platforms )
431 0 f 0 uint <ref> [ clGetPlatformIDs cl-success ] keep uint deref
432 dup void* <c-array> [ f clGetPlatformIDs cl-success ] keep
436 [ platform-devices [ device-info ] { } map-as ] bi
440 : <cl-context> ( devices -- cl-context )
442 [ length ] [ [ id>> ] void*-array{ } map-as ] bi
443 f f 0 int <ref> [ clCreateContext ] keep int deref cl-success
444 cl-context new-disposable swap >>handle ;
446 : <cl-queue> ( context device out-of-order? profiling? -- command-queue )
447 [ [ handle>> ] [ id>> ] bi* ] 2dip
448 [ [ CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE ] [ 0 ] if ]
449 [ [ CL_QUEUE_PROFILING_ENABLE ] [ 0 ] if ] bi* bitor
450 0 int <ref> [ clCreateCommandQueue ] keep int deref cl-success
451 cl-queue new-disposable swap >>handle ;
453 : cl-out-of-order-execution? ( command-queue -- ? )
454 CL_QUEUE_PROPERTIES command-queue-info-ulong
455 CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE bitand 0 = not ; inline
457 : cl-profiling? ( command-queue -- ? )
458 CL_QUEUE_PROPERTIES command-queue-info-ulong
459 CL_QUEUE_PROFILING_ENABLE bitand 0 = not ; inline
461 : <cl-buffer> ( buffer-access-mode size initial-data -- buffer )
462 [ (current-cl-context) ] 3dip
465 [ buffer-access-constant ]
466 [ [ CL_MEM_COPY_HOST_PTR ] [ CL_MEM_ALLOC_HOST_PTR ] if ] tri* bitor
468 0 int <ref> [ clCreateBuffer ] keep int deref cl-success
469 cl-buffer new-disposable swap >>handle ;
471 : cl-read-buffer ( buffer-range -- byte-array )
472 [ (current-cl-queue) handle>> ] dip
473 [ buffer>> handle>> CL_TRUE ]
475 [ size>> dup <byte-array> ] tri
476 [ 0 f f clEnqueueReadBuffer cl-success ] keep ; inline
478 : cl-write-buffer ( buffer-range byte-array -- )
480 [ (current-cl-queue) handle>> ] dip
481 [ buffer>> handle>> CL_TRUE ]
484 ] dip 0 f f clEnqueueWriteBuffer cl-success ; inline
486 : cl-queue-copy-buffer ( src-buffer-ptr dst-buffer-ptr size dependent-events -- event )
490 [ [ buffer>> handle>> ] [ offset>> ] bi ]
491 [ [ buffer>> handle>> ] [ offset>> ] bi ]
493 ] 2dip [ length ] keep [ f ] [ [ handle>> ] void*-array{ } map-as ] if-empty
494 f void* <ref> [ clEnqueueCopyBuffer cl-success ] keep void* deref cl-event
495 new-disposable swap >>handle ;
497 : cl-queue-read-buffer ( buffer-range alien dependent-events -- event )
499 [ (current-cl-queue) handle>> ] dip
500 [ buffer>> handle>> CL_FALSE ] [ offset>> ] [ size>> ] tri
501 ] 2dip [ length ] keep [ f ] [ [ handle>> ] void*-array{ } map-as ] if-empty
502 f void* <ref> [ clEnqueueReadBuffer cl-success ] keep void* <ref> cl-event
503 new-disposable swap >>handle ;
505 : cl-queue-write-buffer ( buffer-range alien dependent-events -- event )
507 [ (current-cl-queue) handle>> ] dip
508 [ buffer>> handle>> CL_FALSE ] [ offset>> ] [ size>> ] tri
509 ] 2dip [ length ] keep [ f ] [ [ handle>> ] void*-array{ } map-as ] if-empty
510 f void* <ref> [ clEnqueueWriteBuffer cl-success ] keep void* deref cl-event
511 new-disposable swap >>handle ;
513 : <cl-sampler> ( normalized-coords? addressing-mode filter-mode -- sampler )
514 [ (current-cl-context) ] 3dip
515 [ [ CL_TRUE ] [ CL_FALSE ] if ]
516 [ addressing-mode-constant ]
517 [ filter-mode-constant ]
518 tri* 0 int <ref> [ clCreateSampler ] keep int deref cl-success
519 cl-sampler new-disposable swap >>handle ;
521 : cl-normalized-coords? ( sampler -- ? )
522 handle>> CL_SAMPLER_NORMALIZED_COORDS sampler-info-bool ; inline
524 : sampler>cl-addressing-mode ( sampler -- addressing-mode )
525 handle>> CL_SAMPLER_ADDRESSING_MODE sampler-info-uint cl_addressing_mode>addressing-mode ; inline
527 : sampler>cl-filter-mode ( sampler -- filter-mode )
528 handle>> CL_SAMPLER_FILTER_MODE sampler-info-uint cl_filter_mode>filter-mode ; inline
530 : <cl-program> ( options strings -- program )
531 [ (current-cl-device) ] 2dip
532 [ (current-cl-context) ] dip
533 (program) -rot (build-program) ;
535 : <cl-kernel> ( program kernel-name -- kernel )
536 [ handle>> ] [ ascii encode 0 suffix ] bi*
537 0 int <ref> [ clCreateKernel ] keep int deref cl-success
538 cl-kernel new-disposable swap >>handle ; inline
540 : cl-kernel-name ( kernel -- string )
541 handle>> CL_KERNEL_FUNCTION_NAME kernel-info-string ;
543 : cl-kernel-arity ( kernel -- arity )
544 handle>> CL_KERNEL_NUM_ARGS kernel-info-uint ;
546 : cl-kernel-local-size ( kernel -- size )
547 (current-cl-device) [ handle>> ] bi@ CL_KERNEL_WORK_GROUP_SIZE kernel-work-group-info-size_t ; inline
549 :: cl-queue-kernel ( kernel args sizes dependent-events -- event )
550 args [| arg idx | kernel idx arg bind-kernel-arg ] each-index
551 (current-cl-queue) handle>>
553 sizes [ length f ] [ [ ] size_t-array{ } map-as f ] bi
554 dependent-events [ length ] [ [ f ] [ [ handle>> ] void*-array{ } map-as ] if-empty ] bi
555 f void* <ref> [ clEnqueueNDRangeKernel cl-success ] keep void* deref
556 cl-event new-disposable swap >>handle ;
558 : cl-event-type ( event -- command-type )
559 handle>> CL_EVENT_COMMAND_TYPE event-info-uint cl_command_type>command-type ; inline
561 : cl-event-status ( event -- execution-status )
562 handle>> CL_EVENT_COMMAND_EXECUTION_STATUS event-info-int cl_int>execution-status ; inline
564 : cl-profile-counters ( event -- queued submitted started finished )
566 [ CL_PROFILING_COMMAND_QUEUED profiling-info-ulong ]
567 [ CL_PROFILING_COMMAND_SUBMIT profiling-info-ulong ]
568 [ CL_PROFILING_COMMAND_START profiling-info-ulong ]
569 [ CL_PROFILING_COMMAND_END profiling-info-ulong ]
572 : cl-barrier-events ( event/events -- )
573 [ (current-cl-queue) handle>> ] dip
574 dup sequence? [ 1array ] unless
575 [ handle>> ] void*-array{ } map-as [ length ] keep clEnqueueWaitForEvents cl-success ; inline
577 : cl-marker ( -- event )
579 f void* <ref> [ clEnqueueMarker cl-success ] keep void* deref cl-event new-disposable
580 swap >>handle ; inline
583 (current-cl-queue) clEnqueueBarrier cl-success ; inline
586 (current-cl-queue) handle>> clFlush cl-success ; inline
588 : cl-wait ( event/events -- )
589 dup sequence? [ 1array ] unless
590 [ handle>> ] void*-array{ } map-as [ length ] keep clWaitForEvents cl-success ; inline
593 (current-cl-queue) handle>> clFinish cl-success ; inline