SWDEV-329687 - update hip docs (#2632)

Change-Id: Iad6e7bf270a6a57256f8f38b67b7f7c504da2c5b

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 <br><sub> Serialize kernel enqueue. </sub> |  0  | 1: Wait for completion before enqueue. <br> 2: Wait for completion after enqueue. <br> 3: Both. |
 | AMD_SERIALIZE_COPY <br><sub> Serialize copies. </sub> |  0  | 1: Wait for completion before enqueue. <br> 2: Wait for completion after enqueue. <br> 3: Both. |
 | HIP_HOST_COHERENT <br><sub> Coherent memory in hipHostMalloc. </sub> |  0  |  0: memory is not coherent between host and GPU. <br> 1: memory is coherent with host. |
-| AMD_DIRECT_DISPATCH <br><sub> Enable direct kernel dispatch. </sub> | 0  | 0: Disable. <br> 1: Enable. |
+| AMD_DIRECT_DISPATCH <br><sub> Enable direct kernel dispatch. </sub> | 1  | 0: Disable. <br> 1: Enable. |
 
 
 ## General Debugging Tips

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.

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".