The OpenCL^™ C++ 1.0 Specification

The half data type must be IEEE 754-2008 compliant. half numbers have 1 sign bit, 5 exponent bits, and 10 mantissa bits. The interpretation of the sign, exponent and mantissa is analogous to IEEE 754 floating-point numbers.

The exponent bias is 15. The half data type must represent finite and normal numbers, denormalized numbers, infinities and NaN. Denormalized numbers for the half data type which may be generated when converting a float to a half using vstore_half and converting a half to a float using vload_half cannot be flushed to zero.

Conversions from float to half correctly round the mantissa to 11 bits of precision.

Conversions from half to float are lossless; all half numbers are exactly representable as float values.

The half data type can only be used to declare a pointer to a buffer that contains half values. All other operations are not allowed if the cl_khr_fp16 extension is not supported.

A few valid examples are given below:

The half scalar data type is required to be supported as a data storage format. Vector data load and store functions (described in the Vector Data Load and Store Functions section) must be supported.

This extension adds support for half scalar and vector types as built-in types that can be used for arithmetic operations, conversions etc. An application that wants to use half and halfn types will need to specify the -cl-fp16-enable compiler option (see the Double and half-precision floating-point options section).

The OpenCL compiler accepts an h and H suffix on floating point literals, indicating the literal is typed as a half

A few valid examples:

Hexadecimal floating point literals are supported in OpenCL C++.

The bool, char, unsigned char, short, unsigned short, int, unsigned int, long, unsigned long, half, float and double vector data types are supported. The vector data type is defined with the type name i.e. bool, char, uchar, short, ushort, int, uint, long, ulong, half, float or double followed by a literal value n that defines the number of elements in the vector. Supported values of n are 2, 3, 4, 8, and 16 for all vector data types.

The built-in vector data types are also declared as appropriate types in the OpenCL API (and header files) that can be used by an application. The following table describes the built-in vector data type in the OpenCL C++ programming language and the corresponding data type available to the application:

The halfn vector data type is required to be supported as a data storage format. Vector data load and store functions (described in the Vector Data Load and Store Functions section) must be supported.

Support for the doublen vector data type is optional.

Vector constructors are defined to initialize a vector data type from a list of scalar or vectors. The forms of the constructors that are available is the set of possible argument lists for which all arguments have the same element type as the result vector, and the total number of elements is equal to the number of elements in the result vector. In addition, a form with a single scalar of the same type as the element type of the vector is available.

For example, the following forms are available for float4:

Operands are evaluated by standard rules for function evaluation, except that implicit scalar-to-vector conversion shall not occur. The order in which the operands are evaluated is undefined. The operands are assigned to their respective positions in the result vector as they appear in memory order. That is, the first element of the first operand is assigned to result.x, the second element of the first operand (or the first element of the second operand if the first operand was a scalar) is assigned to result.y, etc. In the case of the form that has a single scalar operand, the operand is replicated across all lanes of the vector.

Examples:

Many operators that expect operands of arithmetic type cause conversions and yield result types in a similar way. The purpose is to determine a common real type for the operands and result. For the specified operands, each operand is converted, without change of type domain, to a type whose corresponding real type is the common real type. For this purpose, all vector types shall be considered to have higher conversion ranks than scalars. Unless explicitly stated otherwise, the common real type is also the corresponding real type of the result, whose type domain is the type domain of the operands if they are the same, and complex otherwise. This pattern is called the usual arithmetic conversions. If the operands are of more than one vector type, then an error shall occur. Implicit conversions between vector types are not permitted, per the Implicit Type Conversions section.

Otherwise, if there is only a single vector type, and all other operands are scalar types, the scalar types are converted to the type of the vector element, then widened into a new vector containing the same number of elements as the vector, by duplication of the scalar value across the width of the new vector.

The variables are allocated from the global memory pool if they meet the criteria described in the Implicit Storage Classes section for the implicit global storage class or they are declared using explicit global storage class from the standard library (see the global class section).

The global memory objects can be:

The non-trivial constructors and destructors are supported with limitations described in the Memory initialization section.

The constructors of objects in global memory are executed before the first kernel execution in the program. The destructors executed at program release time.

The additional restrictions may apply if the explicit global storage class is used. Please refer to the Restrictions section for more details.

The local variables can be only allocated in a program using the explicit local storage class from the standard library (see the local class section). This type of memory is allocated for each work-group executing the kernel and exist only for the lifetime of the work-group executing the kernel.

The non-trivial constructors and destructors are supported with limitations described in the Memory initialization section.

The constructors of objects in local memory are executed by one work-item before the kernel body execution. The destructors are executed by one work-item after the kernel body execution.

The additional restrictions may apply if the explicit local storage class is used. Please refer to the Restrictions section for more details.

The variables are allocated from the private memory pool if they meet the criteria described in Implicit Storage Classes for the implicit private storage class or they were declared using explicit private storage class from the standard library (see the priv class section).

The non-trivial constructors and destructors are supported.

The additional restrictions may apply if the explicit priv storage class is used. Please refer to the Restrictions section for more details.

The constant variables can be only allocated in a program using the explicit constant storage class from the standard library (see the constant class section). The variables declared using the constant<T> class refer to memory objects allocated from the global memory pool and which are accessed inside a kernel(s) as read-only variables. These read-only variables can be accessed by all (global) work-items of the kernel during its execution.

The constant objects must be constructible at compile time, they cannot have any user defined constructors, destructors, methods and operators. Otherwise behavior is undefined.

The additional restrictions may apply if the explicit constant storage class is used. Please refer to the Restrictions section for more details.

The OpenCL host compiler and the OpenCL C++ kernel language device compiler can have different requirements for i.e. type sizes, data packing and alignment, etc., therefore the kernel parameters must meet the following requirements:

This attribute specifies a minimum alignment (in bytes) for variables of the specified type. For example, the declarations:

force the compiler to insure (as far as it can) that each variable whose type is struct S or more_aligned_int will be allocated and aligned at least on a 8-byte boundary.

