Khronos

WebCL logo

WebCL Specification

Editor's Draft, 10 July 2014

This version:
https://github.com/KhronosGroup/WebCL/blob/master/specs/latest/1.0/
WebIDL: https://github.com/KhronosGroup/WebCL/blob/master/specs/latest/1.0/webcl.idl
Latest version:
https://github.com/KhronosGroup/WebCL/blob/master/specs/latest/1.0/
WebIDL: https://github.com/KhronosGroup/WebCL/blob/master/specs/latest/1.0/webcl.idl
Previous version:
http://www.khronos.org/registry/webcl/specs/1.0.0/
Editors:
Mikaël Bourges-Sévenier (Advanced Micro Devices, Inc.)
Tomi Aarnio (Nokia Research)

Copyright © 2014 Khronos Group


Abstract

This specification defines WebCL (Web Computing Language). WebCL is a JavaScript binding to the Khronos OpenCL standard for heterogeneous parallel computing. It enables web applications to harness GPU and multi-core CPU parallel processing from within a Web browser, enabling significant acceleration of computationally intensive applications, such as image and video processing, data visualization, and 3D games.

Status of this document

This document is an Editor's Draft. Do not cite this document as anything other than work in progress. Public discussion of this document is welcome on the (archived) WebCL mailing list public_webcl@khronos.org (see instructions).

Contents

Introduction

With the popularity of web-centric platforms, there is a need to leverage high-performance parallel processing to accelerate compute-intensive web applications. Such applications include, for example, data visualization, image and video processing, 3D gaming, computational photography, computer vision, and cryptography. By providing application developers with a standardized JavaScript API and a portable kernel programming language, WebCL enables parallel computing on heterogeneous multicore hardware across a breadth of devices, including mobile, desktop, and server. The WebCL API and kernel language are designed to be substantially similar to those of OpenCL (Open Computing Language), and implementable on top of a conformant OpenCL driver.

Conventions

Basic types

The following basic types are used in this document.

typedef boolean             CLboolean;
typedef long                CLint;      // 32-bit signed integer
typedef long long           CLlong;     // 64-bit signed integer
typedef unsigned long       CLuint;     // 32-bit unsigned integer
typedef unsigned long long  CLulong;    // 52-bit unsigned integer
typedef unsigned long       CLenum;     // Used for enumerated types, such as WebCL.DEVICE_TYPE_GPU

The implementation of any WebCL function that accepts numeric values as input must ensure that no undefined behavior will occur if the input value is outside of the specified numeric range. Where necessary, the implementation may fit the incoming JavaScript number into the valid range using a clamp or modulo operator, or alternatively throw an exception. Note that the range of `CLulong` is limited to 52 bits, due to the lack of 64-bit integers in JavaScript.

Exceptions

WebCL methods throw exceptions instead of returning error codes as return values as in OpenCL. Exceptions that may be thrown by each method are listed in that method's documentation. Additionally, almost all methods may throw `OUT_OF_RESOURCES`, `OUT_OF_HOST_MEMORY`, or the WebCL-specific `WEBCL_IMPLEMENTATION_FAILURE`, so these are not listed separately. Furthermore, calling any function (except `release`) on an object that has been released, will throw INVALID_TYPE, where `TYPE` is `CONTEXT`, `COMMAND_QUEUE`, `MEM_OBJECT`, `SAMPLER`, `PROGRAM`, `KERNEL`, or `EVENT`. These exceptions are also not listed separately in the method descriptions.

exception WebCLException : DOMException {
  DOMString name;              // A string representation of the numeric error code, e.g. "INVALID_VALUE"
  DOMString? message;          // An implementation-specific description of what caused the exception
};

Both the `name` and `message` fields are present in any WebCLException thrown by the implementation, but the message may be `null`.

Resource management

Each dynamically created WebCL object has a `release` method that releases the resources consumed by that object. This does not cause the object to be deleted or garbage collected; it will remain in place, but trying to use it will cause an exception. Trying to release an object that has already been released will be silently ignored.

For convenience, the WebCL and WebCLContext classes contain an additional `releaseAll` function, which releases the context(s), as well as all other WebCL objects created from them. The usage and behavior of the release methods are illustrated in the example below.

  var ctx1 = webcl.createContext(...);
  var ctx2 = webcl.createContext(...);
  var ctx3 = webcl.createContext(...);
  var A = ctx1.createBuffer(...);
  var B = ctx2.createBuffer(...);
  ctx1.release();                    // releases ctx1, but not buffer A
  ctx2.releaseAll();                 // releases ctx2 and buffer B
  A.release();                       // releases buffer A 
  B.release();                       // does nothing: B is already released
  var C = ctx1.createBuffer(...);    // EXCEPTION: ctx1 is no longer valid
  webcl.releaseAll();                // releases ctx3

Applications are strongly recommended to explicitly release all WebCL objects as soon as they are no longer needed, instead of relying on the JavaScript garbage collector. This is necessary because garbage collectors typically do not give any guarantees on when (or indeed if) they will reclaim the objects that are no longer in scope.

When the global `document` object goes out of scope, the WebCL implementation must implicitly call `release` on all remaining WebCL objects.

Concurrency

Depending on the implementation, WebCL operations may be running concurrently with JavaScript. In particular, it may be possible for the application to modify an ArrayBuffer while it's being asynchronously copied to/from a WebCLMemoryObject. To avoid corrupting the contents of either buffer, applications should not modify an ArrayBuffer that has been enqueued for async read/write, until the relevant WebCL command queue has finished.

Callbacks

WebCL allows certain long-running functions to be executed either synchronously or asynchronously. The asynchronous mode is used if a user-defined callback function is given as an argument to such functions. The asynchronous mode is strongly recommended, so as to avoid blocking the JavaScript main thread. If a callback function is associated with a WebCL object that is subsequently released, the callback will no longer be invoked. The signature of the callback function is as follows:

callback WebCLCallback = void (optional WebCLEvent event);

To avoid blocking the main thread, the callback function should complete quickly and not call any potentially long-running functions. Certain WebCL functions are explicitly prohibited, and will cause an `INVALID_OPERATION` exception if invoked from a callback. The disallowed functions include `createContext`, `createCommandQueue`, and the blocking forms of `finish`, `waitForEvents`, `build`, and enqueue{Read, Write}{Buffer, BufferRect, Image}.

Interfaces

WebCL

The WebCL API is accessible through a singleton `webcl` object that resides in the global `window` namespace. The `webcl` object implements the WebCL interface, providing functions for creating computing contexts and querying the available WebCL platforms and extensions.

partial interface Window {
  readonly attribute WebCL webcl;
};
interface WebCL {

  sequence<WebCLPlatform> getPlatforms();

  WebCLContext createContext(optional CLenum deviceType = WebCL.DEVICE_TYPE_DEFAULT);

  WebCLContext createContext(WebCLPlatform platform, optional CLenum deviceType = WebCL.DEVICE_TYPE_DEFAULT);

  WebCLContext createContext(WebCLDevice device);

  WebCLContext createContext(sequence<WebCLDevice> devices);

  sequence<DOMString>? getSupportedExtensions();

  CLboolean enableExtension(DOMString extensionName);

  void waitForEvents(sequence<WebCLEvent> eventWaitList,
                     optional WebCLCallback whenFinished);

  void releaseAll();


  /* Error Codes */
  const CLint SUCCESS                                   = 0;
  const CLint DEVICE_NOT_FOUND                          = -1;
  const CLint DEVICE_NOT_AVAILABLE                      = -2;
  const CLint COMPILER_NOT_AVAILABLE                    = -3;
  const CLint MEM_OBJECT_ALLOCATION_FAILURE             = -4;
  const CLint OUT_OF_RESOURCES                          = -5;
  const CLint OUT_OF_HOST_MEMORY                        = -6;
  const CLint PROFILING_INFO_NOT_AVAILABLE              = -7;
  const CLint MEM_COPY_OVERLAP                          = -8;
  const CLint IMAGE_FORMAT_MISMATCH                     = -9;
  const CLint IMAGE_FORMAT_NOT_SUPPORTED                = -10;
  const CLint BUILD_PROGRAM_FAILURE                     = -11;
  const CLint MAP_FAILURE                               = -12;
  const CLint MISALIGNED_SUB_BUFFER_OFFSET              = -13;
  const CLint EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST = -14;
  const CLint INVALID_VALUE                             = -30;
  const CLint INVALID_DEVICE_TYPE                       = -31;
  const CLint INVALID_PLATFORM                          = -32;
  const CLint INVALID_DEVICE                            = -33;
  const CLint INVALID_CONTEXT                           = -34;
  const CLint INVALID_QUEUE_PROPERTIES                  = -35;
  const CLint INVALID_COMMAND_QUEUE                     = -36;
  const CLint INVALID_HOST_PTR                          = -37;
  const CLint INVALID_MEM_OBJECT                        = -38;
  const CLint INVALID_IMAGE_FORMAT_DESCRIPTOR           = -39;
  const CLint INVALID_IMAGE_SIZE                        = -40;
  const CLint INVALID_SAMPLER                           = -41;
  const CLint INVALID_BINARY                            = -42;
  const CLint INVALID_BUILD_OPTIONS                     = -43;
  const CLint INVALID_PROGRAM                           = -44;
  const CLint INVALID_PROGRAM_EXECUTABLE                = -45;
  const CLint INVALID_KERNEL_NAME                       = -46;
  const CLint INVALID_KERNEL_DEFINITION                 = -47;
  const CLint INVALID_KERNEL                            = -48;
  const CLint INVALID_ARG_INDEX                         = -49;
  const CLint INVALID_ARG_VALUE                         = -50;
  const CLint INVALID_ARG_SIZE                          = -51;
  const CLint INVALID_KERNEL_ARGS                       = -52;
  const CLint INVALID_WORK_DIMENSION                    = -53;
  const CLint INVALID_WORK_GROUP_SIZE                   = -54;
  const CLint INVALID_WORK_ITEM_SIZE                    = -55;
  const CLint INVALID_GLOBAL_OFFSET                     = -56;
  const CLint INVALID_EVENT_WAIT_LIST                   = -57;
  const CLint INVALID_EVENT                             = -58;
  const CLint INVALID_OPERATION                         = -59;
  //const CLint INVALID_GL_OBJECT                         = -60; // moved to extension
  const CLint INVALID_BUFFER_SIZE                       = -61;
  //const CLint INVALID_MIP_LEVEL                         = -62; // moved to extension
  const CLint INVALID_GLOBAL_WORK_SIZE                  = -63;
  const CLint INVALID_PROPERTY                          = -64;

