add C++ language support documentation #3512
Conversation
docs/reference/c++_support.rst
Outdated
| The source code is compiled according to the ``C++03``, ``C++11``, ``C++14``, ``C++17``, | ||
| and ``C++20`` standards, along with HIP-specific extensions, but is subject to | ||
| restrictions. The key restriction is the lack of standard library support. This is due to | ||
| the HIP device's Single Instruction/Multiple Data (SIMD) nature, which makes most of the |
There was a problem hiding this comment.
C++ standard libraries are supported for HIP host code. constexpr functions of C++ standard libraries are supported for HIP device code. Most functionality of C++ standard libraries are not supported for HIP device code since non-constexpr C++ functions are host functions, which cannot be called by device functions or kernels in HIP.
docs/reference/c++_support.rst
Outdated
|
|
||
| The C++11 standard introduced many new features. These features are supported in HIP | ||
| device code, with some notable omissions. The most significant is the lack of concurrency | ||
| support on the device. This is because the HIP device concurrency model fundamentally |
There was a problem hiding this comment.
It is important to note that although C++11 language features are supported by HIP for both host and device code, C++11 standard library features are mostly supported for host code only. The implementation of the C++11 standard library does not consider being used by HIP device code. Even though constexpr functions are implicitly host-device functions that can be called by HIP device code, the implementation of certain C++11 standard library features involves using non-constexpr functions or static members, making them unavailable for HIP device code (e.g., std::function). Therefore, C++11 standard library features are not officially supported by HIP for device code and are in a state of "work as-is."
docs/reference/c++_support.rst
Outdated
| C++14 support | ||
| ------------------------------------------------------------------------------- | ||
|
|
||
| The C++14 language features are supported, except for the updates to the memory model. |
There was a problem hiding this comment.
remove 'except for the updates to the memory model'
docs/reference/c++_support.rst
Outdated
| * have a ``constexpr`` specifier | ||
| * have a parameter of type ``std::initializer_list`` or ``va_list`` | ||
| * use an rvalue reference as a parameter. | ||
|
|
There was a problem hiding this comment.
Add
'* use parameters having different sizes in host and device code, e.g. long double arguments, or structs containing long double members.
'* use struct-type arguments which have different layout in host and device code.'
docs/reference/c++_support.rst
Outdated
| in host or device memory and how its memory should be allocated. HIP supports the | ||
| ``__device__``, ``__shared__``, ``__managed__``, and ``__constant__`` specifiers. | ||
|
|
||
| The ``__device__`` and ``__constant__`` specifiers define static variables, which are |
docs/reference/c++_support.rst
Outdated
| The ``__managed__`` variable specifier creates global variables that are initially | ||
| undefined and unaddressed within the global symbol table. The HIP runtime allocates | ||
| managed memory and defines the symbol when it loads the device binary. A managed variable | ||
| can be accessed in both device and host code. To register managed variables, use |
There was a problem hiding this comment.
remove ' To register managed variables, use
__hipRegisterManagedVariable in an initialization function.'
docs/reference/c++_support.rst
Outdated
|
|
||
| It's important to know where a variable is stored because it is only available from | ||
| certain locations. Generally, variables allocated in the host memory are not accessible | ||
| from the device code, while variables allocated in the device memory are not accessible |
There was a problem hiding this comment.
insert 'directly' between 'not accessible'.
Add 'Accessing device variables in host code should be done through kernel execution or HIP API's like hipMemCpyToSymbol.'
docs/reference/c++_support.rst
Outdated
| ------------------------------------------------------------------------------- | ||
|
|
||
| Classes work on both the host and device side, but there are some constraints. The static | ||
| data members need to be ``const`` qualified, and ``static`` member functions can't be |
There was a problem hiding this comment.
remove 'The static
data members need to be const qualified'
docs/reference/c++_support.rst
Outdated
| ``__global__``. ``Virtual`` member functions work, but a ``virtual`` function must not be | ||
| called from the host if the parent object was created on the device, or the other way | ||
| around, because this behavior is undefined. This also means you can't pass an object with | ||
| ``virtual`` functions as a parameter to a kernel. |
There was a problem hiding this comment.
remove ' This also means you can't pass an object with
virtual functions as a parameter to a kernel.'
docs/reference/c++_support.rst
Outdated
|
|
||
| HIP supports Lambdas, which by default work as expected. | ||
|
|
||
| Lambdas inherit the execution space specification from the surrounding context. For |
There was a problem hiding this comment.
A lambda has implicit host device attributes, therefore can be executed by both device and host code. A lambda defined in a host function can be passed to a template kernel as a template argument and called in the kernel. A lambda can access host variables by copy but not by reference. Accessing host variables by reference in device code results in compilation error or undefined behavior.
docs/reference/c++_support.rst
Outdated
| also means that lambdas can't be used as a template argument for kernels unless defined | ||
| in a device function or a kernel. | ||
|
|
||
| To help develop versatile software, HIP supports an extension that makes lambdas even |
There was a problem hiding this comment.
remove this paragraph and add
'To make a lamba callable only by host or device code, users can add __host__ or __device__ attribute.'
f716b8d
to
eb90780
Compare
Co-authored-by: Leo Paoletti <[email protected]>
docs/reference/c++_support.rst
Outdated
| Classes work on both the host and device side, but there are some constraints. The | ||
| ``static`` member functions can't be ``__global__``. ``Virtual`` member functions work, | ||
| but a ``virtual`` function must not be called from the host if the parent object was | ||
| created on the device, or the other way around, because this behavior is undefined. |
There was a problem hiding this comment.
Is there not still a constraint that __device__ global variables must have trivial constructors?
add sentence on device global variables
add documentation c++ language support