Note that the alignment of any given struct or union type is required by the C++ standard to be at least a perfect multiple of the lowest common multiple of the alignments of all of the members of the struct or union in question and must also be a power of two. This means that you can effectively adjust the alignment of a class or union type by attaching an aligned attribute to any one of the members of such a type, but the notation illustrated in the example above is a more obvious, intuitive, and readable way to request the compiler to adjust the alignment of an entire class or union type.

As in the preceding example, you can explicitly specify the alignment (in bytes) that you wish the compiler to use for a given class or union type. Alternatively, you can leave out the alignment factor and just ask the compiler to align a type to the maximum useful alignment for the target machine you are compiling for. For example, you could write:

Whenever you leave out the alignment factor in an aligned attribute specification, the compiler automatically sets the alignment for the type to the largest alignment which is ever used for any data type on the target machine you are compiling for. In the example above, the size of each short is 2 bytes, and therefore the size of the entire struct S type is 6 bytes. The smallest power of two which is greater than or equal to that is 8, so the compiler sets the alignment for the entire struct S type to 8 bytes.

Note that the effectiveness of aligned attributes may be limited by inherent limitations of the OpenCL device and compiler. For some devices, the OpenCL compiler may only be able to arrange for variables to be aligned up to a certain maximum alignment. If the OpenCL compiler is only able to align variables up to a maximum of 8 byte alignment, then specifying aligned(16) will still only provide you with 8 byte alignment. See your platform-specific documentation for further information.

The aligned attribute can only increase the alignment; but you can decrease it by specifying packed as well. See below.

This attribute, attached to class or union type definition, specifies that each member of the structure or union is placed to minimize the memory required. When attached to an enum definition, it indicates that the smallest integral type should be used.

Specifying this attribute for class and union types is equivalent to specifying the packed attribute on each of the structure or union members.

In the following example struct my_packed_struct’s members are packed closely together, but the internal layout of its s member is not packed. To do that, struct my_unpacked_struct would need to be packed, too.

You may only specify this attribute on the definition of an enum, class or union, not on a typedef which does not also define the enumerated type, structure or union.

This attribute specifies a minimum alignment for the variable or class field, measured in bytes. For example, the declaration:

causes the compiler to allocate the global variable x on a 16-byte boundary. The alignment value specified must be a power of two.

You can also specify the alignment of structure fields. For example, to create double-word aligned int pair, you could write:

This is an alternative to creating a union with a double member that forces the union to be double-word aligned.

As in the preceding examples, you can explicitly specify the alignment (in bytes) that you wish the compiler to use for a given variable or structure field. Alternatively, you can leave out the alignment factor and just ask the compiler to align a variable or field to the maximum useful alignment for the target machine you are compiling for. For example, you could write:

Whenever you leave out the alignment factor in an aligned attribute specification, the OpenCL compiler automatically sets the alignment for the declared variable or field to the largest alignment which is ever used for any data type on the target device you are compiling for.

When used on a class, or class member, the aligned attribute can only increase the alignment; in order to decrease it, the packed attribute must be specified as well. When used as part of a typedef, the aligned attribute can both increase and decrease alignment, and specifying the packed attribute will generate a warning.

Note that the effectiveness of aligned attributes may be limited by inherent limitations of the OpenCL device and compiler. For some devices, the OpenCL compiler may only be able to arrange for variables to be aligned up to a certain maximum alignment. If the OpenCL compiler is only able to align variables up to a maximum of 8 byte alignment, then specifying aligned(16) will still only provide you with 8 byte alignment. See your platform-specific documentation for further information.

The packed attribute specifies that a variable or class field should have the smallest possible alignment - one byte for a variable, unless you specify a larger value with the aligned attribute.

Here is a structure in which the field x is packed, so that it immediately follows a:

An attribute list placed at the beginning of a user-defined type applies to the variable of that type and not the type, while attributes following the type body apply to the type.

For example:

The optional [[cl::work_group_size_hint(X, Y, Z)]] is a hint to the compiler and is intended to specify the work-group size that may be used i.e. value most likely to be specified by the local_work_size argument to clEnqueueNDRangeKernel. For example the [[cl::work_group_size_hint(1, 1, 1)]] is a hint to the compiler that the kernel will most likely be executed with a work-group size of 1.

The specialization constants (see the Specialization Constants section) can be used as arguments of cl::work_group_size_hint attribute.

The optional [[cl::required_work_group_size(X, Y, Z)]] is the work-group size that must be used as the local_work_size argument to clEnqueueNDRangeKernel. This allows the compiler to optimize the generated code appropriately for this kernel.

If Z is one, the work_dim argument to clEnqueueNDRangeKernel can be 2 or 3. If Y and Z are one, the work_dim argument to clEnqueueNDRangeKernel can be 1, 2 or 3.

The specialization constants (see the Specialization Constants section) can be used as arguments of cl::required_work_group_size(X, Y, Z) attribute.

The optional [[cl::required_num_sub_groups(X)]] is the number of sub-groups that must be generated by a kernel launch. To ensure that this number is created the queries mapping number of sub-groups to local size may be used. This allows the compiler to optimize the kernel based on the sub-group count and in addition allows the API to enforce correctness of kernel use to the user when concurrency of sub-groups is a requirement.

The specialization constants (see the Specialization Constants section) can be used as argument of cl::required_num_sub_groups attribute.

The optional [[cl::vec_type_hint(<type>)]] is a hint to the compiler and is intended to be a representation of the computational width of the kernel, and should serve as the basis for calculating processor bandwidth utilization when the compiler is looking to autovectorize the code. In the [[cl::vec_type_hint(<type>)]] qualifier <type> is one of the built-in vector types listed in Device built-in vector data types table or the constituent scalar element types. If cl::vec_type_hint(<type>) is not specified, the kernel is assumed to have the [[cl::vec_type_hint(int)]] qualifier.

For example, where the developer specified a width of float4, the compiler should assume that the computation usually uses up to 4 lanes of a float vector, and would decide to merge work-items or possibly even separate one work-item into many threads to better match the hardware capabilities. A conforming implementation is not required to autovectorize code, but shall support the hint. A compiler may autovectorize, even if no hint is provided. If an implementation merges N work-items into one thread, it is responsible for correctly handling cases where the number of global or local work-items in any dimension modulo N is not zero.