  /* cl_bool */
  const CLenum FALSE                                     = 0;
  const CLenum TRUE                                      = 1;

  /* cl_platform_info */
  const CLenum PLATFORM_PROFILE                          = 0x0900;
  const CLenum PLATFORM_VERSION                          = 0x0901;
  const CLenum PLATFORM_NAME                             = 0x0902;
  const CLenum PLATFORM_VENDOR                           = 0x0903;
  const CLenum PLATFORM_EXTENSIONS                       = 0x0904;

  /* cl_device_type - bitfield */
  const CLenum DEVICE_TYPE_DEFAULT                       = 0x1;
  const CLenum DEVICE_TYPE_CPU                           = 0x2;
  const CLenum DEVICE_TYPE_GPU                           = 0x4;
  const CLenum DEVICE_TYPE_ACCELERATOR                   = 0x8;
  const CLenum DEVICE_TYPE_ALL                           = 0xFFFFFFFF;

  /* cl_device_info */
  const CLenum DEVICE_TYPE                               = 0x1000;
  const CLenum DEVICE_VENDOR_ID                          = 0x1001;
  const CLenum DEVICE_MAX_COMPUTE_UNITS                  = 0x1002;
  const CLenum DEVICE_MAX_WORK_ITEM_DIMENSIONS           = 0x1003;
  const CLenum DEVICE_MAX_WORK_GROUP_SIZE                = 0x1004;
  const CLenum DEVICE_MAX_WORK_ITEM_SIZES                = 0x1005;
  const CLenum DEVICE_PREFERRED_VECTOR_WIDTH_CHAR        = 0x1006;
  const CLenum DEVICE_PREFERRED_VECTOR_WIDTH_SHORT       = 0x1007;
  const CLenum DEVICE_PREFERRED_VECTOR_WIDTH_INT         = 0x1008;
  const CLenum DEVICE_PREFERRED_VECTOR_WIDTH_LONG        = 0x1009;
  const CLenum DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT       = 0x100A;
  //const CLenum DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE      = 0x100B; // moved to extension
  const CLenum DEVICE_MAX_CLOCK_FREQUENCY                = 0x100C;
  const CLenum DEVICE_ADDRESS_BITS                       = 0x100D;
  const CLenum DEVICE_MAX_READ_IMAGE_ARGS                = 0x100E;
  const CLenum DEVICE_MAX_WRITE_IMAGE_ARGS               = 0x100F;
  const CLenum DEVICE_MAX_MEM_ALLOC_SIZE                 = 0x1010;
  const CLenum DEVICE_IMAGE2D_MAX_WIDTH                  = 0x1011;
  const CLenum DEVICE_IMAGE2D_MAX_HEIGHT                 = 0x1012;
  const CLenum DEVICE_IMAGE3D_MAX_WIDTH                  = 0x1013;
  const CLenum DEVICE_IMAGE3D_MAX_HEIGHT                 = 0x1014;
  const CLenum DEVICE_IMAGE3D_MAX_DEPTH                  = 0x1015;
  const CLenum DEVICE_IMAGE_SUPPORT                      = 0x1016;
  const CLenum DEVICE_MAX_PARAMETER_SIZE                 = 0x1017;
  const CLenum DEVICE_MAX_SAMPLERS                       = 0x1018;
  const CLenum DEVICE_MEM_BASE_ADDR_ALIGN                = 0x1019;
  //const CLenum DEVICE_MIN_DATA_TYPE_ALIGN_SIZE           = 0x101A; // removed; deprecated in OpenCL 1.2
  const CLenum DEVICE_SINGLE_FP_CONFIG                   = 0x101B;
  const CLenum DEVICE_GLOBAL_MEM_CACHE_TYPE              = 0x101C;
  const CLenum DEVICE_GLOBAL_MEM_CACHELINE_SIZE          = 0x101D;
  const CLenum DEVICE_GLOBAL_MEM_CACHE_SIZE              = 0x101E;
  const CLenum DEVICE_GLOBAL_MEM_SIZE                    = 0x101F;
  const CLenum DEVICE_MAX_CONSTANT_BUFFER_SIZE           = 0x1020;
  const CLenum DEVICE_MAX_CONSTANT_ARGS                  = 0x1021;
  const CLenum DEVICE_LOCAL_MEM_TYPE                     = 0x1022;
  const CLenum DEVICE_LOCAL_MEM_SIZE                     = 0x1023;
  const CLenum DEVICE_ERROR_CORRECTION_SUPPORT           = 0x1024;
  const CLenum DEVICE_PROFILING_TIMER_RESOLUTION         = 0x1025;
  const CLenum DEVICE_ENDIAN_LITTLE                      = 0x1026;
  const CLenum DEVICE_AVAILABLE                          = 0x1027;
  const CLenum DEVICE_COMPILER_AVAILABLE                 = 0x1028;
  const CLenum DEVICE_EXECUTION_CAPABILITIES             = 0x1029;
  const CLenum DEVICE_QUEUE_PROPERTIES                   = 0x102A;
  const CLenum DEVICE_NAME                               = 0x102B;
  const CLenum DEVICE_VENDOR                             = 0x102C;
  const CLenum DRIVER_VERSION                            = 0x102D;
  const CLenum DEVICE_PROFILE                            = 0x102E;
  const CLenum DEVICE_VERSION                            = 0x102F;
  const CLenum DEVICE_EXTENSIONS                         = 0x1030;
  const CLenum DEVICE_PLATFORM                           = 0x1031;
  //const CLenum DEVICE_DOUBLE_FP_CONFIG                   = 0x1032; // moved to extension
  //const CLenum DEVICE_HALF_FP_CONFIG                     = 0x1033; // moved to extension
  //const CLenum DEVICE_PREFERRED_VECTOR_WIDTH_HALF        = 0x1034; // moved to extension
  const CLenum DEVICE_HOST_UNIFIED_MEMORY                = 0x1035;
  const CLenum DEVICE_NATIVE_VECTOR_WIDTH_CHAR           = 0x1036;
  const CLenum DEVICE_NATIVE_VECTOR_WIDTH_SHORT          = 0x1037;
  const CLenum DEVICE_NATIVE_VECTOR_WIDTH_INT            = 0x1038;
  const CLenum DEVICE_NATIVE_VECTOR_WIDTH_LONG           = 0x1039;
  const CLenum DEVICE_NATIVE_VECTOR_WIDTH_FLOAT          = 0x103A;
  //const CLenum DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE         = 0x103B; // moved to extension
  //const CLenum DEVICE_NATIVE_VECTOR_WIDTH_HALF           = 0x103C; // moved to extension
  const CLenum DEVICE_OPENCL_C_VERSION                   = 0x103D;

  /* cl_device_fp_config - bitfield */
  const CLenum FP_DENORM                                 = 0x1;
  const CLenum FP_INF_NAN                                = 0x2;
  const CLenum FP_ROUND_TO_NEAREST                       = 0x4;
  const CLenum FP_ROUND_TO_ZERO                          = 0x8;
  const CLenum FP_ROUND_TO_INF                           = 0x10;
  const CLenum FP_FMA                                    = 0x20;
  const CLenum FP_SOFT_FLOAT                             = 0x40;

  /* cl_device_mem_cache_type */
  const CLenum NONE                                      = 0x0;
  const CLenum READ_ONLY_CACHE                           = 0x1;
  const CLenum READ_WRITE_CACHE                          = 0x2;

  /* cl_device_local_mem_type */
  const CLenum LOCAL                                     = 0x1;
  const CLenum GLOBAL                                    = 0x2;

  /* cl_device_exec_capabilities - bitfield */
  const CLenum EXEC_KERNEL                               = 0x1;
  //const CLenum EXEC_NATIVE_KERNEL                        = 0x2; // disallowed

  /* cl_command_queue_properties - bitfield */
  const CLenum QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE       = 0x1;
  const CLenum QUEUE_PROFILING_ENABLE                    = 0x2;

  /* cl_context_info  */
  //const CLenum CONTEXT_REFERENCE_COUNT                   = 0x1080; // disallowed
  const CLenum CONTEXT_DEVICES                           = 0x1081;
  //const CLenum CONTEXT_PROPERTIES                        = 0x1082; // disallowed; no context properties in WebCL
  const CLenum CONTEXT_NUM_DEVICES                       = 0x1083;

  /* cl_context_properties */
  //const CLenum CONTEXT_PLATFORM                          = 0x1084; // disallowed; no context properties in WebCL

  /* cl_command_queue_info */
  const CLenum QUEUE_CONTEXT                             = 0x1090;
  const CLenum QUEUE_DEVICE                              = 0x1091;
  //const CLenum QUEUE_REFERENCE_COUNT                     = 0x1092; // disallowed
  const CLenum QUEUE_PROPERTIES                          = 0x1093;

  /* cl_mem_flags - bitfield */
  const CLenum MEM_READ_WRITE                            = 0x1;
  const CLenum MEM_WRITE_ONLY                            = 0x2;
  const CLenum MEM_READ_ONLY                             = 0x4;

  /* cl_channel_order */
  const CLenum R                                         = 0x10B0;
  const CLenum A                                         = 0x10B1;
  const CLenum RG                                        = 0x10B2;
  const CLenum RA                                        = 0x10B3;
  const CLenum RGB                                       = 0x10B4;
  const CLenum RGBA                                      = 0x10B5;
  const CLenum BGRA                                      = 0x10B6;
  const CLenum ARGB                                      = 0x10B7;
  const CLenum INTENSITY                                 = 0x10B8;
  const CLenum LUMINANCE                                 = 0x10B9;
  const CLenum Rx                                        = 0x10BA;
  const CLenum RGx                                       = 0x10BB;
  const CLenum RGBx                                      = 0x10BC;

