extra/opencl/opencl.factor

   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 specialized-arrays
   7 variants ;
   8 IN: opencl
   9 SPECIALIZED-ARRAYS: void* char size_t ;
  10
  11 <PRIVATE
  12
  13 ERROR: cl-error err ;
  14
  15 : cl-success ( err -- )
  16     dup CL_SUCCESS = [ drop ] [ cl-error ] if ; inline
  17
  18 : cl-not-null ( err -- )
  19     dup f = [ cl-error ] [ drop ] if ; inline
  20
  21 : info-data-size ( handle name info-quot -- size_t )
  22     [ 0 f 0 size_t <ref> ] dip [ call cl-success ] 2keep drop size_t deref ; inline
  23
  24 : info-data-bytes ( handle name info-quot size -- bytes )
  25     swap [ dup <byte-array> f ] dip [ call cl-success ] 3keep 2drop ; inline
  26
  27 : info ( handle name info-quot lift-quot -- value )
  28     [ 3dup info-data-size info-data-bytes ] dip call ; inline
  29
  30 : 2info-data-size ( handle1 handle2 name info-quot -- size_t )
  31     [ 0 f 0 size_t <ref> ] dip [ call cl-success ] 2keep drop size_t deref ; inline
  32
  33 : 2info-data-bytes ( handle1 handle2 name info-quot size -- bytes )
  34     swap [ dup <byte-array> f ] dip [ call cl-success ] 3keep 2drop ; inline
  35
  36 : 2info ( handle1 handle2 name info_quot lift_quot -- value )
  37     [ 4dup 2info-data-size 2info-data-bytes ] dip call ; inline
  38
  39 : info-bool ( handle name quot -- ? )
  40     [ uint deref CL_TRUE = ] info ; inline
  41
  42 : info-ulong ( handle name quot -- ulong )
  43     [ ulonglong deref ] info ; inline
  44
  45 : info-int ( handle name quot -- int )
  46     [ int deref ] info ; inline
  47
  48 : info-uint ( handle name quot -- uint )
  49     [ uint deref ] info ; inline
  50
  51 : info-size_t ( handle name quot -- size_t )
  52     [ size_t deref ] info ; inline
  53
  54 : 2info-size_t ( handle1 handle2 name quot -- size_t )
  55     [ size_t deref ] 2info ; inline
  56
  57 : info-string ( handle name quot -- string )
  58     [ ascii decode but-last ] info ; inline
  59
  60 : 2info-string ( handle name quot -- string )
  61     [ ascii decode but-last ] 2info ; inline
  62
  63 : info-size_t-array ( handle name quot -- size_t-array )
  64     [ [ length size_t heap-size / ] keep swap size_t <c-direct-array> ] info ; inline
  65
  66 TUPLE: cl-handle < disposable handle ;
  67
  68 PRIVATE>
  69
  70 VARIANT: cl-device-type
  71     cl-device-default cl-device-cpu cl-device-gpu cl-device-accelerator ;
  72
  73 : size_t>cl-device-type ( size_t -- cl-device-type )
  74     {
  75         { CL_DEVICE_TYPE_DEFAULT     [ cl-device-default     ] }
  76         { CL_DEVICE_TYPE_CPU         [ cl-device-cpu         ] }
  77         { CL_DEVICE_TYPE_GPU         [ cl-device-gpu         ] }
  78         { CL_DEVICE_TYPE_ACCELERATOR [ cl-device-accelerator ] }
  79     } case ; inline
  80
  81 VARIANT: cl-fp-feature
  82     cl-denorm cl-inf-and-nan cl-round-to-nearest cl-round-to-zero cl-round-to-inf cl-fma ;
  83
  84 VARIANT: cl-cache-type
  85     cl-no-cache cl-read-only-cache cl-read-write-cache ;
  86
  87 VARIANT: cl-buffer-access-mode
  88     cl-read-access cl-write-access cl-read-write-access ;
  89
  90 VARIANT: cl-image-channel-order
  91     cl-channel-order-r cl-channel-order-a cl-channel-order-rg cl-channel-order-ra
  92     cl-channel-order-rga cl-channel-order-rgba cl-channel-order-bgra cl-channel-order-argb
  93     cl-channel-order-intensity cl-channel-order-luminance ;
  94
  95 VARIANT: cl-image-channel-type
  96     cl-channel-type-snorm-int8 cl-channel-type-snorm-int16 cl-channel-type-unorm-int8
  97     cl-channel-type-unorm-int16 cl-channel-type-unorm-short-565
  98     cl-channel-type-unorm-short-555 cl-channel-type-unorm-int-101010
  99     cl-channel-type-signed-int8 cl-channel-type-signed-int16 cl-channel-type-signed-int32
 100     cl-channel-type-unsigned-int8 cl-channel-type-unsigned-int16
 101     cl-channel-type-unsigned-int32 cl-channel-type-half-float cl-channel-type-float ;
 102
 103 VARIANT: cl-addressing-mode
 104     cl-repeat-addressing cl-clamp-to-edge-addressing cl-clamp-addressing cl-no-addressing ;
 105
 106 VARIANT: cl-filter-mode
 107     cl-filter-nearest cl-filter-linear ;
 108
 109 VARIANT: cl-command-type
 110     cl-ndrange-kernel-command cl-task-command cl-native-kernel-command cl-read-buffer-command
 111     cl-write-buffer-command cl-copy-buffer-command cl-read-image-command cl-write-image-command
 112     cl-copy-image-command cl-copy-buffer-to-image-command cl-copy-image-to-buffer-command
 113     cl-map-buffer-command cl-map-image-command cl-unmap-mem-object-command
 114     cl-marker-command cl-acquire-gl-objects-command cl-release-gl-objects-command ;
 115
 116 VARIANT: cl-execution-status
 117     cl-queued cl-submitted cl-running cl-complete cl-failure ;
 118
 119 TUPLE: cl-platform
 120     id profile version name vendor extensions devices ;
 121
 122 TUPLE: cl-device
 123     id type vendor-id max-compute-units max-work-item-dimensions
 124     max-work-item-sizes max-work-group-size preferred-vector-width-char
 125     preferred-vector-width-short preferred-vector-width-int
 126     preferred-vector-width-long preferred-vector-width-float
 127     preferred-vector-width-double max-clock-frequency address-bits
 128     max-mem-alloc-size image-support max-read-image-args max-write-image-args
 129     image2d-max-width image2d-max-height image3d-max-width image3d-max-height
 130     image3d-max-depth max-samplers max-parameter-size mem-base-addr-align
 131     min-data-type-align-size single-fp-config global-mem-cache-type
 132     global-mem-cacheline-size global-mem-cache-size global-mem-size
 133     max-constant-buffer-size max-constant-args local-mem? local-mem-size
 134     error-correction-support profiling-timer-resolution endian-little
 135     available compiler-available execute-kernels? execute-native-kernels?
 136     out-of-order-exec-available? profiling-available?
 137     name vendor driver-version profile version extensions ;
 138
 139 TUPLE: cl-context < cl-handle ;
 140 TUPLE: cl-queue   < cl-handle ;
 141 TUPLE: cl-buffer  < cl-handle ;
 142 TUPLE: cl-sampler < cl-handle ;
 143 TUPLE: cl-program < cl-handle ;
 144 TUPLE: cl-kernel  < cl-handle ;
 145 TUPLE: cl-event   < cl-handle ;
 146
 147 M: cl-context dispose* handle>> clReleaseContext      cl-success ;
 148 M: cl-queue   dispose* handle>> clReleaseCommandQueue cl-success ;
 149 M: cl-buffer  dispose* handle>> clReleaseMemObject    cl-success ;
 150 M: cl-sampler dispose* handle>> clReleaseSampler      cl-success ;
 151 M: cl-program dispose* handle>> clReleaseProgram      cl-success ;
 152 M: cl-kernel  dispose* handle>> clReleaseKernel       cl-success ;
 153 M: cl-event   dispose* handle>> clReleaseEvent        cl-success ;
 154
 155 TUPLE: cl-buffer-ptr
 156     { buffer cl-buffer read-only }
 157     { offset integer   read-only } ;
 158 C: <cl-buffer-ptr> cl-buffer-ptr
 159
 160 TUPLE: cl-buffer-range
 161     { buffer cl-buffer read-only }
 162     { offset integer   read-only }
 163     { size   integer   read-only } ;
 164 C: <cl-buffer-range> cl-buffer-range
 165
 166 SYMBOLS: cl-current-context cl-current-queue cl-current-device ;
 167
 168 <PRIVATE
 169
 170 : (current-cl-context) ( -- cl-context )
 171     cl-current-context get ; inline
 172
 173 : (current-cl-queue) ( -- cl-queue )
 174     cl-current-queue get ; inline
 175
 176 : (current-cl-device) ( -- cl-device )
 177     cl-current-device get ; inline
 178
 179 GENERIC: buffer-access-constant ( buffer-access-mode -- n )
 180 M: cl-read-write-access buffer-access-constant drop CL_MEM_READ_WRITE ;
 181 M: cl-read-access       buffer-access-constant drop CL_MEM_READ_ONLY ;
 182 M: cl-write-access      buffer-access-constant drop CL_MEM_WRITE_ONLY ;
 183
 184 GENERIC: buffer-map-flags ( buffer-access-mode -- n )
 185 M: cl-read-write-access buffer-map-flags drop CL_MAP_READ CL_MAP_WRITE bitor ;
 186 M: cl-read-access       buffer-map-flags drop CL_MAP_READ ;
 187 M: cl-write-access      buffer-map-flags drop CL_MAP_WRITE ;
 188
 189 GENERIC: addressing-mode-constant ( addressing-mode -- n )
 190 M: cl-repeat-addressing        addressing-mode-constant drop CL_ADDRESS_REPEAT ;
 191 M: cl-clamp-to-edge-addressing addressing-mode-constant drop CL_ADDRESS_CLAMP_TO_EDGE ;
 192 M: cl-clamp-addressing         addressing-mode-constant drop CL_ADDRESS_CLAMP ;
 193 M: cl-no-addressing            addressing-mode-constant drop CL_ADDRESS_NONE ;
 194
 195 GENERIC: filter-mode-constant ( filter-mode -- n )
 196 M: cl-filter-nearest filter-mode-constant drop CL_FILTER_NEAREST ;
 197 M: cl-filter-linear  filter-mode-constant drop CL_FILTER_LINEAR ;
 198
 199 : cl_addressing_mode>addressing-mode ( cl_addressing_mode -- addressing-mode )
 200     {
 201         { CL_ADDRESS_REPEAT        [ cl-repeat-addressing        ] }
 202         { CL_ADDRESS_CLAMP_TO_EDGE [ cl-clamp-to-edge-addressing ] }
 203         { CL_ADDRESS_CLAMP         [ cl-clamp-addressing         ] }
 204         { CL_ADDRESS_NONE          [ cl-no-addressing            ] }
 205     } case ; inline
 206
 207 : cl_filter_mode>filter-mode ( cl_filter_mode -- filter-mode )
 208     {
 209         { CL_FILTER_LINEAR  [ cl-filter-linear  ] }
 210         { CL_FILTER_NEAREST [ cl-filter-nearest ] }
 211     } case ; inline
 212
 213 : platform-info-string ( handle name -- string )
 214     [ clGetPlatformInfo ] info-string ;
 215
 216 : platform-info ( id -- profile version name vendor extensions )
 217     {
 218         [ CL_PLATFORM_PROFILE    platform-info-string ]
 219         [ CL_PLATFORM_VERSION    platform-info-string ]
 220         [ CL_PLATFORM_NAME       platform-info-string ]
 221         [ CL_PLATFORM_VENDOR     platform-info-string ]
 222         [ CL_PLATFORM_EXTENSIONS platform-info-string ]
 223     } cleave ;
 224
 225 : cl_device_fp_config>flags ( ulong -- sequence )
 226     [ {
 227         [ CL_FP_DENORM           bitand 0 = [ f ] [ cl-denorm           ] if ]
 228         [ CL_FP_INF_NAN          bitand 0 = [ f ] [ cl-inf-and-nan      ] if ]
 229         [ CL_FP_ROUND_TO_NEAREST bitand 0 = [ f ] [ cl-round-to-nearest ] if ]
 230         [ CL_FP_ROUND_TO_ZERO    bitand 0 = [ f ] [ cl-round-to-zero    ] if ]
 231         [ CL_FP_ROUND_TO_INF     bitand 0 = [ f ] [ cl-round-to-inf     ] if ]
 232         [ CL_FP_FMA              bitand 0 = [ f ] [ cl-fma              ] if ]
 233     } cleave ] { } output>sequence sift ;
 234
 235 : cl_device_mem_cache_type>cache-type ( uint -- cache-type )
 236     {
 237         { CL_NONE             [ cl-no-cache         ] }
 238         { CL_READ_ONLY_CACHE  [ cl-read-only-cache  ] }
 239         { CL_READ_WRITE_CACHE [ cl-read-write-cache ] }
 240     } case ; inline
 241
 242 : device-info-bool ( handle name -- ? )
 243     [ clGetDeviceInfo ] info-bool ;
 244
 245 : device-info-ulong ( handle name -- ulong )
 246     [ clGetDeviceInfo ] info-ulong ;
 247
 248 : device-info-uint ( handle name -- uint )
 249     [ clGetDeviceInfo ] info-uint ;
 250
 251 : device-info-string ( handle name -- string )
 252     [ clGetDeviceInfo ] info-string ;
 253
 254 : device-info-size_t ( handle name -- size_t )
 255     [ clGetDeviceInfo ] info-size_t ;
 256
 257 : device-info-size_t-array ( handle name -- size_t-array )
 258     [ clGetDeviceInfo ] info-size_t-array ;
 259
 260 : device-info ( device-id -- device )
 261     dup {
 262         [ CL_DEVICE_TYPE                          device-info-size_t size_t>cl-device-type ]
 263         [ CL_DEVICE_VENDOR_ID                     device-info-uint         ]
 264         [ CL_DEVICE_MAX_COMPUTE_UNITS             device-info-uint         ]
 265         [ CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS      device-info-uint         ]
 266         [ CL_DEVICE_MAX_WORK_ITEM_SIZES           device-info-size_t-array ]
 267         [ CL_DEVICE_MAX_WORK_GROUP_SIZE           device-info-size_t       ]
 268         [ CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR   device-info-uint         ]
 269         [ CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT  device-info-uint         ]
 270         [ CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT    device-info-uint         ]
 271         [ CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG   device-info-uint         ]
 272         [ CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT  device-info-uint         ]
 273         [ CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE device-info-uint         ]
 274         [ CL_DEVICE_MAX_CLOCK_FREQUENCY           device-info-uint         ]
 275         [ CL_DEVICE_ADDRESS_BITS                  device-info-uint         ]
 276         [ CL_DEVICE_MAX_MEM_ALLOC_SIZE            device-info-ulong        ]
 277         [ CL_DEVICE_IMAGE_SUPPORT                 device-info-bool         ]
 278         [ CL_DEVICE_MAX_READ_IMAGE_ARGS           device-info-uint         ]
 279         [ CL_DEVICE_MAX_WRITE_IMAGE_ARGS          device-info-uint         ]
 280         [ CL_DEVICE_IMAGE2D_MAX_WIDTH             device-info-size_t       ]
 281         [ CL_DEVICE_IMAGE2D_MAX_HEIGHT            device-info-size_t       ]
 282         [ CL_DEVICE_IMAGE3D_MAX_WIDTH             device-info-size_t       ]
 283         [ CL_DEVICE_IMAGE3D_MAX_HEIGHT            device-info-size_t       ]
 284         [ CL_DEVICE_IMAGE3D_MAX_DEPTH             device-info-size_t       ]
 285         [ CL_DEVICE_MAX_SAMPLERS                  device-info-uint         ]
 286         [ CL_DEVICE_MAX_PARAMETER_SIZE            device-info-size_t       ]
 287         [ CL_DEVICE_MEM_BASE_ADDR_ALIGN           device-info-uint         ]
 288         [ CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE      device-info-uint         ]
 289         [ CL_DEVICE_SINGLE_FP_CONFIG              device-info-ulong cl_device_fp_config>flags           ]
 290         [ CL_DEVICE_GLOBAL_MEM_CACHE_TYPE         device-info-uint  cl_device_mem_cache_type>cache-type ]
 291         [ CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE     device-info-uint         ]
 292         [ CL_DEVICE_GLOBAL_MEM_CACHE_SIZE         device-info-ulong        ]
 293         [ CL_DEVICE_GLOBAL_MEM_SIZE               device-info-ulong        ]
 294         [ CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE      device-info-ulong        ]
 295         [ CL_DEVICE_MAX_CONSTANT_ARGS             device-info-uint         ]
 296         [ CL_DEVICE_LOCAL_MEM_TYPE                device-info-uint CL_LOCAL = ]
 297         [ CL_DEVICE_LOCAL_MEM_SIZE                device-info-ulong        ]
 298         [ CL_DEVICE_ERROR_CORRECTION_SUPPORT      device-info-bool         ]
 299         [ CL_DEVICE_PROFILING_TIMER_RESOLUTION    device-info-size_t       ]
 300         [ CL_DEVICE_ENDIAN_LITTLE                 device-info-bool         ]
 301         [ CL_DEVICE_AVAILABLE                     device-info-bool         ]
 302         [ CL_DEVICE_COMPILER_AVAILABLE            device-info-bool         ]
 303         [ CL_DEVICE_EXECUTION_CAPABILITIES        device-info-ulong CL_EXEC_KERNEL                         bitand 0 = not ]
 304         [ CL_DEVICE_EXECUTION_CAPABILITIES        device-info-ulong CL_EXEC_NATIVE_KERNEL                  bitand 0 = not ]
 305         [ CL_DEVICE_QUEUE_PROPERTIES              device-info-ulong CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE bitand 0 = not ]
 306         [ CL_DEVICE_QUEUE_PROPERTIES              device-info-ulong CL_QUEUE_PROFILING_ENABLE              bitand 0 = not ]
 307         [ CL_DEVICE_NAME                          device-info-string       ]
 308         [ CL_DEVICE_VENDOR                        device-info-string       ]
 309         [ CL_DRIVER_VERSION                       device-info-string       ]
 310         [ CL_DEVICE_PROFILE                       device-info-string       ]
 311         [ CL_DEVICE_VERSION                       device-info-string       ]
 312         [ CL_DEVICE_EXTENSIONS                    device-info-string       ]
 313     } cleave cl-device boa ;
 314
 315 : platform-devices ( platform-id -- devices )
 316     CL_DEVICE_TYPE_ALL [
 317         0 f 0 uint <ref> [ clGetDeviceIDs cl-success ] keep uint deref
 318     ] [
 319         rot dup void* <c-array> [ f clGetDeviceIDs cl-success ] keep
 320     ] 2bi ; inline
 321
 322 : command-queue-info-ulong ( handle name -- ulong )
 323     [ clGetCommandQueueInfo ] info-ulong ;
 324
 325 : sampler-info-bool ( handle name -- ? )
 326     [ clGetSamplerInfo ] info-bool ;
 327
 328 : sampler-info-uint ( handle name -- uint )
 329     [ clGetSamplerInfo ] info-uint ;
 330
 331 : program-build-info-string ( program-handle device-handle name -- string )
 332     [ clGetProgramBuildInfo ] 2info-string ;
 333
 334 : program-build-log ( program-handle device-handle -- string )
 335     CL_PROGRAM_BUILD_LOG program-build-info-string ;
 336
 337 : strings>char*-array ( strings -- char*-array )
 338     [
 339         ascii encode dup length dup malloc
 340         [ cl-not-null ] keep &free [ -rot memcpy ] keep
 341     ] void*-array{ } map-as ;
 342
 343 : (program) ( cl-context sources -- program-handle )
 344     [ handle>> ] dip [
 345         [ length ]
 346         [ strings>char*-array ]
 347         [ [ length ] size_t-array{ } map-as ] tri
 348         0 int <ref> [ clCreateProgramWithSource ] keep int deref cl-success
 349     ] with-destructors ;
 350
 351 :: (build-program) ( program-handle device options -- program )
 352     program-handle 1 device 1array [ id>> ] void*-array{ } map-as
 353     options ascii encode 0 suffix f f clBuildProgram
 354     {
 355         { CL_BUILD_PROGRAM_FAILURE [
 356             program-handle device id>> program-build-log program-handle
 357             clReleaseProgram cl-success cl-error f ] }
 358         { CL_SUCCESS [ cl-program new-disposable program-handle >>handle ] }
 359         [ program-handle clReleaseProgram cl-success cl-success f ]
 360     } case ;
 361
 362 : kernel-info-string ( handle name -- string )
 363     [ clGetKernelInfo ] info-string ;
 364
 365 : kernel-info-uint ( handle name -- uint )
 366     [ clGetKernelInfo ] info-uint ;
 367
 368 : kernel-work-group-info-size_t ( handle1 handle2 name -- size_t )
 369     [ clGetKernelWorkGroupInfo ] 2info-size_t ;
 370
 371 : event-info-uint ( handle name -- uint )
 372     [ clGetEventInfo ] info-uint ;
 373
 374 : event-info-int ( handle name -- int )
 375     [ clGetEventInfo ] info-int ;
 376
 377 : cl_command_type>command-type ( cl_command-type -- command-type )
 378     {
 379         { CL_COMMAND_NDRANGE_KERNEL       [ cl-ndrange-kernel-command       ] }
 380         { CL_COMMAND_TASK                 [ cl-task-command                 ] }
 381         { CL_COMMAND_NATIVE_KERNEL        [ cl-native-kernel-command        ] }
 382         { CL_COMMAND_READ_BUFFER          [ cl-read-buffer-command          ] }
 383         { CL_COMMAND_WRITE_BUFFER         [ cl-write-buffer-command         ] }
 384         { CL_COMMAND_COPY_BUFFER          [ cl-copy-buffer-command          ] }
 385         { CL_COMMAND_READ_IMAGE           [ cl-read-image-command           ] }
 386         { CL_COMMAND_WRITE_IMAGE          [ cl-write-image-command          ] }
 387         { CL_COMMAND_COPY_IMAGE           [ cl-copy-image-command           ] }
 388         { CL_COMMAND_COPY_BUFFER_TO_IMAGE [ cl-copy-buffer-to-image-command ] }
 389         { CL_COMMAND_COPY_IMAGE_TO_BUFFER [ cl-copy-image-to-buffer-command ] }
 390         { CL_COMMAND_MAP_BUFFER           [ cl-map-buffer-command           ] }
 391         { CL_COMMAND_MAP_IMAGE            [ cl-map-image-command            ] }
 392         { CL_COMMAND_UNMAP_MEM_OBJECT     [ cl-unmap-mem-object-command     ] }
 393         { CL_COMMAND_MARKER               [ cl-marker-command               ] }
 394         { CL_COMMAND_ACQUIRE_GL_OBJECTS   [ cl-acquire-gl-objects-command   ] }
 395         { CL_COMMAND_RELEASE_GL_OBJECTS   [ cl-release-gl-objects-command   ] }
 396     } case ;
 397
 398 : cl_int>execution-status ( clint -- execution-status )
 399     {
 400         { CL_QUEUED    [ cl-queued    ] }
 401         { CL_SUBMITTED [ cl-submitted ] }
 402         { CL_RUNNING   [ cl-running   ] }
 403         { CL_COMPLETE  [ cl-complete  ] }
 404         [ drop cl-failure ]
 405     } case ; inline
 406
 407 : profiling-info-ulong ( handle name -- ulong )
 408     [ clGetEventProfilingInfo ] info-ulong ;
 409
 410 : bind-kernel-arg-buffer ( kernel index buffer -- )
 411     [ handle>> ] [ cl_mem heap-size ] [ handle>> void* deref ] tri*
 412     clSetKernelArg cl-success ; inline
 413
 414 : bind-kernel-arg-data ( kernel index byte-array -- )
 415     [ handle>> ] 2dip
 416     [ byte-length ] keep clSetKernelArg cl-success ; inline
 417
 418 GENERIC: bind-kernel-arg ( kernel index data -- )
 419 M: cl-buffer  bind-kernel-arg bind-kernel-arg-buffer ;
 420 M: byte-array bind-kernel-arg bind-kernel-arg-data ;
 421
 422 PRIVATE>
 423
 424 : with-cl-state ( context/f device/f queue/f quot -- )
 425     [
 426         [
 427             [ cl-current-queue   ,, ] when*
 428             [ cl-current-device  ,, ] when*
 429             [ cl-current-context ,, ] when*
 430         ] 3curry H{ } make
 431     ] dip with-variables ; inline
 432
 433 : cl-platforms ( -- platforms )
 434     0 f 0 uint <ref> [ clGetPlatformIDs cl-success ] keep uint deref
 435     dup void* <c-array> [ f clGetPlatformIDs cl-success ] keep
 436     [
 437         dup
 438         [ platform-info ]
 439         [ platform-devices [ device-info ] { } map-as ] bi
 440         cl-platform boa
 441     ] { } map-as ;
 442
 443 : <cl-context> ( devices -- cl-context )
 444     [ f ] dip
 445     [ length ] [ [ id>> ] void*-array{ } map-as ] bi
 446     f f 0 int <ref> [ clCreateContext ] keep int deref cl-success
 447     cl-context new-disposable swap >>handle ;
 448
 449 : <cl-queue> ( context device out-of-order? profiling? -- command-queue )
 450     [ [ handle>> ] [ id>> ] bi* ] 2dip
 451     [ [ CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE ] [ 0 ] if ]
 452     [ [ CL_QUEUE_PROFILING_ENABLE ] [ 0 ] if ] bi* bitor
 453     0 int <ref> [ clCreateCommandQueue ] keep int deref cl-success
 454     cl-queue new-disposable swap >>handle ;
 455
 456 : cl-out-of-order-execution? ( command-queue -- ? )
 457     CL_QUEUE_PROPERTIES command-queue-info-ulong
 458     CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE bitand 0 = not ; inline
 459
 460 : cl-profiling? ( command-queue -- ? )
 461     CL_QUEUE_PROPERTIES command-queue-info-ulong
 462     CL_QUEUE_PROFILING_ENABLE bitand 0 = not ; inline
 463
 464 : <cl-buffer> ( buffer-access-mode size initial-data -- buffer )
 465     [ (current-cl-context) ] 3dip
 466     swap over [
 467         [ handle>> ]
 468         [ buffer-access-constant ]
 469         [ [ CL_MEM_COPY_HOST_PTR ] [ CL_MEM_ALLOC_HOST_PTR ] if ] tri* bitor
 470     ] 2dip
 471     0 int <ref> [ clCreateBuffer ] keep int deref cl-success
 472     cl-buffer new-disposable swap >>handle ;
 473
 474 : cl-read-buffer ( buffer-range -- byte-array )
 475     [ (current-cl-queue) handle>> ] dip
 476     [ buffer>> handle>> CL_TRUE ]
 477     [ offset>> ]
 478     [ size>> dup <byte-array> ] tri
 479     [ 0 f f clEnqueueReadBuffer cl-success ] keep ; inline
 480
 481 : cl-write-buffer ( buffer-range byte-array -- )
 482     [
 483         [ (current-cl-queue) handle>> ] dip
 484         [ buffer>> handle>> CL_TRUE ]
 485         [ offset>> ]
 486         [ size>> ] tri
 487     ] dip 0 f f clEnqueueWriteBuffer cl-success ; inline
 488
 489 : cl-queue-copy-buffer ( src-buffer-ptr dst-buffer-ptr size dependent-events -- event )
 490     [
 491         (current-cl-queue)
 492         [ handle>> ]
 493         [ [ buffer>> handle>> ] [ offset>> ] bi ]
 494         [ [ buffer>> handle>> ] [ offset>> ] bi ]
 495         tri* swapd
 496     ] 2dip [ length ] keep [ f ] [ [ handle>> ] void*-array{ } map-as ] if-empty
 497     f void* <ref> [ clEnqueueCopyBuffer cl-success ] keep void* deref cl-event
 498     new-disposable swap >>handle ;
 499
 500 : cl-queue-read-buffer ( buffer-range alien dependent-events -- event )
 501     [
 502         [ (current-cl-queue) handle>> ] dip
 503         [ buffer>> handle>> CL_FALSE ] [ offset>> ] [ size>> ] tri
 504     ] 2dip [ length ] keep [ f ] [ [ handle>> ] void*-array{ } map-as ] if-empty
 505     f void* <ref> [ clEnqueueReadBuffer cl-success ] keep void* deref cl-event
 506     new-disposable swap >>handle ;
 507
 508 : cl-queue-write-buffer ( buffer-range alien dependent-events -- event )
 509     [
 510         [ (current-cl-queue) handle>> ] dip
 511         [ buffer>> handle>> CL_FALSE ] [ offset>> ] [ size>> ] tri
 512     ] 2dip [ length ] keep [ f ] [ [ handle>> ] void*-array{ } map-as ] if-empty
 513     f void* <ref> [ clEnqueueWriteBuffer cl-success ] keep void* deref cl-event
 514     new-disposable swap >>handle ;
 515
 516 : <cl-sampler> ( normalized-coords? addressing-mode filter-mode -- sampler )
 517     [ (current-cl-context) ] 3dip
 518     [ [ CL_TRUE ] [ CL_FALSE ] if ]
 519     [ addressing-mode-constant ]
 520     [ filter-mode-constant ]
 521     tri* 0 int <ref> [ clCreateSampler ] keep int deref cl-success
 522     cl-sampler new-disposable swap >>handle ;
 523
 524 : cl-normalized-coords? ( sampler -- ? )
 525     handle>> CL_SAMPLER_NORMALIZED_COORDS sampler-info-bool ; inline
 526
 527 : sampler>cl-addressing-mode ( sampler -- addressing-mode )
 528     handle>> CL_SAMPLER_ADDRESSING_MODE sampler-info-uint cl_addressing_mode>addressing-mode ; inline
 529
 530 : sampler>cl-filter-mode ( sampler -- filter-mode )
 531     handle>> CL_SAMPLER_FILTER_MODE sampler-info-uint cl_filter_mode>filter-mode ; inline
 532
 533 : <cl-program> ( options strings -- program )
 534     [ (current-cl-device) ] 2dip
 535     [ (current-cl-context) ] dip
 536     (program) -rot (build-program) ;
 537
 538 : <cl-kernel> ( program kernel-name -- kernel )
 539     [ handle>> ] [ ascii encode 0 suffix ] bi*
 540     0 int <ref> [ clCreateKernel ] keep int deref cl-success
 541     cl-kernel new-disposable swap >>handle ; inline
 542
 543 : cl-kernel-name ( kernel -- string )
 544     handle>> CL_KERNEL_FUNCTION_NAME kernel-info-string ;
 545
 546 : cl-kernel-arity ( kernel -- arity )
 547     handle>> CL_KERNEL_NUM_ARGS kernel-info-uint ;
 548
 549 : cl-kernel-local-size ( kernel -- size )
 550     (current-cl-device) [ handle>> ] bi@ CL_KERNEL_WORK_GROUP_SIZE kernel-work-group-info-size_t ; inline
 551
 552 :: cl-queue-kernel ( kernel args sizes dependent-events -- event )
 553     args [| arg idx | kernel idx arg bind-kernel-arg ] each-index
 554     (current-cl-queue) handle>>
 555     kernel handle>>
 556     sizes [ length f ] [ [ ] size_t-array{ } map-as f ] bi
 557     dependent-events [ length ] [ [ f ] [ [ handle>> ] void*-array{ } map-as ] if-empty ] bi
 558     f void* <ref> [ clEnqueueNDRangeKernel cl-success ] keep void* deref
 559     cl-event new-disposable swap >>handle ;
 560
 561 : cl-event-type ( event -- command-type )
 562     handle>> CL_EVENT_COMMAND_TYPE event-info-uint cl_command_type>command-type ; inline
 563
 564 : cl-event-status ( event -- execution-status )
 565     handle>> CL_EVENT_COMMAND_EXECUTION_STATUS event-info-int cl_int>execution-status ; inline
 566
 567 : cl-profile-counters ( event -- queued submitted started finished )
 568     handle>> {
 569         [ CL_PROFILING_COMMAND_QUEUED profiling-info-ulong ]
 570         [ CL_PROFILING_COMMAND_SUBMIT profiling-info-ulong ]
 571         [ CL_PROFILING_COMMAND_START  profiling-info-ulong ]
 572         [ CL_PROFILING_COMMAND_END    profiling-info-ulong ]
 573     } cleave ; inline
 574
 575 : cl-barrier-events ( event/events -- )
 576     [ (current-cl-queue) handle>> ] dip
 577     dup sequence? [ 1array ] unless
 578     [ handle>> ] void*-array{ } map-as [ length ] keep clEnqueueWaitForEvents cl-success ; inline
 579
 580 : cl-marker ( -- event )
 581     (current-cl-queue)
 582     f void* <ref> [ clEnqueueMarker cl-success ] keep void* deref cl-event new-disposable
 583     swap >>handle ; inline
 584
 585 : cl-barrier ( -- )
 586     (current-cl-queue) clEnqueueBarrier cl-success ; inline
 587
 588 : cl-flush ( -- )
 589     (current-cl-queue) handle>> clFlush cl-success ; inline
 590
 591 : cl-wait ( event/events -- )
 592     dup sequence? [ 1array ] unless
 593     [ handle>> ] void*-array{ } map-as [ length ] keep clWaitForEvents cl-success ; inline
 594
 595 : cl-finish ( -- )
 596     (current-cl-queue) handle>> clFinish cl-success ; inline