-
Notifications
You must be signed in to change notification settings - Fork 526
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
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.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
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.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
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.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
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.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
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 |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
static -> global
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.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
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.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
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.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
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.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
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.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
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.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
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]>
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Pretty good. Thanks!
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.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Is there not still a constraint that __device__
global variables must have trivial constructors?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
add sentence on device global variables
add documentation c++ language support