Skip to content
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

[HIP] Document func ptr and virtual func #68126

Merged
merged 2 commits into from
Oct 18, 2023
Merged

Conversation

yxsamliu
Copy link
Collaborator

@yxsamliu yxsamliu commented Oct 3, 2023

Document clang support for function pointers and virtual functions with HIP

@llvmbot llvmbot added the clang Clang issues not falling into any other category label Oct 3, 2023
@llvmbot
Copy link
Member

llvmbot commented Oct 3, 2023

@llvm/pr-subscribers-clang

Changes

Document clang support for function pointers and virtual functions with HIP


Full diff: https://github.com/llvm/llvm-project/pull/68126.diff

1 Files Affected:

  • (modified) clang/docs/HIPSupport.rst (+67)
diff --git a/clang/docs/HIPSupport.rst b/clang/docs/HIPSupport.rst
index 8b4649733a9c777..7a4db10789f2c90 100644
--- a/clang/docs/HIPSupport.rst
+++ b/clang/docs/HIPSupport.rst
@@ -176,3 +176,70 @@ Predefined Macros
    * - ``HIP_API_PER_THREAD_DEFAULT_STREAM``
      - Alias to ``__HIP_API_PER_THREAD_DEFAULT_STREAM__``. Deprecated.
 
+Function Pointers Support in Clang with HIP
+===========================================
+
+Function pointers' support varies with the usage mode in Clang with HIP. The following table provides an overview of the support status across different use-cases and modes.
+
+.. list-table:: Function Pointers Support Overview
+   :widths: 25 25 25
+   :header-rows: 1
+
+   * - Use Case
+     - ``-fno-gpu-rdc`` Mode (default)
+     - ``-fgpu-rdc`` Mode
+   * - Defined and used in the same TU
+     - Supported
+     - Supported
+   * - Defined in one TU and used in another TU
+     - Not Supported
+     - Supported
+
+In the ``-fno-gpu-rdc`` mode, the compiler calculates the resource usage of kernels based only on functions present within the same Translation Unit (TU). This mode does not support the use of function pointers defined in a different TU due to the possibility of incorrect resource usage calculations, leading to undefined behavior. 
+
+On the other hand, the ``-fgpu-rdc`` mode allows the definition and use of function pointers across different TUs, as resource usage calculations can accommodate functions from disparate TUs.
+
+Virtual Function Support in Clang with HIP
+==========================================
+
+In Clang with HIP, support for calling virtual functions of an object in device or host code is contingent on where the object is constructed. 
+
+- **Constructed in Device Code**: Virtual functions of an object can be called in device code if the object is constructed in device code.
+- **Constructed in Host Code**: Virtual functions of an object can be called in host code if the object is constructed in host code.
+
+In other scenarios, calling virtual functions is not allowed.
+
+Explanation
+-----------
+
+An object constructed on the device side contains a pointer to the virtual function table on the device side, which is not accessible in host code, and vice versa. Thus, trying to invoke virtual functions from a context different from where the object was constructed will be disallowed because the appropriate virtual table cannot be accessed.
+
+Example Usage
+-------------
+
+.. code-block:: c++
+
+   class Base {
+   public:
+      __device__ virtual void virtualFunction() {
+         // Base virtual function implementation
+      }
+   };
+
+   class Derived : public Base {
+   public:
+      __device__ void virtualFunction() override {
+         // Derived virtual function implementation
+      }
+   };
+
+   __global__ void kernel() {
+      Derived obj;
+      Base* basePtr = &obj;
+      basePtr->virtualFunction(); // Allowed since obj is constructed in device code
+   }
+
+Note
+----
+
+Ensure to construct objects in the appropriate context (host or device) depending on where you intend to call their virtual functions to avoid runtime errors.

Note
----

Ensure to construct objects in the appropriate context (host or device) depending on where you intend to call their virtual functions to avoid runtime errors.
Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Not sure that this note is needed given the preceding information. If kept, replace "Ensure to construct" with "Ensure that objects are constructed"...

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

will remove

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I added a separate section about -fgpu-rdc and -fno-gpu-rdc mode

Copy link

@b-sumner b-sumner left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This looks fine to me.

- Not Supported
- Supported

In the ``-fno-gpu-rdc`` mode, the compiler calculates the resource usage of kernels based only on functions present within the same Translation Unit (TU). This mode does not support the use of function pointers defined in a different TU due to the possibility of incorrect resource usage calculations, leading to undefined behavior.
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I'd mention that in -fno-gpu-rdc each TU compiles to a fully linked GPU executable, vs an object file with -fgpu-rdc. This should give more context for reasoning about accessibility of various data bits.

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

will do

clang/docs/HIPSupport.rst Show resolved Hide resolved

An object constructed on the device side contains a pointer to the virtual function table on the device side, which is not accessible in host code, and vice versa. Thus, trying to invoke virtual functions from a context different from where the object was constructed will be disallowed because the appropriate virtual table cannot be accessed. The virtual function tables for offloading devices with different architecures are different, therefore trying to invoke virtual functions from an offloading device with a different architecture than where the object is constructed is also disallowed.

A possible way to alleviate the current limitation of virtual function support in HIP is through the use of a "composite vtable". This involves creating a vtable that combines those from the host and all offloading device architectures, storing it in memory accessible by both. A dedicated registration function is introduced to populate this composite vtable. This function is invoked during global initialization to ensure the vtable is ready before any virtual function calls are made. For every virtual function call, irrespective of context, the system refers to this composite vtable to determine the correct function execution.
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Could you clarify whether the purpose of this pragraph is to provide a potential workaround to users or a solution in a future compiler?

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It is a possible solution. However, there may be better solutions that I am open to discussion.

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I would avoid speculating in this documentation. Maybe we should remove this paragraph?

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

removed the paragraph

@yxsamliu
Copy link
Collaborator Author

ping

Document clang support for function pointers and virtual functions with HIP
Document clang support for function pointers and virtual functions with HIP
Copy link

@b-sumner b-sumner left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Looks good to me.

@scchan
Copy link
Contributor

scchan commented Oct 18, 2023

LGTM thanks

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
clang Clang issues not falling into any other category
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

5 participants