  /* cl_channel_type */
  const CLenum SNORM_INT8                                = 0x10D0;
  const CLenum SNORM_INT16                               = 0x10D1;
  const CLenum UNORM_INT8                                = 0x10D2;
  const CLenum UNORM_INT16                               = 0x10D3;
  const CLenum UNORM_SHORT_565                           = 0x10D4;
  const CLenum UNORM_SHORT_555                           = 0x10D5;
  const CLenum UNORM_INT_101010                          = 0x10D6;
  const CLenum SIGNED_INT8                               = 0x10D7;
  const CLenum SIGNED_INT16                              = 0x10D8;
  const CLenum SIGNED_INT32                              = 0x10D9;
  const CLenum UNSIGNED_INT8                             = 0x10DA;
  const CLenum UNSIGNED_INT16                            = 0x10DB;
  const CLenum UNSIGNED_INT32                            = 0x10DC;
  const CLenum HALF_FLOAT                                = 0x10DD;
  const CLenum FLOAT                                     = 0x10DE;

  /* cl_mem_object_type */
  const CLenum MEM_OBJECT_BUFFER                         = 0x10F0;
  const CLenum MEM_OBJECT_IMAGE2D                        = 0x10F1;
  //const CLenum MEM_OBJECT_IMAGE3D                        = 0x10F2; // moved to extension (TBD)

  /* cl_mem_info */
  const CLenum MEM_TYPE                                  = 0x1100;
  const CLenum MEM_FLAGS                                 = 0x1101;
  const CLenum MEM_SIZE                                  = 0x1102;
  //const CLenum MEM_HOST_PTR                              = 0x1103; // disallowed
  //const CLenum MEM_MAP_COUNT                             = 0x1104; // disallowed
  //const CLenum MEM_REFERENCE_COUNT                       = 0x1105; // disallowed
  const CLenum MEM_CONTEXT                               = 0x1106;
  const CLenum MEM_ASSOCIATED_MEMOBJECT                  = 0x1107;
  const CLenum MEM_OFFSET                                = 0x1108;

  /* cl_image_info */
  const CLenum IMAGE_FORMAT                              = 0x1110;
  const CLenum IMAGE_ELEMENT_SIZE                        = 0x1111;
  const CLenum IMAGE_ROW_PITCH                           = 0x1112;
  const CLenum IMAGE_WIDTH                               = 0x1114;
  const CLenum IMAGE_HEIGHT                              = 0x1115;

  /* cl_addressing_mode */
  //const CLenum ADDRESS_NONE                              = 0x1130; // disallowed
  const CLenum ADDRESS_CLAMP_TO_EDGE                     = 0x1131;
  const CLenum ADDRESS_CLAMP                             = 0x1132;
  const CLenum ADDRESS_REPEAT                            = 0x1133;
  const CLenum ADDRESS_MIRRORED_REPEAT                   = 0x1134;

  /* cl_filter_mode */
  const CLenum FILTER_NEAREST                            = 0x1140;
  const CLenum FILTER_LINEAR                             = 0x1141;

  /* cl_sampler_info */
  //const CLenum SAMPLER_REFERENCE_COUNT                   = 0x1150; // disallowed
  const CLenum SAMPLER_CONTEXT                           = 0x1151;
  const CLenum SAMPLER_NORMALIZED_COORDS                 = 0x1152;
  const CLenum SAMPLER_ADDRESSING_MODE                   = 0x1153;
  const CLenum SAMPLER_FILTER_MODE                       = 0x1154;

  /* cl_map_flags - bitfield */
  //const CLenum MAP_READ                                  = 0x1; // disallowed
  //const CLenum MAP_WRITE                                 = 0x2; // disallowed

  /* cl_program_info */
  //const CLenum PROGRAM_REFERENCE_COUNT                   = 0x1160; // disallowed
  const CLenum PROGRAM_CONTEXT                           = 0x1161;
  const CLenum PROGRAM_NUM_DEVICES                       = 0x1162;
  const CLenum PROGRAM_DEVICES                           = 0x1163;
  const CLenum PROGRAM_SOURCE                            = 0x1164;
  //const CLenum PROGRAM_BINARY_SIZES                      = 0x1165; // disallowed
  //const CLenum PROGRAM_BINARIES                          = 0x1166; // disallowed

  /* cl_program_build_info */
  const CLenum PROGRAM_BUILD_STATUS                      = 0x1181;
  const CLenum PROGRAM_BUILD_OPTIONS                     = 0x1182;
  const CLenum PROGRAM_BUILD_LOG                         = 0x1183;

  /* cl_build_status */
  const CLint BUILD_SUCCESS                             = 0;
  const CLint BUILD_NONE                                = -1;
  const CLint BUILD_ERROR                               = -2;
  const CLint BUILD_IN_PROGRESS                         = -3;

  /* cl_kernel_info */
  const CLenum KERNEL_FUNCTION_NAME                      = 0x1190;
  const CLenum KERNEL_NUM_ARGS                           = 0x1191;
  //const CLenum KERNEL_REFERENCE_COUNT                    = 0x1192; // disallowed
  const CLenum KERNEL_CONTEXT                            = 0x1193;
  const CLenum KERNEL_PROGRAM                            = 0x1194;

  /* cl_kernel_work_group_info */
  const CLenum KERNEL_WORK_GROUP_SIZE                    = 0x11B0;
  const CLenum KERNEL_COMPILE_WORK_GROUP_SIZE            = 0x11B1;
  const CLenum KERNEL_LOCAL_MEM_SIZE                     = 0x11B2;
  const CLenum KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE = 0x11B3;
  const CLenum KERNEL_PRIVATE_MEM_SIZE                   = 0x11B4;

  /* cl_event_info  */
  const CLenum EVENT_COMMAND_QUEUE                       = 0x11D0;
  const CLenum EVENT_COMMAND_TYPE                        = 0x11D1;
  //const CLenum EVENT_REFERENCE_COUNT                     = 0x11D2; // disallowed
  const CLenum EVENT_COMMAND_EXECUTION_STATUS            = 0x11D3;
  const CLenum EVENT_CONTEXT                             = 0x11D4;

  /* cl_command_type */
  const CLenum COMMAND_NDRANGE_KERNEL                    = 0x11F0;
  const CLenum COMMAND_TASK                              = 0x11F1;
  //const CLenum COMMAND_NATIVE_KERNEL                     = 0x11F2; // disallowed
  const CLenum COMMAND_READ_BUFFER                       = 0x11F3;
  const CLenum COMMAND_WRITE_BUFFER                      = 0x11F4;
  const CLenum COMMAND_COPY_BUFFER                       = 0x11F5;
  const CLenum COMMAND_READ_IMAGE                        = 0x11F6;
  const CLenum COMMAND_WRITE_IMAGE                       = 0x11F7;
  const CLenum COMMAND_COPY_IMAGE                        = 0x11F8;
  const CLenum COMMAND_COPY_IMAGE_TO_BUFFER              = 0x11F9;
  const CLenum COMMAND_COPY_BUFFER_TO_IMAGE              = 0x11FA;
  //const CLenum COMMAND_MAP_BUFFER                        = 0x11FB; // disallowed
  //const CLenum COMMAND_MAP_IMAGE                         = 0x11FC; // disallowed
  //const CLenum COMMAND_UNMAP_MEM_OBJECT                  = 0x11FD; // disallowed
  const CLenum COMMAND_MARKER                            = 0x11FE;
  //const CLenum COMMAND_ACQUIRE_GL_OBJECTS                = 0x11FF; // moved to extension
  //const CLenum COMMAND_RELEASE_GL_OBJECTS                = 0x1200; // moved to extension
  const CLenum COMMAND_READ_BUFFER_RECT                  = 0x1201;
  const CLenum COMMAND_WRITE_BUFFER_RECT                 = 0x1202;
  const CLenum COMMAND_COPY_BUFFER_RECT                  = 0x1203;
  const CLenum COMMAND_USER                              = 0x1204;

  /* command execution status */
  const CLenum COMPLETE                                  = 0x0;
  const CLenum RUNNING                                   = 0x1;
  const CLenum SUBMITTED                                 = 0x2;
  const CLenum QUEUED                                    = 0x3;

