AttributeReference.rst

A compute unit (CU) is responsible for executing the wavefronts of a work-group.
It is composed of one or more execution units (EU), which are responsible for
executing the wavefronts. An EU can have enough resources to maintain the state
of more than one executing wavefront. This allows an EU to hide latency by
switching between wavefronts in a similar way to symmetric multithreading on a
CPU. In order to allow the state for multiple wavefronts to fit on an EU, the
resources used by a single wavefront have to be limited. For example, the number
of SGPRs and VGPRs. Limiting such resources can allow greater latency hiding,
but can result in having to spill some register state to memory.

Clang supports the ``__attribute__((amdgpu_waves_per_eu(<min>[, <max>])))``
attribute for the AMDGPU target. This attribute may be attached to a kernel
function definition and is an optimization hint.

``<min>`` parameter specifies the requested minimum number of waves per EU, and
*optional* ``<max>`` parameter specifies the requested maximum number of waves
per EU (must be greater than ``<min>`` if specified). If ``<max>`` is omitted,
then there is no restriction on the maximum number of waves per EU other than
the one dictated by the hardware for which the kernel is compiled. Passing
``0, 0`` as ``<min>, <max>`` implies the default behavior (no limits).

If specified, this attribute allows an advanced developer to tune the number of
wavefronts that are capable of fitting within the resources of an EU. The AMDGPU
target backend can use this information to limit resources, such as number of
SGPRs, number of VGPRs, size of available group and private memory segments, in
such a way that guarantees that at least ``<min>`` wavefronts and at most
``<max>`` wavefronts are able to fit within the resources of an EU. Requesting
more wavefronts can hide memory latency but limits available registers which
can result in spilling. Requesting fewer wavefronts can help reduce cache
thrashing, but can reduce memory latency hiding.

This attribute controls the machine code generated by the AMDGPU target backend
to ensure it is capable of meeting the requested values. However, when the
kernel is executed, there may be other reasons that prevent meeting the request,
for example, there may be wavefronts from other kernels executing on the EU.

An error will be given if:
  - Specified values violate subtarget specifications;
  - Specified values are not compatible with values provided through other
    attributes;
  - The AMDGPU target backend is unable to create machine code that can meet the
    request.


Consumed Annotation Checking
============================
Clang supports additional attributes for checking basic resource management
properties, specifically for unique objects that have a single owning reference.
The following attributes are currently supported, although **the implementation
for these annotations is currently in development and are subject to change.**

callable_when (clang::callable_when)
------------------------------------
.. csv-table:: Supported Syntaxes
   :header: "GNU", "C++11", "C2x", "__declspec", "Keyword", "Pragma", "Pragma clang attribute"

   "X","X","","","", "", "X"

Use ``__attribute__((callable_when(...)))`` to indicate what states a method
may be called in.  Valid states are unconsumed, consumed, or unknown.  Each
argument to this attribute must be a quoted string.  E.g.:

``__attribute__((callable_when("unconsumed", "unknown")))``


consumable (clang::consumable)
------------------------------
.. csv-table:: Supported Syntaxes
   :header: "GNU", "C++11", "C2x", "__declspec", "Keyword", "Pragma", "Pragma clang attribute"

   "X","X","","","", "", "X"

Each ``class`` that uses any of the typestate annotations must first be marked
using the ``consumable`` attribute.  Failure to do so will result in a warning.

This attribute accepts a single parameter that must be one of the following:
``unknown``, ``consumed``, or ``unconsumed``.


param_typestate (clang::param_typestate)
----------------------------------------
.. csv-table:: Supported Syntaxes
   :header: "GNU", "C++11", "C2x", "__declspec", "Keyword", "Pragma", "Pragma clang attribute"

   "X","X","","","", "", "X"

This attribute specifies expectations about function parameters.  Calls to an
function with annotated parameters will issue a warning if the corresponding
argument isn't in the expected state.  The attribute is also used to set the
initial state of the parameter when analyzing the function's body.


