From 346117774ce7e3d80883ad9db63d60fc19ec4012 Mon Sep 17 00:00:00 2001
From: Julia Jiang <56359287+jujiang-del@users.noreply.github.com>
Date: Tue, 14 Sep 2021 04:05:51 -0400
Subject: [PATCH] SWDEV-291894 - Update the launch bounds error message and
 version note in FAQ (#2335)

Change-Id: I5b876e80c69d3e04fec92bc3eb995fcd3e94edcd
---
 docs/markdown/hip_faq.md             | 2 ++
 docs/markdown/hip_kernel_language.md | 4 ++++
 2 files changed, 6 insertions(+)

diff --git a/docs/markdown/hip_faq.md b/docs/markdown/hip_faq.md
index a0309612c3..d1d816123c 100644
--- a/docs/markdown/hip_faq.md
+++ b/docs/markdown/hip_faq.md
@@ -271,3 +271,5 @@ HIP version can be queried from HIP API call,
 hipRuntimeGetVersion(&runtimeVersion);
 
 The version returned will always be greater than the versions in previous ROCm releases.
+
+Note: The version definition of HIP runtime is different from CUDA. On AMD platform, the function returns HIP runtime version, while on NVIDIA platform, it returns CUDA runtime version. And there is no mapping/correlation between HIP version and CUDA version.
diff --git a/docs/markdown/hip_kernel_language.md b/docs/markdown/hip_kernel_language.md
index 8ab3ec99e1..03bcd9810e 100644
--- a/docs/markdown/hip_kernel_language.md
+++ b/docs/markdown/hip_kernel_language.md
@@ -734,6 +734,10 @@ __launch_bounds__ supports two parameters:
 The threads-per-block is the product of (hipBlockDim_x * hipBlockDim_y * hipBlockDim_z).
 - MIN_WARPS_PER_EU - directs the compiler to minimize resource usage so that the requested number of warps can be simultaneously active on a multi-processor. Since active warps compete for the same fixed pool of resources, the compiler must reduce resources required by each warp(primarily registers). MIN_WARPS_PER_EU is optional and defaults to 1 if not specified. Specifying a MIN_WARPS_PER_EU greater than the default 1 effectively constrains the compiler's resource usage.
 
+When launch kernel with HIP APIs, for example, hipModuleLaunchKernel(), HIP will do validation to make sure input kernel dimension size is not larger than specified launch_bounds.
+In case exceeded, HIP would return launch failure, if AMD_LOG_LEVEL is set with proper value (for details, please refer to docs/markdown/hip_logging.md), detail information will be shown in the error log message, including
+launch parameters of kernel dim size, launch bounds, and the name of the faulting kernel. It's helpful to figure out which is the faulting kernel, besides, the kernel dim size and launch bounds values will also assist in debugging such failures.
+
 ### Compiler Impact
 The compiler uses these parameters as follows:
 - The compiler uses the hints only to manage register usage, and does not automatically reduce shared memory or other resources.