  /* cl_profiling_info  */
  const CLenum PROFILING_COMMAND_QUEUED                  = 0x1280;
  const CLenum PROFILING_COMMAND_SUBMIT                  = 0x1281;
  const CLenum PROFILING_COMMAND_START                   = 0x1282;
  const CLenum PROFILING_COMMAND_END                     = 0x1283;
};
sequence<WebCLPlatform> getPlatforms() (OpenCL 1.1 §4.1, clGetPlatformIDs)
Retrieves all WebCL platforms that are available in this system.
WebCLContext createContext(optional CLenum deviceType = WebCL.DEVICE_TYPE_DEFAULT) (OpenCL 1.1 §4.3, clCreateContextFromType)
Returns a newly created WebCL context for the given type of device, if available. If there are multiple devices of the given type, the implementation may select any one of them, a subset of them, or all of them. The available values for `deviceType` are listed in the table below.
deviceTypeDescription
DEVICE_TYPE_CPUA single-core or multi-core CPU, typically the host processor.
DEVICE_TYPE_GPUA graphics processing unit, typically also used by WebGL.
DEVICE_TYPE_ACCELERATORA dedicated OpenCL accelerator.
DEVICE_TYPE_DEFAULTThe default device(s) on the system.
Exceptions:
WebCLContext createContext(WebCLPlatform platform, optional CLenum deviceType = WebCL.DEVICE_TYPE_DEFAULT) (OpenCL 1.1 §4.3, clCreateContextFromType)
Returns a newly created WebCL context for the given type of device, if available on the given platform. If there are multiple devices of the given type, the implementation may select any one of them, a subset of them, or all of them. The available values for `deviceType` are listed in the table below.
deviceTypeDescription
DEVICE_TYPE_CPUA single-core or multi-core CPU, typically the host processor.
DEVICE_TYPE_GPUA graphics processing unit, typically also used by WebGL.
DEVICE_TYPE_ACCELERATORA dedicated OpenCL accelerator.
DEVICE_TYPE_DEFAULTThe default device(s) on the given platform.
DEVICE_TYPE_ALLAll devices available on the given platform.
Exceptions:
WebCLContext createContext(WebCLDevice device) (OpenCL 1.1 §4.3, clCreateContext)
Returns a newly created WebCL context for the given device.
Exceptions:
WebCLContext createContext(sequence<WebCLDevice> devices) (OpenCL 1.1 §4.3, clCreateContext)
Returns a newly created WebCL context spanning the given list of devices.
Exceptions:
sequence<DOMString>? getSupportedExtensions() (OpenCL 1.1 §9)
Returns an array of extension names that are supported by all WebCLPlatforms and WebCLDevices in this implementation. Any string in this list, when passed to `enableExtension`, must enable the corresponding extension.
CLboolean enableExtension(DOMString extensionName)
Enables the given extension on all WebCLPlatforms and WebCLDevices. Returns `true` if the extension is successfully enabled, or `false` if not. The available extension names can be queried by `getSupportedExtensions`. Note that enabling an extension does not take effect retroactively, i.e., contexts that were created before enabling the extension will continue to not have the extended capabilities.
void waitForEvents(sequence<WebCLEvent> eventWaitList, optional WebCLCallback whenFinished) (OpenCL 1.1 §5.9, clWaitForEvents)
Invokes the given callback when all events in `eventWaitList` have completed. If a callback is not provided, the JavaScript main thread will be blocked until the events have completed. Applications are strongly advised to provide a callback; the blocking mode mainly exists to ease the porting of existing OpenCL code.
Exceptions:
void releaseAll() (OpenCL 1.1 §4.3, clReleaseContext)
Releases the resources held up by all WebCLContext instances, as well as their descendant objects, if any.
// Valid examples of createContext. All of these will create a valid
// context, assuming that a suitable device is available.

ctx = webcl.createContext();
ctx = webcl.createContext(WebCL.DEVICE_TYPE_CPU);
ctx = webcl.createContext(aPlatform);
ctx = webcl.createContext(aPlatform, WebCL.DEVICE_TYPE_GPU);
ctx = webcl.createContext(aPlatform, WebCL.DEVICE_TYPE_ALL);
ctx = webcl.createContext(aDevice);
ctx = webcl.createContext([aDevice, anotherDeviceOnTheSamePlatform]);

// Invalid examples of createContext. These will throw an exception.

ctx = webcl.createContext(0x1234);
ctx = webcl.createContext(WebCL.DEVICE_TYPE_ALL);
ctx = webcl.createContext({});
ctx = webcl.createContext([]);
ctx = webcl.createContext(aPlatform, 0x1234);
ctx = webcl.createContext([aValidDevice, null]);
ctx = webcl.createContext([aDeviceOnPlatformA, aDeviceOnPlatformB]);

WebCLPlatform

interface WebCLPlatform {
  any getInfo(CLenum name);
  sequence<WebCLDevice> getDevices(optional CLenum deviceType = WebCL.DEVICE_TYPE_ALL);
  sequence<DOMString>? getSupportedExtensions();
  CLboolean enableExtension(DOMString extensionName);
};
sequence<WebCLDevice> getDevices(optional CLenum deviceType = WebCL.DEVICE_TYPE_ALL) (OpenCL 1.1 §4.2, clGetDeviceIDs)

Retrieves the WebCLDevices that are available on this WebCLPlatform and match the given device type. The valid device types are listed in the table below. Omitting the device type is equivalent to specifying `DEVICE_TYPE_ALL`.

deviceTypedescription
DEVICE_TYPE_CPUA single-core or multi-core CPU, typically the host processor.
DEVICE_TYPE_GPUA graphics processing unit, typically also used by WebGL.
DEVICE_TYPE_ACCELERATORA dedicated OpenCL accelerator.
DEVICE_TYPE_DEFAULTThe default device on this platform.
DEVICE_TYPE_ALLAll devices available on this platform.
Exceptions:
any getInfo(CLenum name) (OpenCL 1.1 §4.1, clGetPlatformInfo)

Retrieves information about this WebCLPlatform. The available query parameters are listed in the table below. Note that depending on the privacy policy of the implementation, some of the return values may be empty or otherwise not reflective of the underlying OpenCL platform. Implementations that choose to mask the underlying values are recommended to make them available for privileged applications (such as browser extensions and development tools) via the System Info extension.

namereturn typeexpected return value
PLATFORM_PROFILEDOMString"WEBCL_PROFILE" + optional vendor string
PLATFORM_VERSIONDOMString"WebCL 1.0" + optional vendor string
PLATFORM_NAMEDOMStringVendor specific
PLATFORM_VENDORDOMStringVendor specific
PLATFORM_EXTENSIONSDOMStringVendor specific
Exceptions:
sequence<DOMString>? getSupportedExtensions() (OpenCL 1.1 §9, clGetPlatformInfo)
Returns an array of extension names that are supported by all WebCLDevices on this WebCLPlatform. Any string in this list, when passed to `enableExtension` on this platform, or any device on this platform, must enable the corresponding extension.
CLboolean enableExtension(DOMString extensionName)
Enables the given WebCL extension on this WebCLPlatform. Returns `true` if the extension is successfully enabled, or `false` if not. The available extension names can be queried by `getSupportedExtensions`. Note that enabling an extension does not take effect retroactively, i.e., contexts that were created before enabling the extension will continue to not have the extended capabilities.

WebCLDevice

interface WebCLDevice {
  any getInfo(CLenum name);
  sequence<DOMString>? getSupportedExtensions();
  CLboolean enableExtension(DOMString extensionName);
};
any getInfo(CLenum name) (OpenCL 1.1 §4.2, clGetDeviceInfo)

Retrieves information about this WebCLDevice. The available query parameters are listed in the table below. Note that depending on the privacy policy of the implementation, some of the return values may be empty or otherwise not reflective of the underlying OpenCL device. Implementations that choose to mask the underlying values are recommended to make them available for privileged applications (such as browser extensions and development tools) via the System Info extension.

namereturn typeexpected return value
DEVICE_PROFILEDOMString"WEBCL_PROFILE" + optional vendor string
DEVICE_VERSIONDOMString"WebCL 1.0" + optional vendor string
DEVICE_OPENCL_C_VERSIONDOMString"WebCL C 1.0" + optional vendor string
DEVICE_NAMEDOMStringVendor specific
DEVICE_VENDORDOMStringVendor specific
DEVICE_VENDOR_IDCLuintVendor specific
DRIVER_VERSIONDOMStringVendor specific
DEVICE_EXTENSIONSDOMStringVendor specific
DEVICE_AVAILABLECLbooleantrue
DEVICE_COMPILER_AVAILABLECLbooleantrue
DEVICE_TYPECLenum One or more of { DEVICE_TYPE_CPU, DEVIDE_TYPE_GPU, DEVICE_TYPE_ACCELERATOR }
DEVICE_MAX_COMPUTE_UNITSCLuint>= 1
DEVICE_MAX_WORK_ITEM_DIMENSIONSCLuint>= 3
DEVICE_MAX_WORK_ITEM_SIZESCLuint>= (1, 1, 1)
DEVICE_MAX_WORK_GROUP_SIZECLuint>= 1
DEVICE_PREFERRED_VECTOR_WIDTH_{CHAR, SHORT, INT, LONG, FLOAT}CLuint>= 1
DEVICE_NATIVE_VECTOR_WIDTH_{CHAR, SHORT, INT, LONG, FLOAT}CLuint>= 1
DEVICE_MAX_CLOCK_FREQUENCYCLuint>= 0
DEVICE_ADDRESS_BITSCLuint32 or 64
DEVICE_MAX_MEM_ALLOC_SIZECLulong>= max(DEVICE_GLOBAL_MEM_SIZE/4, 1024*1024)
DEVICE_IMAGE_SUPPORTCLbooleantrue
DEVICE_MAX_READ_IMAGE_ARGSCLuint>= 8
DEVICE_MAX_WRITE_IMAGE_ARGSCLuint>= 1
DEVICE_IMAGE2D_MAX_WIDTHCLuint>= 2048
DEVICE_IMAGE2D_MAX_HEIGHTCLuint>= 2048
DEVICE_IMAGE3D_MAX_WIDTHCLuint>= 0
DEVICE_IMAGE3D_MAX_HEIGHTCLuint>= 0
DEVICE_IMAGE3D_MAX_DEPTHCLuint>= 0
DEVICE_MAX_SAMPLERSCLuint>= 8
DEVICE_MAX_PARAMETER_SIZECLuint>= 256 bytes
DEVICE_MEM_BASE_ADDR_ALIGNCLuint>= numBits(float16) = 16*32 = 512
DEVICE_SINGLE_FP_CONFIGCLenum A combination of one or more of { FP_ROUND_TO_NEAREST, FP_ROUND_TO_ZERO } and zero or more of { FP_DENORM, FP_INF_NAN, FP_ROUND_TO_INF, FP_FMA, FP_SOFT_FLOAT }.
DEVICE_GLOBAL_MEM_CACHE_TYPECLenumExactly one of { NONE, READ_ONLY_CACHE, READ_WRITE_CACHE }
DEVICE_GLOBAL_MEM_CACHELINE_SIZECLuint>= 0 bytes
DEVICE_GLOBAL_MEM_CACHE_SIZECLulong>= 0 bytes
DEVICE_GLOBAL_MEM_SIZECLulong>= 1024*1024 bytes
DEVICE_MAX_CONSTANT_BUFFER_SIZECLulong>= 1024 bytes
DEVICE_MAX_CONSTANT_ARGSCLuint>= 4
DEVICE_LOCAL_MEM_TYPECLenumExactly one of { LOCAL, GLOBAL }
DEVICE_LOCAL_MEM_SIZECLulong>= 1024 bytes
DEVICE_ERROR_CORRECTION_SUPPORTCLbooleantrue/false
DEVICE_HOST_UNIFIED_MEMORYCLbooleantrue/false
DEVICE_PROFILING_TIMER_RESOLUTIONCLuint>= 0
DEVICE_ENDIAN_LITTLECLbooleantrue/false
DEVICE_EXECUTION_CAPABILITIESCLenumEXEC_KERNEL
DEVICE_QUEUE_PROPERTIESCLenumZero or more of { QUEUE_PROFILING_ENABLE, QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE }
DEVICE_PLATFORMWebCLPlatformObject
Exceptions:
sequence<DOMString>? getSupportedExtensions() (OpenCL 1.1 §9, clGetDeviceInfo)
Returns an array of extension names that are supported by this WebCLDevice. Any string in this list, when passed to `enableExtension`, must enable the corresponding extension.
CLboolean enableExtension(DOMString extensionName)
Enables the given WebCL extension on this WebCLDevice. Returns `true` if the extension is successfully enabled, or `false` if not. The available extension names can be queried by `getSupportedExtensions`. Note that enabling an extension does not take effect retroactively, i.e., contexts that were created before enabling the extension will continue to not have the extended capabilities.

