From 7c551c94819def450423083d2b9a72f7e2343110 Mon Sep 17 00:00:00 2001
From: Istvan Kiss <neon60@gmail.com>
Date: Fri, 31 May 2024 15:21:46 +0200
Subject: [PATCH] Update C++ language extensions

- Add clr surface functions
- Update based on review comments
- Fix clocks
---
 docs/reference/kernel_language.rst | 40 +++++++++++++++++++++++++++---
 include/hip/hip_runtime_api.h      |  5 ++--
 2 files changed, 39 insertions(+), 6 deletions(-)

diff --git a/docs/reference/kernel_language.rst b/docs/reference/kernel_language.rst
index 6d2d0ccdb9..e6d3df266b 100644
--- a/docs/reference/kernel_language.rst
+++ b/docs/reference/kernel_language.rst
@@ -51,8 +51,7 @@ Supported ``__global__`` functions are:
   * Run on the device
   * Called (launched) from the host
 
-HIP ``__global__`` functions must have a ``void`` return type. The first parameter in a HIP ``__global__``
-function must have the type ``hipLaunchParm``. Refer to :ref:`kernel-launch-example` to see usage.
+HIP ``__global__`` functions must have a ``void`` return type.
 
 HIP doesn't support dynamic-parallelism, which means that you can't call ``__global__`` functions from
 the device.
@@ -105,7 +104,7 @@ You can include your kernel arguments after these parameters.
 .. code-block:: cpp
 
   // Example hipLaunchKernelGGL pseudocode:
-  __global__ MyKernel(hipLaunchParm lp, float *A, float *B, float *C, size_t N)
+  __global__ MyKernel(float *A, float *B, float *C, size_t N)
   {
   ...
   }
@@ -1441,7 +1440,38 @@ code.
 Surface functions
 ===============================================
 
-Surface functions are not supported.
+The following surface functions are supported in HIP:
+
+.. doxygengroup:: Surface
+   :content-only:
+
+.. doxygenfunction:: surf1Dread
+
+.. doxygenfunction:: surf1DWrite
+
+.. doxygenfunction:: surf2Dread
+
+.. doxygenfunction:: surf2DWrite
+
+.. doxygenfunction:: surf3Dread
+
+.. doxygenfunction:: surf3Dwrite
+
+.. doxygenfunction:: surf1DLayeredread
+
+.. doxygenfunction:: surf1DLayeredWrite
+
+.. doxygenfunction:: surf2DLayeredread
+
+.. doxygenfunction:: surf2DLayeredWrite
+
+.. doxygenfunction:: surfCubemapread
+
+.. doxygenfunction:: surfCubemapwrite
+
+.. doxygenfunction:: surfCubemapLayeredread
+
+.. doxygenfunction:: surfCubemapLayeredwrite
 
 Timer functions
 ===============================================
@@ -1473,6 +1503,8 @@ To read a high-resolution timer from the device, HIP provides the following buil
 
   Where ``hipDeviceAttributeWallClockRate`` is a device attribute. Note that wall clock frequency is a
   per-device attribute.
+  
+  Note that ``clock()`` and ``clock64()`` do not work properly on AMD RDNA3 (GFX11) graphic processors.
 
 Atomic functions
 ===============================================
diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h
index 38438f85a2..ca92ee5639 100644
--- a/include/hip/hip_runtime_api.h
+++ b/include/hip/hip_runtime_api.h
@@ -3176,7 +3176,8 @@ hipError_t hipMemAllocHost(void** ptr, size_t size);
  *  @param[out] ptr Pointer to the allocated host pinned memory
  *  @param[in]  size Requested memory size in bytes
  *  If size is 0, no memory is allocated, *ptr returns nullptr, and hipSuccess is returned.
- *  @param[in]  flags Type of host memory allocation
+ *  @param[in]  flags Type of host memory allocation. See the description of flags in
+ *  hipSetDeviceFlags.
  *
  *  If no input for flags, it will be the default pinned memory allocation on the host.
  *
@@ -3803,7 +3804,7 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr);
  *
  *  After registering the memory, use #hipHostGetDevicePointer to obtain the mapped device pointer.
  *  On many systems, the mapped device pointer will have a different value than the mapped host
- * pointer.  Applications must use the device pointer in device code, and the host pointer in device
+ * pointer.  Applications must use the device pointer in device code, and the host pointer in host
  * code.
  *
  *  On some systems, registered memory is pinned.  On some systems, registered memory may not be