The OpenCL^™ C Specification

The following table describes the list of built-in scalar data types.

If the double-precision floating-point extension {cl_khr_fp64} or the __opencl_c_fp64 feature is not supported, implementations may implicitly cast double-precision floating-point literals to single-precision literals. The use of double-precision literals without double-precision support should result in a diagnostic.

Most built-in scalar 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 scalar data type in the OpenCL C programming language and the corresponding data type available to the application:

The char, unsigned char, short, unsigned short, int, unsigned int, long, unsigned long, float and double vector data types are supported. ^[6] The vector data type is defined with the type name, i.e. char, uchar, short, ushort, int, uint, long, ulong, 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 following table describes the list of built-in 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 following table describes the list of additional data types supported by OpenCL.

The C99 derived types (arrays, structs, unions, functions, and pointers), constructed from the built-in scalar, vector, and other data types are supported, with specified restrictions.

The following tables describe the other built-in data types in OpenCL described in Other Built-in Data Types and the corresponding data type available to the application:

The data type names described in the following table are reserved and cannot be used by applications as type names. The vector data type names defined in Built-in Vector Data Types, but where n is any value other than 2, 3, 4, 8 and 16, are also reserved.

A data item declared to be a data type in memory is always aligned to the size of the data type in bytes. For example, a float4 variable will be aligned to a 16-byte boundary, a char2 variable will be aligned to a 2-byte boundary.

For 3-component vector data types, the size of the data type is 4 * sizeof(component). This means that a 3-component vector data type will be aligned to a 4 * sizeof(component) boundary. The vload3 and vstore3 built-in functions can be used to read and write, respectively, 3-component vector data types from an array of packed scalar data type.

A built-in data type that is not a power of two bytes in size must be aligned to the next larger power of two. This rule applies to built-in types only, not structs or unions.

The OpenCL compiler is responsible for aligning data items to the appropriate alignment as required by the data type. For arguments to a __kernel function declared to be a pointer to a data type, the OpenCL compiler can assume that the pointee is always appropriately aligned as required by the data type. The behavior of an unaligned load or store is undefined, except for the vector data load and store functions vloadn, vload_halfn, vstoren, and vstore_halfn. The vector load functions can read a vector from an address aligned to the element type of the vector. The vector store functions can write a vector to an address aligned to the element type of the vector.

The arithmetic operators add (+), subtract (-), multiply (*) and divide (/) operate on built-in integer and floating-point scalar, and vector data types. The arithmetic operator remainder (%) operates on built-in integer scalar and integer vector data types. All arithmetic operators return result of the same built-in type (integer or floating-point) as the type of the operands, after operand type conversion. After conversion, the following cases are valid:

All other cases of implicit conversions are illegal. Division on integer types which results in a value that lies outside of the range bounded by the maximum and minimum representable values of the integer type will not cause an exception but will result in an unspecified value. A divide by zero with integer types does not cause an exception but will result in an unspecified value. Division by zero for floating-point types will result in ±∞ or NaN as prescribed by the IEEE-754 standard. Use the built-in functions dot and cross to get, respectively, the vector dot product and the vector cross product.

The arithmetic unary operators (+ and -) operate on built-in scalar and vector types.

The arithmetic post- and pre-increment and decrement operators (-- and ++) operate on built-in scalar and vector types except the built-in scalar and vector float types ^[18]. All unary operators work component-wise on their operands. These result with the same type they operated on. For post- and pre-increment and decrement, the expression must be one that could be assigned to (an l-value). Pre-increment and pre-decrement add or subtract 1 to the contents of the expression they operate on, and the value of the pre-increment or pre-decrement expression is the resulting value of that modification. Post-increment and post-decrement expressions add or subtract 1 to the contents of the expression they operate on, but the resulting expression has the expression’s value before the post-increment or post-decrement was executed.

The relational operators greater than (>), less than (<), greater than or equal (>=), and less than or equal (<=) operate on scalar and vector types ^[19]. All relational operators result in an integer type. After operand type conversion, the following cases are valid:

All other cases of implicit conversions are illegal.

The result is a scalar signed integer of type int if the source operands are scalar and a vector signed integer type of the same size as the source operands if the source operands are vector types. Vector source operands of type charn and ucharn return a charn result; vector source operands of type shortn and ushortn return a shortn result; vector source operands of type intn, uintn and floatn return an intn result; vector source operands of type longn, ulongn and doublen return a longn result. For scalar types, the relational operators shall return 0 if the specified relation is false and 1 if the specified relation is true. For vector types, the relational operators shall return 0 if the specified relation is false and -1 (i.e. all bits set) if the specified relation is true. The relational operators always return 0 if either argument is not a number (NaN).