WebCLContext

interface WebCLContext {

  WebCLBuffer createBuffer(CLenum memFlags, CLuint sizeInBytes, optional ArrayBufferView hostPtr);

  WebCLCommandQueue createCommandQueue(optional WebCLDevice? device = null, optional CLenum properties = 0);

  WebCLImage createImage(CLenum memFlags,
                         WebCLImageDescriptor descriptor,
                         optional ArrayBufferView hostPtr);

  WebCLProgram createProgram(DOMString source);

  WebCLSampler createSampler(CLboolean normalizedCoords,
                             CLenum addressingMode,
                             CLenum filterMode);

  WebCLUserEvent createUserEvent();

  any getInfo(CLenum name);

  sequence<WebCLImageDescriptor>? getSupportedImageFormats(optional CLenum memFlags = WebCL.MEM_READ_WRITE);

  void release();

  void releaseAll();
};
WebCLBuffer createBuffer(CLenum memFlags, CLuint sizeInBytes, optional ArrayBufferView hostPtr) (OpenCL 1.1 §5.2.1, clCreateBuffer)
Creates a new WebCLBuffer with a capacity of `sizeInBytes`, and optionally fills it with data from the given ArrayBufferView.
Exceptions:
WebCLCommandQueue createCommandQueue(optional WebCLDevice? device = null, optional CLenum properties = 0) (OpenCL 1.1 §5.1, clCreateCommandQueue)
Creates a new command queue for the given device. If `device` is `null`, the WebCL implementation will select any WebCLDevice that matches the given `properties` and is covered by this WebCLContext. If `properties` is omitted, the command queue is created with out-of-order execution disabled and profiling disabled.
Exceptions:
WebCLImage createImage(CLenum memFlags, WebCLImageDescriptor descriptor, optional ArrayBufferView hostPtr) (OpenCL 1.1 §5.3.1, clCreateImage2D)
Creates a new WebCLImage with the width and height specified in `descriptor`, and optionally fills it with data from the given ArrayBufferView.
Exceptions:
WebCLProgram createProgram(DOMString source) (OpenCL 1.1 §5.6.1, clCreateProgramWithSource)
Creates a WebCLProgram from a UTF-8 string.
Exceptions:
WebCLSampler createSampler(CLboolean normalizedCoords, CLenum addressingMode, CLenum filterMode) (OpenCL 1.1 §5.5.1, clCreateSampler)
Creates a new WebCLSampler with the given coordinate mode, addressing mode, and filtering mode. The `ADDRESS_NONE` mode is not available, due to its lack of bounds checking.
Exceptions:
WebCLUserEvent createUserEvent() (OpenCL 1.1 §5.9, clCreateUserEvent)
any getInfo(CLenum name) (OpenCL 1.1 §4.3, clGetContextInfo)

Returns the value for the passed name. The type returned is the natural type for the requested name, as given in the table below.

namereturn typedescription
CONTEXT_NUM_DEVICESCLuintthe number of devices in this context
CONTEXT_DEVICESsequence<WebCLDevice>the devices in this context
Exceptions:
sequence<WebCLImageDescriptor>? getSupportedImageFormats(optional CLenum memFlags = WebCL.MEM_READ_WRITE) (OpenCL 1.1 §5.3.2, clGetSupportedImageFormats)

Returns a list of image formats that are supported by this WebCLContext with the given `memFlags`. For example, passing in `MEM_READ_WRITE` will return a list of image formats that are available for both reading and writing in kernel code. The returned WebCLImageDescriptor objects will only have the `channelOrder` and `channelType` fields filled in; the other fields will have their default value (zero).

The minimum set of image formats that must be supported (for both reading and writing) by all devices is defined in the table below. Note that the set of required formats is the same as on OpenCL 1.1 Embedded Profile.

channelOrderchannelType
RGBA UNORM_INT8
UNORM_INT16
SIGNED_INT8
SIGNED_INT16
SIGNED_INT32
UNSIGNED_INT8
UNSIGNED_INT16
UNSIGNED_INT32
HALF_FLOAT
FLOAT
Exceptions:
void release() (OpenCL 1.1 §4.3, clReleaseContext)
Releases the resources held up by this object.
void releaseAll() (OpenCL 1.1 §4.3, clReleaseContext)
Releases the resources held up by this WebCLContext and all objects created from it, if any.

WebCLImageDescriptor

dictionary WebCLImageDescriptor {
  CLenum channelOrder = WebCL.RGBA;            
  CLenum channelType = WebCL.UNORM_INT8;
  CLuint width = 0;
  CLuint height = 0;
  CLuint rowPitch = 0;
};

WebCLCommandQueue

interface WebCLCommandQueue {

  ////////////////////////////////////////////////////////////////////////////
  //
  // Copying: Buffer <-> Buffer, Image <-> Image, Buffer <-> Image
  //

  void enqueueCopyBuffer(
                    WebCLBuffer                           srcBuffer,
                    WebCLBuffer                           dstBuffer,
                    CLuint                                srcOffset,
                    CLuint                                dstOffset,
                    CLuint                                numBytes,
                    optional sequence<WebCLEvent>?        eventWaitList = null,
                    optional WebCLEvent?                  event = null);

  void enqueueCopyBufferRect(
                    WebCLBuffer                           srcBuffer,
                    WebCLBuffer                           dstBuffer,
                    sequence<CLuint>                      srcOrigin,
                    sequence<CLuint>                      dstOrigin,
                    sequence<CLuint>                      region,
                    CLuint                                srcRowPitch,
                    CLuint                                srcSlicePitch,
                    CLuint                                dstRowPitch,
                    CLuint                                dstSlicePitch,
                    optional sequence<WebCLEvent>?        eventWaitList = null,
                    optional WebCLEvent?                  event = null);

  void enqueueCopyImage(
                    WebCLImage                            srcImage,
                    WebCLImage                            dstImage,
                    sequence<CLuint>                      srcOrigin,
                    sequence<CLuint>                      dstOrigin,
                    sequence<CLuint>                      region,
                    optional sequence<WebCLEvent>?        eventWaitList = null,
                    optional WebCLEvent?                  event = null);

  void enqueueCopyImageToBuffer(
                    WebCLImage                            srcImage,
                    WebCLBuffer                           dstBuffer,
                    sequence<CLuint>                      srcOrigin,
                    sequence<CLuint>                      srcRegion,
                    CLuint                                dstOffset,
                    optional sequence<WebCLEvent>?        eventWaitList = null,
                    optional WebCLEvent?                  event = null);

  void enqueueCopyBufferToImage(
                    WebCLBuffer                           srcBuffer,
                    WebCLImage                            dstImage,
                    CLuint                                srcOffset,
                    sequence<CLuint>                      dstOrigin,
                    sequence<CLuint>                      dstRegion,
                    optional sequence<WebCLEvent>?        eventWaitList = null,
                    optional WebCLEvent?                  event = null);

  ////////////////////////////////////////////////////////////////////////////
  //
  // Reading: Buffer -> Host, Image -> Host
  //

  void enqueueReadBuffer(
                    WebCLBuffer                           buffer,
                    CLboolean                             blockingRead,
                    CLuint                                bufferOffset,
                    CLuint                                numBytes,
                    ArrayBufferView                       hostPtr,
                    optional sequence<WebCLEvent>?        eventWaitList = null,
                    optional WebCLEvent?                  event = null);

  void enqueueReadBufferRect(
                    WebCLBuffer                           buffer,
                    CLboolean                             blockingRead,
                    sequence<CLuint>                      bufferOrigin,
                    sequence<CLuint>                      hostOrigin,
                    sequence<CLuint>                      region,
                    CLuint                                bufferRowPitch,
                    CLuint                                bufferSlicePitch,
                    CLuint                                hostRowPitch,
                    CLuint                                hostSlicePitch,
                    ArrayBufferView                       hostPtr,
                    optional sequence<WebCLEvent>?        eventWaitList = null,
                    optional WebCLEvent?                  event = null);

  void enqueueReadImage(
                    WebCLImage                            image,
                    CLboolean                             blockingRead,
                    sequence<CLuint>                      origin,
                    sequence<CLuint>                      region,
                    CLuint                                hostRowPitch,
                    ArrayBufferView                       hostPtr,
                    optional sequence<WebCLEvent>?        eventWaitList = null,
                    optional WebCLEvent?                  event = null);

  ////////////////////////////////////////////////////////////////////////////
  //
  // Writing: Host -> Buffer, Host -> Image
  //

  void enqueueWriteBuffer(
                    WebCLBuffer                           buffer,
                    CLboolean                             blockingWrite,
                    CLuint                                bufferOffset,
                    CLuint                                numBytes,
                    ArrayBufferView                       hostPtr,
                    optional sequence<WebCLEvent>?        eventWaitList = null,
                    optional WebCLEvent?                  event = null);

  void enqueueWriteBufferRect(
                    WebCLBuffer                           buffer,
                    CLboolean                             blockingWrite,
                    sequence<CLuint>                      bufferOrigin,
                    sequence<CLuint>                      hostOrigin,
                    sequence<CLuint>                      region,
                    CLuint                                bufferRowPitch,
                    CLuint                                bufferSlicePitch,
                    CLuint                                hostRowPitch,
                    CLuint                                hostSlicePitch,
                    ArrayBufferView                       hostPtr,
                    optional sequence<WebCLEvent>?        eventWaitList = null,
                    optional WebCLEvent?                  event = null);

