[clang] 10079a2 - [HIP] Document func ptr and virtual func (#68126)

[via cfe-commits]

Author: Yaxun (Sam) Liu
Date: 2023-10-18T14:14:04-04:00
New Revision: 10079a23c5b6becfdaaa34d8ba72c91706ad4a6d

URL: https://github.com/llvm/llvm-project/commit/10079a23c5b6becfdaaa34d8ba72c91706ad4a6d
DIFF: https://github.com/llvm/llvm-project/commit/10079a23c5b6becfdaaa34d8ba72c91706ad4a6d.diff

LOG: [HIP] Document func ptr and virtual func (#68126)

Document clang support for function pointers and virtual functions with
HIP

Added: 
    

Modified: 
    clang/docs/HIPSupport.rst

Removed: 
    


################################################################################
diff  --git a/clang/docs/HIPSupport.rst b/clang/docs/HIPSupport.rst
index 8b4649733a9c777..84cee45e83ba3c8 100644
--- a/clang/docs/HIPSupport.rst
+++ b/clang/docs/HIPSupport.rst
@@ -176,3 +176,93 @@ Predefined Macros
    * - ``HIP_API_PER_THREAD_DEFAULT_STREAM``
      - Alias to ``__HIP_API_PER_THREAD_DEFAULT_STREAM__``. Deprecated.
 
+Compilation Modes
+=================
+
+Each HIP source file contains intertwined device and host code. Depending on the chosen compilation mode by the compiler options ``-fno-gpu-rdc`` and ``-fgpu-rdc``, these portions of code are compiled 
diff erently.
+
+Device Code Compilation
+-----------------------
+
+**``-fno-gpu-rdc`` Mode (default)**:
+
+- Compiles to a self-contained, fully linked offloading device binary for each offloading device architecture.
+- Device code within a Translation Unit (TU) cannot call functions located in another TU.
+
+**``-fgpu-rdc`` Mode**:
+
+- Compiles to a bitcode for each GPU architecture.
+- For each offloading device architecture, the bitcode from 
diff erent TUs are linked together to create a single offloading device binary.
+- Device code in one TU can call functions located in another TU.
+
+Host Code Compilation
+---------------------
+
+**Both Modes**:
+
+- Compiles to a relocatable object for each TU.
+- These relocatable objects are then linked together.
+- Host code within a TU can call host functions and launch kernels from another TU.
+
+Function Pointers Support
+=========================
+
+Function pointers' support varies with the usage mode in Clang with HIP. The following table provides an overview of the support status across 
diff erent 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 TU. This mode does not support the use of function pointers defined in a 
diff erent 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 
diff erent TUs, as resource usage calculations can accommodate functions from disparate TUs.
+
+Virtual Function Support
+========================
+
+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 on a specific offloading device if the object is constructed in device code on an offloading device with the same architecture.
+- **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 
diff erent 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 
diff erent architecures are 
diff erent, therefore trying to invoke virtual functions from an offloading device with a 
diff erent architecture than where the object is constructed is also disallowed.
+
+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
+   }