The equality operators equal (==) and not equal (!=) operate on built-in scalar and vector types ^[20]. All equality operators result in an integer type. After operand type conversion, the following cases are valid:

All other cases of implicit conversions are illegal.

The result is a scalar signed integer of type int if the source operands are scalar and a vector signed integer type of the same size as the source operands if the source operands are vector types. Vector source operands of type charn and ucharn return a charn result; vector source operands of type shortn and ushortn return a shortn result; vector source operands of type intn, uintn and floatn return an intn result; vector source operands of type longn, ulongn and doublen return a longn result.

For scalar types, the equality operators return 0 if the specified relation is false and return 1 if the specified relation is true. For vector types, the equality operators shall return 0 if the specified relation is false and -1 (i.e. all bits set) if the specified relation is true. The equality operator equal (==) returns 0 if one or both arguments are not a number (NaN). The equality operator not equal (!=) returns 1 (for scalar source operands) or -1 (for vector source operands) if one or both arguments are not a number (NaN).

The bitwise operators and (&), or (|), exclusive or (^), and not (~) operate on all scalar and vector built-in types except the built-in scalar and vector float types. For vector built-in types, the operators are applied component-wise. If one operand is a scalar and the other is a vector, the scalar may be subject to the usual arithmetic conversion to the element type used by the vector operand. The scalar type is then widened to a vector that has the same number of components as the vector operand. The operation is done component-wise resulting in the same size vector.

The logical operators and (&&) and or (||) operate on all scalar and vector built-in types. For scalar built-in types only, and (&&) will only evaluate the right hand operand if the left hand operand compares unequal to 0. For scalar built-in types only, or (||) will only evaluate the right hand operand if the left hand operand compares equal to 0. For built-in vector types, both operands are evaluated and the operators are applied component-wise. If one operand is a scalar and the other is a vector, the scalar may be subject to the usual arithmetic conversion to the element type used by the vector operand. The scalar type is then widened to a vector that has the same number of components as the vector operand. The operation is done component-wise resulting in the same size vector.

The logical operator exclusive or (^^) is reserved.

The result is a scalar signed integer of type int if the source operands are scalar and a vector signed integer type of the same size as the source operands if the source operands are vector types. Vector source operands of type charn and ucharn return a charn result; vector source operands of type shortn and ushortn return a shortn result; vector source operands of type intn, uintn and floatn return an intn result; vector source operands of type longn, ulongn and doublen return a longn result.

For scalar types, the logical operators shall return 0 if the result of the operation is false and 1 if the result is true. For vector types, the logical operators shall return 0 if the result of the operation is false and -1 (i.e. all bits set) if the result is true.

The logical unary operator not (!) operates on all scalar and vector built-in types. For built-in vector types, the operators are applied component-wise.

The result is a scalar signed integer of type int if the source operands are scalar and a vector signed integer type of the same size as the source operands if the source operands are vector types. Vector source operands of type charn and ucharn return a charn result; vector source operands of type shortn and ushortn return a shortn result; vector source operands of type intn, uintn and floatn return an intn result; vector source operands of type longn, ulongn and doublen return a longn result.

For scalar types, the result of the logical unary operator is 0 if the value of its operand compares unequal to 0, and 1 if the value of its operand compares equal to 0. For vector types, the unary operator shall return a 0 if the value of its operand compares unequal to 0, and -1 (i.e. all bits set) if the value of its operand compares equal to 0.

The ternary selection operator (?:) operates on three expressions (exp1 ? exp2 : exp3). This operator evaluates the first expression exp1, which can be a scalar or vector result except float. If all three expressions are scalar values, the C99 rules for ternary operator are followed. If the result is a vector value, then this is equivalent to calling select(exp3, exp2, exp1). The select function is described in Built-in Scalar and Vector Relational Functions. The second and third expressions can be any type, as long their types match, or there is an implicit conversion that can be applied to one of the expressions to make their types match, or one is a vector and the other is a scalar and the scalar may be subject to the usual arithmetic conversion to the element type used by the vector operand and widened to the same type as the vector type. This resulting matching type is the type of the entire expression.

The operators right-shift (>>), left-shift (<<) operate on all scalar and vector built-in types except the built-in scalar and vector float types. For built-in vector types, the operators are applied component-wise. For the right-shift (>>), left-shift (<<) operators, the rightmost operand must be a scalar if the first operand is a scalar, and the rightmost operand can be a vector or scalar if the first operand is a vector.