return_typestate (clang::return_typestate)
------------------------------------------
.. csv-table:: Supported Syntaxes
   :header: "GNU", "C++11", "C2x", "__declspec", "Keyword", "Pragma", "Pragma clang attribute"

   "X","X","","","", "", "X"

The ``return_typestate`` attribute can be applied to functions or parameters.
When applied to a function the attribute specifies the state of the returned
value.  The function's body is checked to ensure that it always returns a value
in the specified state.  On the caller side, values returned by the annotated
function are initialized to the given state.

When applied to a function parameter it modifies the state of an argument after
a call to the function returns.  The function's body is checked to ensure that
the parameter is in the expected state before returning.


set_typestate (clang::set_typestate)
------------------------------------
.. csv-table:: Supported Syntaxes
   :header: "GNU", "C++11", "C2x", "__declspec", "Keyword", "Pragma", "Pragma clang attribute"

   "X","X","","","", "", "X"

Annotate methods that transition an object into a new state with
``__attribute__((set_typestate(new_state)))``.  The new state must be
unconsumed, consumed, or unknown.


test_typestate (clang::test_typestate)
--------------------------------------
.. csv-table:: Supported Syntaxes
   :header: "GNU", "C++11", "C2x", "__declspec", "Keyword", "Pragma", "Pragma clang attribute"

   "X","X","","","", "", "X"

Use ``__attribute__((test_typestate(tested_state)))`` to indicate that a method
returns true if the object is in the specified state..


Type Safety Checking
====================
Clang supports additional attributes to enable checking type safety properties
that can't be enforced by the C type system. To see warnings produced by these
checks, ensure that -Wtype-safety is enabled. Use cases include:

* MPI library implementations, where these attributes enable checking that
  the buffer type matches the passed ``MPI_Datatype``;
* for HDF5 library there is a similar use case to MPI;
* checking types of variadic functions' arguments for functions like
  ``fcntl()`` and ``ioctl()``.

You can detect support for these attributes with ``__has_attribute()``.  For
example:

.. code-block:: c++

  #if defined(__has_attribute)
  #  if __has_attribute(argument_with_type_tag) && \
        __has_attribute(pointer_with_type_tag) && \
        __has_attribute(type_tag_for_datatype)
  #    define ATTR_MPI_PWT(buffer_idx, type_idx) __attribute__((pointer_with_type_tag(mpi,buffer_idx,type_idx)))
  /* ... other macros ...  */
  #  endif
  #endif

  #if !defined(ATTR_MPI_PWT)
  # define ATTR_MPI_PWT(buffer_idx, type_idx)
  #endif

  int MPI_Send(void *buf, int count, MPI_Datatype datatype /*, other args omitted */)
      ATTR_MPI_PWT(1,3);

argument_with_type_tag
----------------------
.. csv-table:: Supported Syntaxes
   :header: "GNU", "C++11", "C2x", "__declspec", "Keyword", "Pragma", "Pragma clang attribute"

   "X","","","","", "", ""

Use ``__attribute__((argument_with_type_tag(arg_kind, arg_idx,
type_tag_idx)))`` on a function declaration to specify that the function
accepts a type tag that determines the type of some other argument.

This attribute is primarily useful for checking arguments of variadic functions
(``pointer_with_type_tag`` can be used in most non-variadic cases).

In the attribute prototype above:
  * ``arg_kind`` is an identifier that should be used when annotating all
    applicable type tags.
  * ``arg_idx`` provides the position of a function argument. The expected type of
    this function argument will be determined by the function argument specified
    by ``type_tag_idx``. In the code example below, "3" means that the type of the
    function's third argument will be determined by ``type_tag_idx``.
  * ``type_tag_idx`` provides the position of a function argument. This function
    argument will be a type tag. The type tag will determine the expected type of
    the argument specified by ``arg_idx``. In the code example below, "2" means
    that the type tag associated with the function's second argument should agree
    with the type of the argument specified by ``arg_idx``.

For example:

.. code-block:: c++

  int fcntl(int fd, int cmd, ...)
      __attribute__(( argument_with_type_tag(fcntl,3,2) ));
  // The function's second argument will be a type tag; this type tag will
  // determine the expected type of the function's third argument.


pointer_with_type_tag
---------------------
.. csv-table:: Supported Syntaxes
   :header: "GNU", "C++11", "C2x", "__declspec", "Keyword", "Pragma", "Pragma clang attribute"

   "X","","","","", "", ""

Use ``__attribute__((pointer_with_type_tag(ptr_kind, ptr_idx, type_tag_idx)))``
on a function declaration to specify that the function accepts a type tag that
determines the pointee type of some other pointer argument.

In the attribute prototype above:
  * ``ptr_kind`` is an identifier that should be used when annotating all
    applicable type tags.
  * ``ptr_idx`` provides the position of a function argument; this function
    argument will have a pointer type. The expected pointee type of this pointer
    type will be determined by the function argument specified by
    ``type_tag_idx``. In the code example below, "1" means that the pointee type
    of the function's first argument will be determined by ``type_tag_idx``.
  * ``type_tag_idx`` provides the position of a function argument; this function
    argument will be a type tag. The type tag will determine the expected pointee
    type of the pointer argument specified by ``ptr_idx``. In the code example
    below, "3" means that the type tag associated with the function's third
    argument should agree with the pointee type of the pointer argument specified
    by ``ptr_idx``.

For example:

.. code-block:: c++

  typedef int MPI_Datatype;
  int MPI_Send(void *buf, int count, MPI_Datatype datatype /*, other args omitted */)
      __attribute__(( pointer_with_type_tag(mpi,1,3) ));
  // The function's 3rd argument will be a type tag; this type tag will
  // determine the expected pointee type of the function's 1st argument.


type_tag_for_datatype
---------------------
.. csv-table:: Supported Syntaxes
   :header: "GNU", "C++11", "C2x", "__declspec", "Keyword", "Pragma", "Pragma clang attribute"

   "X","","","","", "", ""

When declaring a variable, use
``__attribute__((type_tag_for_datatype(kind, type)))`` to create a type tag that
is tied to the ``type`` argument given to the attribute.

In the attribute prototype above:
  * ``kind`` is an identifier that should be used when annotating all applicable
    type tags.
  * ``type`` indicates the name of the type.

Clang supports annotating type tags of two forms.

  * **Type tag that is a reference to a declared identifier.**
    Use ``__attribute__((type_tag_for_datatype(kind, type)))`` when declaring that
    identifier:

    .. code-block:: c++

      typedef int MPI_Datatype;
      extern struct mpi_datatype mpi_datatype_int
          __attribute__(( type_tag_for_datatype(mpi,int) ));
      #define MPI_INT ((MPI_Datatype) &mpi_datatype_int)
      // &mpi_datatype_int is a type tag. It is tied to type "int".

  * **Type tag that is an integral literal.**
    Declare a ``static const`` variable with an initializer value and attach
    ``__attribute__((type_tag_for_datatype(kind, type)))`` on that declaration:

    .. code-block:: c++

      typedef int MPI_Datatype;
      static const MPI_Datatype mpi_datatype_int
          __attribute__(( type_tag_for_datatype(mpi,int) )) = 42;
      #define MPI_INT ((MPI_Datatype) 42)
      // The number 42 is a type tag. It is tied to type "int".