  void enqueueWriteImage(
                    WebCLImage                            image,
                    CLboolean                             blockingWrite,
                    sequence<CLuint>                      origin,
                    sequence<CLuint>                      region,
                    CLuint                                hostRowPitch,
                    ArrayBufferView                       hostPtr,
                    optional sequence<WebCLEvent>?        eventWaitList = null,
                    optional WebCLEvent?                  event = null);

  ////////////////////////////////////////////////////////////////////////////
  //
  // Executing kernels
  //

  void enqueueNDRangeKernel(
                    WebCLKernel                           kernel,
                    CLuint                                workDim,
                    sequence<CLuint>?                     globalWorkOffset = null, 
                    sequence<CLuint>                      globalWorkSize,
                    optional sequence<CLuint>?            localWorkSize = null,
                    optional sequence<WebCLEvent>?        eventWaitList = null,
                    optional WebCLEvent?                  event = null);

  ////////////////////////////////////////////////////////////////////////////
  //
  // Synchronization
  //

  void enqueueMarker(WebCLEvent event);

  void enqueueBarrier();

  void enqueueWaitForEvents (sequence<WebCLEvent> eventWaitList);
  
  void finish(optional WebCLCallback whenFinished);

  void flush();

  ////////////////////////////////////////////////////////////////////////////
  //
  // Querying command queue information
  //

  any getInfo(CLenum name);

  void release();
};
void enqueueCopyBuffer( WebCLBuffer srcBuffer, WebCLBuffer dstBuffer, CLuint srcOffset, CLuint dstOffset, CLuint numBytes, optional sequence<WebCLEvent>? eventWaitList = null, optional WebCLEvent? event = null) (OpenCL 1.1 §5.2.2, clEnqueueCopyBuffer)
Exceptions:
void enqueueCopyBufferRect( WebCLBuffer srcBuffer, WebCLBuffer dstBuffer, sequence<CLuint> srcOrigin, sequence<CLuint> dstOrigin, sequence<CLuint> region, CLuint srcRowPitch, CLuint srcSlicePitch, CLuint dstRowPitch, CLuint dstSlicePitch, optional sequence<WebCLEvent>? eventWaitList = null, optional WebCLEvent? event = null) (OpenCL 1.1 §5.2.2, clEnqueueCopyBufferRect)
Exceptions:
void enqueueCopyImage( WebCLImage srcImage, WebCLImage dstImage, sequence<CLuint> srcOrigin, sequence<CLuint> dstOrigin, sequence<CLuint> region, optional sequence<WebCLEvent>? eventWaitList = null, optional WebCLEvent? event = null) (OpenCL 1.1 §5.2.2, clEnqueueCopyImage)
Exceptions:
void enqueueCopyBufferToImage( WebCLBuffer srcBuffer, WebCLImage dstImage, CLuint srcOffset, sequence<CLuint> dstOrigin, sequence<CLuint> dstRegion, optional sequence<WebCLEvent>? eventWaitList = null, optional WebCLEvent? event = null) (OpenCL 1.1 §5.3.4, clEnqueueCopyBufferToImage)
Exceptions:
void enqueueCopyImageToBuffer( WebCLImage srcImage, WebCLBuffer dstBuffer, sequence<CLuint> srcOrigin, sequence<CLuint> srcRegion, CLuint dstOffset, optional sequence<WebCLEvent>? eventWaitList = null, optional WebCLEvent? event = null) (OpenCL 1.1 §5.3.4, clEnqueueCopyImageToBuffer)
Exceptions:
void enqueueReadBuffer( WebCLBuffer buffer, CLboolean blockingRead, CLuint bufferOffset, CLuint numBytes, ArrayBufferView hostPtr, optional sequence<WebCLEvent>? eventWaitList = null, optional WebCLEvent? event = null) (OpenCL 1.1 §5.2.2, clEnqueueReadBuffer)
Exceptions:
void enqueueReadBufferRect( WebCLBuffer buffer, CLboolean blockingRead, sequence<CLuint> bufferOrigin, sequence<CLuint> hostOrigin, sequence<CLuint> region, CLuint bufferRowPitch, CLuint bufferSlicePitch, CLuint hostRowPitch, CLuint hostSlicePitch, ArrayBufferView hostPtr, optional sequence<WebCLEvent>? eventWaitList = null, optional WebCLEvent? event = null) (OpenCL 1.1 §5.2.2, clEnqueueReadBufferRect)
Exceptions:
void enqueueReadImage( WebCLImage image, CLboolean blockingRead, sequence<CLuint> origin, sequence<CLuint> region, CLuint hostRowPitch, ArrayBufferView hostPtr, optional sequence<WebCLEvent>? eventWaitList = null, optional WebCLEvent? event = null) (OpenCL 1.1 §5.2.2, clEnqueueReadImage)
Exceptions:
void enqueueWriteBuffer( WebCLBuffer buffer, CLboolean blockingWrite, CLuint bufferOffset, CLuint numBytes, ArrayBufferView hostPtr, optional sequence<WebCLEvent>? eventWaitList = null, optional WebCLEvent? event = null) (OpenCL 1.1 §5.2.2, clEnqueueWriteBuffer)
Exceptions:
void enqueueWriteBufferRect( WebCLBuffer buffer, CLboolean blockingWrite, sequence<CLuint> bufferOrigin, sequence<CLuint> hostOrigin, sequence<CLuint> region, CLuint bufferRowPitch, CLuint bufferSlicePitch, CLuint hostRowPitch, CLuint hostSlicePitch, ArrayBufferView hostPtr, optional sequence<WebCLEvent>? eventWaitList = null, optional WebCLEvent? event = null) (OpenCL 1.1 §5.2.2, clEnqueueWriteBufferRect)
Exceptions:
void enqueueWriteImage( WebCLImage image, CLboolean blockingWrite, sequence<CLuint> origin, sequence<CLuint> region, CLuint hostRowPitch, ArrayBufferView hostPtr, optional sequence<WebCLEvent>? eventWaitList = null, optional WebCLEvent? event = null) (OpenCL 1.1 §5.2.2, clEnqueueWriteImage)
Exceptions:
void enqueueNDRangeKernel( WebCLKernel kernel, CLuint workDim, sequence<CLuint>? globalWorkOffset = null, sequence<CLuint> globalWorkSize, optional sequence<CLuint>? localWorkSize = null, optional sequence<WebCLEvent>? eventWaitList = null, optional WebCLEvent? event = null) (OpenCL 1.1 §5.8, clEnqueueNDRangeKernel)
Exceptions:
void enqueueMarker(WebCLEvent event) (OpenCL 1.1 §5.10, clEnqueueMarker)
Exceptions:
void enqueueBarrier() (OpenCL 1.1 §5.10, clEnqueueBarrier)
void enqueueWaitForEvents (sequence<WebCLEvent> eventWaitList); (OpenCL 1.1 §5.10, clEnqueueWaitForEvents)
Exceptions:
void finish(optional WebCLCallback whenFinished) (OpenCL 1.1 §5.13, clFinish)
Invokes the given callback when all commands in this queue have completed. If a callback is not provided, the JavaScript main thread will be blocked until the commands have completed. Applications are strongly advised to provide a callback; the blocking mode mainly exists to ease the porting of existing OpenCL code.
Exceptions:
void flush() (OpenCL 1.1 §5.13, clFlush)
any getInfo(CLenum name) (OpenCL 1.1 §5.1, clGetCommandQueueInfo)
namereturn type
QUEUE_CONTEXTWebCLContext
QUEUE_DEVICEWebCLDevice
QUEUE_PROPERTIESCLenum
Exceptions:
void release() (OpenCL 1.1 §5.1, clReleaseCommandQueue)
Releases the resources held up by this object.
// This example shows how to wait for two kernels to complete before a
// third kernel executes

var eventWaitList = [ new WebCLEvent(), new WebCLEvent() ];
queue.enqueueNDRangeKernel(kernel1, dim, null, globals, null, null, eventWaitList[0]);
queue.enqueueNDRangeKernel(kernel2, dim, null, globals, null, null, eventWaitList[1]);
queue.enqueueNDRangeKernel(kernel3, dim, null, globals, null, eventWaitList);
// This example reads the entire contents of a WebCL buffer object
// into a newly created Uint8Array in host memory.

function readBufferToHost(srcBuffer) {

 // Query the number of bytes in the source buffer, create a new
 // Uint8Array of that size, then fill it with a blocking read.

 var numBytes = srcBuffer.getInfo(WebCL.MEM_SIZE);
 var dstArray = new Uint8Array(numBytes);
 queue.enqueueReadBuffer(srcBuffer, true, 0, numBytes, dstArray);

 return dstArray;
};
// This example fills a WebCL image with pixels from a WebCL buffer.

function copyBufferToImage(srcBuffer, dstImage) {

  // Extract the dimensions of the image.

  var imgWidth = dstImage.getInfo(WebCL.IMAGE_WIDTH);
  var imgHeight = dstImage.getInfo(WebCL.IMAGE_HEIGHT);
  queue.enqueueCopyBufferToImage(srcBuffer, dstImage, 0, [0,0], [imgWidth, imgHeight]);
};
// This example copies a rectangular region of memory from a
// buffer object to another.  The buffer objects are assumed
// to be the same size. The origin, width, height, and pitch
// (a.k.a. row stride) are given in bytes (not pixels, because
// these are buffers rather than images).

function copyBufferRect(srcBuffer, dstBuffer, srcX, srcY, dstX, dstY, w, h, pitch) {

  // The row pitch could be different for source and destination,
  // but is assumed to be the same in this example.

  queue.enqueueCopyBufferRect(srcBuffer, dstBuffer, [srcX, srcY], [dstX, dstY], 
                              [w, h], pitch, 0, pitch, 0);
};

WebCLMemoryObject