Examples:

This attribute can be provided with a kernel argument of type constant_ptr<T>, constant<T>*, constant<T>&, local_ptr<T>, local<T>*, local<T>&. The value of the attribute specifies the maximum size in bytes of the corresponding memory object. This size cannot exceed the limits supported by the device:

The specialization constants (see the Specialization Constants section) can be used as argument of cl::max_size attribute.

Examples:

The and [[cl::unroll_hint(n)]] attribute qualifiers can be used to specify that a loop (for, while and do loops) can be unrolled. This attribute qualifier can be used to specify full unrolling or partial unrolling by a specified amount. This is a compiler hint and the compiler may ignore this directive.

n is the loop unrolling factor and must be a positive integral compile time constant expression. An unroll factor of 1 disables unrolling. If n is not specified, the compiler determines the unrolling factor for the loop.

Examples:

This tells the compiler to unroll the above while loop by a factor of 2.

In the example above, the compiler will determine how much to unroll the loop.

The above is an example where the loop should not be unrolled.

Below are some examples of invalid usage of [[cl::unroll_hint(n)]].

The above example is an invalid usage of the loop unroll factor as the loop unroll factor is negative.

The above example is invalid because the unroll_hint attribute qualifier is used on a non-loop construct.

The above example is invalid because the loop unroll factor is not a compile-time constant expression.

The [[cl::ivdep]] (ignore vector dependencies) attribute qualifier is a hint to the compiler and may appear in loops to indicate that the compiler may assume there are no memory dependencies across loop iterations in order to autovectorize consecutive iterations of the loop. This attribute qualifier may appear in one of the following forms:

If the parameter len is specified, it is used to specify the maximum number of consecutive iterations without loop-carried dependencies. len is a lower bound on the distance of any loop-carried dependence, and it applies to arbitrary alignment. For example, any 4 consecutive iterations can be vectorized with cl::ivdep(4). The len parameter must be a positive integer. The final decision whether to autovectorize the complete loop may be subject to other compiler heuristics as well as flags e.g., -cl-fast-relaxed-math to ignore non-associated operations.

Examples:

In the example above, assuming that A and B are not restricted pointers, it is unknown if C aliases A or B. Placing the [[cl::ivdep]] attribute before the loop lets the compiler assume there are no memory dependencies across the loop iterations.

In the example above, buffer A is read from and written to in the loop iterations. In each iteration, the read and write to A are to different indices. In this case it is not safe to vectorize the loop to a vector length greater than K, so the len parameter is specified with a value that is known to be not greater than any value that K may take during the execution of loop. In this example we are guaranteed (by len) that K will always be greater than or equal to 8.

Below is an example of invalid usage of [[cl::ivdep]].

The above example is an invalid usage of the attribute qualifier as len is negative.

Examples of casting between two vector types with saturation.

Examples of casting from float to integer vector type with saturation and rounding mode specified.

Examples of casting from integer to float vector type.

Examples of reinterpreting data types using as_type<> function.

The variables declared using global<T> class refer to memory objects allocated from the global memory pool (see the Global Memory Pool section). The global storage class can only be used to declare variables at program, function and class scope. The variables at function and class scope must be declared with static specifier.

If T is a fundamental or an array type, the global class should meet the following requirements:

If T is a class type, the global class should provide the following interface:

The variables declared using local<T> class refer to memory objects allocated from the local memory pool (see the Local Memory Pool section). The local storage class can only be used to declare variables at program, kernel and class scope. The variables at class scope must be declared with static specifier.

If T is a fundamental or an array type, the local class should meet the following requirements:

If T is a class type, the local class should provide the following interface:

The variables declared using the priv<T> class refer to memory objects allocated from the private memory pool.

The priv storage class cannot be used to declare variables in the program scope, with static specifier or extern specifier.

If T is a fundamental or an array type, the priv class should meet the following requirements:

If T is a class type, the priv class should provide the following interface:

The variables declared using the constant<T> class refer to memory objects allocated from the global memory pool and which are accessed inside a kernel(s) as read-only variables. The constant storage class can only be used to declare variables at program, kernel and class scope. The variables at class scope must be declared with static specifier.

The T type must meet the following requirements:

If T is a fundamental, array or class type, the constant class should meet the following requirements:

Constructs an object which points to nothing.

Constructs an object which points to p.

Copy constructor.

Move constructor.

Constructs an object initialized with nullptr.

Copy assignment operator

Move assignment operator

Assigns r pointer to the stored pointer

Assigns nullptr to the stored pointer

Returns *get(). It is only defined in single object version of the explicit address space pointer class. The result of this operator is undefined if get() == nullptr.

Returns get(). It is only defined in single object version of the explicit address space pointer class. The result of this operator is undefined if get() == nullptr.

Returns get()[pos]. The subscript operator is only defined in specialized global_ptr<T[]>, local_ptr<T[]>, private_ptr<T[]> and constant_ptr<T[]> version for array types. The result of this operator is undefined if pos >= the number of elements in the array to which the stored pointer points.

Returns the stored pointer.

Returns get() != nullptr.

Assigns nullptr to the stored pointer and returns the value get() had at the start of the call to release.

Assigned p to the stored pointer. It is only defined in single object version of the explicit address space pointer class

Assigned p to the stored pointer. It is only defined in specialized global_ptr<T[]>, local_ptr<T[]>, private_ptr<T[]> and constant_ptr<T[]> version for array types.

Equivalent to reset(pointer()). It is only defined in specialized global_ptr<T[]>, local_ptr<T[]>, private_ptr<T[]> and constant_ptr<T[]> version for array types.

Invokes swap on the stored pointers.

Prefix increment operator. Increments the stored pointer by one.

Postfix increment operator. Increments the stored pointer by one.

Prefix decrement operator. Decrements the stored pointer by one.

Postfix decrement operator. Decrements the stored pointer by one.

Adds r to the stored pointer and returns *this.

Subtracts r to the stored pointer and returns *this.

Adds r to the stored pointer and returns the value *this has at the start of operator+.

