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