The ``type_tag_for_datatype`` attribute also accepts an optional third argument
that determines how the type of the function argument specified by either
``arg_idx`` or ``ptr_idx`` is compared against the type associated with the type
tag. (Recall that for the ``argument_with_type_tag`` attribute, the type of the
function argument specified by ``arg_idx`` is compared against the type
associated with the type tag. Also recall that for the ``pointer_with_type_tag``
attribute, the pointee type of the function argument specified by ``ptr_idx`` is
compared against the type associated with the type tag.) There are two supported
values for this optional third argument:

  * ``layout_compatible`` will cause types to be compared according to
    layout-compatibility rules (In C++11 [class.mem] p 17, 18, see the
    layout-compatibility rules for two standard-layout struct types and for two
    standard-layout union types). This is useful when creating a type tag
    associated with a struct or union type. For example:

    .. code-block:: c++

      /* In mpi.h */
      typedef int MPI_Datatype;
      struct internal_mpi_double_int { double d; int i; };
      extern struct mpi_datatype mpi_datatype_double_int
          __attribute__(( type_tag_for_datatype(mpi,
                          struct internal_mpi_double_int, layout_compatible) ));

      #define MPI_DOUBLE_INT ((MPI_Datatype) &mpi_datatype_double_int)

      int MPI_Send(void *buf, int count, MPI_Datatype datatype, ...)
          __attribute__(( pointer_with_type_tag(mpi,1,3) ));

      /* In user code */
      struct my_pair { double a; int b; };
      struct my_pair *buffer;
      MPI_Send(buffer, 1, MPI_DOUBLE_INT /*, ...  */); // no warning because the
                                                       // layout of my_pair is
                                                       // compatible with that of
                                                       // internal_mpi_double_int

      struct my_int_pair { int a; int b; }
      struct my_int_pair *buffer2;
      MPI_Send(buffer2, 1, MPI_DOUBLE_INT /*, ...  */); // warning because the
                                                        // layout of my_int_pair
                                                        // does not match that of
                                                        // internal_mpi_double_int

  * ``must_be_null`` specifies that the function argument specified by either
    ``arg_idx`` (for the ``argument_with_type_tag`` attribute) or ``ptr_idx`` (for
    the ``pointer_with_type_tag`` attribute) should be a null pointer constant.
    The second argument to the ``type_tag_for_datatype`` attribute is ignored. For
    example:

    .. code-block:: c++

      /* In mpi.h */
      typedef int MPI_Datatype;
      extern struct mpi_datatype mpi_datatype_null
          __attribute__(( type_tag_for_datatype(mpi, void, must_be_null) ));

      #define MPI_DATATYPE_NULL ((MPI_Datatype) &mpi_datatype_null)
      int MPI_Send(void *buf, int count, MPI_Datatype datatype, ...)
          __attribute__(( pointer_with_type_tag(mpi,1,3) ));

      /* In user code */
      struct my_pair { double a; int b; };
      struct my_pair *buffer;
      MPI_Send(buffer, 1, MPI_DATATYPE_NULL /*, ...  */); // warning: MPI_DATATYPE_NULL
                                                          // was specified but buffer
                                                          // is not a null pointer


OpenCL Address Spaces
=====================
The address space qualifier may be used to specify the region of memory that is
used to allocate the object. OpenCL supports the following address spaces:
__generic(generic), __global(global), __local(local), __private(private),
__constant(constant).

  .. code-block:: c

    __constant int c = ...;

    __generic int* foo(global int* g) {
      __local int* l;
      private int p;
      ...
      return l;
    }

More details can be found in the OpenCL C language Spec v2.0, Section 6.5.

constant (__constant)
---------------------
.. csv-table:: Supported Syntaxes
   :header: "GNU", "C++11", "C2x", "__declspec", "Keyword", "Pragma", "Pragma clang attribute"

   "","","","","X", "", ""

The constant address space attribute signals that an object is located in
a constant (non-modifiable) memory region. It is available to all work items.
Any type can be annotated with the constant address space attribute. Objects
with the constant address space qualifier can be declared in any scope and must
have an initializer.


generic (__generic)
-------------------
.. csv-table:: Supported Syntaxes
   :header: "GNU", "C++11", "C2x", "__declspec", "Keyword", "Pragma", "Pragma clang attribute"

   "","","","","X", "", ""

The generic address space attribute is only available with OpenCL v2.0 and later.
It can be used with pointer types. Variables in global and local scope and
function parameters in non-kernel functions can have the generic address space
type attribute. It is intended to be a placeholder for any other address space
except for '__constant' in OpenCL code which can be used with multiple address
spaces.