Subtracts r to the stored pointer and returns the value *this has at the start of operator-.

Comparison operator== for the explicit address space pointer classes.

Comparison operator== for the explicit address space pointer classes with a nullptr_t.

Comparison operator!= for the explicit address space pointer classes.

Comparison operator!= for the explicit address space pointer classes with a nullptr_t.

Comparison operator< for the explicit address space pointer classes.

Comparison operator< for the explicit address space pointer classes with a nullptr_t.

Comparison operator> for the explicit address space pointer classes.

Comparison operator> for the explicit address space pointer classes with a nullptr_t.

Comparison operator<= for the explicit address space pointer classes.

Comparison operator<= for the explicit address space pointer classes with a nullptr_t.

Comparison operator>= for the explicit address space pointer classes.

Comparison operator>= for the explicit address space pointer classes with a nullptr_t.

Calls a.swap(b)

Returns the mem_fence value for ptr. ptr must be the generic pointer and it cannot be the explicit address space pointer (global_ptr<>, local_ptr<>, private_ptr<> and constant_ptr<>) or pointer to address space storage class (global<>*, local<>*, priv<>* and constant<>*).

Returns a pointer that points to a region in the address space pointer class specified in T if dynamic_asptr_cast can cast ptr to the specified address space. Otherwise it returns nullptr. Only global_ptr<U>, local_ptr<U> and private_ptr<U> are valid T template arguments. ptr must be the generic pointer and it cannot be the explicit address space pointer (global_ptr<>, local_ptr<>, private_ptr<> and constant_ptr<>) or pointer to address space storage class (global<>*, local<>*, priv<>* and constant<>*).

The expression static_cast(r.get()) shall be well formed.

The expression reinterpret_cast(r.get()) shall be well formed.

Example of passing an explicit address space storage object to a kernel.

Example of passing an explicit address space pointer object to a kernel.

Example of casting a generic pointer to an explicit address space pointer object. This is the runtime operation and the dynamic_asptr_cast can fail.

Example of using an array with an explicit address space storage class.

Example of using a fundamental type with an explicit address space storage class.

Constructor of spec_constant class. The value parameter is a default value of the specialization constant that will be used if a value is not set by the host API. It must be a literal value.

Return a value of specialization constant. If an object is not specialized from the host, the default value will be returned.

Template parameter T in spec_constant class template denotes the data type of specialization constant. The type T must be integral or floating point type.

Template parameter ID in spec_constant class template denotes an unique ID of the specialization constant that can be used to set a value from the host API. The value of ID must be unique within this compilation unit and across any other SPIR-V modules that it is linked with.

Example of using the specialization constant in the kernel.

Example of specializing one of the dimensions in cl::required_work_group_size attribute.

Example of using built-in vector iterators.

Example of iterating though built-in vector channels and using range library.

The examples of invalid use of marker types.

The examples of how to use is_marker_type trait.

Reads a color value from the non-multisample image using sampler and floating point coordinates.

Reads a color value from non-multisample image using sampler and integer coordinates.

A sampler must use filter mode set to filtering_mode::nearest, normalized coordinates and addressing mode set to addressing_mode::clamp_to_edge, addressing_mode::clamp, addressing_mode::none, otherwise the values returned are undefined.

Reads a color value from non-multisample image using sampler and floating point coordinates in the mip-level specified by lod.

Method is present for non-multisample images if the cl_khr_mipmap_image extension is enabled.

Use the gradients to compute the lod and coordinate coord to do an element lookup in the mip-level specified by the computed lod.

Method is present if the cl_khr_mipmap_image extension is enabled.

Based on the parameters with which image was created on host side the function will return different ranges of values

Values returned by image::sample where T is a floating-point type for image objects with image_channel_type values not specified in the description above are undefined.

The image::sample functions that take an image object where T is a signed integer type can only be used with image objects created with:

If the image_channel_type is not one of the above values, the values returned by image::sample are undefined.

The image::sample functions that take an image object where T is an unsigned integer type can only be used with image objects created with:

If the image_channel_type is not one of the above values, the values returned by image::sample are undefined.

Reads a color value from non-multisample image without sampler and integral coordinates. If the cl_khr_mipmap_image extension is present, may also perform reads from mipmap layer 0.

The image read function behaves exactly as the corresponding image sample function using a sampler that has filter mode set to filtering_mode::nearest, normalized coordinates set to normalized_coordinates::unnormalized and addressing mode set to addressing_mode::none. There is one execption for cases where the image sample type is image_channel_type::fp32. In this exceptional case, when channel data values are denormalized, the image read function may return the denormalized data, while the sample function may flush denormalized channel data values to zero. The coordinates must be between 0 and image size in that dimension non inclusive.

Based on the parameters with which image was created on host side the function will return different ranges of values

Values returned by image::read where T is a floating-point type for image objects with image_channel_type values not specified in the description above are undefined.

The image::read functions that take an image object where T is a signed integer type can only be used with image objects created with:

If the image_channel_type is not one of the above values, the values returned by image::read are undefined.

The image::read functions that take an image object where T is an unsigned integer type can only be used with image objects created with image_channel_type set to one of the following values:

If the image_channel_type is not one of the above values, the values returned by image::read are undefined.

Use the coordinate and sample to do an element lookup in the image object. Method is only available in the MSAA image types, and if the cl_khr_gl_msaa_sharing and cl_khr_gl_depth_images extensions are enabled.

When a multisample image is accessed in a kernel, the access takes one vector of integers describing which pixel to fetch and an integer corresponding to the sample numbers describing which sample within the pixel to fetch. sample identifies the sample position in the multi-sample image.

For best performance, we recommend that sample be a literal value so it is known at compile time and the OpenCL compiler can perform appropriate optimizations for multisample reads on the device.

No standard sampling instructions are allowed on the multisample image. Accessing a coordinate outside the image and/or a sample that is outside the number of samples associated with each pixel in the image is undefined.

Writes a color value to location specified by coordinates from non-multisample image. If the cl_khr_mipmap_image_writes extension is present, may also perform writes to mipmap layer 0. The coordinates must be between 0 and image size in that dimension non inclusive.

