Skip to content

Commit

Permalink
Update C++ language extensions
Browse files Browse the repository at this point in the history
- Add clr surface functions
- Update based on review comments
- Fix clocks
  • Loading branch information
neon60 committed May 31, 2024
1 parent b8189b7 commit 1494532
Show file tree
Hide file tree
Showing 2 changed files with 39 additions and 6 deletions.
40 changes: 36 additions & 4 deletions docs/reference/kernel_language.rst
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand Down Expand Up @@ -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)
{
...
}
Expand Down Expand Up @@ -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
===============================================
Expand Down Expand Up @@ -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
===============================================
Expand Down
5 changes: 3 additions & 2 deletions include/hip/hip_runtime_api.h
Original file line number Diff line number Diff line change
Expand Up @@ -3191,7 +3191,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.
*
Expand Down Expand Up @@ -3818,7 +3819,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
Expand Down

0 comments on commit 1494532

Please sign in to comment.