global (__global)
-----------------
.. csv-table:: Supported Syntaxes
   :header: "GNU", "C++11", "C2x", "__declspec", "Keyword", "Pragma", "Pragma clang attribute"

   "","","","","X", "", ""

The global address space attribute specifies that an object is allocated in
global memory, which is accessible by all work items. The content stored in this
memory area persists between kernel executions. Pointer types to the global
address space are allowed as function parameters or local variables. Starting
with OpenCL v2.0, the global address space can be used with global (program
scope) variables and static local variable as well.


local (__local)
---------------
.. csv-table:: Supported Syntaxes
   :header: "GNU", "C++11", "C2x", "__declspec", "Keyword", "Pragma", "Pragma clang attribute"

   "","","","","X", "", ""

The local address space specifies that an object is allocated in the local (work
group) memory area, which is accessible to all work items in the same work
group. The content stored in this memory region is not accessible after
the kernel execution ends. In a kernel function scope, any variable can be in
the local address space. In other scopes, only pointer types to the local address
space are allowed. Local address space variables cannot have an initializer.


private (__private)
-------------------
.. csv-table:: Supported Syntaxes
   :header: "GNU", "C++11", "C2x", "__declspec", "Keyword", "Pragma", "Pragma clang attribute"

   "","","","","X", "", ""

The private address space specifies that an object is allocated in the private
(work item) memory. Other work items cannot access the same memory area and its
content is destroyed after work item execution ends. Local variables can be
declared in the private address space. Function arguments are always in the
private address space. Kernel function arguments of a pointer or an array type
cannot point to the private address space.


Nullability Attributes
======================
Whether a particular pointer may be "null" is an important concern when working with pointers in the C family of languages. The various nullability attributes indicate whether a particular pointer can be null or not, which makes APIs more expressive and can help static analysis tools identify bugs involving null pointers. Clang supports several kinds of nullability attributes: the ``nonnull`` and ``returns_nonnull`` attributes indicate which function or method parameters and result types can never be null, while nullability type qualifiers indicate which pointer types can be null (``_Nullable``) or cannot be null (``_Nonnull``).

The nullability (type) qualifiers express whether a value of a given pointer type can be null (the ``_Nullable`` qualifier), doesn't have a defined meaning for null (the ``_Nonnull`` qualifier), or for which the purpose of null is unclear (the ``_Null_unspecified`` qualifier). Because nullability qualifiers are expressed within the type system, they are more general than the ``nonnull`` and ``returns_nonnull`` attributes, allowing one to express (for example) a nullable pointer to an array of nonnull pointers. Nullability qualifiers are written to the right of the pointer to which they apply. For example:

  .. code-block:: c

    // No meaningful result when 'ptr' is null (here, it happens to be undefined behavior).
    int fetch(int * _Nonnull ptr) { return *ptr; }

    // 'ptr' may be null.
    int fetch_or_zero(int * _Nullable ptr) {
      return ptr ? *ptr : 0;
    }

    // A nullable pointer to non-null pointers to const characters.
    const char *join_strings(const char * _Nonnull * _Nullable strings, unsigned n);

In Objective-C, there is an alternate spelling for the nullability qualifiers that can be used in Objective-C methods and properties using context-sensitive, non-underscored keywords. For example:

  .. code-block:: objective-c

    @interface NSView : NSResponder
      - (nullable NSView *)ancestorSharedWithView:(nonnull NSView *)aView;
      @property (assign, nullable) NSView *superview;
      @property (readonly, nonnull) NSArray *subviews;
    @end

_Nonnull
--------
.. csv-table:: Supported Syntaxes
   :header: "GNU", "C++11", "C2x", "__declspec", "Keyword", "Pragma", "Pragma clang attribute"

   "","","","","X", "", ""

The ``_Nonnull`` nullability qualifier indicates that null is not a meaningful value for a value of the ``_Nonnull`` pointer type. For example, given a declaration such as:

  .. code-block:: c

    int fetch(int * _Nonnull ptr);