Based on the parameters with which image was created on host side the function will perform appropriate data format conversions before writing a color value.

Writes a color value to location specified by coordinates and lod from mipmap image. The coordinates must be between 0 and image size in that dimension non inclusive.

Method is present if the cl_khr_mipmap_image extension is enabled.

Based on the parameters with which image was created on host side the function will perform appropriate data format conversions before writing a color value.

The image::write functions that take an image object where T is a floating-point type can only be used with image objects created with image_channel_type set to one of the pre-defined packed formats or set to:

The image::write functions that take an image object where T is a signed integer type can only be used with image objects created with:

The image::write functions that take an image object where T is an unsigned integer type can only be used with image objects created with:

The behavior of image::write for image objects created with image_channel_type values not specified in the description above is undefined.

Creates a pixel which can be used to read or/and write operation(s). It depends on image_access specified in the image.

Reads a color value from non-multisample image without sampler and integral coordinates specified in pixel. If the cl_khr_mipmap_image extension is present, may also perform reads from mipmap layer 0.

This function is similar to image::read method. Please refer to description of this method for more details.

Writes a color value to location specified by coordinates in pixel from non-multisample image. If the cl_khr_mipmap_image_writes extension is present, may also perform writes to mipmap layer 0. The coordinates specified in pixel must be between 0 and image size in that dimension non inclusive.

Based on the parameters with which image was created on host side the function will perform appropriate data format conversions before writing a color value.

This function is similar to image::write method. Please refer to description of this method for more details.

Returns width of the image.

Returns width of the mip-level specified by lod.

Method is present in the non-multisample image types if the cl_khr_mipmap_image extension is enabled.

Returns height of the image.

Returns height of the mip-level specified by lod.

Method is present in the non-multisample image types if the cl_khr_mipmap_image extension is enabled.

Returns depth of the image.

Returns depth of the mip-level specified by lod.

Method is present in the non-multisample image types if the cl_khr_mipmap_image extension is enabled.

Returns size of the image array.

Returns size of the image array specified by lod.

Method is present in the non-multisample image types if the cl_khr_mipmap_image extension is enabled.

Returns appropriately sized vector, or scalar for 1 dimensional images, containing all image dimensions followed by array size.

Returns format of the image as specified upon its creation on host side.

Returns channel order of the image as specified upon its creation on host side.

Returns number of mipmaps of image. Method is present if the cl_khr_mipmap_image or cl_khr_mipmap_image_writes extension is enabled.

The example how to use an image object with sampler-less reads.

The example how to use an image object with image_access::read_write access and atomic_fence function.

The example how to use an image object with sampler passed by a kernel argument.

The example how to use an image object with sampler declared at program scope.

The example how to use an image object with sampler declared at non-program scope.

Read packet from pipe into ref.

Returns true if read is successful and false if the pipe is empty.

Write packet specified by ref to pipe. Returns true if write is successful and false if the pipe is full.

Reserve num_packets entries for reading/writing from/to pipe. Returns a valid reservation if the reservation is successful.

The reserved pipe entries are referred to by indices that go from 0 …​ num_packets - 1.

Reserve num_packets entries for reading/writing from/to pipe. Returns a valid reservation if the reservation is successful.

The reserved pipe entries are referred to by indices that go from 0 …​ num_packets - 1.

Reserve num_packets entries for reading/writing from/to pipe. Returns a valid reservation if the reservation is successful.

The reserved pipe entries are referred to by indices that go from 0 …​ num_packets - 1.

Returns the current number of packets that have been written to the pipe, but have not yet been read from the pipe. The number of available entries in a pipe is a dynamic value. The value returned should be considered immediately stale.

Returns the maximum number of packets specified when pipe was created.

Read packet from the reserved area of the pipe referred to by index into ref.

The reserved pipe entries are referred to by indices that go from 0 …​ num_packets - 1.

Returns true if read is successful and false otherwise.

Write packet specified by ref to the reserved area of the pipe referred to by index.

The reserved pipe entries are referred to by indices that go from 0 …​ num_packets - 1.

Returns true if write is successful and false otherwise.

Indicates that all reads/writes to num_packets associated with reservation are completed.

Return true if reservation is a valid reservation and false otherwise.

Return true if reservation is a valid reservation and false otherwise.

Constructs a read only or write only pipe from pipe_storage object. One kernel can have only one pipe accessor associated with one pipe_storage object.

Constructs a read only or write only pipe from pipe_storage object. One kernel can have only one pipe accessor associated with one pipe_storage object.

Template parameter T in pipe and pipe_storage class template denotes the data type stored in pipe. The type T must be a POD type i.e. satisfy is_pod<T>::value == true.

All work-group specific functions must be encountered by all work items in a work-group executing the kernel with the same argument values, otherwise the behavior is undefined.

All sub-group specific functions must be encountered by all work items in a sub-group executing the kernel with the same argument values, otherwise the behavior is undefined.

The following behavior is undefined:

The following behavior is undefined:

Example of reading from a pipe object.

Example of writing to a pipe object.

Example of reading from a pipe object using reservations.

Example of using a pipe_storage object and how to create the pipe objects/accessors from it.

Example of using more than one pipe_storage object.

This method allows to enqueue functor or lambda fun on the device with specified policy over the specified ndrange.

args are the arguments that will be passed to fun when kernel will be enqueued with the exception for local_ptr parameters. For local pointers user must supply the size of local memory that will be allocated using local_ptr<T>::size_type\{num elements}. Please see examples how to use enqueue_kernel are in the Examples section.

This method enqueues functor or lambda fun in the same way as the overload above with the exception for the passed event list. If an event is returned, enqueue_kernel performs an implicit retain on the returned event.

This method enqueues a marker to device queue. The marker command waits for a list of events specified by event_wait_list to complete before the marker completes. event_ret must not be nullptr as otherwise this is a no-op.

If an event is returned, enqueue_marker performs an implicit retain on the returned event.

Returns true if event object is a valid event. Otherwise returns false.

Returns true if event object is a valid event. Otherwise returns false.

Increments the event reference count. Event must be an event returned by enqueue_kernel or enqueue_marker or a user event.

