SWDEV-243933-update HIP README document

Change-Id: Ie303698ae92c7fa1b1202c1ee06beda1005ba065
2020-07-09 16:36:30 -04:00
@@ -4,11 +4,11 @@

 Key features include:

-* HIP is very thin and has little or no performance impact over coding directly in CUDA or hcc "HC" mode.
+* HIP is very thin and has little or no performance impact over coding directly in CUDA mode.
 * HIP allows coding in a single-source C++ programming language including features such as templates, C++11 lambdas, classes, namespaces, and more.
 * HIP allows developers to use the "best" development environment and tools on each target platform.
 * The [HIPIFY](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/README.md) tools automatically convert source from CUDA to HIP.
-* Developers can specialize for the platform (CUDA or hcc) to tune for performance or handle tricky cases 
+* Developers can specialize for the platform (CUDA or AMD) to tune for performance or handle tricky cases.

 New projects can be developed directly in the portable HIP C++ language and can run on either NVIDIA or AMD platforms.  Additionally, HIP provides porting tools which make it easy to port existing CUDA codes to the HIP layer, with no loss of performance as compared to the original CUDA application.  HIP is not intended to be a drop-in replacement for CUDA, and developers should expect to do some manual coding and performance tuning work to complete the port.

@@ -27,16 +27,15 @@ HIP releases are typically of two types. The tag naming convention is different
 * preview_x.yy.zzzz: These denote pre-release code and are based on the developer-preview branch. This type of release is typically made once a week.

 ## More Info:
- [Installation](INSTALL.md) 
+- [Installation](INSTALL.md)
 - [HIP FAQ](docs/markdown/hip_faq.md)
 - [HIP Kernel Language](docs/markdown/hip_kernel_language.md)
 - [HIP Runtime API (Doxygen)](http://rocm-developer-tools.github.io/HIP)
 - [HIP Porting Guide](docs/markdown/hip_porting_guide.md)
 - [HIP Porting Driver Guide](docs/markdown/hip_porting_driver_api.md)
 - [HIP Programming Guide](docs/markdown/hip_programming_guide.md)
- [HIP Profiling ](docs/markdown/hip_profiling.md)
- [HIP Debugging](docs/markdown/hip_debugging.md)
- [HIP Terminology](docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenCL)
+- [HIP Logging ](docs/markdown/hip_logging.md)
+- [HIP Terminology](docs/markdown/hip_terms2.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/OpenCL)
 - [HIPIFY](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/README.md)
 - Supported CUDA APIs:
  * [Runtime API](docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md)
@@ -55,7 +54,7 @@ HIP releases are typically of two types. The tag naming convention is different
 See the [Installation](INSTALL.md) notes.

 ## Simple Example
-The HIP API includes functions such as hipMalloc, hipMemcpy, and hipFree.  
+The HIP API includes functions such as hipMalloc, hipMemcpy, and hipFree.
 Programmers familiar with CUDA will also be able to quickly learn and start coding with the HIP API.
 Compute kernels are launched with the "hipLaunchKernel" macro call.    Here is simple example showing a
 snippet of HIP API code:
@@ -76,11 +75,10 @@ hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost);
 ```


-The HIP kernel language defines builtins for determining grid and block coordinates, math functions, short vectors, 
-atomics, and timer functions. It also specifies additional defines and keywords for function types, address spaces, and 
-optimization controls.  (See the [HIP Kernel Language](docs/markdown/hip_kernel_language.md) for a full description).
-Here's an example of defining a simple 'vector_square' kernel.  
-
+The HIP kernel language defines builtins for determining grid and block coordinates, math functions, short vectors,
+atomics, and timer functions.
+It also specifies additional defines and keywords for function types, address spaces, and optimization controls (See the [HIP Kernel Language](docs/markdown/hip_kernel_language.md) for a full description).
+Here's an example of defining a simple 'vector_square' kernel.


 ```cpp
@@ -100,14 +98,14 @@ vector_square(T *C_d, const T *A_d, size_t N)
 The HIP Runtime API code and compute kernel definition can exist in the same source file - HIP takes care of generating host and device code appropriately.

 ## HIP Portability and Compiler Technology
-HIP C++ code can be compiled with either :
+HIP C++ code can be compiled with either,
 - On the NVIDIA CUDA platform, HIP provides header file which translate from the HIP runtime APIs to CUDA runtime APIs.  The header file contains mostly inlined
  functions and thus has very low overhead - developers coding in HIP should expect the same performance as coding in native CUDA.  The code is then
  compiled with nvcc, the standard C++ compiler provided with the CUDA SDK.  Developers can use any tools supported by the CUDA SDK including the CUDA
  profiler and debugger.
- On the AMD ROCm platform, HIP provides a header and runtime library built on top of hcc compiler.  The HIP runtime implements HIP streams, events, and memory APIs,
+- On the AMD ROCm platform, HIP provides a header and runtime library built on top of HIP-Clang compiler.  The HIP runtime implements HIP streams, events, and memory APIs,
  and is a object library that is linked with the application.  The source code for all headers and the library implementation is available on GitHub.
-  HIP developers on ROCm can use AMD's CodeXL for debugging and profiling.
+  HIP developers on ROCm can use AMD's ROCgdb (https://github.com/ROCm-Developer-Tools/ROCgdb) for debugging and profiling.

 Thus HIP source code can be compiled to run on either platform.  Platform-specific features can be isolated to a specific platform using conditional compilation.  Thus HIP
 provides source portability to either platform.   HIP provides the _hipcc_ compiler driver which will call the appropriate toolchain depending on the desired platform.
@@ -117,7 +115,7 @@ provides source portability to either platform.   HIP provides the _hipcc_ compi

 * A sample and [blog](http://gpuopen.com/hip-to-be-squared-an-introductory-hip-tutorial) that uses any of [HIPIFY](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/README.md) tools to convert a simple app from CUDA to HIP:

- 
+
 ```shell
 cd samples/01_Intro/square
 # follow README / blog steps to hipify the application.
@@ -142,11 +140,11 @@ The README with the procedures and tips the team used during this porting effort
    * **hip_runtime.h** : Includes everything in hip_runtime_api.h PLUS hipLaunchKernel and syntax for writing device kernels and device functions.  hip_runtime.h can only be compiled with hcc.
    * **hcc_detail/**** , **nvcc_detail/**** : Implementation details for specific platforms. HIP applications should not include these files directly.
    * **hcc.h** : Includes interop APIs for HIP and HCC
-    
+
 * **bin**: Tools and scripts to help with hip porting
    * **hipify-perl** : Script based tool to convert CUDA code to portable CPP. Converts CUDA APIs and kernel builtins.
-    * **hipcc** : Compiler driver that can be used to replace nvcc in existing CUDA code. hipcc will call nvcc or hcc depending on platform and include appropriate platform-specific headers and libraries.
-    * **hipconfig** : Print HIP configuration (HIP_PATH, HIP_PLATFORM, CXX config flags, etc.)
+    * **hipcc** : Compiler driver that can be used to replace nvcc in existing CUDA code. hipcc will call nvcc or HIP-Clang depending on platform and include appropriate platform-specific headers and libraries.
+    * **hipconfig** : Print HIP configuration (HIP_PATH, HIP_PLATFORM, HIP_COMPILER, HIP_RUNTIME, CXX config flags, etc.)
    * **hipexamine-perl.sh** : Script to scan the directory, find all code, and report statistics on how much can be ported with HIP (and identify likely features not yet supported).
    * **hipconvertinplace-perl.sh** : Script to scan the directory, find all code, and convert the found CUDA code to HIP reporting all unconverted things.