interface WebCLMemoryObject {
  any getInfo(CLenum name);
  void release();
};
any getInfo(CLenum name) (OpenCL 1.1 §5.4.1, clGetMemObjectInfo)
namereturn typereturn value
`MEM_TYPE``CLenum``MEM_OBJECT_BUFFER` or `MEM_OBJECT_IMAGE2D`
`MEM_FLAGS``CLenum`The `memFlags` as specified at construction time
`MEM_SIZE``CLuint`The size of this memory object in bytes
`MEM_CONTEXT``WebCLContext`the WebCL context of this memory object
`MEM_ASSOCIATED_MEMOBJECT``WebCLBuffer`The buffer object that this buffer was created from, or `null` if this buffer was not created using `createSubBuffer`
`MEM_OFFSET``CLuint`The offset given to `createSubBuffer`, or zero if this buffer was not created using `createSubBuffer`
Exceptions:
void release() (OpenCL 1.1 §5.4.1, clReleaseMemObject)
Releases the resources held up by this object.

WebCLBuffer

interface WebCLBuffer : WebCLMemoryObject {
  WebCLBuffer createSubBuffer(CLenum memFlags, CLuint origin, CLuint sizeInBytes);
};
WebCLBuffer createSubBuffer(CLenum memFlags, CLuint origin, CLuint sizeInBytes) (OpenCL 1.1 §5.2.1, clCreateSubBuffer)
Creates a new WebCLBuffer that represents a sub-region of this WebCLBuffer. The two buffers reference the same area of global memory, so changes made into one are immediately reflected in the other.
Exceptions:

WebCLImage

interface WebCLImage : WebCLMemoryObject {
  WebCLImageDescriptor getInfo();
};
WebCLImageDescriptor getInfo() (OpenCL 1.1 §5.3.6, clGetImageInfo)
Returns a new WebCLImageDescriptor containing information about the pixel format and dimensions of this WebCLImage.

WebCLSampler

interface WebCLSampler {
  any getInfo(CLenum name);
  void release();
};
any getInfo(CLenum name) (OpenCL 1.1 §5.5.2, clGetSamplerInfo)
namereturn type
SAMPLER_CONTEXTWebCLContext
SAMPLER_NORMALIZED_COORDSCLboolean
SAMPLER_ADDRESSING_MODECLenum
SAMPLER_FILTER_MODECLenum
Exceptions:
void release() (OpenCL 1.1 §5.5.1, clReleaseSampler)
Releases the resources held up by this object.

WebCLProgram

interface WebCLProgram {
  any getInfo(CLenum name);

  any getBuildInfo(WebCLDevice device, CLenum name);

  void build(optional sequence<WebCLDevice>? devices = null,
             optional DOMString? options = null,
             optional WebCLCallback whenFinished);

  WebCLKernel createKernel(DOMString kernelName);

  sequence<WebCLKernel> createKernelsInProgram();

  void release();
};
any getInfo(CLenum name) (OpenCL 1.1 §5.6.5, clGetProgramInfo)
namereturn type
PROGRAM_CONTEXTWebCLContext
PROGRAM_NUM_DEVICESCLuint
PROGRAM_DEVICESsequence<WebCLDevice>
PROGRAM_SOURCEDOMString
Exceptions:
any getBuildInfo(WebCLDevice device, CLenum name) (OpenCL 1.1 §5.6.5, clGetProgramBuildInfo)
namereturn type
PROGRAM_BUILD_STATUSCLint
PROGRAM_BUILD_OPTIONSDOMString
PROGRAM_BUILD_LOGDOMString
Exceptions:
void build(optional sequence<WebCLDevice>? devices = null, optional DOMString? options = null, optional WebCLCallback whenFinished) (OpenCL 1.1 §5.6.2, clBuildProgram)

Compiles this WebCLProgram for the given list of devices, or in absence of the list, for all devices associated with the WebCLContext that this WebCLProgram was created from. A string of compiler options and an asynchronous callback function to invoke when the build is completed can also be provided. Applications are strongly advised to provide a callback, to avoid blocking the JavaScript main thread. The available compiler options are listed in the table below; any other option is considered invalid.

Build optionDescription
-D nameEquivalent to #define name in OpenCL C.
-D name=definitionEquivalent to #define name definition in OpenCL C.
-cl-opt-disableDisable all optimizations.
-cl-single-precision-constantTreat double-precision constants as single-precision.
-cl-denorms-are-zeroAllow denormalized numbers to be flushed to zero.
-cl-mad-enableAllow a*b+c to be computed with potentially reduced accuracy.
-cl-no-signed-zerosAllow optimizations that ignore the signedness of zero.
-cl-unsafe-math-optimizationsAllow optimizations that may violate IEEE 754 and OpenCL numerical compliance requirements. This option includes -cl-no-signed-zeros and cl-mad-enable.
-cl-finite-math-onlyAllow optimizations that ignore NaNs and infinities.
-cl-fast-relaxed-mathEquivalent to -cl-finite-math-only and -cl-unsafe-math-optimizations.
-wInhibit all warning messages.
-WerrorMake all warnings into errors.
Exceptions:
WebCLKernel createKernel(DOMString kernelName) (OpenCL 1.1 §5.7.1, clCreateKernel)
Exceptions:
sequence<WebCLKernel> createKernelsInProgram() (OpenCL 1.1 §5.7.1, clCreateKernelsInProgram)
Exceptions:
void release() (OpenCL 1.1 §5.6.1, clReleaseProgram)
Releases the resources held up by this object.

WebCLKernel

The following methods are available for setting kernel arguments and querying kernel-specific information.

interface WebCLKernel {
  any getInfo(CLenum name);
  any getWorkGroupInfo(WebCLDevice? device, CLenum name);
  WebCLKernelArgInfo getArgInfo(CLuint index);
  void setArg(CLuint index, WebCLBuffer buffer);
  void setArg(CLuint index, WebCLImage image);
  void setArg(CLuint index, WebCLSampler value);
  void setArg(CLuint index, ArrayBufferView value);
  void release();
};
any getInfo(CLenum name) (OpenCL 1.1 §5.7.3, clGetKernelInfo)
namereturn type
KERNEL_FUNCTION_NAMEDOMString
KERNEL_NUM_ARGSCLuint
KERNEL_CONTEXTWebCLContext
KERNEL_PROGRAMWebCLProgram
Exceptions:
any getWorkGroupInfo(WebCLDevice? device, CLenum name) (OpenCL 1.1 §5.7.3, clGetKernelWorkGroupInfo)

Returns the value corresponding to the given `device` and `name`. If there is only a single WebCLDevice associated with this WebCLKernel, `null` may be passed for `device`.

namereturn type
KERNEL_WORK_GROUP_SIZECLuint
KERNEL_COMPILE_WORK_GROUP_SIZEsequence<CLuint>
KERNEL_LOCAL_MEM_SIZECLuint
KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLECLuint
KERNEL_PRIVATE_MEM_SIZECLuint
Exceptions:
WebCLKernelArgInfo getArgInfo(CLuint index) (OpenCL 1.2 §5.7.3, clGetKernelArgInfo)
Retrieves the name, type, and qualífiers of the kernel argument at the given `index`. Legal values for `index` range from zero to `N-1`, where `N` is the total number of arguments declared by the kernel.
Exceptions:
void setArg(CLuint index, WebCLBuffer buffer) (OpenCL 1.1 §5.7.2, clSetKernelArg)

Sets the given `buffer` to the kernel argument at the given `index`. Legal values for `index` range from zero to `N-1`, where `N` is the total number of arguments declared by the kernel.

Exceptions:
void setArg(CLuint index, WebCLImage image) (OpenCL 1.1 §5.7.2, clSetKernelArg)

Sets the given `image` to the kernel argument at the given `index`. Legal values for `index` range from zero to `N-1`, where `N` is the total number of arguments declared by the kernel.

Exceptions:
void setArg(CLuint index, WebCLSampler sampler) (OpenCL 1.1 §5.7.2, clSetKernelArg)

Sets the given `sampler` to the kernel argument at the given `index`. Legal values for `index` range from zero to `N-1`, where `N` is the total number of arguments declared by the kernel.

Exceptions:
void setArg(CLuint index, ArrayBufferView value) (OpenCL 1.1 §5.7.2, clSetKernelArg)

Sets the given `value` to the kernel argument at the given `index`. Legal values for `index` range from zero to `N-1`, where `N` is the total number of arguments declared by the kernel.

The type and length of `value` must match the base type and vector width of the target kernel argument. For example, if the kernel argument is of type `float4`, then `value` must be a Float32Array of length 4. Arguments with the `local` address space qualifier must be set using a Uint32Array of length 1, specifying the number of bytes to be allocated.

Since there is no 64-bit integer variant of ArrayBufferView, 64-bit integers must be represented as pairs of 32-bit unsigned integers. The low-order 32 bits are stored in the first element of each pair, and the high-order 32 bits in the second element.

Exceptions:
void release() (OpenCL 1.1 §5.7.1, clReleaseKernel)
Releases the resources held up by this object.
myKernel.setArg(0, 3.14159);                                  // ERROR: Passing in a Number is not allowed
myKernel.setArg(0, new Float32Array([3.14159]));              // cast 3.14159 to `float`, then pass to kernel as arg #0
myKernel.setArg(1, new Uint32Array([1.23, 2.34]));            // cast the numbers to `uint2`, then pass to kernel as arg #1
myKernel.setArg(2, new Uint32Array([512]));                   // reserve 512 bytes of local memory for arg #2
myKernel.setArg(3, myImage);                                  // pass `myImage` to kernel as arg #3

WebCLKernelArgInfo

dictionary WebCLKernelArgInfo {
  DOMString name;
  DOMString typeName;         // 'char', 'float', 'uint4', 'image2d_t', 'sampler_t', etc.
  DOMString addressQualifier; // 'global', 'local', 'constant', or 'private'
  DOMString accessQualifier;  // 'read_only', 'write_only', or 'none'
};
DOMString name
The name of the kernel argument; for example, srcPixels or foobarCoefficient.
DOMString typeName (OpenCL 1.1 §6.1, Data types)
The data type of the kernel argument in normalized form; for example, uint4 or image2d_t.
DOMString addressQualifier (OpenCL 1.1 §6.5, Address space qualifiers)
The address space qualifier of the kernel argument in normalized form (that is, not prefixed with underscores). The possible values are global, local, constant, and private.
DOMString accessQualifier (OpenCL 1.1 §6.6, Access qualifiers)
The access qualifier of the kernel argument in normalized form (that is, not prefixed with underscores). If `typeName` is image2d_t, the possible values are read_only and write_only. Otherwise, the value is none.

