-unportable bindings
\ No newline at end of file
+bindings\r
-unportable bindings\r
+bindings\r
-unportable bindings
+bindings
--- /dev/null
+! Copyright (C) 2010 Samuel Tardieu.
+! See http://factorcode.org/license.txt for BSD license.
+USING: help.markup help.syntax ;
+IN: astar
+
+{ find-path <astar> considered } related-words
+
+HELP: <astar>
+{ $values
+ { "neighbours" "a quotation with stack effect ( node -- seq )" }
+ { "cost" "a quotation with stack effect ( from to -- cost )" }
+ { "heuristic" "a quotation with stack effect ( pos target -- cost )" }
+ { "astar" "a astar tuple" }
+}
+{ $description "Build an astar object from the given quotations. The "
+ { $snippet "neighbours" } " one builds the list of neighbours. The "
+ { $snippet "cost" } " and " { $snippet "heuristic" } " ones represent "
+ "respectively the cost for transitioning from a node to one of its neighbour, "
+ "and the underestimated cost for going from a node to the target."
+} ;
+
+HELP: find-path
+{ $values
+ { "start" "a node" }
+ { "target" "a node" }
+ { "astar" "a astar tuple" }
+ { "path/f" "an optimal path from " { $snippet "start" } " to " { $snippet "target" }
+ ", or f if no such path exists" }
+}
+{ $description "Find a path between " { $snippet "start" } " and " { $snippet "target" }
+ " using the A* algorithm. The " { $snippet "astar" } " tuple must have been previously "
+ " built using " { $link <astar> } "."
+} ;
+
+HELP: considered
+{ $values
+ { "astar" "a astar tuple" }
+ { "considered" "a sequence" }
+}
+{ $description "When called after a call to " { $link find-path } ", return a list of nodes "
+ "which have been examined during the A* exploration."
+} ;
--- /dev/null
+! Copyright (C) 2010 Samuel Tardieu.
+! See http://factorcode.org/license.txt for BSD license.
+USING: arrays assocs astar combinators hashtables kernel literals math math.functions
+math.vectors sequences sorting splitting strings tools.test ;
+IN: astar.tests
+
+<<
+
+! Use a 10x9 maze (see below) to try to go from s to e, f or g.
+! X means that a position is unreachable.
+! The costs model is:
+! - going up costs 5 points
+! - going down costs 1 point
+! - going left or right costs 2 points
+
+: reachable? ( pos -- ? )
+ first2 [ 2 * 5 + ] [ 2 + ] bi* $[
+" 0 1 2 3 4 5 6 7 8 9
+
+ 0 X X X X X X X X X X
+ 1 X s f X X
+ 2 X X X X X X X X X
+ 3 X X X X X X X X X
+ 4 X X X X X X
+ 5 X X X X X
+ 6 X X X X X X e X
+ 7 X g X X
+ 8 X X X X X X X X X X"
+ "\n" split ] nth nth CHAR: X = not ;
+
+: neighbours ( pos -- neighbours )
+ first2
+ { [ 1 + 2array ] [ 1 - 2array ] [ [ 1 + ] dip 2array ] [ [ 1 - ] dip 2array ] } 2cleave
+ 4array
+ [ reachable? ] filter ;
+
+: heuristic ( from to -- cost )
+ v- [ abs ] [ + ] map-reduce ;
+
+: cost ( from to -- cost )
+ 2dup [ first ] bi@ = [ [ second ] bi@ > 1 5 ? ] [ 2drop 2 ] if ;
+
+: test1 ( to -- path considered )
+ { 1 1 } swap [ neighbours ] [ cost ] [ heuristic ] <astar> [ find-path ] [ considered ] bi ;
+>>
+
+! Existing path from s to f
+[
+ {
+ { 1 1 }
+ { 2 1 }
+ { 3 1 }
+ { 4 1 }
+ { 4 2 }
+ { 4 3 }
+ { 4 4 }
+ { 4 5 }
+ { 4 6 }
+ { 4 7 }
+ { 5 7 }
+ { 6 7 }
+ { 7 7 }
+ { 8 7 }
+ { 8 6 }
+ }
+] [
+ { 8 6 } test1 drop
+] unit-test
+
+! Check that only the right positions have been considered in the s to f path
+[ 7 ] [ { 7 1 } test1 nip length ] unit-test
+
+! Non-existing path from s to g -- all positions must have been considered
+[ f 26 ] [ { 1 7 } test1 length ] unit-test
+
+<<
+
+! Look for a path between A and C. The best path is A --> D --> C. C will be placed
+! in the open set early because B will be examined first. This checks that the evaluation
+! of C is correctly replaced in the open set.
+!
+! We use no heuristic here and always return 0.
+!
+! (5)
+! B ---> C <--------
+! \ (2)
+! ^ ^ |
+! | | |
+! (1) | | (2) |
+! | | |
+!
+! A ---> D ---------> E ---> F
+! (2) (1) (1)
+
+: n ( pos -- neighbours )
+ $[ { "ABD" "BC" "C" "DCE" "ECF" } [ unclip swap 2array ] map >hashtable ] at ;
+
+: c ( from to -- cost )
+ "" 2sequence H{ { "AB" 1 } { "AD" 2 } { "BC" 5 } { "DC" 2 } { "DE" 1 } { "EC" 2 } { "EF" 1 } } at ;
+
+: test2 ( fromto -- path considered )
+ first2 [ n ] [ c ] [ 2drop 0 ] <astar> [ find-path ] [ considered natural-sort >string ] bi ;
+>>
+
+! Check path from A to C -- all nodes but F must have been examined
+[ "ADC" "ABCDE" ] [ "AC" test2 [ >string ] dip ] unit-test
+
+! No path from D to B -- all nodes reachable from D must have been examined
+[ f "CDEF" ] [ "DB" test2 ] unit-test
--- /dev/null
+! Copyright (C) 2010 Samuel Tardieu.
+! See http://factorcode.org/license.txt for BSD license.
+USING: accessors assocs heaps kernel math sequences sets shuffle ;
+IN: astar
+
+! This implements the A* algorithm. See http://en.wikipedia.org/wiki/A*
+
+<PRIVATE
+
+TUPLE: astar neighbours heuristic cost
+ goal g origin in-open-set in-closed-set open-set ;
+
+: (add-to-open-set) ( h node astar -- )
+ 2dup in-open-set>> at* [ over open-set>> heap-delete ] [ drop ] if
+ [ swapd open-set>> heap-push* ] [ in-open-set>> set-at ] 2bi ;
+
+: add-to-open-set ( node astar -- )
+ [ g>> at ] 2keep
+ [ [ goal>> ] [ heuristic>> call( n1 n2 -- c ) ] bi + ] 2keep
+ (add-to-open-set) ;
+
+: ?add-to-open-set ( node astar -- )
+ 2dup in-closed-set>> key? [ 2drop ] [ add-to-open-set ] if ;
+
+: move-to-closed-set ( node astar -- )
+ [ in-closed-set>> conjoin ] [ in-open-set>> delete-at ] 2bi ;
+
+: get-first ( astar -- node )
+ [ open-set>> heap-pop drop dup ] [ move-to-closed-set ] bi ;
+
+: set-g ( origin g node astar -- )
+ [ [ origin>> set-at ] [ g>> set-at ] bi-curry bi-curry bi* ] [ ?add-to-open-set ] 2bi ;
+
+: cost-through ( origin node astar -- cost )
+ [ cost>> call( n1 n2 -- c ) ] [ nip g>> at ] 3bi + ;
+
+: ?set-g ( origin node astar -- )
+ [ cost-through ] 3keep [ swap ] 2dip
+ 3dup g>> at [ 1/0. ] unless* > [ 4drop ] [ set-g ] if ;
+
+: build-path ( target astar -- path )
+ [ over ] [ over [ [ origin>> at ] keep ] dip ] produce 2nip reverse ;
+
+: handle ( node astar -- )
+ dupd [ neighbours>> call( node -- neighbours ) ] keep [ ?set-g ] curry with each ;
+
+: (find-path) ( astar -- path/f )
+ dup open-set>> heap-empty? [
+ drop f
+ ] [
+ [ get-first ] keep 2dup goal>> = [ build-path ] [ [ handle ] [ (find-path) ] bi ] if
+ ] if ;
+
+: (init) ( from to astar -- )
+ swap >>goal
+ H{ } clone >>g
+ H{ } clone >>origin
+ H{ } clone >>in-open-set
+ H{ } clone >>in-closed-set
+ <min-heap> >>open-set
+ [ 0 ] 2dip [ (add-to-open-set) ] [ g>> set-at ] 3bi ;
+
+PRIVATE>
+
+: find-path ( start target astar -- path/f )
+ [ (init) ] [ (find-path) ] bi ;
+
+: <astar> ( neighbours cost heuristic -- astar )
+ astar new swap >>heuristic swap >>cost swap >>neighbours ;
+
+: considered ( astar -- considered )
+ in-closed-set>> keys ;
--- /dev/null
+Samuel Tardieu
--- /dev/null
+A* path-finding algorithm
--- /dev/null
+Erik Charlebois
--- /dev/null
+Erik Charlebois
--- /dev/null
+! Copyright (C) 2010 Erik Charlebois.
+! See http://factorcode.org/license.txt for BSD license.
+USING: tools.test opencl.ffi multiline locals kernel io.encodings.ascii
+io.encodings.string sequences libc alien.c-types destructors math specialized-arrays
+math.order alien ;
+FROM: alien.c-types => float ;
+SPECIALIZED-ARRAYS: float void* ;
+IN: opencl.ffi.tests
+
+STRING: kernel-source
+__kernel void square(
+ __global float* input,
+ __global float* output,
+ const unsigned int count)
+{
+ int i = get_global_id(0);
+ if (i < count)
+ output[i] = input[i] * input[i];
+}
+;
+
+ERROR: cl-error err ;
+: cl-success ( err -- )
+ dup CL_SUCCESS = [ drop ] [ cl-error ] if ;
+
+:: cl-string-array ( str -- alien )
+ str ascii encode 0 suffix :> str-buffer
+ str-buffer length malloc &free :> str-alien
+ str-alien str-buffer dup length memcpy str-alien ;
+
+:: opencl-square ( in -- out )
+ 0 f 0 <uint> [ clGetPlatformIDs cl-success ] keep *uint
+ dup <void*-array> [ f clGetPlatformIDs cl-success ] keep first
+ CL_DEVICE_TYPE_DEFAULT 1 f <void*> [ f clGetDeviceIDs cl-success ] keep *void* :> device-id
+ f 1 device-id <void*> f f 0 <int> [ clCreateContext ] keep *int cl-success :> context
+ context device-id 0 0 <int> [ clCreateCommandQueue ] keep *int cl-success :> queue
+
+ [
+ context 1 kernel-source cl-string-array <void*>
+ f 0 <int> [ clCreateProgramWithSource ] keep *int cl-success
+ [ 0 f f f f clBuildProgram cl-success ]
+ [ "square" cl-string-array 0 <int> [ clCreateKernel ] keep *int cl-success ]
+ [ ] tri
+ ] with-destructors :> ( kernel program )
+
+ context CL_MEM_READ_ONLY in byte-length f
+ 0 <int> [ clCreateBuffer ] keep *int cl-success :> input
+
+ context CL_MEM_WRITE_ONLY in byte-length f
+ 0 <int> [ clCreateBuffer ] keep *int cl-success :> output
+
+ queue input CL_TRUE 0 in byte-length in 0 f f clEnqueueWriteBuffer cl-success
+
+ kernel 0 cl_mem heap-size input <void*> clSetKernelArg cl-success
+ kernel 1 cl_mem heap-size output <void*> clSetKernelArg cl-success
+ kernel 2 uint heap-size in length <uint> clSetKernelArg cl-success
+
+ queue kernel 1 f in length <ulonglong> f
+ 0 f f clEnqueueNDRangeKernel cl-success
+
+ queue clFinish cl-success
+
+ queue output CL_TRUE 0 in byte-length in length <float-array>
+ [ 0 f f clEnqueueReadBuffer cl-success ] keep
+
+ input clReleaseMemObject cl-success
+ output clReleaseMemObject cl-success
+ program clReleaseProgram cl-success
+ kernel clReleaseKernel cl-success
+ queue clReleaseCommandQueue cl-success
+ context clReleaseContext cl-success ;
+
+[ float-array{ 1.0 4.0 9.0 16.0 100.0 } ]
+[ float-array{ 1.0 2.0 3.0 4.0 10.0 } opencl-square ] unit-test
--- /dev/null
+! Copyright (C) 2010 Erik Charlebois.
+! See http://factorcode.org/license.txt for BSD license.
+USING: alien.c-types alien.libraries alien.syntax classes.struct
+combinators system alien.accessors byte-arrays kernel ;
+IN: opencl.ffi
+
+<< "opencl" {
+ { [ os windows? ] [ "OpenCL.dll" ] }
+ { [ os macosx? ] [ "/System/Library/Frameworks/OpenCL.framework/OpenCL" ] }
+ { [ os unix? ] [ "libOpenCL.so" ] }
+ } cond "stdcall" add-library >>
+LIBRARY: opencl
+
+! cl_platform.h
+TYPEDEF: char cl_char
+TYPEDEF: uchar cl_uchar
+TYPEDEF: short cl_short
+TYPEDEF: ushort cl_ushort
+TYPEDEF: int cl_int
+TYPEDEF: uint cl_uint
+TYPEDEF: longlong cl_long
+TYPEDEF: ulonglong cl_ulong
+TYPEDEF: ushort cl_half
+TYPEDEF: float cl_float
+TYPEDEF: double cl_double
+
+CONSTANT: CL_CHAR_BIT 8
+CONSTANT: CL_SCHAR_MAX 127
+CONSTANT: CL_SCHAR_MIN -128
+CONSTANT: CL_CHAR_MAX 127
+CONSTANT: CL_CHAR_MIN -128
+CONSTANT: CL_UCHAR_MAX 255
+CONSTANT: CL_SHRT_MAX 32767
+CONSTANT: CL_SHRT_MIN -32768
+CONSTANT: CL_USHRT_MAX 65535
+CONSTANT: CL_INT_MAX 2147483647
+CONSTANT: CL_INT_MIN -2147483648
+CONSTANT: CL_UINT_MAX HEX: ffffffff
+CONSTANT: CL_LONG_MAX HEX: 7FFFFFFFFFFFFFFF
+CONSTANT: CL_LONG_MIN HEX: 8000000000000000
+CONSTANT: CL_ULONG_MAX HEX: FFFFFFFFFFFFFFFF
+
+CONSTANT: CL_FLT_DIG 6
+CONSTANT: CL_FLT_MANT_DIG 24
+CONSTANT: CL_FLT_MAX_10_EXP 38
+CONSTANT: CL_FLT_MAX_EXP 128
+CONSTANT: CL_FLT_MIN_10_EXP -37
+CONSTANT: CL_FLT_MIN_EXP -125
+CONSTANT: CL_FLT_RADIX 2
+CONSTANT: CL_FLT_MAX 340282346638528859811704183484516925440.0
+CONSTANT: CL_FLT_MIN 1.175494350822287507969e-38
+CONSTANT: CL_FLT_EPSILON HEX: 1.0p-23
+
+CONSTANT: CL_DBL_DIG 15
+CONSTANT: CL_DBL_MANT_DIG 53
+CONSTANT: CL_DBL_MAX_10_EXP 308
+CONSTANT: CL_DBL_MAX_EXP 1024
+CONSTANT: CL_DBL_MIN_10_EXP -307
+CONSTANT: CL_DBL_MIN_EXP -1021
+CONSTANT: CL_DBL_RADIX 2
+CONSTANT: CL_DBL_MAX 179769313486231570814527423731704356798070567525844996598917476803157260780028538760589558632766878171540458953514382464234321326889464182768467546703537516986049910576551282076245490090389328944075868508455133942304583236903222948165808559332123348274797826204144723168738177180919299881250404026184124858368.0
+CONSTANT: CL_DBL_MIN 2.225073858507201383090e-308
+CONSTANT: CL_DBL_EPSILON 2.220446049250313080847e-16
+
+CONSTANT: CL_NAN NAN: 0
+CONSTANT: CL_HUGE_VALF 1.0e50
+CONSTANT: CL_HUGE_VAL 1.0e500
+CONSTANT: CL_MAXFLOAT 340282346638528859811704183484516925440.0
+CONSTANT: CL_INFINITY 1.0e50
+
+TYPEDEF: uint cl_GLuint
+TYPEDEF: int cl_GLint
+TYPEDEF: uint cl_GLenum
+
+! cl.h
+C-TYPE: _cl_platform_id
+C-TYPE: _cl_device_id
+C-TYPE: _cl_context
+C-TYPE: _cl_command_queue
+C-TYPE: _cl_mem
+C-TYPE: _cl_program
+C-TYPE: _cl_kernel
+C-TYPE: _cl_event
+C-TYPE: _cl_sampler
+
+TYPEDEF: _cl_platform_id* cl_platform_id
+TYPEDEF: _cl_device_id* cl_device_id
+TYPEDEF: _cl_context* cl_context
+TYPEDEF: _cl_command_queue* cl_command_queue
+TYPEDEF: _cl_mem* cl_mem
+TYPEDEF: _cl_program* cl_program
+TYPEDEF: _cl_kernel* cl_kernel
+TYPEDEF: _cl_event* cl_event
+TYPEDEF: _cl_sampler* cl_sampler
+
+TYPEDEF: cl_uint cl_bool
+TYPEDEF: cl_ulong cl_bitfield
+TYPEDEF: cl_bitfield cl_device_type
+TYPEDEF: cl_uint cl_platform_info
+TYPEDEF: cl_uint cl_device_info
+TYPEDEF: cl_bitfield cl_device_address_info
+TYPEDEF: cl_bitfield cl_device_fp_config
+TYPEDEF: cl_uint cl_device_mem_cache_type
+TYPEDEF: cl_uint cl_device_local_mem_type
+TYPEDEF: cl_bitfield cl_device_exec_capabilities
+TYPEDEF: cl_bitfield cl_command_queue_properties
+
+TYPEDEF: intptr_t cl_context_properties
+TYPEDEF: cl_uint cl_context_info
+TYPEDEF: cl_uint cl_command_queue_info
+TYPEDEF: cl_uint cl_channel_order
+TYPEDEF: cl_uint cl_channel_type
+TYPEDEF: cl_bitfield cl_mem_flags
+TYPEDEF: cl_uint cl_mem_object_type
+TYPEDEF: cl_uint cl_mem_info
+TYPEDEF: cl_uint cl_image_info
+TYPEDEF: cl_uint cl_addressing_mode
+TYPEDEF: cl_uint cl_filter_mode
+TYPEDEF: cl_uint cl_sampler_info
+TYPEDEF: cl_bitfield cl_map_flags
+TYPEDEF: cl_uint cl_program_info
+TYPEDEF: cl_uint cl_program_build_info
+TYPEDEF: cl_int cl_build_status
+TYPEDEF: cl_uint cl_kernel_info
+TYPEDEF: cl_uint cl_kernel_work_group_info
+TYPEDEF: cl_uint cl_event_info
+TYPEDEF: cl_uint cl_command_type
+TYPEDEF: cl_uint cl_profiling_info
+
+STRUCT: cl_image_format
+ { image_channel_order cl_channel_order }
+ { image_channel_data_type cl_channel_type } ;
+
+CONSTANT: CL_SUCCESS 0
+CONSTANT: CL_DEVICE_NOT_FOUND -1
+CONSTANT: CL_DEVICE_NOT_AVAILABLE -2
+CONSTANT: CL_COMPILER_NOT_AVAILABLE -3
+CONSTANT: CL_MEM_OBJECT_ALLOCATION_FAILURE -4
+CONSTANT: CL_OUT_OF_RESOURCES -5
+CONSTANT: CL_OUT_OF_HOST_MEMORY -6
+CONSTANT: CL_PROFILING_INFO_NOT_AVAILABLE -7
+CONSTANT: CL_MEM_COPY_OVERLAP -8
+CONSTANT: CL_IMAGE_FORMAT_MISMATCH -9
+CONSTANT: CL_IMAGE_FORMAT_NOT_SUPPORTED -10
+CONSTANT: CL_BUILD_PROGRAM_FAILURE -11
+CONSTANT: CL_MAP_FAILURE -12
+
+CONSTANT: CL_INVALID_VALUE -30
+CONSTANT: CL_INVALID_DEVICE_TYPE -31
+CONSTANT: CL_INVALID_PLATFORM -32
+CONSTANT: CL_INVALID_DEVICE -33
+CONSTANT: CL_INVALID_CONTEXT -34
+CONSTANT: CL_INVALID_QUEUE_PROPERTIES -35
+CONSTANT: CL_INVALID_COMMAND_QUEUE -36
+CONSTANT: CL_INVALID_HOST_PTR -37
+CONSTANT: CL_INVALID_MEM_OBJECT -38
+CONSTANT: CL_INVALID_IMAGE_FORMAT_DESCRIPTOR -39
+CONSTANT: CL_INVALID_IMAGE_SIZE -40
+CONSTANT: CL_INVALID_SAMPLER -41
+CONSTANT: CL_INVALID_BINARY -42
+CONSTANT: CL_INVALID_BUILD_OPTIONS -43
+CONSTANT: CL_INVALID_PROGRAM -44
+CONSTANT: CL_INVALID_PROGRAM_EXECUTABLE -45
+CONSTANT: CL_INVALID_KERNEL_NAME -46
+CONSTANT: CL_INVALID_KERNEL_DEFINITION -47
+CONSTANT: CL_INVALID_KERNEL -48
+CONSTANT: CL_INVALID_ARG_INDEX -49
+CONSTANT: CL_INVALID_ARG_VALUE -50
+CONSTANT: CL_INVALID_ARG_SIZE -51
+CONSTANT: CL_INVALID_KERNEL_ARGS -52
+CONSTANT: CL_INVALID_WORK_DIMENSION -53
+CONSTANT: CL_INVALID_WORK_GROUP_SIZE -54
+CONSTANT: CL_INVALID_WORK_ITEM_SIZE -55
+CONSTANT: CL_INVALID_GLOBAL_OFFSET -56
+CONSTANT: CL_INVALID_EVENT_WAIT_LIST -57
+CONSTANT: CL_INVALID_EVENT -58
+CONSTANT: CL_INVALID_OPERATION -59
+CONSTANT: CL_INVALID_GL_OBJECT -60
+CONSTANT: CL_INVALID_BUFFER_SIZE -61
+CONSTANT: CL_INVALID_MIP_LEVEL -62
+CONSTANT: CL_INVALID_GLOBAL_WORK_SIZE -63
+
+CONSTANT: CL_VERSION_1_0 1
+
+CONSTANT: CL_FALSE 0
+CONSTANT: CL_TRUE 1
+
+CONSTANT: CL_PLATFORM_PROFILE HEX: 0900
+CONSTANT: CL_PLATFORM_VERSION HEX: 0901
+CONSTANT: CL_PLATFORM_NAME HEX: 0902
+CONSTANT: CL_PLATFORM_VENDOR HEX: 0903
+CONSTANT: CL_PLATFORM_EXTENSIONS HEX: 0904
+
+CONSTANT: CL_DEVICE_TYPE_DEFAULT 1
+CONSTANT: CL_DEVICE_TYPE_CPU 2
+CONSTANT: CL_DEVICE_TYPE_GPU 4
+CONSTANT: CL_DEVICE_TYPE_ACCELERATOR 8
+CONSTANT: CL_DEVICE_TYPE_ALL HEX: FFFFFFFF
+
+CONSTANT: CL_DEVICE_TYPE HEX: 1000
+CONSTANT: CL_DEVICE_VENDOR_ID HEX: 1001
+CONSTANT: CL_DEVICE_MAX_COMPUTE_UNITS HEX: 1002
+CONSTANT: CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS HEX: 1003
+CONSTANT: CL_DEVICE_MAX_WORK_GROUP_SIZE HEX: 1004
+CONSTANT: CL_DEVICE_MAX_WORK_ITEM_SIZES HEX: 1005
+CONSTANT: CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR HEX: 1006
+CONSTANT: CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT HEX: 1007
+CONSTANT: CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT HEX: 1008
+CONSTANT: CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG HEX: 1009
+CONSTANT: CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT HEX: 100A
+CONSTANT: CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE HEX: 100B
+CONSTANT: CL_DEVICE_MAX_CLOCK_FREQUENCY HEX: 100C
+CONSTANT: CL_DEVICE_ADDRESS_BITS HEX: 100D
+CONSTANT: CL_DEVICE_MAX_READ_IMAGE_ARGS HEX: 100E
+CONSTANT: CL_DEVICE_MAX_WRITE_IMAGE_ARGS HEX: 100F
+CONSTANT: CL_DEVICE_MAX_MEM_ALLOC_SIZE HEX: 1010
+CONSTANT: CL_DEVICE_IMAGE2D_MAX_WIDTH HEX: 1011
+CONSTANT: CL_DEVICE_IMAGE2D_MAX_HEIGHT HEX: 1012
+CONSTANT: CL_DEVICE_IMAGE3D_MAX_WIDTH HEX: 1013
+CONSTANT: CL_DEVICE_IMAGE3D_MAX_HEIGHT HEX: 1014
+CONSTANT: CL_DEVICE_IMAGE3D_MAX_DEPTH HEX: 1015
+CONSTANT: CL_DEVICE_IMAGE_SUPPORT HEX: 1016
+CONSTANT: CL_DEVICE_MAX_PARAMETER_SIZE HEX: 1017
+CONSTANT: CL_DEVICE_MAX_SAMPLERS HEX: 1018
+CONSTANT: CL_DEVICE_MEM_BASE_ADDR_ALIGN HEX: 1019
+CONSTANT: CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE HEX: 101A
+CONSTANT: CL_DEVICE_SINGLE_FP_CONFIG HEX: 101B
+CONSTANT: CL_DEVICE_GLOBAL_MEM_CACHE_TYPE HEX: 101C
+CONSTANT: CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE HEX: 101D
+CONSTANT: CL_DEVICE_GLOBAL_MEM_CACHE_SIZE HEX: 101E
+CONSTANT: CL_DEVICE_GLOBAL_MEM_SIZE HEX: 101F
+CONSTANT: CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE HEX: 1020
+CONSTANT: CL_DEVICE_MAX_CONSTANT_ARGS HEX: 1021
+CONSTANT: CL_DEVICE_LOCAL_MEM_TYPE HEX: 1022
+CONSTANT: CL_DEVICE_LOCAL_MEM_SIZE HEX: 1023
+CONSTANT: CL_DEVICE_ERROR_CORRECTION_SUPPORT HEX: 1024
+CONSTANT: CL_DEVICE_PROFILING_TIMER_RESOLUTION HEX: 1025
+CONSTANT: CL_DEVICE_ENDIAN_LITTLE HEX: 1026
+CONSTANT: CL_DEVICE_AVAILABLE HEX: 1027
+CONSTANT: CL_DEVICE_COMPILER_AVAILABLE HEX: 1028
+CONSTANT: CL_DEVICE_EXECUTION_CAPABILITIES HEX: 1029
+CONSTANT: CL_DEVICE_QUEUE_PROPERTIES HEX: 102A
+CONSTANT: CL_DEVICE_NAME HEX: 102B
+CONSTANT: CL_DEVICE_VENDOR HEX: 102C
+CONSTANT: CL_DRIVER_VERSION HEX: 102D
+CONSTANT: CL_DEVICE_PROFILE HEX: 102E
+CONSTANT: CL_DEVICE_VERSION HEX: 102F
+CONSTANT: CL_DEVICE_EXTENSIONS HEX: 1030
+CONSTANT: CL_DEVICE_PLATFORM HEX: 1031
+
+CONSTANT: CL_FP_DENORM 1
+CONSTANT: CL_FP_INF_NAN 2
+CONSTANT: CL_FP_ROUND_TO_NEAREST 4
+CONSTANT: CL_FP_ROUND_TO_ZERO 8
+CONSTANT: CL_FP_ROUND_TO_INF 16
+CONSTANT: CL_FP_FMA 32
+
+CONSTANT: CL_NONE 0
+CONSTANT: CL_READ_ONLY_CACHE 1
+CONSTANT: CL_READ_WRITE_CACHE 2
+
+CONSTANT: CL_LOCAL 1
+CONSTANT: CL_GLOBAL 2
+
+CONSTANT: CL_EXEC_KERNEL 1
+CONSTANT: CL_EXEC_NATIVE_KERNEL 2
+
+CONSTANT: CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE 1
+CONSTANT: CL_QUEUE_PROFILING_ENABLE 2
+
+CONSTANT: CL_CONTEXT_REFERENCE_COUNT HEX: 1080
+CONSTANT: CL_CONTEXT_DEVICES HEX: 1081
+CONSTANT: CL_CONTEXT_PROPERTIES HEX: 1082
+
+CONSTANT: CL_CONTEXT_PLATFORM HEX: 1084
+
+CONSTANT: CL_QUEUE_CONTEXT HEX: 1090
+CONSTANT: CL_QUEUE_DEVICE HEX: 1091
+CONSTANT: CL_QUEUE_REFERENCE_COUNT HEX: 1092
+CONSTANT: CL_QUEUE_PROPERTIES HEX: 1093
+
+CONSTANT: CL_MEM_READ_WRITE 1
+CONSTANT: CL_MEM_WRITE_ONLY 2
+CONSTANT: CL_MEM_READ_ONLY 4
+CONSTANT: CL_MEM_USE_HOST_PTR 8
+CONSTANT: CL_MEM_ALLOC_HOST_PTR 16
+CONSTANT: CL_MEM_COPY_HOST_PTR 32
+
+CONSTANT: CL_R HEX: 10B0
+CONSTANT: CL_A HEX: 10B1
+CONSTANT: CL_RG HEX: 10B2
+CONSTANT: CL_RA HEX: 10B3
+CONSTANT: CL_RGB HEX: 10B4
+CONSTANT: CL_RGBA HEX: 10B5
+CONSTANT: CL_BGRA HEX: 10B6
+CONSTANT: CL_ARGB HEX: 10B7
+CONSTANT: CL_INTENSITY HEX: 10B8
+CONSTANT: CL_LUMINANCE HEX: 10B9
+
+CONSTANT: CL_SNORM_INT8 HEX: 10D0
+CONSTANT: CL_SNORM_INT16 HEX: 10D1
+CONSTANT: CL_UNORM_INT8 HEX: 10D2
+CONSTANT: CL_UNORM_INT16 HEX: 10D3
+CONSTANT: CL_UNORM_SHORT_565 HEX: 10D4
+CONSTANT: CL_UNORM_SHORT_555 HEX: 10D5
+CONSTANT: CL_UNORM_INT_101010 HEX: 10D6
+CONSTANT: CL_SIGNED_INT8 HEX: 10D7
+CONSTANT: CL_SIGNED_INT16 HEX: 10D8
+CONSTANT: CL_SIGNED_INT32 HEX: 10D9
+CONSTANT: CL_UNSIGNED_INT8 HEX: 10DA
+CONSTANT: CL_UNSIGNED_INT16 HEX: 10DB
+CONSTANT: CL_UNSIGNED_INT32 HEX: 10DC
+CONSTANT: CL_HALF_FLOAT HEX: 10DD
+CONSTANT: CL_FLOAT HEX: 10DE
+
+CONSTANT: CL_MEM_OBJECT_BUFFER HEX: 10F0
+CONSTANT: CL_MEM_OBJECT_IMAGE2D HEX: 10F1
+CONSTANT: CL_MEM_OBJECT_IMAGE3D HEX: 10F2
+
+CONSTANT: CL_MEM_TYPE HEX: 1100
+CONSTANT: CL_MEM_FLAGS HEX: 1101
+CONSTANT: CL_MEM_SIZE HEX: 1102
+CONSTANT: CL_MEM_HOST_PTR HEX: 1103
+CONSTANT: CL_MEM_MAP_COUNT HEX: 1104
+CONSTANT: CL_MEM_REFERENCE_COUNT HEX: 1105
+CONSTANT: CL_MEM_CONTEXT HEX: 1106
+
+CONSTANT: CL_IMAGE_FORMAT HEX: 1110
+CONSTANT: CL_IMAGE_ELEMENT_SIZE HEX: 1111
+CONSTANT: CL_IMAGE_ROW_PITCH HEX: 1112
+CONSTANT: CL_IMAGE_SLICE_PITCH HEX: 1113
+CONSTANT: CL_IMAGE_WIDTH HEX: 1114
+CONSTANT: CL_IMAGE_HEIGHT HEX: 1115
+CONSTANT: CL_IMAGE_DEPTH HEX: 1116
+
+CONSTANT: CL_ADDRESS_NONE HEX: 1130
+CONSTANT: CL_ADDRESS_CLAMP_TO_EDGE HEX: 1131
+CONSTANT: CL_ADDRESS_CLAMP HEX: 1132
+CONSTANT: CL_ADDRESS_REPEAT HEX: 1133
+
+CONSTANT: CL_FILTER_NEAREST HEX: 1140
+CONSTANT: CL_FILTER_LINEAR HEX: 1141
+
+CONSTANT: CL_SAMPLER_REFERENCE_COUNT HEX: 1150
+CONSTANT: CL_SAMPLER_CONTEXT HEX: 1151
+CONSTANT: CL_SAMPLER_NORMALIZED_COORDS HEX: 1152
+CONSTANT: CL_SAMPLER_ADDRESSING_MODE HEX: 1153
+CONSTANT: CL_SAMPLER_FILTER_MODE HEX: 1154
+
+CONSTANT: CL_MAP_READ 1
+CONSTANT: CL_MAP_WRITE 2
+
+CONSTANT: CL_PROGRAM_REFERENCE_COUNT HEX: 1160
+CONSTANT: CL_PROGRAM_CONTEXT HEX: 1161
+CONSTANT: CL_PROGRAM_NUM_DEVICES HEX: 1162
+CONSTANT: CL_PROGRAM_DEVICES HEX: 1163
+CONSTANT: CL_PROGRAM_SOURCE HEX: 1164
+CONSTANT: CL_PROGRAM_BINARY_SIZES HEX: 1165
+CONSTANT: CL_PROGRAM_BINARIES HEX: 1166
+
+CONSTANT: CL_PROGRAM_BUILD_STATUS HEX: 1181
+CONSTANT: CL_PROGRAM_BUILD_OPTIONS HEX: 1182
+CONSTANT: CL_PROGRAM_BUILD_LOG HEX: 1183
+
+CONSTANT: CL_BUILD_SUCCESS 0
+CONSTANT: CL_BUILD_NONE -1
+CONSTANT: CL_BUILD_ERROR -2
+CONSTANT: CL_BUILD_IN_PROGRESS -3
+
+CONSTANT: CL_KERNEL_FUNCTION_NAME HEX: 1190
+CONSTANT: CL_KERNEL_NUM_ARGS HEX: 1191
+CONSTANT: CL_KERNEL_REFERENCE_COUNT HEX: 1192
+CONSTANT: CL_KERNEL_CONTEXT HEX: 1193
+CONSTANT: CL_KERNEL_PROGRAM HEX: 1194
+
+CONSTANT: CL_KERNEL_WORK_GROUP_SIZE HEX: 11B0
+CONSTANT: CL_KERNEL_COMPILE_WORK_GROUP_SIZE HEX: 11B1
+CONSTANT: CL_KERNEL_LOCAL_MEM_SIZE HEX: 11B2
+
+CONSTANT: CL_EVENT_COMMAND_QUEUE HEX: 11D0
+CONSTANT: CL_EVENT_COMMAND_TYPE HEX: 11D1
+CONSTANT: CL_EVENT_REFERENCE_COUNT HEX: 11D2
+CONSTANT: CL_EVENT_COMMAND_EXECUTION_STATUS HEX: 11D3
+
+CONSTANT: CL_COMMAND_NDRANGE_KERNEL HEX: 11F0
+CONSTANT: CL_COMMAND_TASK HEX: 11F1
+CONSTANT: CL_COMMAND_NATIVE_KERNEL HEX: 11F2
+CONSTANT: CL_COMMAND_READ_BUFFER HEX: 11F3
+CONSTANT: CL_COMMAND_WRITE_BUFFER HEX: 11F4
+CONSTANT: CL_COMMAND_COPY_BUFFER HEX: 11F5
+CONSTANT: CL_COMMAND_READ_IMAGE HEX: 11F6
+CONSTANT: CL_COMMAND_WRITE_IMAGE HEX: 11F7
+CONSTANT: CL_COMMAND_COPY_IMAGE HEX: 11F8
+CONSTANT: CL_COMMAND_COPY_IMAGE_TO_BUFFER HEX: 11F9
+CONSTANT: CL_COMMAND_COPY_BUFFER_TO_IMAGE HEX: 11FA
+CONSTANT: CL_COMMAND_MAP_BUFFER HEX: 11FB
+CONSTANT: CL_COMMAND_MAP_IMAGE HEX: 11FC
+CONSTANT: CL_COMMAND_UNMAP_MEM_OBJECT HEX: 11FD
+CONSTANT: CL_COMMAND_MARKER HEX: 11FE
+CONSTANT: CL_COMMAND_ACQUIRE_GL_OBJECTS HEX: 11FF
+CONSTANT: CL_COMMAND_RELEASE_GL_OBJECTS HEX: 1200
+
+CONSTANT: CL_COMPLETE HEX: 0
+CONSTANT: CL_RUNNING HEX: 1
+CONSTANT: CL_SUBMITTED HEX: 2
+CONSTANT: CL_QUEUED HEX: 3
+
+CONSTANT: CL_PROFILING_COMMAND_QUEUED HEX: 1280
+CONSTANT: CL_PROFILING_COMMAND_SUBMIT HEX: 1281
+CONSTANT: CL_PROFILING_COMMAND_START HEX: 1282
+CONSTANT: CL_PROFILING_COMMAND_END HEX: 1283
+
+FUNCTION: cl_int clGetPlatformIDs ( cl_uint num_entries, cl_platform_id* platforms, cl_uint* num_platforms ) ;
+FUNCTION: cl_int clGetPlatformInfo ( cl_platform_id platform, cl_platform_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret ) ;
+FUNCTION: cl_int clGetDeviceIDs ( cl_platform_id platform, cl_device_type device_type, cl_uint num_entries, cl_device_id* devices, cl_uint* num_devices ) ;
+FUNCTION: cl_int clGetDeviceInfo ( cl_device_id device, cl_device_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret ) ;
+CALLBACK: void cl_create_context_cb ( char* a, void* b, size_t s, void* c ) ;
+FUNCTION: cl_context clCreateContext ( cl_context_properties* properties, cl_uint num_devices, cl_device_id* devices, cl_create_context_cb pfn_notify, void* user_data, cl_int* errcode_ret ) ;
+FUNCTION: cl_context clCreateContextFromType ( cl_context_properties* properties, cl_device_type device_type, cl_create_context_cb pfn_notify, void* user_data, cl_int* errcode_ret ) ;
+FUNCTION: cl_int clRetainContext ( cl_context context ) ;
+FUNCTION: cl_int clReleaseContext ( cl_context context ) ;
+FUNCTION: cl_int clGetContextInfo ( cl_context context, cl_context_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret ) ;
+FUNCTION: cl_command_queue clCreateCommandQueue ( cl_context context, cl_device_id device, cl_command_queue_properties properties, cl_int* errcode_ret ) ;
+FUNCTION: cl_int clRetainCommandQueue ( cl_command_queue command_queue ) ;
+FUNCTION: cl_int clReleaseCommandQueue ( cl_command_queue command_queue ) ;
+FUNCTION: cl_int clGetCommandQueueInfo ( cl_command_queue command_queue, cl_command_queue_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret ) ;
+FUNCTION: cl_int clSetCommandQueueProperty ( cl_command_queue command_queue, cl_command_queue_properties properties, cl_bool enable, cl_command_queue_properties* old_properties ) ;
+FUNCTION: cl_mem clCreateBuffer ( cl_context context, cl_mem_flags flags, size_t size, void* host_ptr, cl_int* errcode_ret ) ;
+FUNCTION: cl_mem clCreateImage2D ( cl_context context, cl_mem_flags flags, cl_image_format* image_format, size_t image_width, size_t image_height, size_t image_row_pitch, void* host_ptr, cl_int* errcode_ret ) ;
+FUNCTION: cl_mem clCreateImage3D ( cl_context context, cl_mem_flags flags, cl_image_format* image_format, size_t image_width, size_t image_height, size_t image_depth, size_t image_row_pitch, size_t image_slice_pitch, void* host_ptr, cl_int* errcode_ret ) ;
+FUNCTION: cl_int clRetainMemObject ( cl_mem memobj ) ;
+FUNCTION: cl_int clReleaseMemObject ( cl_mem memobj ) ;
+FUNCTION: cl_int clGetSupportedImageFormats ( cl_context context, cl_mem_flags flags, cl_mem_object_type image_type, cl_uint num_entries, cl_image_format* image_formats, cl_uint* num_image_formats ) ;
+FUNCTION: cl_int clGetMemObjectInfo ( cl_mem memobj, cl_mem_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret ) ;
+FUNCTION: cl_int clGetImageInfo ( cl_mem image, cl_image_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret ) ;
+FUNCTION: cl_sampler clCreateSampler ( cl_context context, cl_bool normalized_coords, cl_addressing_mode addressing_mode, cl_filter_mode filter_mode, cl_int* errcode_ret ) ;
+FUNCTION: cl_int clRetainSampler ( cl_sampler sampler ) ;
+FUNCTION: cl_int clReleaseSampler ( cl_sampler sampler ) ;
+FUNCTION: cl_int clGetSamplerInfo ( cl_sampler sampler, cl_sampler_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret ) ;
+FUNCTION: cl_program clCreateProgramWithSource ( cl_context context, cl_uint count, char** strings, size_t* lengths, cl_int* errcode_ret ) ;
+FUNCTION: cl_program clCreateProgramWithBinary ( cl_context context, cl_uint num_devices, cl_device_id* device_list, size_t* lengths, char** binaries, cl_int* binary_status, cl_int* errcode_ret ) ;
+FUNCTION: cl_int clRetainProgram ( cl_program program ) ;
+FUNCTION: cl_int clReleaseProgram ( cl_program program ) ;
+CALLBACK: void cl_build_program_cb ( cl_program program, void* user_data ) ;
+FUNCTION: cl_int clBuildProgram ( cl_program program, cl_uint num_devices, cl_device_id* device_list, char* options, cl_build_program_cb pfn_notify, void* user_data ) ;
+FUNCTION: cl_int clUnloadCompiler ( ) ;
+FUNCTION: cl_int clGetProgramInfo ( cl_program program, cl_program_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret ) ;
+FUNCTION: cl_int clGetProgramBuildInfo ( cl_program program, cl_device_id device, cl_program_build_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret ) ;
+FUNCTION: cl_kernel clCreateKernel ( cl_program program, char* kernel_name, cl_int* errcode_ret ) ;
+FUNCTION: cl_int clCreateKernelsInProgram ( cl_program program, cl_uint num_kernels, cl_kernel* kernels, cl_uint* num_kernels_ret ) ;
+FUNCTION: cl_int clRetainKernel ( cl_kernel kernel ) ;
+FUNCTION: cl_int clReleaseKernel ( cl_kernel kernel ) ;
+FUNCTION: cl_int clSetKernelArg ( cl_kernel kernel, cl_uint arg_index, size_t arg_size, void* arg_value ) ;
+FUNCTION: cl_int clGetKernelInfo ( cl_kernel kernel, cl_kernel_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret ) ;
+FUNCTION: cl_int clGetKernelWorkGroupInfo ( cl_kernel kernel, cl_device_id device, cl_kernel_work_group_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret ) ;
+FUNCTION: cl_int clWaitForEvents ( cl_uint num_events, cl_event* event_list ) ;
+FUNCTION: cl_int clGetEventInfo ( cl_event event, cl_event_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret ) ;
+FUNCTION: cl_int clRetainEvent ( cl_event event ) ;
+FUNCTION: cl_int clReleaseEvent ( cl_event event ) ;
+FUNCTION: cl_int clGetEventProfilingInfo ( cl_event event, cl_profiling_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret ) ;
+FUNCTION: cl_int clFlush ( cl_command_queue command_queue ) ;
+FUNCTION: cl_int clFinish ( cl_command_queue command_queue ) ;
+FUNCTION: cl_int clEnqueueReadBuffer ( cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_read, size_t offset, size_t cb, void* ptr, cl_uint num_events_in_wait_list, cl_event* event_wait_list, cl_event* event ) ;
+FUNCTION: cl_int clEnqueueWriteBuffer ( cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_write, size_t offset, size_t cb, void* ptr, cl_uint num_events_in_wait_list, cl_event* event_wait_list, cl_event* event ) ;
+FUNCTION: cl_int clEnqueueCopyBuffer ( cl_command_queue command_queue, cl_mem src_buffer, cl_mem dst_buffer, size_t src_offset, size_t dst_offset, size_t cb, cl_uint num_events_in_wait_list, cl_event* event_wait_list, cl_event* event ) ;
+FUNCTION: cl_int clEnqueueReadImage ( cl_command_queue command_queue, cl_mem image, cl_bool blocking_read, size_t** origin, size_t** region, size_t row_pitch, size_t slice_pitch, void* ptr, cl_uint num_events_in_wait_list, cl_event* event_wait_list, cl_event* event ) ;
+FUNCTION: cl_int clEnqueueWriteImage ( cl_command_queue command_queue, cl_mem image, cl_bool blocking_write, size_t** origin, size_t** region, size_t input_row_pitch, size_t input_slice_pitch, void* ptr, cl_uint num_events_in_wait_list, cl_event* event_wait_list, cl_event* event ) ;
+FUNCTION: cl_int clEnqueueCopyImage ( cl_command_queue command_queue, cl_mem src_image, cl_mem dst_image, size_t** src_origin, size_t** dst_origin, size_t** region, cl_uint num_events_in_wait_list, cl_event* event_wait_list, cl_event* event ) ;
+FUNCTION: cl_int clEnqueueCopyImageToBuffer ( cl_command_queue command_queue, cl_mem src_image, cl_mem dst_buffer, size_t** src_origin, size_t** region, size_t dst_offset, cl_uint num_events_in_wait_list, cl_event* event_wait_list, cl_event* event ) ;
+FUNCTION: cl_int clEnqueueCopyBufferToImage ( cl_command_queue command_queue, cl_mem src_buffer, cl_mem dst_image, size_t src_offset, size_t** dst_origin, size_t** region, cl_uint num_events_in_wait_list, cl_event* event_wait_list, cl_event* event ) ;
+FUNCTION: void* clEnqueueMapBuffer ( cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_map, cl_map_flags map_flags, size_t offset, size_t cb, cl_uint num_events_in_wait_list, cl_event* event_wait_list, cl_event* event, cl_int* errcode_ret ) ;
+FUNCTION: void* clEnqueueMapImage ( cl_command_queue command_queue, cl_mem image, cl_bool blocking_map, cl_map_flags map_flags, size_t** origin, size_t** region, size_t* image_row_pitch, size_t* image_slice_pitch, cl_uint num_events_in_wait_list, cl_event* event_wait_list, cl_event* event, cl_int* errcode_ret ) ;
+FUNCTION: cl_int clEnqueueUnmapMemObject ( cl_command_queue command_queue, cl_mem memobj, void* mapped_ptr, cl_uint num_events_in_wait_list, cl_event* event_wait_list, cl_event* event ) ;
+FUNCTION: cl_int clEnqueueNDRangeKernel ( cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim, size_t* global_work_offset, size_t* global_work_size, size_t* local_work_size, cl_uint num_events_in_wait_list, cl_event* event_wait_list, cl_event* event ) ;
+CALLBACK: void cl_enqueue_task_cb ( void* args ) ;
+FUNCTION: cl_int clEnqueueTask ( cl_command_queue command_queue, cl_kernel kernel, cl_uint num_events_in_wait_list, cl_event* event_wait_list, cl_event* event ) ;
+FUNCTION: cl_int clEnqueueNativeKernel ( cl_command_queue command_queue, cl_enqueue_task_cb user_func, void* args, size_t cb_args, cl_uint num_mem_objects, cl_mem* mem_list, void** args_mem_loc, cl_uint num_events_in_wait_list, cl_event* event_wait_list, cl_event* event ) ;
+FUNCTION: cl_int clEnqueueMarker ( cl_command_queue command_queue, cl_event* event ) ;
+FUNCTION: cl_int clEnqueueWaitForEvents ( cl_command_queue command_queue, cl_uint num_events, cl_event* event_list ) ;
+FUNCTION: cl_int clEnqueueBarrier ( cl_command_queue command_queue ) ;
+FUNCTION: void* clGetExtensionFunctionAddress ( char* func_name ) ;
+
+! cl_ext.h
+CONSTANT: CL_DEVICE_DOUBLE_FP_CONFIG HEX: 1032
+CONSTANT: CL_DEVICE_HALF_FP_CONFIG HEX: 1033
+
+! cl_khr_icd.txt
+CONSTANT: CL_PLATFORM_ICD_SUFFIX_KHR HEX: 0920
+CONSTANT: CL_PLATFORM_NOT_FOUND_KHR -1001
+
+FUNCTION: cl_int clIcdGetPlatformIDsKHR ( cl_uint num_entries, cl_platform_id* platforms, cl_uint* num_platforms ) ;
+
+! cl_gl.h
+TYPEDEF: cl_uint cl_gl_object_type
+TYPEDEF: cl_uint cl_gl_texture_info
+TYPEDEF: cl_uint cl_gl_platform_info
+
+CONSTANT: CL_GL_OBJECT_BUFFER HEX: 2000
+CONSTANT: CL_GL_OBJECT_TEXTURE2D HEX: 2001
+CONSTANT: CL_GL_OBJECT_TEXTURE3D HEX: 2002
+CONSTANT: CL_GL_OBJECT_RENDERBUFFER HEX: 2003
+CONSTANT: CL_GL_TEXTURE_TARGET HEX: 2004
+CONSTANT: CL_GL_MIPMAP_LEVEL HEX: 2005
+
+FUNCTION: cl_mem clCreateFromGLBuffer ( cl_context context, cl_mem_flags flags, cl_GLuint bufobj, int* errcode_ret ) ;
+FUNCTION: cl_mem clCreateFromGLTexture2D ( cl_context context, cl_mem_flags flags, cl_GLenum target, cl_GLint miplevel, cl_GLuint texture, cl_int* errcode_ret ) ;
+FUNCTION: cl_mem clCreateFromGLTexture3D ( cl_context context, cl_mem_flags flags, cl_GLenum target, cl_GLint miplevel, cl_GLuint texture, cl_int* errcode_ret ) ;
+FUNCTION: cl_mem clCreateFromGLRenderbuffer ( cl_context context, cl_mem_flags flags, cl_GLuint renderbuffer, cl_int* errcode_ret ) ;
+FUNCTION: cl_int clGetGLObjectInfo ( cl_mem memobj, cl_gl_object_type* gl_object_type, cl_GLuint* gl_object_name ) ;
+FUNCTION: cl_int clGetGLTextureInfo ( cl_mem memobj, cl_gl_texture_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret ) ;
+FUNCTION: cl_int clEnqueueAcquireGLObjects ( cl_command_queue command_queue, cl_uint num_objects, cl_mem* mem_objects, cl_uint num_events_in_wait_list, cl_event* event_wait_list, cl_event* event ) ;
+FUNCTION: cl_int clEnqueueReleaseGLObjects ( cl_command_queue command_queue, cl_uint num_objects, cl_mem* mem_objects, cl_uint num_events_in_wait_list, cl_event* event_wait_list, cl_event* event ) ;
+
+! cl_khr_gl_sharing.txt
+TYPEDEF: cl_uint cl_gl_context_info
+
+CONSTANT: CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR -1000
+CONSTANT: CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR HEX: 2006
+CONSTANT: CL_DEVICES_FOR_GL_CONTEXT_KHR HEX: 2007
+CONSTANT: CL_GL_CONTEXT_KHR HEX: 2008
+CONSTANT: CL_EGL_DISPLAY_KHR HEX: 2009
+CONSTANT: CL_GLX_DISPLAY_KHR HEX: 200A
+CONSTANT: CL_WGL_HDC_KHR HEX: 200B
+CONSTANT: CL_CGL_SHAREGROUP_KHR HEX: 200C
+
+FUNCTION: cl_int clGetGLContextInfoKHR ( cl_context_properties* properties, cl_gl_context_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret ) ;
+
+! cl_nv_d3d9_sharing.txt
+CONSTANT: CL_D3D9_DEVICE_NV HEX: 4022
+CONSTANT: CL_D3D9_ADAPTER_NAME_NV HEX: 4023
+CONSTANT: CL_PREFERRED_DEVICES_FOR_D3D9_NV HEX: 4024
+CONSTANT: CL_ALL_DEVICES_FOR_D3D9_NV HEX: 4025
+CONSTANT: CL_CONTEXT_D3D9_DEVICE_NV HEX: 4026
+CONSTANT: CL_MEM_D3D9_RESOURCE_NV HEX: 4027
+CONSTANT: CL_IMAGE_D3D9_FACE_NV HEX: 4028
+CONSTANT: CL_IMAGE_D3D9_LEVEL_NV HEX: 4029
+CONSTANT: CL_COMMAND_ACQUIRE_D3D9_OBJECTS_NV HEX: 402A
+CONSTANT: CL_COMMAND_RELEASE_D3D9_OBJECTS_NV HEX: 402B
+CONSTANT: CL_INVALID_D3D9_DEVICE_NV -1010
+CONSTANT: CL_INVALID_D3D9_RESOURCE_NV -1011
+CONSTANT: CL_D3D9_RESOURCE_ALREADY_ACQUIRED_NV -1012
+CONSTANT: CL_D3D9_RESOURCE_NOT_ACQUIRED_NV -1013
+
+TYPEDEF: void* cl_d3d9_device_source_nv
+TYPEDEF: void* cl_d3d9_device_set_nv
+
+FUNCTION: cl_int clGetDeviceIDsFromD3D9NV ( cl_platform_id platform, cl_d3d9_device_source_nv d3d_device_source, void* d3d_object, cl_d3d9_device_set_nv d3d_device_set, cl_uint num_entries, cl_device_id* devices, cl_uint* num_devices ) ;
+FUNCTION: cl_mem clCreateFromD3D9VertexBufferNV ( cl_context context, cl_mem_flags flags, void* id3dvb9_resource, cl_int* errcode_ret ) ;
+FUNCTION: cl_mem clCreateFromD3D9IndexBufferNV ( cl_context context, cl_mem_flags flags, void* id3dib9_resource, cl_int* errcode_ret ) ;
+FUNCTION: cl_mem clCreateFromD3D9SurfaceNV ( cl_context context, cl_mem_flags flags, void* id3dsurface9_resource, cl_int* errcode_ret ) ;
+FUNCTION: cl_mem clCreateFromD3D9TextureNV ( cl_context context, cl_mem_flags flags, void* id3dtexture9_resource, uint miplevel, cl_int* errcode_ret ) ;
+FUNCTION: cl_mem clCreateFromD3D9CubeTextureNV ( cl_context context, cl_mem_flags flags, void* id3dct9_resource, int facetype, uint miplevel, cl_int* errcode_ret ) ;
+FUNCTION: cl_mem clCreateFromD3D9VolumeTextureNV ( cl_context context, cl_mem_flags flags, void* id3dvt9-resource, uint miplevel, cl_int* errcode_ret ) ;
+FUNCTION: cl_int clEnqueueAcquireD3D9ObjectsNV ( cl_command_queue command_queue, cl_uint num_objects, cl_mem* mem_objects, cl_uint num_events_in_wait_list, cl_event* event_wait_list, cl_event* event ) ;
+FUNCTION: cl_int clEnqueueReleaseD3D9ObjectsNV ( cl_command_queue command_queue, cl_uint num_objects, cl_mem* mem_objects, cl_uint num_events_in_wait_list, cl_event* event_wait_list, cl_event* event ) ;
+
+! cl_nv_d3d10_sharing.txt
+CONSTANT: CL_D3D10_DEVICE_NV HEX: 4010
+CONSTANT: CL_D3D10_DXGI_ADAPTER_NV HEX: 4011
+CONSTANT: CL_PREFERRED_DEVICES_FOR_D3D10_NV HEX: 4012
+CONSTANT: CL_ALL_DEVICES_FOR_D3D10_NV HEX: 4013
+CONSTANT: CL_CONTEXT_D3D10_DEVICE_NV HEX: 4014
+CONSTANT: CL_MEM_D3D10_RESOURCE_NV HEX: 4015
+CONSTANT: CL_IMAGE_D3D10_SUBRESOURCE_NV HEX: 4016
+CONSTANT: CL_COMMAND_ACQUIRE_D3D10_OBJECTS_NV HEX: 4017
+CONSTANT: CL_COMMAND_RELEASE_D3D10_OBJECTS_NV HEX: 4018
+CONSTANT: CL_INVALID_D3D10_DEVICE_NV -1002
+CONSTANT: CL_INVALID_D3D10_RESOURCE_NV -1003
+CONSTANT: CL_D3D10_RESOURCE_ALREADY_ACQUIRED_NV -1004
+CONSTANT: CL_D3D10_RESOURCE_NOT_ACQUIRED_NV -1005
+
+TYPEDEF: void* cl_d3d10_device_source_nv
+TYPEDEF: void* cl_d3d10_device_set_nv
+
+FUNCTION: cl_int clGetDeviceIDsFromD3D10NV ( cl_platform_id platform, cl_d3d10_device_source_nv d3d_device_source, void* d3d_object, cl_d3d10_device_set_nv d3d_device_set, cl_uint num_entries, cl_device_id* devices, cl_uint* num_devices ) ;
+FUNCTION: cl_mem clCreateFromD3D10BufferNV ( cl_context context, cl_mem_flags flags, void* id3d10buffer_resource, cl_int* errcode_ret ) ;
+FUNCTION: cl_mem clCreateFromD3D10Texture2DNV ( cl_context context, cl_mem_flags flags, void* id3d10texture2d_resource, uint subresource, cl_int* errcode_ret ) ;
+FUNCTION: cl_mem clCreateFromD3D10Texture3DNV ( cl_context context, cl_mem_flags flags, void* id3d10texture3d_resource, uint subresource, cl_int* errcode_ret ) ;
+FUNCTION: cl_int clEnqueueAcquireD3D10ObjectsNV ( cl_command_queue command_queue, cl_uint num_objects, cl_mem* mem_objects, cl_uint num_events_in_wait_list, cl_event* event_wait_list, cl_event* event ) ;
+FUNCTION: cl_int clEnqueueReleaseD3D10ObjectsNV ( cl_command_queue command_queue, cl_uint num_objects, cl_mem* mem_objects, cl_uint num_events_in_wait_list, cl_event* event_wait_list, cl_event* event ) ;
+
+! cl_nv_d3d11_sharing.txt
+CONSTANT: CL_D3D11_DEVICE_NV HEX: 4019
+CONSTANT: CL_D3D11_DXGI_ADAPTER_NV HEX: 401A
+CONSTANT: CL_PREFERRED_DEVICES_FOR_D3D11_NV HEX: 401B
+CONSTANT: CL_ALL_DEVICES_FOR_D3D11_NV HEX: 401C
+CONSTANT: CL_CONTEXT_D3D11_DEVICE_NV HEX: 401D
+CONSTANT: CL_MEM_D3D11_RESOURCE_NV HEX: 401E
+CONSTANT: CL_IMAGE_D3D11_SUBRESOURCE_NV HEX: 401F
+CONSTANT: CL_COMMAND_ACQUIRE_D3D11_OBJECTS_NV HEX: 4020
+CONSTANT: CL_COMMAND_RELEASE_D3D11_OBJECTS_NV HEX: 4021
+CONSTANT: CL_INVALID_D3D11_DEVICE_NV -1006
+CONSTANT: CL_INVALID_D3D11_RESOURCE_NV -1007
+CONSTANT: CL_D3D11_RESOURCE_ALREADY_ACQUIRED_NV -1008
+CONSTANT: CL_D3D11_RESOURCE_NOT_ACQUIRED_NV -1009
+
+TYPEDEF: void* cl_d3d11_device_source_nv
+TYPEDEF: void* cl_d3d11_device_set_nv
+
+FUNCTION: cl_int clGetDeviceIDsFromD3D11NV ( cl_platform_id platform, cl_d3d11_device_source_nv d3d_device_source, void* d3d_object, cl_d3d11_device_set_nv d3d_device_set, cl_uint num_entries, cl_device_id* devices, cl_uint* num_devices ) ;
+FUNCTION: cl_mem clCreateFromD3D11BufferNV ( cl_context context, cl_mem_flags flags, void* id3d11buffer_resource, cl_int* errcode_ret ) ;
+FUNCTION: cl_mem clCreateFromD3D11Texture2DNV ( cl_context context, cl_mem_flags flags, void* id3d11texture2d_resource, uint subresource, cl_int* errcode_ret ) ;
+FUNCTION: cl_mem clCreateFromD3D11Texture3DNV ( cl_context context, cl_mem_flags flags, void* id3dtexture3d_resource, uint subresource, cl_int* errcode_ret ) ;
+FUNCTION: cl_int clEnqueueAcquireD3D11ObjectsNV ( cl_command_queue command_queue, cl_uint num_objects, cl_mem* mem_objects, cl_uint num_events_in_wait_list, cl_event* event_wait_list, cl_event* event ) ;
+FUNCTION: cl_int clEnqueueReleaseD3D11ObjectsNV ( cl_command_queue command_queue, cl_uint num_objects, cl_mem* mem_objects, cl_uint num_events_in_wait_list, cl_event* event_wait_list, cl_event* event ) ;
+
+! Utility words needed for working with the API
+: *size_t ( c-ptr -- value )
+ size_t heap-size {
+ { 4 [ 0 alien-unsigned-4 ] }
+ { 8 [ 0 alien-unsigned-8 ] }
+ } case ; inline
+
+: <size_t> ( value -- c-ptr )
+ size_t heap-size [ (byte-array) ] keep {
+ { 4 [ [ 0 set-alien-unsigned-4 ] keep ] }
+ { 8 [ [ 0 set-alien-unsigned-8 ] keep ] }
+ } case ; inline
--- /dev/null
+Bindings to OpenCL
--- /dev/null
+bindings
+untested
--- /dev/null
+! Copyright (C) 2010 Erik Charlebois.
+! See http://factorcode.org/license.txt for BSD license.
+USING: help.markup help.syntax kernel quotations strings opencl.private
+math byte-arrays alien ;
+IN: opencl
+
+HELP: cl-addressing-mode
+{ $values
+ { "sampler" cl-sampler }
+ { "addressing-mode" cl-addressing-mode }
+}
+{ $description "Returns the addressing mode of the given sampler." } ;
+
+HELP: cl-barrier
+{ $description "Insert a synchronization barrier into the current command queue." } ;
+
+HELP: cl-barrier-events
+{ $values
+ { "event/events" "a single event or sequence of events" }
+}
+{ $description "Insert a synchronization barrier for the specified events into the current command queue." } ;
+
+HELP: cl-buffer
+{ $var-description "Tuple wrapper which will release the memory object handle when disposed." } ;
+
+HELP: cl-buffer-ptr
+{ $var-description "A buffer and offset pair for specifying a starting point for a copy." } ;
+
+HELP: cl-buffer-range
+{ $var-description "A buffer, offset and size triplet for specifying copy ranges." } ;
+
+HELP: cl-context
+{ $var-description "Tuple wrapper which will release the context handle when disposed." } ;
+
+HELP: cl-current-context
+{ $var-description "Symbol for the current cl-context tuple." } ;
+
+HELP: cl-current-device
+{ $var-description "Symbol for the current cl-device tuple." } ;
+
+HELP: cl-current-queue
+{ $var-description "Symbol for the current cl-queue tuple." } ;
+
+HELP: cl-device
+{ $var-description "Tuple wrapper which will release the device handle when disposed." } ;
+
+HELP: cl-event
+{ $var-description "Tuple wrapper which will release the event handle when disposed." } ;
+
+HELP: cl-event-status
+{ $values
+ { "event" cl-event }
+ { "execution-status" cl-execution-status }
+}
+{ $description "Returns the current execution status of the operation represented by the event." } ;
+
+HELP: cl-event-type
+{ $values
+ { "event" cl-event }
+ { "command-type" cl-execution-status }
+}
+{ $description "Returns the type of operation that created the event." } ;
+
+HELP: cl-filter-mode
+{ $values
+ { "sampler" cl-sampler }
+ { "filter-mode" cl-filter-mode }
+}
+{ $description "Returns the filter mode of the sampler object." } ;
+
+HELP: cl-finish
+{ $description "Flush the current command queue and wait till all operations are completed." } ;
+
+HELP: cl-flush
+{ $description "Flush the current command queue to kick off pending operations." } ;
+
+HELP: cl-kernel
+{ $var-description "Tuple wrapper which will release the kernel handle when disposed." } ;
+
+HELP: cl-kernel-arity
+{ $values
+ { "kernel" cl-kernel }
+ { "arity" integer }
+}
+{ $description "Returns the number of inputs that this kernel function accepts." } ;
+
+HELP: cl-kernel-local-size
+{ $values
+ { "kernel" cl-kernel }
+ { "size" integer }
+}
+{ $description "Returns the maximum size of a local work group for this kernel." } ;
+
+HELP: cl-kernel-name
+{ $values
+ { "kernel" cl-kernel }
+ { "string" string }
+}
+{ $description "Returns the name of the kernel function." } ;
+
+HELP: cl-marker
+{ $values
+
+ { "event" cl-event }
+}
+{ $description "Inserts a marker into the current command queue." } ;
+
+HELP: cl-normalized-coords?
+{ $values
+ { "sampler" cl-sampler }
+ { "?" boolean }
+}
+{ $description "Returns whether the sampler uses normalized coords or not." } ;
+
+HELP: cl-out-of-order-execution?
+{ $values
+ { "command-queue" cl-queue }
+ { "?" boolean }
+}
+{ $description "Returns whether the given command queue allows out of order execution or not." } ;
+
+HELP: cl-platform
+{ $var-description "Tuple summarizing the capabilities and devices of an OpenCL platform." } ;
+
+HELP: cl-platforms
+{ $values
+
+ { "platforms" "sequence of cl-platform"}
+}
+{ $description "Returns the platforms available for OpenCL computation on this hardware." } ;
+
+HELP: cl-profile-counters
+{ $values
+ { "event" cl-event }
+ { "queued" integer } { "submitted" integer } { "started" integer } { "finished" integer }
+}
+{ $description "Returns the profiling counters for the operation represented by event." } ;
+
+HELP: cl-profiling?
+{ $values
+ { "command-queue" cl-queue }
+ { "?" boolean }
+}
+{ $description "Returns true if the command queue allows profiling." } ;
+
+HELP: cl-program
+{ $var-description "Tuple wrapper which will release the program handle when disposed." } ;
+
+HELP: cl-queue
+{ $var-description "Tuple wrapper which will release the command queue handle when disposed." } ;
+
+HELP: cl-read-buffer
+{ $values
+ { "buffer-range" cl-buffer-range }
+ { "byte-array" byte-array }
+}
+{ $description "Synchronously read a byte-array from the specified buffer location." } ;
+
+HELP: cl-sampler
+{ $var-description "Tuple wrapper which will release the sampler handle when disposed." } ;
+
+HELP: cl-queue-copy-buffer
+{ $values
+ { "src-buffer-ptr" cl-buffer-ptr } { "dst-buffer-ptr" cl-buffer-ptr } { "size" integer } { "dependent-events" "sequence of events" }
+ { "event" cl-event }
+}
+{ $description "Queue a copy operation from " { $snippet "src-buffer-ptr" } " to " { $snippet "dst-buffer-ptr" } ". Dependent events can be passed to order the operation relative to other operations." } ;
+
+HELP: cl-queue-kernel
+{ $values
+ { "kernel" cl-kernel } { "args" "sequence of cl-buffer or byte-array" } { "sizes" "sequence of integers" } { "dependent-events" "sequence of events" }
+ { "event" cl-event }
+}
+{ $description "Queue a kernel for execution with the given arguments. The " { $snippet "sizes" } " argument specifies input array sizes for each dimension. Dependent events can be passed to order the operation relative to other operations." } ;
+
+HELP: cl-queue-read-buffer
+{ $values
+ { "buffer-range" cl-buffer-range } { "alien" alien } { "dependent-events" "a sequence of events" }
+ { "event" cl-event }
+}
+{ $description "Queue a read operation from " { $snippet "buffer-range" } " to " { $snippet "alien" } ". Dependent events can be passed to order the operation relative to other operations." } ;
+
+HELP: cl-queue-write-buffer
+{ $values
+ { "buffer-range" cl-buffer-range } { "alien" alien } { "dependent-events" "a sequence of events" }
+ { "event" cl-event }
+}
+{ $description "Queue a write operation from " { $snippet "alien" } " to " { $snippet "buffer-range" } ". Dependent events can be passed to order the operation relative to other operations." } ;
+
+HELP: cl-wait
+{ $values
+ { "event/events" "a single event or sequence of events" }
+}
+{ $description "Synchronously wait for the events to complete." } ;
+
+HELP: cl-write-buffer
+{ $values
+ { "buffer-range" cl-buffer-range } { "byte-array" byte-array }
+}
+{ $description "Synchronously write a byte-array to the specified buffer location." } ;
+
+HELP: <cl-program>
+{ $values
+ { "options" string } { "strings" "sequence of source code strings" }
+ { "program" "compiled cl-program" }
+}
+{ $description "Compile the given source code and return a program object. A " { $link cl-error } " is thrown in the event of a compile error." } ;
+
+HELP: with-cl-state
+{ $values
+ { "context/f" { $maybe cl-context } } { "device/f" { $maybe cl-device } } { "queue/f" { $maybe cl-queue } } { "quot" quotation }
+}
+{ $description "Run the specified quotation with the given context, device and command queue. False arguments are not bound." } ;
+
+ARTICLE: "opencl" "OpenCL"
+"The " { $vocab-link "opencl" } " vocabulary provides high-level words for using OpenCL."
+{ $subsections
+ cl-platforms
+ <cl-queue>
+ with-cl-state
+}
+"Memory Objects:"
+{ $subsections
+ <cl-buffer>
+ cl-queue-copy-buffer
+ cl-read-buffer
+ cl-queue-read-buffer
+ cl-write-buffer
+ cl-queue-write-buffer
+}
+"Programs and Kernels:"
+{ $subsections
+ <cl-program>
+ <cl-kernel>
+}
+
+"Running and Waiting for Completion:"
+{ $subsections
+ cl-queue-kernel
+ cl-wait
+ cl-flush
+ cl-finish
+}
+;
+
+ABOUT: "opencl"
--- /dev/null
+! Copyright (C) 2010 Erik Charlebois.
+! See http://factorcode.org/license.txt for BSD license.
+USING: multiline locals io.encodings.ascii io.encodings.string sequences
+math specialized-arrays alien.c-types math.order alien opencl tools.test
+accessors arrays destructors kernel namespaces ;
+FROM: alien.c-types => float ;
+SPECIALIZED-ARRAY: float
+IN: opencl.tests
+
+STRING: kernel-source
+__kernel void square(
+ __global float* input,
+ __global float* output,
+ const unsigned int count)
+{
+ int i = get_global_id(0);
+ if (i < count)
+ output[i] = input[i] * input[i];
+}
+;
+
+:: opencl-square ( in -- out )
+ [
+ in byte-length :> num-bytes
+ in length :> num-floats
+ cl-platforms first devices>> first :> device
+ device 1array <cl-context> &dispose :> context
+ context device f f <cl-queue> &dispose :> queue
+
+ context device queue [
+ "" kernel-source 1array <cl-program> &dispose "square" <cl-kernel> &dispose :> kernel
+ cl-read-access num-bytes in <cl-buffer> &dispose :> in-buffer
+ cl-write-access num-bytes f <cl-buffer> &dispose :> out-buffer
+
+ kernel in-buffer out-buffer num-floats <uint> 3array
+ { num-floats } [ ] cl-queue-kernel &dispose drop
+
+ cl-finish
+ out-buffer 0 num-bytes <cl-buffer-range> cl-read-buffer num-floats <direct-float-array>
+ ] with-cl-state
+ ] with-destructors ;
+
+[ float-array{ 1.0 4.0 9.0 16.0 100.0 } ]
+[ float-array{ 1.0 2.0 3.0 4.0 10.0 } opencl-square ] unit-test
--- /dev/null
+! Copyright (C) 2010 Erik Charlebois.
+! See http://factorcode.org/license.txt for BSD license.
+USING: accessors alien alien.accessors alien.c-types arrays
+byte-arrays combinators combinators.smart continuations destructors
+fry io.encodings.ascii io.encodings.string kernel libc locals macros
+math math.order multiline opencl.ffi prettyprint sequences
+specialized-arrays typed variants namespaces ;
+IN: opencl
+SPECIALIZED-ARRAYS: void* char size_t ;
+
+<PRIVATE
+ERROR: cl-error err ;
+
+: cl-success ( err -- )
+ dup CL_SUCCESS = [ drop ] [ cl-error ] if ; inline
+
+: cl-not-null ( err -- )
+ dup f = [ cl-error ] [ drop ] if ; inline
+
+MACRO: info ( info-quot lift-quot -- quot )
+ [ dup ] dip '[ 2dup 0 f 0 <size_t> _ '[ _ call cl-success ] keep
+ *size_t dup <byte-array> _ '[ f _ call cl-success ] keep
+ _ call ] ;
+
+MACRO: 2info ( info-quot lift-quot -- quot )
+ [ dup ] dip '[ 3dup 0 f 0 <size_t> _ '[ _ call cl-success ] keep
+ *size_t dup <byte-array> _ '[ f _ call cl-success ] keep
+ _ call ] ;
+
+: info-bool ( handle name quot -- ? )
+ [ *uint CL_TRUE = ] info ; inline
+
+: info-ulong ( handle name quot -- ulong )
+ [ *ulonglong ] info ; inline
+
+: info-int ( handle name quot -- int )
+ [ *int ] info ; inline
+
+: info-uint ( handle name quot -- uint )
+ [ *uint ] info ; inline
+
+: info-size_t ( handle name quot -- size_t )
+ [ *size_t ] info ; inline
+
+: 2info-size_t ( handle1 handle2 name quot -- size_t )
+ [ *size_t ] 2info ; inline
+
+: info-string ( handle name quot -- string )
+ [ ascii decode 1 head* ] info ; inline
+
+: 2info-string ( handle name quot -- string )
+ [ ascii decode 1 head* ] 2info ; inline
+
+: info-size_t-array ( handle name quot -- size_t-array )
+ [ [ length size_t heap-size / ] keep swap <direct-size_t-array> ] info ; inline
+
+TUPLE: cl-handle < disposable handle ;
+PRIVATE>
+
+VARIANT: cl-device-type
+ cl-device-default cl-device-cpu cl-device-gpu cl-device-accelerator ;
+
+: size_t>cl-device-type ( size_t -- cl-device-type )
+ {
+ { CL_DEVICE_TYPE_DEFAULT [ cl-device-default ] }
+ { CL_DEVICE_TYPE_CPU [ cl-device-cpu ] }
+ { CL_DEVICE_TYPE_GPU [ cl-device-gpu ] }
+ { CL_DEVICE_TYPE_ACCELERATOR [ cl-device-accelerator ] }
+ } case ; inline
+
+VARIANT: cl-fp-feature
+ cl-denorm cl-inf-and-nan cl-round-to-nearest cl-round-to-zero cl-round-to-inf cl-fma ;
+
+VARIANT: cl-cache-type
+ cl-no-cache cl-read-only-cache cl-read-write-cache ;
+
+VARIANT: cl-buffer-access-mode
+ cl-read-access cl-write-access cl-read-write-access ;
+
+VARIANT: cl-image-channel-order
+ cl-channel-order-r cl-channel-order-a cl-channel-order-rg cl-channel-order-ra
+ cl-channel-order-rga cl-channel-order-rgba cl-channel-order-bgra cl-channel-order-argb
+ cl-channel-order-intensity cl-channel-order-luminance ;
+
+VARIANT: cl-image-channel-type
+ cl-channel-type-snorm-int8 cl-channel-type-snorm-int16 cl-channel-type-unorm-int8
+ cl-channel-type-unorm-int16 cl-channel-type-unorm-short-565
+ cl-channel-type-unorm-short-555 cl-channel-type-unorm-int-101010
+ cl-channel-type-signed-int8 cl-channel-type-signed-int16 cl-channel-type-signed-int32
+ cl-channel-type-unsigned-int8 cl-channel-type-unsigned-int16
+ cl-channel-type-unsigned-int32 cl-channel-type-half-float cl-channel-type-float ;
+
+VARIANT: cl-addressing-mode
+ cl-repeat-addressing cl-clamp-to-edge-addressing cl-clamp-addressing cl-no-addressing ;
+
+VARIANT: cl-filter-mode
+ cl-filter-nearest cl-filter-linear ;
+
+VARIANT: cl-command-type
+ cl-ndrange-kernel-command cl-task-command cl-native-kernel-command cl-read-buffer-command
+ cl-write-buffer-command cl-copy-buffer-command cl-read-image-command cl-write-image-command
+ cl-copy-image-command cl-copy-buffer-to-image-command cl-copy-image-to-buffer-command
+ cl-map-buffer-command cl-map-image-command cl-unmap-mem-object-command
+ cl-marker-command cl-acquire-gl-objects-command cl-release-gl-objects-command ;
+
+VARIANT: cl-execution-status
+ cl-queued cl-submitted cl-running cl-complete cl-failure ;
+
+TUPLE: cl-platform
+ id profile version name vendor extensions devices ;
+
+TUPLE: cl-device
+ id type vendor-id max-compute-units max-work-item-dimensions
+ max-work-item-sizes max-work-group-size preferred-vector-width-char
+ preferred-vector-width-short preferred-vector-width-int
+ preferred-vector-width-long preferred-vector-width-float
+ preferred-vector-width-double max-clock-frequency address-bits
+ max-mem-alloc-size image-support max-read-image-args max-write-image-args
+ image2d-max-width image2d-max-height image3d-max-width image3d-max-height
+ image3d-max-depth max-samplers max-parameter-size mem-base-addr-align
+ min-data-type-align-size single-fp-config global-mem-cache-type
+ global-mem-cacheline-size global-mem-cache-size global-mem-size
+ max-constant-buffer-size max-constant-args local-mem? local-mem-size
+ error-correction-support profiling-timer-resolution endian-little
+ available compiler-available execute-kernels? execute-native-kernels?
+ out-of-order-exec-available? profiling-available?
+ name vendor driver-version profile version extensions ;
+
+TUPLE: cl-context < cl-handle ;
+TUPLE: cl-queue < cl-handle ;
+TUPLE: cl-buffer < cl-handle ;
+TUPLE: cl-sampler < cl-handle ;
+TUPLE: cl-program < cl-handle ;
+TUPLE: cl-kernel < cl-handle ;
+TUPLE: cl-event < cl-handle ;
+
+M: cl-context dispose* handle>> clReleaseContext cl-success ;
+M: cl-queue dispose* handle>> clReleaseCommandQueue cl-success ;
+M: cl-buffer dispose* handle>> clReleaseMemObject cl-success ;
+M: cl-sampler dispose* handle>> clReleaseSampler cl-success ;
+M: cl-program dispose* handle>> clReleaseProgram cl-success ;
+M: cl-kernel dispose* handle>> clReleaseKernel cl-success ;
+M: cl-event dispose* handle>> clReleaseEvent cl-success ;
+
+TUPLE: cl-buffer-ptr
+ { buffer cl-buffer read-only }
+ { offset integer read-only } ;
+C: <cl-buffer-ptr> cl-buffer-ptr
+
+TUPLE: cl-buffer-range
+ { buffer cl-buffer read-only }
+ { offset integer read-only }
+ { size integer read-only } ;
+C: <cl-buffer-range> cl-buffer-range
+
+SYMBOLS: cl-current-context cl-current-queue cl-current-device ;
+
+<PRIVATE
+: (current-cl-context) ( -- cl-context )
+ cl-current-context get ; inline
+
+: (current-cl-queue) ( -- cl-queue )
+ cl-current-queue get ; inline
+
+: (current-cl-device) ( -- cl-device )
+ cl-current-device get ; inline
+
+GENERIC: buffer-access-constant ( buffer-access-mode -- n )
+M: cl-read-write-access buffer-access-constant drop CL_MEM_READ_WRITE ;
+M: cl-read-access buffer-access-constant drop CL_MEM_READ_ONLY ;
+M: cl-write-access buffer-access-constant drop CL_MEM_WRITE_ONLY ;
+
+GENERIC: buffer-map-flags ( buffer-access-mode -- n )
+M: cl-read-write-access buffer-map-flags drop CL_MAP_READ CL_MAP_WRITE bitor ;
+M: cl-read-access buffer-map-flags drop CL_MAP_READ ;
+M: cl-write-access buffer-map-flags drop CL_MAP_WRITE ;
+
+GENERIC: addressing-mode-constant ( addressing-mode -- n )
+M: cl-repeat-addressing addressing-mode-constant drop CL_ADDRESS_REPEAT ;
+M: cl-clamp-to-edge-addressing addressing-mode-constant drop CL_ADDRESS_CLAMP_TO_EDGE ;
+M: cl-clamp-addressing addressing-mode-constant drop CL_ADDRESS_CLAMP ;
+M: cl-no-addressing addressing-mode-constant drop CL_ADDRESS_NONE ;
+
+GENERIC: filter-mode-constant ( filter-mode -- n )
+M: cl-filter-nearest filter-mode-constant drop CL_FILTER_NEAREST ;
+M: cl-filter-linear filter-mode-constant drop CL_FILTER_LINEAR ;
+
+: cl_addressing_mode>addressing-mode ( cl_addressing_mode -- addressing-mode )
+ {
+ { CL_ADDRESS_REPEAT [ cl-repeat-addressing ] }
+ { CL_ADDRESS_CLAMP_TO_EDGE [ cl-clamp-to-edge-addressing ] }
+ { CL_ADDRESS_CLAMP [ cl-clamp-addressing ] }
+ { CL_ADDRESS_NONE [ cl-no-addressing ] }
+ } case ; inline
+
+: cl_filter_mode>filter-mode ( cl_filter_mode -- filter-mode )
+ {
+ { CL_FILTER_LINEAR [ cl-filter-linear ] }
+ { CL_FILTER_NEAREST [ cl-filter-nearest ] }
+ } case ; inline
+
+: platform-info-string ( handle name -- string )
+ [ clGetPlatformInfo ] info-string ; inline
+
+: platform-info ( id -- profile version name vendor extensions )
+ {
+ [ CL_PLATFORM_PROFILE platform-info-string ]
+ [ CL_PLATFORM_VERSION platform-info-string ]
+ [ CL_PLATFORM_NAME platform-info-string ]
+ [ CL_PLATFORM_VENDOR platform-info-string ]
+ [ CL_PLATFORM_EXTENSIONS platform-info-string ]
+ } cleave ;
+
+: cl_device_fp_config>flags ( ulong -- sequence )
+ [ {
+ [ CL_FP_DENORM bitand 0 = [ f ] [ cl-denorm ] if ]
+ [ CL_FP_INF_NAN bitand 0 = [ f ] [ cl-inf-and-nan ] if ]
+ [ CL_FP_ROUND_TO_NEAREST bitand 0 = [ f ] [ cl-round-to-nearest ] if ]
+ [ CL_FP_ROUND_TO_ZERO bitand 0 = [ f ] [ cl-round-to-zero ] if ]
+ [ CL_FP_ROUND_TO_INF bitand 0 = [ f ] [ cl-round-to-inf ] if ]
+ [ CL_FP_FMA bitand 0 = [ f ] [ cl-fma ] if ]
+ } cleave ] { } output>sequence sift ;
+
+: cl_device_mem_cache_type>cache-type ( uint -- cache-type )
+ {
+ { CL_NONE [ cl-no-cache ] }
+ { CL_READ_ONLY_CACHE [ cl-read-only-cache ] }
+ { CL_READ_WRITE_CACHE [ cl-read-write-cache ] }
+ } case ; inline
+
+: device-info-bool ( handle name -- ? )
+ [ clGetDeviceInfo ] info-bool ; inline
+
+: device-info-ulong ( handle name -- ulong )
+ [ clGetDeviceInfo ] info-ulong ; inline
+
+: device-info-uint ( handle name -- uint )
+ [ clGetDeviceInfo ] info-uint ; inline
+
+: device-info-string ( handle name -- string )
+ [ clGetDeviceInfo ] info-string ; inline
+
+: device-info-size_t ( handle name -- size_t )
+ [ clGetDeviceInfo ] info-size_t ; inline
+
+: device-info-size_t-array ( handle name -- size_t-array )
+ [ clGetDeviceInfo ] info-size_t-array ; inline
+
+: device-info ( device-id -- device )
+ dup {
+ [ CL_DEVICE_TYPE device-info-size_t size_t>cl-device-type ]
+ [ CL_DEVICE_VENDOR_ID device-info-uint ]
+ [ CL_DEVICE_MAX_COMPUTE_UNITS device-info-uint ]
+ [ CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS device-info-uint ]
+ [ CL_DEVICE_MAX_WORK_ITEM_SIZES device-info-size_t-array ]
+ [ CL_DEVICE_MAX_WORK_GROUP_SIZE device-info-size_t ]
+ [ CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR device-info-uint ]
+ [ CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT device-info-uint ]
+ [ CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT device-info-uint ]
+ [ CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG device-info-uint ]
+ [ CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT device-info-uint ]
+ [ CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE device-info-uint ]
+ [ CL_DEVICE_MAX_CLOCK_FREQUENCY device-info-uint ]
+ [ CL_DEVICE_ADDRESS_BITS device-info-uint ]
+ [ CL_DEVICE_MAX_MEM_ALLOC_SIZE device-info-ulong ]
+ [ CL_DEVICE_IMAGE_SUPPORT device-info-bool ]
+ [ CL_DEVICE_MAX_READ_IMAGE_ARGS device-info-uint ]
+ [ CL_DEVICE_MAX_WRITE_IMAGE_ARGS device-info-uint ]
+ [ CL_DEVICE_IMAGE2D_MAX_WIDTH device-info-size_t ]
+ [ CL_DEVICE_IMAGE2D_MAX_HEIGHT device-info-size_t ]
+ [ CL_DEVICE_IMAGE3D_MAX_WIDTH device-info-size_t ]
+ [ CL_DEVICE_IMAGE3D_MAX_HEIGHT device-info-size_t ]
+ [ CL_DEVICE_IMAGE3D_MAX_DEPTH device-info-size_t ]
+ [ CL_DEVICE_MAX_SAMPLERS device-info-uint ]
+ [ CL_DEVICE_MAX_PARAMETER_SIZE device-info-size_t ]
+ [ CL_DEVICE_MEM_BASE_ADDR_ALIGN device-info-uint ]
+ [ CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE device-info-uint ]
+ [ CL_DEVICE_SINGLE_FP_CONFIG device-info-ulong cl_device_fp_config>flags ]
+ [ CL_DEVICE_GLOBAL_MEM_CACHE_TYPE device-info-uint cl_device_mem_cache_type>cache-type ]
+ [ CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE device-info-uint ]
+ [ CL_DEVICE_GLOBAL_MEM_CACHE_SIZE device-info-ulong ]
+ [ CL_DEVICE_GLOBAL_MEM_SIZE device-info-ulong ]
+ [ CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE device-info-ulong ]
+ [ CL_DEVICE_MAX_CONSTANT_ARGS device-info-uint ]
+ [ CL_DEVICE_LOCAL_MEM_TYPE device-info-uint CL_LOCAL = ]
+ [ CL_DEVICE_LOCAL_MEM_SIZE device-info-ulong ]
+ [ CL_DEVICE_ERROR_CORRECTION_SUPPORT device-info-bool ]
+ [ CL_DEVICE_PROFILING_TIMER_RESOLUTION device-info-size_t ]
+ [ CL_DEVICE_ENDIAN_LITTLE device-info-bool ]
+ [ CL_DEVICE_AVAILABLE device-info-bool ]
+ [ CL_DEVICE_COMPILER_AVAILABLE device-info-bool ]
+ [ CL_DEVICE_EXECUTION_CAPABILITIES device-info-ulong CL_EXEC_KERNEL bitand 0 = not ]
+ [ CL_DEVICE_EXECUTION_CAPABILITIES device-info-ulong CL_EXEC_NATIVE_KERNEL bitand 0 = not ]
+ [ CL_DEVICE_QUEUE_PROPERTIES device-info-ulong CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE bitand 0 = not ]
+ [ CL_DEVICE_QUEUE_PROPERTIES device-info-ulong CL_QUEUE_PROFILING_ENABLE bitand 0 = not ]
+ [ CL_DEVICE_NAME device-info-string ]
+ [ CL_DEVICE_VENDOR device-info-string ]
+ [ CL_DRIVER_VERSION device-info-string ]
+ [ CL_DEVICE_PROFILE device-info-string ]
+ [ CL_DEVICE_VERSION device-info-string ]
+ [ CL_DEVICE_EXTENSIONS device-info-string ]
+ } cleave cl-device boa ;
+
+: platform-devices ( platform-id -- devices )
+ CL_DEVICE_TYPE_ALL [
+ 0 f 0 <uint> [ clGetDeviceIDs cl-success ] keep *uint
+ ] [
+ rot dup <void*-array> [ f clGetDeviceIDs cl-success ] keep
+ ] 2bi ; inline
+
+: command-queue-info-ulong ( handle name -- ulong )
+ [ clGetCommandQueueInfo ] info-ulong ; inline
+
+: sampler-info-bool ( handle name -- ? )
+ [ clGetSamplerInfo ] info-bool ; inline
+
+: sampler-info-uint ( handle name -- uint )
+ [ clGetSamplerInfo ] info-uint ; inline
+
+: program-build-info-string ( program-handle device-handle name -- string )
+ [ clGetProgramBuildInfo ] 2info-string ; inline
+
+: program-build-log ( program-handle device-handle -- string )
+ CL_PROGRAM_BUILD_LOG program-build-info-string ; inline
+
+: strings>char*-array ( strings -- char*-array )
+ [ ascii encode dup length dup malloc [ cl-not-null ]
+ keep &free [ -rot memcpy ] keep ] void*-array{ } map-as ; inline
+
+: (program) ( cl-context sources -- program-handle )
+ [ handle>> ] dip [
+ [ length ]
+ [ strings>char*-array ]
+ [ [ length ] size_t-array{ } map-as ] tri
+ 0 <int> [ clCreateProgramWithSource ] keep *int cl-success
+ ] with-destructors ;
+
+:: (build-program) ( program-handle device options -- program )
+ program-handle 1 device 1array [ id>> ] void*-array{ } map-as
+ options ascii encode 0 suffix f f clBuildProgram
+ {
+ { CL_BUILD_PROGRAM_FAILURE [
+ program-handle device id>> program-build-log program-handle
+ clReleaseProgram cl-success cl-error f ] }
+ { CL_SUCCESS [ cl-program new-disposable program-handle >>handle ] }
+ [ program-handle clReleaseProgram cl-success cl-success f ]
+ } case ;
+
+: kernel-info-string ( handle name -- string )
+ [ clGetKernelInfo ] info-string ; inline
+
+: kernel-info-uint ( handle name -- uint )
+ [ clGetKernelInfo ] info-uint ; inline
+
+: kernel-work-group-info-size_t ( handle1 handle2 name -- size_t )
+ [ clGetKernelWorkGroupInfo ] 2info-size_t ; inline
+
+: event-info-uint ( handle name -- uint )
+ [ clGetEventInfo ] info-uint ; inline
+
+: event-info-int ( handle name -- int )
+ [ clGetEventInfo ] info-int ; inline
+
+: cl_command_type>command-type ( cl_command-type -- command-type )
+ {
+ { CL_COMMAND_NDRANGE_KERNEL [ cl-ndrange-kernel-command ] }
+ { CL_COMMAND_TASK [ cl-task-command ] }
+ { CL_COMMAND_NATIVE_KERNEL [ cl-native-kernel-command ] }
+ { CL_COMMAND_READ_BUFFER [ cl-read-buffer-command ] }
+ { CL_COMMAND_WRITE_BUFFER [ cl-write-buffer-command ] }
+ { CL_COMMAND_COPY_BUFFER [ cl-copy-buffer-command ] }
+ { CL_COMMAND_READ_IMAGE [ cl-read-image-command ] }
+ { CL_COMMAND_WRITE_IMAGE [ cl-write-image-command ] }
+ { CL_COMMAND_COPY_IMAGE [ cl-copy-image-command ] }
+ { CL_COMMAND_COPY_BUFFER_TO_IMAGE [ cl-copy-buffer-to-image-command ] }
+ { CL_COMMAND_COPY_IMAGE_TO_BUFFER [ cl-copy-image-to-buffer-command ] }
+ { CL_COMMAND_MAP_BUFFER [ cl-map-buffer-command ] }
+ { CL_COMMAND_MAP_IMAGE [ cl-map-image-command ] }
+ { CL_COMMAND_UNMAP_MEM_OBJECT [ cl-unmap-mem-object-command ] }
+ { CL_COMMAND_MARKER [ cl-marker-command ] }
+ { CL_COMMAND_ACQUIRE_GL_OBJECTS [ cl-acquire-gl-objects-command ] }
+ { CL_COMMAND_RELEASE_GL_OBJECTS [ cl-release-gl-objects-command ] }
+ } case ;
+
+: cl_int>execution-status ( clint -- execution-status )
+ {
+ { CL_QUEUED [ cl-queued ] }
+ { CL_SUBMITTED [ cl-submitted ] }
+ { CL_RUNNING [ cl-running ] }
+ { CL_COMPLETE [ cl-complete ] }
+ [ drop cl-failure ]
+ } case ; inline
+
+: profiling-info-ulong ( handle name -- ulong )
+ [ clGetEventProfilingInfo ] info-ulong ; inline
+
+
+: bind-kernel-arg-buffer ( kernel index buffer -- )
+ [ handle>> ] [ cl_mem heap-size ] [ handle>> <void*> ] tri*
+ clSetKernelArg cl-success ; inline
+
+: bind-kernel-arg-data ( kernel index byte-array -- )
+ [ handle>> ] 2dip
+ [ byte-length ] keep clSetKernelArg cl-success ; inline
+
+GENERIC: bind-kernel-arg ( kernel index data -- )
+M: cl-buffer bind-kernel-arg bind-kernel-arg-buffer ;
+M: byte-array bind-kernel-arg bind-kernel-arg-data ;
+PRIVATE>
+
+: with-cl-state ( context/f device/f queue/f quot -- )
+ [
+ [
+ [ cl-current-queue set ] when*
+ [ cl-current-device set ] when*
+ [ cl-current-context set ] when*
+ ] 3curry H{ } make-assoc
+ ] dip bind ; inline
+
+: cl-platforms ( -- platforms )
+ 0 f 0 <uint> [ clGetPlatformIDs cl-success ] keep *uint
+ dup <void*-array> [ f clGetPlatformIDs cl-success ] keep
+ [
+ dup
+ [ platform-info ]
+ [ platform-devices [ device-info ] { } map-as ] bi
+ cl-platform boa
+ ] { } map-as ;
+
+: <cl-context> ( devices -- cl-context )
+ [ f ] dip
+ [ length ] [ [ id>> ] void*-array{ } map-as ] bi
+ f f 0 <int> [ clCreateContext ] keep *int cl-success
+ cl-context new-disposable swap >>handle ;
+
+: <cl-queue> ( context device out-of-order? profiling? -- command-queue )
+ [ [ handle>> ] [ id>> ] bi* ] 2dip
+ [ [ CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE ] [ 0 ] if ]
+ [ [ CL_QUEUE_PROFILING_ENABLE ] [ 0 ] if ] bi* bitor
+ 0 <int> [ clCreateCommandQueue ] keep *int cl-success
+ cl-queue new-disposable swap >>handle ;
+
+: cl-out-of-order-execution? ( command-queue -- ? )
+ CL_QUEUE_PROPERTIES command-queue-info-ulong
+ CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE bitand 0 = not ; inline
+
+: cl-profiling? ( command-queue -- ? )
+ CL_QUEUE_PROPERTIES command-queue-info-ulong
+ CL_QUEUE_PROFILING_ENABLE bitand 0 = not ; inline
+
+: <cl-buffer> ( buffer-access-mode size initial-data -- buffer )
+ [ (current-cl-context) ] 3dip
+ swap over [
+ [ handle>> ]
+ [ buffer-access-constant ]
+ [ [ CL_MEM_COPY_HOST_PTR ] [ CL_MEM_ALLOC_HOST_PTR ] if ] tri* bitor
+ ] 2dip
+ 0 <int> [ clCreateBuffer ] keep *int cl-success
+ cl-buffer new-disposable swap >>handle ;
+
+: cl-read-buffer ( buffer-range -- byte-array )
+ [ (current-cl-queue) handle>> ] dip
+ [ buffer>> handle>> CL_TRUE ]
+ [ offset>> ]
+ [ size>> dup <byte-array> ] tri
+ [ 0 f f clEnqueueReadBuffer cl-success ] keep ; inline
+
+: cl-write-buffer ( buffer-range byte-array -- )
+ [
+ [ (current-cl-queue) handle>> ] dip
+ [ buffer>> handle>> CL_TRUE ]
+ [ offset>> ]
+ [ size>> ] tri
+ ] dip 0 f f clEnqueueWriteBuffer cl-success ; inline
+
+: cl-queue-copy-buffer ( src-buffer-ptr dst-buffer-ptr size dependent-events -- event )
+ [
+ (current-cl-queue)
+ [ handle>> ]
+ [ [ buffer>> handle>> ] [ offset>> ] bi ]
+ [ [ buffer>> handle>> ] [ offset>> ] bi ]
+ tri* swapd
+ ] 2dip [ length ] keep [ f ] [ [ handle>> ] void*-array{ } map-as ] if-empty
+ f <void*> [ clEnqueueCopyBuffer cl-success ] keep *void* cl-event
+ new-disposable swap >>handle ;
+
+: cl-queue-read-buffer ( buffer-range alien dependent-events -- event )
+ [
+ [ (current-cl-queue) handle>> ] dip
+ [ buffer>> handle>> CL_FALSE ] [ offset>> ] [ size>> ] tri
+ ] 2dip [ length ] keep [ f ] [ [ handle>> ] void*-array{ } map-as ] if-empty
+ f <void*> [ clEnqueueReadBuffer cl-success ] keep *void* cl-event
+ new-disposable swap >>handle ;
+
+: cl-queue-write-buffer ( buffer-range alien dependent-events -- event )
+ [
+ [ (current-cl-queue) handle>> ] dip
+ [ buffer>> handle>> CL_FALSE ] [ offset>> ] [ size>> ] tri
+ ] 2dip [ length ] keep [ f ] [ [ handle>> ] void*-array{ } map-as ] if-empty
+ f <void*> [ clEnqueueWriteBuffer cl-success ] keep *void* cl-event
+ new-disposable swap >>handle ;
+
+: <cl-sampler> ( normalized-coords? addressing-mode filter-mode -- sampler )
+ [ (current-cl-context) ] 3dip
+ [ [ CL_TRUE ] [ CL_FALSE ] if ]
+ [ addressing-mode-constant ]
+ [ filter-mode-constant ]
+ tri* 0 <int> [ clCreateSampler ] keep *int cl-success
+ cl-sampler new-disposable swap >>handle ;
+
+: cl-normalized-coords? ( sampler -- ? )
+ handle>> CL_SAMPLER_NORMALIZED_COORDS sampler-info-bool ; inline
+
+: cl-addressing-mode ( sampler -- addressing-mode )
+ handle>> CL_SAMPLER_ADDRESSING_MODE sampler-info-uint cl_addressing_mode>addressing-mode ; inline
+
+: cl-filter-mode ( sampler -- filter-mode )
+ handle>> CL_SAMPLER_FILTER_MODE sampler-info-uint cl_filter_mode>filter-mode ; inline
+
+: <cl-program> ( options strings -- program )
+ [ (current-cl-device) ] 2dip
+ [ (current-cl-context) ] dip
+ (program) -rot (build-program) ;
+
+: <cl-kernel> ( program kernel-name -- kernel )
+ [ handle>> ] [ ascii encode 0 suffix ] bi*
+ 0 <int> [ clCreateKernel ] keep *int cl-success
+ cl-kernel new-disposable swap >>handle ; inline
+
+: cl-kernel-name ( kernel -- string )
+ handle>> CL_KERNEL_FUNCTION_NAME kernel-info-string ; inline
+
+: cl-kernel-arity ( kernel -- arity )
+ handle>> CL_KERNEL_NUM_ARGS kernel-info-uint ; inline
+
+: cl-kernel-local-size ( kernel -- size )
+ (current-cl-device) [ handle>> ] bi@ CL_KERNEL_WORK_GROUP_SIZE kernel-work-group-info-size_t ; inline
+
+:: cl-queue-kernel ( kernel args sizes dependent-events -- event )
+ args [| arg idx | kernel idx arg bind-kernel-arg ] each-index
+ (current-cl-queue) handle>>
+ kernel handle>>
+ sizes [ length f ] [ [ ] size_t-array{ } map-as f ] bi
+ dependent-events [ length ] [ [ f ] [ [ handle>> ] void*-array{ } map-as ] if-empty ] bi
+ f <void*> [ clEnqueueNDRangeKernel cl-success ] keep *void*
+ cl-event new-disposable swap >>handle ;
+
+: cl-event-type ( event -- command-type )
+ handle>> CL_EVENT_COMMAND_TYPE event-info-uint cl_command_type>command-type ; inline
+
+: cl-event-status ( event -- execution-status )
+ handle>> CL_EVENT_COMMAND_EXECUTION_STATUS event-info-int cl_int>execution-status ; inline
+
+: cl-profile-counters ( event -- queued submitted started finished )
+ handle>> {
+ [ CL_PROFILING_COMMAND_QUEUED profiling-info-ulong ]
+ [ CL_PROFILING_COMMAND_SUBMIT profiling-info-ulong ]
+ [ CL_PROFILING_COMMAND_START profiling-info-ulong ]
+ [ CL_PROFILING_COMMAND_END profiling-info-ulong ]
+ } cleave ; inline
+
+: cl-barrier-events ( event/events -- )
+ [ (current-cl-queue) handle>> ] dip
+ dup sequence? [ 1array ] unless
+ [ handle>> ] void*-array{ } map-as [ length ] keep clEnqueueWaitForEvents cl-success ; inline
+
+: cl-marker ( -- event )
+ (current-cl-queue)
+ f <void*> [ clEnqueueMarker cl-success ] keep *void* cl-event new-disposable
+ swap >>handle ; inline
+
+: cl-barrier ( -- )
+ (current-cl-queue) clEnqueueBarrier cl-success ; inline
+
+: cl-flush ( -- )
+ (current-cl-queue) handle>> clFlush cl-success ; inline
+
+: cl-wait ( event/events -- )
+ dup sequence? [ 1array ] unless
+ [ handle>> ] void*-array{ } map-as [ length ] keep clWaitForEvents cl-success ; inline
+
+: cl-finish ( -- )
+ (current-cl-queue) handle>> clFinish cl-success ; inline
--- /dev/null
+High-level vocabulary for using OpenCL
--- /dev/null
+Erik Charlebois
--- /dev/null
+! Copyright (C) 2010 Erik Charlebois.
+! See http://factorcode.org/license.txt for BSD license.
+USING: classes.parser classes.singleton classes.union kernel lexer
+sequences ;
+IN: opencl.syntax
+
+SYNTAX: SINGLETONS-UNION:
+ CREATE-CLASS ";" parse-tokens [ create-class-in [ define-singleton-class ] keep ] map define-union-class ;
--- /dev/null
+bindings
+untested
"LIBRARY:"\r
"M:" "M::" "MACRO:" "MACRO::" "MAIN:" "MATH:"\r
"MEMO:" "MEMO:" "METHOD:" "MIXIN:"\r
+ "NAN:"\r
"OCT:"\r
"POSTPONE:" "PREDICATE:" "PRIMITIVE:" "PRIVATE>" "PROVIDE:"\r
"QUALIFIED-WITH:" "QUALIFIED:"\r
"SINGLETON:" "SINGLETONS:" "SLOT:" "SPECIALIZED-ARRAY:" "SPECIALIZED-ARRAYS:" "STRING:" "STRUCT:" "SYMBOL:" "SYMBOLS:" "SYNTAX:"\r
"TUPLE:" "t" "t?" "TYPEDEF:" "TYPED:" "TYPED::"\r
"UNIFORM-TUPLE:" "UNION:" "UNION-STRUCT:" "USE:" "USING:"\r
- "VARS:" "VERTEX-FORMAT:"))\r
+ "VARIANT:" "VERTEX-FORMAT:"))\r
\r
(defconst fuel-syntax--parsing-words-regex\r
(regexp-opt fuel-syntax--parsing-words 'words))\r
"\\_<-?[0-9]+\\_>")\r
\r
(defconst fuel-syntax--raw-float-regex\r
- "[0-9]*\\.[0-9]*\\([eE][+-]?[0-9]+\\)?")\r
+ "[0-9]*\\.[0-9]*\\([eEpP][+-]?[0-9]+\\)?")\r
\r
(defconst fuel-syntax--float-regex\r
(format "\\_<-?%s\\_>" fuel-syntax--raw-float-regex))\r
'("IN:" "USE:" "FROM:" "EXCLUDE:" "QUALIFIED:" "QUALIFIED-WITH:")))\r
\r
(defconst fuel-syntax--int-constant-def-regex\r
- (fuel-syntax--second-word-regex '("ALIEN:" "CHAR:" "BIN:" "HEX:" "OCT:")))\r
+ (fuel-syntax--second-word-regex '("ALIEN:" "CHAR:" "BIN:" "HEX:" "NAN:" "OCT:")))\r
\r
(defconst fuel-syntax--type-definition-regex\r
(fuel-syntax--second-word-regex\r
"MEMO" "MEMO:" "METHOD"\r
"SYNTAX"\r
"PREDICATE" "PRIMITIVE"\r
- "STRUCT" "TAG" "TUPLE"\r
+ "SINGLETONS"\r
+ "STRUCT" "SYMBOLS" "TAG" "TUPLE"\r
"TYPED" "TYPED:"\r
"UNIFORM-TUPLE"\r
"UNION-STRUCT" "UNION"\r
- "VERTEX-FORMAT"))\r
+ "VARIANT" "VERTEX-FORMAT"))\r
\r
(defconst fuel-syntax--no-indent-def-starts '("ARTICLE"\r
"HELP"\r
- "SINGLETONS"\r
- "SPECIALIZED-ARRAYS"\r
- "SYMBOLS"\r
- "VARS"))\r
+ "SPECIALIZED-ARRAYS"))\r
\r
(defconst fuel-syntax--indent-def-start-regex\r
(format "^\\(%s:\\)\\( \\|\n\\)" (regexp-opt fuel-syntax--indent-def-starts)))\r
"IN:" "INSTANCE:"\r
"LIBRARY:"\r
"MAIN:" "MATH:" "MIXIN:"\r
+ "NAN:"\r
"OCT:"\r
"POSTPONE:" "PRIVATE>" "<PRIVATE"\r
"QUALIFIED-WITH:" "QUALIFIED:"\r
("\\_<C-ENUM:\\( \\|\n\\)" (1 "<b"))\r
("\\_<TUPLE: +\\w+? +< +\\w+? *\\( \\|\n\\)\\([^;]\\|$\\)" (1 "<b"))\r
("\\_<TUPLE: +\\w+? *\\( \\|\n\\)\\([^;<\n]\\|\\_>\\)" (1 "<b"))\r
- ("\\_<\\(SYMBOLS\\|VARS\\|SPECIALIZED-ARRAYS\\|SINGLETONS\\): *?\\( \\|\n\\)\\([^;\n]\\|\\_>\\)"\r
+ ("\\_<\\(SYMBOLS\\|SPECIALIZED-ARRAYS\\|SINGLETONS\\|VARIANT\\): *?\\( \\|\n\\)\\([^;\n]\\|\\_>\\)"\r
(2 "<b"))\r
("\\(\n\\| \\);\\_>" (1 ">b"))\r
;; Let and lambda:\r
projects / factor.git / commitdiff