Decrements the event reference count. The event object is deleted once the event reference count is zero, the specific command identified by this event has completed (or terminated) and there are no commands in any device command queue that require a wait for this event to complete. Event must be an event returned by enqueue_kernel, enqueue_marker or a user event.

Sets the execution status of a user event. Event must be a user event. status can be either event_status::complete or event_status::error value indicating an error.

Captures the profiling information for functions that are enqueued as commands. The specific function being referred to is: enqueue_kernel. These enqueued commands are identified by unique event objects. The profiling information will be available in value once the command identified by event has completed. Event must be an event returned by enqueue_kernel.

name identifies which profiling information is to be queried and can be:

Returns the default device queue. If a default device queue has not been created, device_queue::is_valid() will return false.

Creates a user event. Returns the user event. The execution status of the user event created is set to event_status::submitted.

This provides a mechanism to query the maximum work-group size that can be used to execute a functor on the current device.

fun specifies the functor representing the kernel code that would be enqueued.

args are the arguments that will be passed to fun when kernel will be enqueued with the exception for local_ptr parameters. For local pointers user must supply the size of local memory that will be allocated.

Returns the preferred multiple of work-group size for launch. This is a performance hint. Specifying a work-group size that is not a multiple of the value returned by this query as the value of the local work size argument to enqueue will not fail to enqueue the functor for execution unless the work-group size specified is larger than the device maximum.

fun specifies the functor representing the kernel code that would be enqueued.

args are the arguments that will be passed to fun when kernel will be enqueued with the exception for local_ptr parameters. For local pointers user must supply the size of local memory that will be allocated.

Returns the number of sub-groups in each work-group of the dispatch (except for the last in cases where the global size does not divide cleanly into work-groups) given the combination of the passed ndrange and functor.

fun specifies the functor representing the kernel code that would be enqueued.

args are the arguments that will be passed to fun when kernel will be enqueued with the exception for local_ptr parameters. For local pointers user must supply the size of local memory that will be allocated.

Returns the maximum sub-group size for a functor.

fun specifies the functor representing the kernel code that would be enqueued.

args are the arguments that will be passed to fun when kernel will be enqueued with the exception for local_ptr parameters. For local pointers user must supply the size of local memory that will be allocated.

Returns a valid local size that would produce the requested number of sub-groups such that each sub-group is complete with no partial sub-groups.

fun specifies the functor representing the kernel code that would be enqueued.

args are the arguments that will be passed to fun when kernel will be enqueued with the exception for local_ptr parameters. For local pointers user must supply the size of local memory that will be allocated.

Provides a mechanism to query the maximum number of sub-groups that can be used to execute the passed functor on the current device.

fun specifies the functor representing the kernel code that would be enqueued.

args are the arguments that will be passed to fun when kernel will be enqueued with the exception for local_ptr parameters. For local pointers user must supply the size of local memory that will be allocated.

Functors that are enqueued cannot have any reference and pointer fields, nor can have fields of marker types. If object containing such fields is enqueued the behavior is undefined.

The same restrictions apply to capture-list variables in lambda.

Events can be used to identify commands enqueued to a command-queue from the host. These events created by the OpenCL runtime can only be used on the host i.e. as events passed in event_wait_list argument to various clEnqueue* APIs or runtime APIs that take events as arguments such as clRetainEvent, clReleaseEvent, clGetEventProfilingInfo.

Similarly, events can be used to identify commands enqueued to a device queue (from a kernel). These event objects cannot be passed to the host or used by OpenCL runtime APIs such as the clEnqueue* APIs or runtime APIs that take event arguments.

clRetainEvent and clReleaseEvent will return CL_INVALID_OPERATION if event specified is an event that refers to any kernel enqueued to a device queue using enqueue_kernel or enqueue_marker or is a user event created by make_user_event.

Similarly, clSetUserEventStatus can only be used to set the execution status of events created using clCreateUserEvent. User events created on the device can be set using event::set_status method.

Example of using the explicit local pointer class with enqueue_kernel method.

Example of enqueuing a functor to a device_queue object.

Returns the number of dimensions in use. This is the value given to the work_dim argument specified in clEnqueueNDRangeKernel.

Returns the number of global work-items specified for dimension identified by dimindx. This value is given by the global_work_size argument to clEnqueueNDRangeKernel. Valid values of dimindx are 0 to get_work_dim()-1. For other values of dimindx, get_global_size() returns 1.

Returns the unique global work-item ID value for dimension identified by dimindx. The global work-item ID specifies the work-item ID based on the number of global work-items specified to execute the kernel. Valid values of dimindx are 0 to get_work_dim()-1. For other values of dimindx, get_global_id() returns 0.

Returns the number of local work-items specified in dimension identified by dimindx. This value is at most the value given by the local_work_size argument to clEnqueueNDRangeKernel if local_work_size is not NULL; otherwise the OpenCL implementation chooses an appropriate local_work_size value which is returned by this function. If the kernel is executed with a non-uniform work-group size [19], calls to this built-in from some work-groups may return different values than calls to this built-in from other work-groups.

Valid values of dimindx are 0 to get_work_dim()-1. For other values of dimindx, get_local_size() returns 1.

Returns the same value as that returned by get_local_size(dimindx) if the kernel is executed with a uniform work-group size.

If the kernel is executed with a non-uniform work-group size, returns the number of local work-items in each of the work-groups that make up the uniform region of the global range in the dimension identified by dimindx. If the local_work_size argument to clEnqueueNDRangeKernel is not NULL, this value will match the value specified in local_work_size[dimindx]. If local_work_size is NULL, this value will match the local size that the implementation determined would be most efficient at implementing the uniform region of the global range.

Valid values of dimindx are 0 to get_work_dim()-1. For other values of dimindx, get_enqueued_local_size() returns 1.

Returns the unique local work-item ID i.e. a work-item within a specific work-group for dimension identified by dimindx. Valid values of dimindx are 0 to get_work_dim()-1. For other values of dimindx, get_local_id() returns 0.

Returns the number of work-groups that will execute a kernel for dimension identified by dimindx. Valid values of dimindx are 0 to get_work_dim()-1. For other values of dimindx, get_num_groups() returns 1.