The result of E1 << E2 is E1 left-shifted by log₂(N) least significant bits in E2 viewed as an unsigned integer value, where N is the number of bits used to represent the data type of E1 after integer promotion ^[21], if E1 is a scalar, or the number of bits used to represent the type of E1 elements, if E1 is a vector. The vacated bits are filled with zeros.

The result of E1 >> E2 is E1 right-shifted by log₂(N) least significant bits in E2 viewed as an unsigned integer value, where N is the number of bits used to represent the data type of E1 after integer promotion, if E1 is a scalar, or the number of bits used to represent the type of E1 elements, if E1 is a vector. If E1 has an unsigned type or if E1 has a signed type and a nonnegative value, the vacated bits are filled with zeros. If E1 has a signed type and a negative value, the vacated bits are filled with ones.

The sizeof operator yields the size (in bytes) of its operand, including any padding bytes needed for alignment, which may be an expression or the parenthesized name of a type. The size is determined from the type of the operand. The result is of type size_t. If the type of the operand is a variable length array ^[22] type, the operand is evaluated; otherwise, the operand is not evaluated and the result is an integer constant.

When applied to an operand that has type char or uchar, the result is 1. When applied to an operand that has type short, ushort, or half the result is 2. When applied to an operand that has type int, uint or float, the result is 4. When applied to an operand that has type long, ulong or double, the result is 8. When applied to an operand that is a vector type, the result is the number of components times the size of each scalar component ^[23]. When applied to an operand that has array type, the result is the total number of bytes in the array. When applied to an operand that has structure or union type, the result is the total number of bytes in such an object, including internal and trailing padding. The sizeof operator shall not be applied to an expression that has function type or an incomplete type, to the parenthesized name of such a type, or to an expression that designates a bit-field struct member ^[24].

The behavior of applying the sizeof operator to the bool, image2d_t, image3d_t, image2d_array_t, image1d_t, image1d_buffer_t, image1d_array_t, image2d_depth_t, image2d_array_depth_t, sampler_t, queue_t, ndrange_t, clk_event_t, reserve_id_t, and event_t types is implementation-defined. Additionally, the behavior of applying the sizeof operator to a pipe object (a type with the pipe type specifier keyword) is implementation-defined.

The comma (,) operator operates on expressions by returning the type and value of the right-most expression in a comma separated list of expressions. All expressions are evaluated, in order, from left to right.

The unary (*) operator denotes indirection. If the operand points to an object, the result is an l-value designating the object. If the operand has type "pointer to type", the result has type "type". If an invalid value has been assigned to the pointer, the behavior of the unary * operator is undefined ^[25].

The unary (&) operator returns the address of its operand. If the operand has type "type", the result has type "pointer to type". If the operand is the result of a unary * operator, neither that operator nor the & operator is evaluated and the result is as if both were omitted, except that the constraints on the operators still apply and the result is not an l-value. Similarly, if the operand is the result of a [] operator, neither the & operator nor the unary * that is implied by the [] is evaluated and the result is as if the & operator were removed and the [] operator were changed to a + operator. Otherwise, the result is a pointer to the object designated by its operand ^[26].

Assignments of values to variable names are done with the assignment operator (=), like

The assignment operator stores the value of expression into lvalue. The expression and lvalue must have the same type, or the expression must have a type in Built-in Scalar Data Types, in which case an implicit conversion will be done on the expression before the assignment is done.

If expression is a scalar type and lvalue is a vector type, the scalar is converted to the element type used by the vector operand. The scalar type is then widened to a vector that has the same number of components as the vector operand. The operation is done component-wise resulting in the same size vector.

Any other desired type-conversions must be specified explicitly. L-values must be writable. Variables that are built-in types, entire structures or arrays, structure fields, l-values with the field selector (.) applied to select components or swizzles without repeated fields, l-values within parentheses, and l-values dereferenced with the array subscript operator ([]) are all l-values. Other binary or unary expressions, function names, swizzles with repeated fields, and constants cannot be l-values. The ternary operator (?:) is also not allowed as an l-value.

The order of evaluation of the operands is unspecified. If an attempt is made to modify the result of an assignment operator or to access it after the next sequence point, the behavior is undefined. Other assignment operators are the assignments add into (+=), subtract from (-=), multiply into (=), divide into (/=), modulus into (%=), left shift by (<<=), right shift by (>>=), and into (&=), inclusive or into (|=), and exclusive or into (^=).