WebCLEvent

[Constructor]
interface WebCLEvent {
  any getInfo(CLenum name);
  CLulong getProfilingInfo(CLenum name);
  void setCallback(CLenum commandExecCallbackType, WebCLCallback notify);
  void release();
};
WebCLEvent()

Creates an empty WebCLEvent instance, to be populated by one of the `enqueue` methods. Each WebCLEvent instance can only be used once; the `enqueue` methods will throw an `INVALID_EVENT` exception if a previously populated event is passed in. User events are considered to be populated by `createUserEvent`, and as such cannot be re-populated.

any getInfo(CLenum name) (OpenCL 1.1 §5.9, clGetEventInfo)
namereturn type
EVENT_COMMAND_QUEUEWebCLCommandQueue
EVENT_CONTEXTWebCLContext
EVENT_COMMAND_TYPECLenum
EVENT_COMMAND_EXECUTION_STATUSCLint
Exceptions:
any getProfilingInfo(CLenum name) (OpenCL 1.1 §5.12, clGetEventProfilingInfo)
namereturn type
PROFILING_COMMAND_QUEUEDCLulong
PROFILING_COMMAND_SUBMITCLulong
PROFILING_COMMAND_STARTCLulong
PROFILING_COMMAND_ENDCLulong
Exceptions:
void setCallback(CLenum commandExecCallbackType, WebCLCallback notify) (OpenCL 1.1 §5.9, clSetEventCallback)
Sets a WebCLCallback function `notify` to be called when this WebCLEvent reaches the given `commandExecCallbackType` status (`COMPLETE`, `RUNNING`, `SUBMITTED`, or `QUEUED`). In this version of WebCL, only the `COMPLETE` status is supported.
Exceptions:
void release() (OpenCL 1.1 §5.9, clReleaseEvent)
Releases the resources held up by this object.

WebCLUserEvent

interface WebCLUserEvent : WebCLEvent {
  void  setStatus(CLint executionStatus);
};
void setStatus(CLint executionStatus) (OpenCL 1.1 §5.9, clSetUserEventStatus)
Exceptions:

Extensions

WebCL implementations may expose a number of extensions to the core functionality. The extensions are listed and documented in the WebCL Extension Registry. Each extension, regardless of its standardization status, must be explicitly enabled by the application before it can be used. This policy is aimed to minimize inadvertant usage of features that are not universally available.

The set of available extensions may vary from one WebCLDevice to another, and one WebCLPlatform to another, even within the same system. For example, the WebGL Resource Sharing extension may be supported by the GPU but not the CPU. The names of extensions that are supported system-wide can be queried with `webcl.getSupportedExtensions()`. Similarly, calling `getSupportedExtensions` on a particular WebCLPlatform returns the set of extensions that are supported by all devices on that platform. Finally, `getSupportedExtensions()` on an individual WebCLDevice returns the set of extensions that are supported by that device.

WebCL extensions can be enabled by calling `enableExtension` with the name of the desired extension as a string parameter. This returns `true` if the extension is successfully enabled, and `false` if not. This is illustrated in the following examples.

// Iterates over all WebCL devices and returns the ones that
// support the given extension, with that extension enabled.
//
function getDevicesWithExtensionEnabled(extensionName) {
  var devices = [];
  webcl.getPlatforms().forEach(function(platform) {
    platform.getDevices().forEach(function(device) {
      var isSupported = device.enableExtension(extensionName);
      if (isSupported) devices.push(device);
    });
  });
  return devices;
}
// Use the above function find the set of devices
// that support both KHR_fp64 and KHR_gl_sharing.
//
var withFP64 = getDevicesWithExtensionEnabled("KHR_fp64");
var withGL = getDevicesWithExtensionEnabled("KHR_gl_sharing");
var withBoth = withGL.filter(function(v) { return (withFP64.indexOf(v) !== -1) });

Security and Robustness

WebCL has been designed with security as a primary concern. This section defines the main security requirements that a conformant WebCL implementation must fulfill.

Out-of-Bounds Memory Access

WebCL kernels must not be able to access unauthorized areas of memory, regardless of address space (private, local, global, or constant). If detected during compilation, out-of-bounds (OOB) accesses must generate a compiler error. At runtime, OOB reads must return zero and writes must be discarded. Alternatively, the kernel may be terminated and an exception be thrown.

For purposes of bounds checking, the implementation may treat all `private` variables as one contiguous block of memory, and similarly for `local` variables, instead of enforcing the bounds of each variable separately. For example, in a kernel program containing two `private` arrays, an OOB read from the first array is allowed to return any value from either the first or the second array, rather than zero.

The WebCL Validator open source project was initiated by the WebCL working group for prevention of out-of-bounds memory accesses through instrumentation, analysis and validation of WebCL kernels. The Validator also enforces initialization of local and private memory (Khronos WebCL Validator Project).

Memory Initialization

To ensure that applications cannot inspect data left behind by previous applications, the WebCL implementation must initialize all buffers and variables to zero before giving the application read access to them. This requirement applies regardless of address space (private, local, global, or constant), and regardless of whether the allocation is made in kernel code or host code.

Where available, the OpenCL 1.2 extension `cl_khr_initialize_memory` allows WebCL implementations to automatically and efficiently initialize local and private memory before a kernel begins execution (OpenCL 1.2 Extensions Specification, §9.15).

Denial of Service

Long-running and/or computationally intensive kernels (or other commands in a command queue) may cause the system to become unresponsive by disproportionately consuming device resources. It is generally not possible to guard against this problem at the browser/WebCL level; the necessary mechanisms, such as watchdog timers and pre-emptive multitasking, must be provided by the OpenCL driver and the operating system. On systems where the necessary facilities are in place, WebCL implementations are strongly encouraged to make use of them in order to:

  1. Detect offending kernels and other commands. A command is considered offending if it runs for an excessively long time, or disproportionately consumes system resources.
  2. Terminate the contexts associated with the offending commands, before they render the OpenCL device unresponsive, and result in potential device reset.

Where available, the OpenCL 1.2 extension `cl_khr_terminate_context` can be used for fast termination of a context, if, for example, one or more of the kernels associated with the context has been running too long, or if the program terminates due to an exception (OpenCL 1.2 Extensions Specification, §9.16).

Implementation notes

Although WebCL is based on OpenCL 1.1 (Embedded Profile), it can also be implemented on top of OpenCL 1.2, or any future version that remains backwards compatible with OpenCL 1.1. However, an implementation running on top of OpenCL 1.2 or later must make sure to not allow applications to use any features that are not supported in WebCL. This implies, in particular, that implementations must not pass kernel code to the underlying OpenCL compiler without first checking that the kernel code is valid according to WebCL. For example, usage of `printf()` must be treated as a compilation error, because `printf` is not supported in WebCL.

The WebCL specification and conformance tests are written with the assumption that implementations will be based on conformant OpenCL drivers. Implementations are strongly recommended to support dynamic black-listing of malfunctioning or vulnerable OpenCL drivers in order to quickly eliminate any security or stability issues.

WebCL objects are regular JavaScript objects, and must behave accordingly. For example, the application must be able attach arbitrary properties and functions to them, and if the same object is retrieved twice, the instances returned must be equal according to the `===` operator. For example, the assertions in the following code example must be true:

  var ctx = webcl.createContext();
  ctx.name = "myContext";
  var queue = ctx.createCommandQueue();
  var ctx2 = queue.getInfo(WebCL.QUEUE_CONTEXT);
  ASSERT(ctx2 === ctx);
  ASSERT(ctx2.name === "myContext");

Differences between WebCL and OpenCL 1.1

This section describes changes made to the WebCL API and the WebCL C kernel programming language, relative to OpenCL 1.1 Embedded Profile and OpenCL C. The differences are as follows:

References

Normative references

[OPENCL11]
OpenCL 1.1 Specification, A. Munshi.
[WEBIDL]
Web IDL: W3C Candidate Recommendation, C. McCormack.
[WEBGL]
WebGL 1.0 Specification, C. Marrin.

Acknowledgments

This specification is produced by the Khronos WebCL Working Group.

Main Contributors (in alphabetical order):
Tomi Aarnio (Editor)
Alex Bourd
Mikaël Bourges-Sévenier (Editor)
Tasneem Brutch (Working Group Chair)
Fabien Cellier
Steven Eliuk
Simon Gibbs
Jeff Gilbert
Antonio Gomes
Won Jeon
Tim Johansson
Sharath Kamath
David Ligon
Laurent Morichetti
Igor Oliveira
Amit Rao
Ofer Rosenberg
Dave Shreiner
Shilpa Shri
Neil Trevett
Additional Contributors and Participants (in alphabetical order):
Remi Arnaud
Andrew Brownsword
Robert Elliott
Stephen Frye
Rick Hudson
Adam Lake
Thierry Lepley
Bill Licea-Kane
Mikael Lepisto
Anthony Liot
Aaftab Munshi
Jari Nikara
Janne Pietiäinen
Andrew Richards
Kenneth Russell
Dirk Schulze
Micah Villmow
Rami Ylimaki
Gongyuan Zhuang
Special thanks to the following participants (in alphabetical order):
Bastiaan Aarts
Ben Ashbaugh
Mark Callow
Sean Ellis
Anwar Ghuloum
Pyry Haulos
Dean Jackson
Benoit Jacob
Seungwon Lee
Erik Moller
Matthew Newport
Eisaku Ohbuchi
Rob Quill
Kalle Raita
Antti Rasmus
Lars Remes
Larry Seiler
Tatiana Shpeisman
Kevin Smith
Gowtham Tammana
Thierry Vuillaume
Todd Weybrew
Dong-Hoon Yoo