Description

  1. The use of pointers is somewhat restricted. The following rules apply:

    • Arguments to kernel functions declared in a program that are pointers must be declared with the __global, __constant or __local qualifier.

    • A pointer declared with the __constant qualifier can only be assigned to a pointer declared with the __constant qualifier respectively.

    • Pointers to functions are not allowed.

    • Arguments to kernel functions in a program cannot be declared as a pointer to a pointer(s). Variables inside a function or arguments to non-kernel functions in a program can be declared as a pointer to a pointer(s). This restriction only applies to OpenCL C 1.2 or below.

  2. An image type (image2d_t, image3d_t, image2d_array_t, image1d_t, image1d_buffer_t or image1d_array_t) can only be used as the type of a function argument. An image function argument cannot be modified. Elements of an image can only be accessed using the built-in image read and write functions.

    An image type cannot be used to declare a variable, a structure or union field, an array of images, a pointer to an image, or the return type of a function. An image type cannot be used with the __global, __private, __local and __constant address space qualifiers.

    The sampler type (sampler_t) can only be used as the type of a function argument or a variable declared in the program scope or the outermost scope of a kernel function. The behavior of a sampler variable declared in a non-outermost scope of a kernel function is implementation-defined. A sampler argument or variable cannot be modified.

    The sampler type cannot be used to declare a structure or union field, an array of samplers, a pointer to a sampler, or the return type of a function. The sampler type cannot be used with the __local and __global address space qualifiers.

  3. Bit-field struct members are currently not supported.

  4. Variable length arrays and structures with flexible (or unsized) arrays are not supported.

  5. Variadic functions are not supported, with the exception of printf and enqueue_kernel.

  6. Variadic macros are not supported. This restriction only applies to OpenCL C 2.0 or below.

  7. If a list of parameters in a function declaration is empty, the function takes no arguments. This is due to the above restriction on variadic functions.

  8. Unless defined in the OpenCL specification, the library functions, macros, types, and constants defined in the C99 standard headers assert.h, ctype.h, complex.h, errno.h, fenv.h, float.h, inttypes.h, limits.h, locale.h, setjmp.h, signal.h, stdarg.h, stdio.h, stdlib.h, string.h, tgmath.h, time.h, wchar.h and wctype.h are not available and cannot be included by a program.

  9. The auto and register storage-class specifiers are not supported.

  10. Predefined identifiers are not supported. This restriction only applies to OpenCL C 1.1 or below.

  11. Recursion is not supported.

  12. The return type of a kernel function must be void.

  13. Arguments to kernel functions in a program cannot be declared with the built-in scalar types bool, size_t, ptrdiff_t, intptr_t, and uintptr_t or a struct and/or union that contain fields declared to be one of these built-in scalar types.

  14. half is not supported as half can be used as a storage format [1] only and is not a data type on which floating-point arithmetic can be performed.

  15. Whether or not irreducible control flow is illegal is implementation defined.

  16. The following restriction only applies to OpenCL C 1.0, and only if the cl_khr_byte_addressable_store extension macro is not supported:
    Built-in types that are less than 32-bits in size, i.e. char, uchar, char2, uchar2, short, ushort, and half, have the following restriction:

    • Writes to a pointer (or arrays) of type char, uchar, char2, uchar2, short, ushort, and half or to elements of a struct that are of type char, uchar, char2, uchar2, short and ushort are not supported. Refer to section 9.9 for additional information.

      The kernel example below shows what memory operations are not supported on built-in types less than 32-bits in size.

      kernel void
do_proc (__global char *pA, short b,
         __global short *pB)
{
    char x[100];
    __private char *px = x;
    int id = (int)get_global_id(0);
    short f;

    f = pB[id] + b; // is allowed
    px[1] = pA[1]; // error. px cannot be written.
    pB[id] = b; // error. pB cannot be written
}

  17. The type qualifiers const, restrict and volatile as defined by the C99 specification are supported. These qualifiers cannot be used with image2d_t, image3d_t, image2d_array_t, image2d_depth_t, image2d_array_depth_t, image1d_t, image1d_buffer_t and image1d_array_t types. Types other than pointer types shall not use the restrict qualifier.

  18. The event type (event_t) cannot be used as the type of a kernel function argument. The event type cannot be used to declare a program scope variable. The event type cannot be used to declare a structure or union field. The event type cannot be used with the __local, __constant and __global address space qualifiers.

  19. The clk_event_t, ndrange_t and reserve_id_t types cannot be used as arguments to kernel functions that get enqueued from the host. The clk_event_t and reserve_id_t types cannot be declared in program scope.

  20. Kernels enqueued by the host must continue to have their arguments that are a pointer to a type declared to point to a named address space.

  21. A function in an OpenCL program cannot be called main.

  22. Implicit function declaration is not supported.

  23. Program scope variables can be defined with any valid OpenCL C data type except for those in Other Built-in Data Types. Such program scope variables may be of any user-defined type, or a pointer to a user-defined type.

    In the presence of shared virtual memory, these pointers or pointer members should work as expected as long as they are shared virtual memory pointers and the referenced storage has been mapped appropriately. Program scope variables can be declared with __constant address space qualifiers or if __opencl_c_program_scope_global_variables feature is supported with __global address space qualifier.

See Also

No cross-references are available

Document Notes

For more information, see the OpenCL C Specification

This page is extracted from the OpenCL C Specification. Fixes and changes should be made to the Specification, not directly.

Copyright

Copyright 2014-2025 The Khronos Group Inc.

SPDX-License-Identifier: CC-BY-4.0

1. Unless the cl_khr_fp16 extension is supported and has been enabled.