The expression

is equivalent to

and the lvalue and expression must satisfy the requirements for both operator op and assignment (=).

The __global or global address space name is used to refer to memory objects (buffer or image objects) allocated from the global memory pool.

A buffer memory object can be declared as a pointer to a scalar, vector or user-defined struct. This allows the kernel to read and/or write any location in the buffer.

The actual size of the memory object is determined when the memory object is allocated via appropriate API calls in the host code.

Examples:

As image objects are always allocated from the global address space, the __global or global qualifier should not be specified for image types. The elements of an image object cannot be directly accessed. Built-in functions to read from and write to an image object are provided.

Variables at program scope or static or extern variables inside functions can be declared in global address space if the __opencl_c_program_scope_global_variables feature is supported. These variables in the global address space have the same lifetime as the program, and their values persist between calls to any of the kernels in the program. They are not shared across devices and have distinct storage.

The __local or local address space name is used to describe variables that are allocated in local memory and shared by all work-items in a work-group.

Examples:

The __constant or constant address space name is used to describe read-only variables that are accessible globally. They may be declared in program scope or in the outermost kernel scope or inside functions with a static or extern storage class specifier. Such variables can be accessed by all work-items or by different kernels during the program execution.

It is illegal to write to a variable in the constant address space and will result in a compilation error.

Example:

Implementations are not required to aggregate these declarations into the fewest number of constant arguments. This behavior is implementation defined.

Thus portable code must conservatively assume that each variable declared inside a function or in program scope allocated in the __constant address space counts as a separate constant argument.

The private address space is a memory segment that can only be accessed by one work item. Variables that are not shareable among work items are allocated in the private address space, and it is the default address space for most variables, in particular variables with automatic storage duration.

Example:

The generic address space requires support for OpenCL C 2.0 or OpenCL C 3.0 with the __opencl_c_generic_address_space feature. It can be used with pointer types and it represents a placeholder for any of the named address spaces - global, local or private. It signals that a pointer points to an object in one of these concrete named address spaces. The exact address space resolution can occur dynamically during the kernel execution.

This section describes use of address space qualifiers with respect to declaration scopes or variable types.

Local variables inside functions can be qualified by the private address space qualifier.

Variables declared in the outermost compound statement inside the body of the kernel function can be qualified by the local or constant address spaces.

Examples:

Program scope variables or variables with a extern or static storage class specifier:

Examples:

Kernel function arguments declared to be a pointer or an array of a type must point to one of the named address spaces __global, __local or __constant.

Examples:

Program scope and static variables in the __global address space are zero initialized by default. A constant expression may be given as an initializer.

Variables allocated in the __local address space inside a kernel function cannot be initialized.

Variables allocated in the __constant address space are required to be initialized and the values used to initialize these variables must be a compile time constant.

Private address space objects are not initialized by default; any initializer is allowed to be given.

Examples:

Address space qualifiers are not required in many cases. If they are not specified explicitly the default address space will be inferred depending on the declaration scope and the object type.

There is no syntax to provide address space in the source for some situations, therefore only the default address space is applicable.

For OpenCL C 2.0 or with the __opencl_c_program_scope_global_variables feature, the address space for a variable at program scope or a static or extern variable inside a function are inferred to be __global.

If the generic address space is supported i.e. for OpenCL C 2.0 or OpenCL C 3.0 with __opencl_c_generic_address_space feature, pointers that are declared without pointing to a named address space point to the generic address space.

All string literal storage shall be in the __constant address space.

For all other cases that are not listed above the address space is inferred to __private. This includes:

Examples:

OpenCL implements the address space nesting model for pointers from Embedded C, section 5.1.3 as follows:

Following section 5.3 of the Embedded C Specification, it is only allowed to convert pointers implicitly, i.e. in assignments, function parameters, operations, if the original pointer points to an object qualified by an address space enclosed into the address space pointed by the destination pointer.

In contrast to the Embedded C Specification, explicitly converting i.e. casting between pointers to non-overlapping address spaces is illegal in OpenCL.

Considering the above, the following applies to conversions of pointers pointing to different address spaces:

Examples:

This is the canonical example. In this example, function foo is declared with an argument that is a pointer with the unnamed generic address space address space qualifier.

In the example below, var is a pointer to the unnamed generic address space. A pointer to the global or local address space may be assigned to var depending on the result of a conditional expression.