a caller of ``fetch`` should not provide a null value, and the compiler will produce a warning if it sees a literal null value passed to ``fetch``. Note that, unlike the declaration attribute ``nonnull``, the presence of ``_Nonnull`` does not imply that passing null is undefined behavior: ``fetch`` is free to consider null undefined behavior or (perhaps for backward-compatibility reasons) defensively handle null.


_Null_unspecified
-----------------
.. csv-table:: Supported Syntaxes
   :header: "GNU", "C++11", "C2x", "__declspec", "Keyword", "Pragma", "Pragma clang attribute"

   "","","","","X", "", ""

The ``_Null_unspecified`` nullability qualifier indicates that neither the ``_Nonnull`` nor ``_Nullable`` qualifiers make sense for a particular pointer type. It is used primarily to indicate that the role of null with specific pointers in a nullability-annotated header is unclear, e.g., due to overly-complex implementations or historical factors with a long-lived API.


_Nullable
---------
.. csv-table:: Supported Syntaxes
   :header: "GNU", "C++11", "C2x", "__declspec", "Keyword", "Pragma", "Pragma clang attribute"

   "","","","","X", "", ""

The ``_Nullable`` nullability qualifier indicates that a value of the ``_Nullable`` pointer type can be null. For example, given:

  .. code-block:: c

    int fetch_or_zero(int * _Nullable ptr);

a caller of ``fetch_or_zero`` can provide null.


nonnull (gnu::nonnull)
----------------------
.. csv-table:: Supported Syntaxes
   :header: "GNU", "C++11", "C2x", "__declspec", "Keyword", "Pragma", "Pragma clang attribute"

   "X","X","","","", "", ""

The ``nonnull`` attribute indicates that some function parameters must not be null, and can be used in several different ways. It's original usage (`from GCC <https://gcc.gnu.org/onlinedocs/gcc/Common-Function-Attributes.html#Common-Function-Attributes>`_) is as a function (or Objective-C method) attribute that specifies which parameters of the function are nonnull in a comma-separated list. For example:

  .. code-block:: c

    extern void * my_memcpy (void *dest, const void *src, size_t len)
                    __attribute__((nonnull (1, 2)));

Here, the ``nonnull`` attribute indicates that parameters 1 and 2
cannot have a null value. Omitting the parenthesized list of parameter indices means that all parameters of pointer type cannot be null:

  .. code-block:: c

    extern void * my_memcpy (void *dest, const void *src, size_t len)
                    __attribute__((nonnull));

Clang also allows the ``nonnull`` attribute to be placed directly on a function (or Objective-C method) parameter, eliminating the need to specify the parameter index ahead of type. For example:

  .. code-block:: c

    extern void * my_memcpy (void *dest __attribute__((nonnull)),
                             const void *src __attribute__((nonnull)), size_t len);

Note that the ``nonnull`` attribute indicates that passing null to a non-null parameter is undefined behavior, which the optimizer may take advantage of to, e.g., remove null checks. The ``_Nonnull`` type qualifier indicates that a pointer cannot be null in a more general manner (because it is part of the type system) and does not imply undefined behavior, making it more widely applicable.


returns_nonnull (gnu::returns_nonnull)
--------------------------------------
.. csv-table:: Supported Syntaxes
   :header: "GNU", "C++11", "C2x", "__declspec", "Keyword", "Pragma", "Pragma clang attribute"

   "X","X","","","", "", "X"

The ``returns_nonnull`` attribute indicates that a particular function (or Objective-C method) always returns a non-null pointer. For example, a particular system ``malloc`` might be defined to terminate a process when memory is not available rather than returning a null pointer:

  .. code-block:: c

    extern void * malloc (size_t size) __attribute__((returns_nonnull));

The ``returns_nonnull`` attribute implies that returning a null pointer is undefined behavior, which the optimizer may take advantage of. The ``_Nonnull`` type qualifier indicates that a pointer cannot be null in a more general manner (because it is part of the type system) and does not imply undefined behavior, making it more widely applicable