From faccab9d44560b524b25ba983e07df7ddc576ddf Mon Sep 17 00:00:00 2001 From: Julia Jiang Date: Tue, 4 Apr 2023 15:07:41 -0400 Subject: [PATCH] SWDEV-389297 - Update Kernel language for HIP programming guide Change-Id: I29a81ee92524a5d1be981ef828b204c033ebd1f9 --- docs/markdown/hip_kernel_language.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/docs/markdown/hip_kernel_language.md b/docs/markdown/hip_kernel_language.md index 5f8e9f8f95..5e32e923fa 100644 --- a/docs/markdown/hip_kernel_language.md +++ b/docs/markdown/hip_kernel_language.md @@ -186,10 +186,10 @@ Built-ins determine the coordinate of the active work item in the execution grid In HIP, built-ins coordinate variable definitions are the same as in Cuda, for instance: 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. +Coordinates builtins are implemented as structures for better performance. When used with printf, they needs to be casted to integer types explicitly. ### warpSize -The warpSize variable is of type int and contains the warp size (in threads) for the target device. Note that all current Nvidia devices return 32 for this variable, and all current AMD devices return 64. Device code should use the warpSize built-in to develop portable wave-aware code. - +The warpSize variable is of type int and contains the warp size (in threads) for the target device. Note that all current Nvidia devices return 32 for this variable, and current AMD devices return 64 for gfx9 and 32 for gfx10 and above. The warpSize variable should only be used in device functions. Device code should use the warpSize built-in to develop portable wave-aware code. ## Vector Types