In the example below, the same pointer to the unnamed generic address space is used to point to objects allocated in different named address spaces. A pointer to the unnamed generic address space may point to objects in the global, local, and private address spaces, but it is not legal for a pointer to the unnamed generic address to point to an object in the constant address space.

In the example below, pointers to named address spaces are assigned to a pointer to the unnamed generic address space. It is legal to assign a pointer to the global, local, and private address spaces to a pointer to the unnamed generic address space without an explicit cast. It is not legal to assign a pointer to the constant address space to a pointer to the unnamed generic address space. It is also not legal to assign a pointer to the unnamed generic address space to a pointer to a named address space without a cast.

The example below illustrates the implicit conversion between named address spaces.

The example below demonstrates explicit conversions for pointers pointing to different address spaces.

For nested pointers, implicit conversions between address spaces are disallowed. Explicitly casting between different address spaces in nested pointers is allowed but the use of such pointers can lead to incorrect behavior such as accessing invalid memory locations.

Various clarifications and examples illustrating how changes to ISO/IEC 9899:1999 detailed in Embedded C, section 5.3 apply to OpenCL C with the generic address space.

Clause 6.2.5 - Types:

If address space qualifier on type T is omitted refer to Inference.

Clause 6.3.2.3 - Pointers

Conversions between disjoint address spaces are disallowed in OpenCL (Address space conversions).

Clause 6.5.8 - Relational operators:

Examples:

Clause 6.5.9 - Equality operators:

Examples:

Consider the following example:

The behavior of callee is undefined as gptr and pptr are in different address spaces. The example above would have the same undefined behavior if the equality operator is replaced with a relational operator.

Examples:

Clause 6.5.15 - Conditional operator:

Examples:

Clause 6.5.16.1 - Simple assignment:

Examples:

Clause 6.7.3 - Type qualifiers

The type of an object with automatic storage duration are in private address space and therefore can be qualified with private/__private.

The __kernel (or kernel) qualifier declares a function to be a kernel that can be executed by an application on an OpenCL device(s). The following rules apply to functions that are declared with this qualifier:

The __kernel and kernel names are reserved for use as functions qualifiers and shall not be used otherwise.

The __kernel qualifier can be used with the keyword attribute to declare additional information about the kernel function as described below.

The optional __attribute__((vec_type_hint(<type>))) ^[27] 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 __attribute__((vec_type_hint(<type>))) qualifier <type> is one of the built-in vector types listed in Built-in Vector Data Types or the constituent scalar element types. If vec_type_hint (<type>) is not specified, the kernel is assumed to have the __attribute__((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:

If for example, a __kernel function is declared with

(meaning that most operations in the __kernel function are explicitly vectorized using float4) and the kernel is running using Intel^® Advanced Vector Instructions (Intel^® AVX) which implements a 8-float-wide vector unit, the autovectorizer might choose to merge two work-items to one thread, running a second work-item in the high half of the 256-bit AVX register.

As another example, a Power4 machine has two scalar double precision floating-point units with an 6-cycle deep pipe. An autovectorizer for the Power4 machine might choose to interleave six kernels declared with the __attribute__(( vec_type_hint (double2))) qualifier into one hardware thread, to ensure that there is always 12-way parallelism available to saturate the FPUs. It might also choose to merge 4 or 8 work-items (or some other number) if it concludes that these are better choices, due to resource utilization concerns or some preference for divisibility by 2.

The optional __attribute__((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 __attribute__((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 optional __attribute__((reqd_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 keyword __attribute__ allows you to specify special attributes of enum, struct and union types when you define such types. This keyword is followed by an attribute specification inside double parentheses. Two attributes are currently defined for types: aligned, and packed.

You may specify type attributes in an enum, struct or union type declaration or definition, or for other types in a typedef declaration.

For an enum, struct or union type, you may specify attributes either between the enum, struct or union tag and the name of the type, or just past the closing curly brace of the definition. The former syntax is preferred.

The keyword __attribute__ allows you to specify special attributes of variables or structure fields. This keyword is followed by an attribute specification inside double parentheses. The following attribute qualifiers are currently defined:

For basic blocks and control-flow-statements the attribute is placed before the structure in question, for example:

Here attr1 applies to the block in braces and attr2 and attr3 apply to the loop’s control construct and body, respectively.

No attribute qualifiers for blocks and control-flow-statements are currently defined.

The __attribute__((opencl_unroll_hint)) and __attribute__((opencl_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:

The 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 __attribute__((opencl_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 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.