diff --git a/docs/markdown/hip_debugging.md b/docs/markdown/hip_debugging.md index 9127085a4d..b26915c50e 100644 --- a/docs/markdown/hip_debugging.md +++ b/docs/markdown/hip_debugging.md @@ -261,7 +261,7 @@ The following is the summary of the most useful environment variables in HIP. | AMD_SERIALIZE_KERNEL
Serialize kernel enqueue. | 0 | 1: Wait for completion before enqueue.
2: Wait for completion after enqueue.
3: Both. | | AMD_SERIALIZE_COPY
Serialize copies. | 0 | 1: Wait for completion before enqueue.
2: Wait for completion after enqueue.
3: Both. | | HIP_HOST_COHERENT
Coherent memory in hipHostMalloc. | 0 | 0: memory is not coherent between host and GPU.
1: memory is coherent with host. | -| AMD_DIRECT_DISPATCH
Enable direct kernel dispatch. | 0 | 0: Disable.
1: Enable. | +| AMD_DIRECT_DISPATCH
Enable direct kernel dispatch. | 1 | 0: Disable.
1: Enable. | ## General Debugging Tips diff --git a/docs/markdown/hip_faq.md b/docs/markdown/hip_faq.md index a26597ca67..71dd1084ca 100644 --- a/docs/markdown/hip_faq.md +++ b/docs/markdown/hip_faq.md @@ -57,7 +57,6 @@ The HIP API documentation describes each API and its limitations, if any, compar At a high-level, the following features are not supported: - Textures (partial support available) - Dynamic parallelism (CUDA 5.0) -- Managed memory (CUDA 6.5) - Graphics interoperability with OpenGL or Direct3D - CUDA IPC Functions (Under Development) - CUDA array, mipmappedArray and pitched memory @@ -70,8 +69,7 @@ See the [API Support Table](CUDA_Runtime_API_functions_supported_by_HIP.md) for - Virtual functions, indirect functions and try/catch (CUDA 4.0) - `__prof_trigger` - PTX assembly (CUDA 4.0). HIP-Clang supports inline GCN assembly. -- Several kernel features are under development. See the [HIP Kernel Language](hip_kernel_language.md) for more information. These include: - - printf +- Several kernel features are under development. See the [HIP Kernel Language](hip_kernel_language.md) for more information. ### Is HIP a drop-in replacement for CUDA? @@ -233,6 +231,7 @@ See the [HIP Logging](hip_logging.md) for more information. ### What is maximum limit of kernel launching parameter? Product of block.x, block.y, and block.z should be less than 1024. +Please note, HIP does not support kernel launch with total work items defined in dimension with size gridDim x blockDim >= 2^32, so gridDim.x * blockDim.x, gridDim.y * blockDim.y and gridDim.z * blockDim.z are always less than 2^32. ### Are __shfl_*_sync functions supported on HIP platform? __shfl_*_sync is not supported on HIP but for nvcc path CUDA 9.0 and above all shuffle calls get redirected to it's sync version. diff --git a/docs/markdown/hip_kernel_language.md b/docs/markdown/hip_kernel_language.md index 9af6a8ec3a..81b3e076be 100644 --- a/docs/markdown/hip_kernel_language.md +++ b/docs/markdown/hip_kernel_language.md @@ -184,7 +184,7 @@ The `__restrict__` keyword tells the compiler that the associated memory pointer ### Coordinate Built-Ins Built-ins determine the coordinate of the active work item in the execution grid. They are defined in amd_hip_runtime.h (rather than being implicitly defined by the compiler). In HIP, built-ins coordinate variable definitions are the same as in Cuda, for instance: -threadIdx.x, blockIdx.y, gridDim.y, etc. +threadIdx.x, blockIdx.y, gridDim.y, etc. The products gridDim.x * blockDim.x, gridDim.y * blockDim.y and gridDim.z * blockDim.z are always less than 2^32. ### warpSize @@ -696,7 +696,7 @@ void assert(int input) ``` There are two kinds of implementations for assert functions depending on the use sceneries, -- One is for the host version of assert, which is defined in assert.h, +- One is for the host version of assert, which is defined in assert.h, - Another is the device version of assert, which is implemented in hip/hip_runtime.h. Users need to include assert.h to use assert. For assert to work in both device and host functions, users need to include "hip/hip_runtime.h".