get_group_id returns the work-group ID which is a number from 0 …​ get_num_groups(dimindx)-1. Valid values of dimindx are 0 to get_work_dim()-1. For other values, get_group_id() returns 0.

get_global_offset returns the offset values specified in global_work_offset argument to clEnqueueNDRangeKernel. Valid values of dimindx are 0 to get_work_dim()-1. For other values, get_global_offset() returns 0.

Returns the work-items 1-dimensional global ID.

For 1D work-groups, it is computed as:

For 2D work-groups, it is computed as:

For 3D work-groups, it is computed as:

Returns the work-items 1-dimensional local ID.

For 1D work-groups, it is the same value as:

For 2D work-groups, it is computed as:

For 3D work-groups, it is computed as:

Returns the number of work-items in the sub-group. This value is no more than the maximum sub-group size and is implementation-defined based on a combination of the compiled kernel and the dispatch dimensions. This will be a constant value for the lifetime of the sub-group.

Returns the maximum size of a sub-group within the dispatch. This value will be invariant for a given set of dispatch dimensions and a kernel object compiled for a given device.

Returns the number of sub-groups that the current work-group is divided into.

This number will be constant for the duration of a work-group’s execution. If the kernel is executed with a non-uniform work-group size [17] values for any dimension, calls to this built-in from some work-groups may return different values than calls to this built-in from other work-groups.

Returns the same value as that returned by get_num_sub_groups() if the kernel is executed with a uniform work-group size.

If the kernel is executed with a non-uniform work-group size, returns the number of sub groups in each of the work groups that make up the uniform region of the global range.

get_sub_group_id() returns the sub-group ID which is a number from 0 …​ get_num_sub_groups()-1.

For clEnqueueTask, this returns 0.

Returns the unique work-item ID within the current sub-group. The mapping from get_local_id(dimindx) to get_sub_group_local_id() will be invariant for the lifetime of the work-group.

Evaluates predicate for all work-items in the work-group and returns true if predicate evaluates to true for all work-items in the work-group.

Evaluates predicate for all work-items in the work-group and returns true if predicate evaluates to true for any work-items in the work-group.

Evaluates predicate for all work-items in the sub-group and returns true value if predicate evaluates to true for all work-items in the sub-group.

Evaluates predicate for all work-items in the sub-group and returns true value if predicate evaluates to true for any work-items in the sub-group.

Example:

In this case work_group_all would return true for all work-items in work-group if all elements in p, in range specified by work-group’s size, are true.

One could achieve similar result by using analogical call to work_group_any:

Broadcast the value of a for work-item identified by local_id to all work-items in the work-group.

local_id must be the same value for all work-items in the work-group.

Broadcast the value of a for work-item identified by sub_group_local_id (value returned by get_sub_group_local_id) to all work-items in the sub-group.

sub_group_local_id must be the same value for all work-items in the sub-group.

Example:

Here we are broadcasting value passed to work_group_broadcast function by work-item with local_id = 0 (which is p[0]). This function will return p[0] for all callers. Please note that local_id must be the same for all work-items, therefore something like this is invalid:

Return result of reduction operation specified by op for all values of x specified by work-items in a work-group.

Do an exclusive scan operation specified by op of all values specified by work-items in the work-group. The scan results are returned for each work-item.

The scan order is defined by increasing 1D linear global ID within the work-group.

Do an inclusive scan operation specified by op of all values specified by work-items in the work-group. The scan results are returned for each work-item.

The scan order is defined by increasing 1D linear global ID within the work-group.

Return result of reduction operation specified by op for all values of x specified by work-items in a sub-group.

Do an exclusive scan operation specified by op of all values specified by work-items in a sub-group. The scan results are returned for each work-item.

The scan order is defined by increasing 1D linear global ID within the sub-group.

Do an inclusive scan operation specified by op of all values specified by work-items in a sub-group. The scan results are returned for each work-item.

The scan order is defined by increasing 1D linear global ID within the sub-group.

The inclusive scan operation takes a binary operator op with an identity I and n (where n is the size of the work-group) elements [a₀, a₁, …​ a_n-1] and returns [a₀, (a₀ op a₁), …​ (a₀ op a₁ op …​ op a_n-1)]. If op is work_group_op::add, the identity I is 0. If op is work_group_op::min, the identity I is INT_MAX, UINT_MAX, LONG_MAX, ULONG_MAX, for int, uint, long, ulong types and is +INF for floating-point types. Similarly if op is work_group_op::max, the identity I is INT_MIN, 0, LONG_MIN, 0 and -INF.

Consider the following example:

For the example above, let’s assume that the work-group size is 8 and p points to the following elements [3 1 7 0 4 1 6 3]. Work-item 0 calls work_group_scan_inclusive<work_group_op::add> with 3 and returns 3. Work-item 1 calls work_group_scan_inclusive<work_group_op::add> with 1 and returns 4. The full set of values returned by work_group_scan_inclusive<work_group_op::add> for work-items 0 …​ 7 is [3 4 11 11 15 16 22 25].

The exclusive scan operation takes a binary associative operator op with an identity I and n (where n is the size of the work-group) elements [a₀, a₁, …​ a_n-1] and returns [I, a₀, (a₀ op a₁), …​ (a₀ op a₁ op …​ op a_n-2)]. For the example above, the exclusive scan add operation on the ordered set [3 1 7 0 4 1 6 3] would return [0 3 4 11 11 15 16 22].

All work-items in a work-group executing the kernel on a processor must execute this function before any are allowed to continue execution beyond the work_group_barrier. This function must be encountered by all work-items in a work-group executing the kernel. These rules apply to ND-ranges implemented with uniform and non-uniform work-groups.

If work_group_barrier is inside a conditional statement, then all work-items must enter the conditional if any work-item enters the conditional statement and executes the work_group_barrier.

If work_group_barrier is inside a loop, all work-items must execute the work_group_barrier for each iteration of the loop before any are allowed to continue execution beyond the work_group_barrier.

The scope argument specifies whether the memory accesses of work-items in the work-group to memory address space(s) identified by flags become visible to all work-items in the work-group, the device or all SVM devices.

The work_group_barrier function can also be used to specify which memory operations i.e. to global memory, local memory or images become visible to the appropriate memory scope identified by scope. The flags argument specifies the memory address spaces. This is a bitfield and can be set to 0 or a combination of the following values ORed together. When these flags are ORed together the work_group_barrier acts as a combined barrier for all address spaces specified by the flags ordering memory accesses both within and across the specified address spaces.

mem_fence::image cannot be used together with mem_fence::local or mem_fence::global.

The values of flags and scope must be the same for all work-items in the work-group.

All work-items in a sub-group executing the kernel on a processor must execute this function before any are allowed to continue execution beyond the sub-group barrier. This function must be encountered by all work-items in a sub-group executing the kernel. These rules apply to ND-ranges implemented with uniform and non-uniform work-groups.

If sub_group_barrier is inside a conditional statement, then all work-items within the sub-group must enter the conditional if any work-item in the sub-group enters the conditional statement and executes the sub_group_barrier.

If sub_group_barrier is inside a loop, all work-items within the sub-group must execute the sub_group_barrier for each iteration of the loop before any are allowed to continue execution beyond the sub_group_barrier.

The sub_group_barrier function also queues a memory fence (reads and writes) to ensure correct ordering of memory operations to local or global memory.

The flags argument specifies the memory address spaces. This is a bitfield and can be set to 0 or a combination of the following values ORed together. When these flags are ORed together the sub_group_barrier acts as a combined barrier for all address spaces specified by the flags ordering memory accesses both within and across the specified address spaces.

mem_fence::local - The sub_group_barrier function will ensure that all local memory accesses become visible to all work-items in the sub-group. Note that the value of scope is ignored as the memory scope is always memory_scope_work_group.

mem_fence::global - The sub_group_barrier function ensure that all global memory accesses become visible to the appropriate scope as given by scope.

mem_fence::image - The sub_group_barrier function will ensure that all image memory accesses become visible to the appropriate scope as given by scope. The value of scope must be memory_scope_work_group or memory_scope_device.

mem_fence::image cannot be used together with mem_fence::local or mem_fence::global.

The values of flags and scope must be the same for all work-items in the sub-group.

Initialize a new named barrier object to synchronize sub_group_count sub-groups in the current work-group. Construction of a named-barrier object is a work-group operation and hence must be called uniformly across the work-group. sub_group_count must be uniform across the work-group.

Named barrier objects can be reconstructed and assigned to underlying entities by the compiler, or reused. Reused barriers will always be the same size and act in phases such that when each participating sub-group has waited the wait count is set to 0 and the entire process can start again. The internal wait count will cycle through the range from 0 to sub_group_count in each phase of use of the barrier.

Named barrier objects can only be constructed within kernels, not within arbitrary functions.

All work-items in a sub-group executing the kernel on a processor must execute this method before any are allowed to continue execution beyond the barrier. This function must be encountered by all work-items in a sub-group executing the kernel.

These rules apply to ND-ranges implemented with uniform and non-uniform work-groups.

If wait is called inside a conditional statement, then all work-items within the sub-group must enter the conditional if any work-item in the sub-group enters the conditional statement and executes the call to wait.

If wait is called inside a loop, all work-items within the sub-group must execute the wait operation for each iteration of the loop before any are allowed to continue execution beyond the call to wait. The wait function causes the entire sub-group to wait until sub_group_count sub-groups have waited on the named barrier, where sub_group_count is the initialization value passed to the call to the constructor of the named barrier. Once the wait count equals sub_group_count, any sub-groups waiting at the named barrier will be released and the barrier’s wait count reset to 0.

The flags argument specifies the memory address spaces. This is a bitfield and can be set to 0 or a combination of the following values ORed together. When these flags are ORed together the wait acts as a combined wait operation for all address spaces specified by the flags ordering memory accesses both within and across the specified address spaces.

mem_fence::local - The wait function will ensure that all local memory accesses become visible to all work-items in the sub-group. Note that the value of scope is ignored as the memory scope is always memory_scope_work_group.

mem_fence::global - The wait function ensure that all global memory accesses become visible to the appropriate scope as given by scope.

mem_fence::image cannot be specified as a flag for this function.

The values of flags and scope must be the same for all work-items in the sub-group.

The below example shows how to use the named barriers:

Returns fmin(fmax(x, minval), maxval).

Results are undefined if minval > maxval.

Converts radians to degrees, i.e. (180 / π) * radians.

Returns y if x < y, otherwise it returns x. If x or y are infinite or NaN, the return values are undefined.

Returns y if y < x, otherwise it returns x. If x or y are infinite or NaN, the return values are undefined.

Returns the linear blend of x and y implemented as:

x + (y - x) * a

a must be a value in the range 0.0 …​ 1.0. If a is not in the range 0.0 …​ 1.0, the return values are undefined.

Converts degrees to radians, i.e. (π / 180) * degrees.

Returns 0.0 if x < edge, otherwise it returns 1.0.

Returns 0.0 if x <= edge0 and 1.0 if x >= edge1 and performs smooth Hermite interpolation between 0 and 1 when edge0 < x < edge1. This is useful in cases where you would want a threshold function with a smooth transition.

This is equivalent to:

Results are undefined if edge0 >= edge1 or if x, edge0 or edge1 is a NaN.

Returns 1.0 if x > 0, -0.0 if x = -0.0, +0.0 if x = 0.0, or -1.0 if x < 0. Returns 0.0 if x is a NaN.

Returns the cross product of p0.xyz and p1.xyz. The w component of float4 result returned will be 0.0.

Compute dot product.

Returns the distance between p0 and p1.

This is calculated as length(p0 - p1).

Return the length of vector p, i.e., \(\sqrt{p.x^2 + p.y^2 + \ldots}\)

Returns a vector in the same direction as p but with a length of 1.

Returns fast_length(p0 - p1).

Returns the length of vector p computed as:

Returns a vector in the same direction as p but with a length of 1. fast_normalize is computed as:

The result shall be within 8192 ulps error from the infinitely precise result of

with the following exceptions: