From e3873dec5aba93f9d82b189fc804ee5e0159d225 Mon Sep 17 00:00:00 2001 From: Julia Jiang Date: Mon, 26 Apr 2021 19:00:03 -0400 Subject: [PATCH] SWDEV-282350 - update HIP documents Change-Id: Iff3a49e75279b1ecee0672a6d35f6b84a8256296 [ROCm/hip commit: ba1c5a955a3e32d66964dabfb539a3de1ac16191] --- projects/hip/README.md | 18 ++++++++++-------- .../hip/docs/markdown/hip_kernel_language.md | 10 ++++++---- .../hip/docs/markdown/hip_porting_guide.md | 8 +++----- 3 files changed, 19 insertions(+), 17 deletions(-) diff --git a/projects/hip/README.md b/projects/hip/README.md index b6bec6efd0..4644e9a716 100644 --- a/projects/hip/README.md +++ b/projects/hip/README.md @@ -41,18 +41,20 @@ HIP releases are typically naming convention for each ROCM release to help diffe - [HIP Porting Driver Guide](docs/markdown/hip_porting_driver_api.md) - [HIP Programming Guide](docs/markdown/hip_programming_guide.md) - [HIP Logging ](docs/markdown/hip_logging.md) +- [HIP Debugging ](docs/markdown/hip_debugging.md) - [Code Object tooling ] (docs/markdown/obj_tooling.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) - * [Driver API](docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md) - * [cuComplex API](docs/markdown/cuComplex_API_supported_by_HIP.md) - * [cuBLAS](docs/markdown/CUBLAS_API_supported_by_HIP.md) - * [cuRAND](docs/markdown/CURAND_API_supported_by_HIP.md) - * [cuDNN](docs/markdown/CUDNN_API_supported_by_HIP.md) - * [cuFFT](docs/markdown/CUFFT_API_supported_by_HIP.md) - * [cuSPARSE](docs/markdown/CUSPARSE_API_supported_by_HIP.md) + * [Runtime API](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/doc/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md) + * [Driver API](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/doc/markdown/CUDA_Driver_API_functions_supported_by_HIP.md) + * [cuComplex API](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/doc/markdown/cuComplex_API_supported_by_HIP.md) + * [Device API](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/doc/markdown/CUDA_Device_API_supported_by_HIP.md) + * [cuBLAS](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/doc/markdown/CUBLAS_API_supported_by_HIP.md) + * [cuRAND](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/doc/markdown/CURAND_API_supported_by_HIP.md) + * [cuDNN](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/doc/markdown/CUDNN_API_supported_by_HIP.md) + * [cuFFT](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/doc/markdown/CUFFT_API_supported_by_HIP.md) + * [cuSPARSE](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/doc/markdown/CUSPARSE_API_supported_by_HIP.md) - [Developer/CONTRIBUTING Info](CONTRIBUTING.md) - [Release Notes](RELEASE.md) diff --git a/projects/hip/docs/markdown/hip_kernel_language.md b/projects/hip/docs/markdown/hip_kernel_language.md index f721ca068f..84873a1317 100644 --- a/projects/hip/docs/markdown/hip_kernel_language.md +++ b/projects/hip/docs/markdown/hip_kernel_language.md @@ -119,12 +119,12 @@ __global__ MyKernel(hipLaunchParm lp, float *A, float *B, float *C, size_t N) } MyKernel<<>> (a,b,c,n); -// Alternatively, kernel can be launched by +// Alternatively, kernel can be launched by // hipLaunchKernel(MyKernel, dim3(gridDim), dim3(groupDim), 0/*dynamicShared*/, 0/*stream), a, b, c, n); ``` -The hipLaunchKernel macro always starts with the five parameters specified above, followed by the kernel arguments. HIPIFY tools optionally convert Cuda launch syntax to hipLaunchKernel, including conversion of optional arguments in <<< >>> to the five required hipLaunchKernel parameters. The dim3 constructor accepts zero to three arguments and will by default initialize unspecified dimensions to 1. See [dim3](#dim3). The kernel uses the coordinate built-ins (hipThread*, hipBlock*, hipGrid*) to determine coordinate index and coordinate bounds of the work item that’s currently executing. See [Coordinate Built-Ins](#coordinate-builtins). +The hipLaunchKernel macro always starts with the five parameters specified above, followed by the kernel arguments. HIPIFY tools optionally convert Cuda launch syntax to hipLaunchKernel, including conversion of optional arguments in <<< >>> to the five required hipLaunchKernel parameters. The dim3 constructor accepts zero to three arguments and will by default initialize unspecified dimensions to 1. See [dim3](#dim3). The kernel uses the coordinate built-ins (thread*, block*, grid*) to determine coordinate index and coordinate bounds of the work item that’s currently executing. See [Coordinate Built-Ins](#coordinate-builtins). ## Kernel-Launch Example @@ -298,7 +298,7 @@ Following is the list of supported single precision mathematical functions. | float log10f ( float x )
Calculate the base 10 logarithm of the input argument. | ✓ | ✓ | | float log1pf ( float x )
Calculate the value of loge( 1 + x ). | ✓ | ✓ | | float logbf ( float x )
Calculate the floating point representation of the exponent of the input argument. | ✓ | ✓ | -| float log2f ( float x )
Calculate the base 2 logarithm of the input argument. | ✓ | ✓ | +| float log2f ( float x )
Calculate the base 2 logarithm of the input argument. | ✓ | ✓ | | float logf ( float x )
Calculate the natural logarithm of the input argument. | ✓ | ✓ | | float modff ( float x, float* iptr )
Break down the input argument into fractional and integral parts. | ✓ | ✗ | | float nanf ( const char* tagp )
Returns "Not a Number" value. | ✗ | ✓ | @@ -379,7 +379,7 @@ Following is the list of supported double precision mathematical functions. | double exp2 ( double x )
Calculate the base 2 exponential of the input argument. | ✓ | ✓ | | double expm1 ( double x )
Calculate the base e exponential of the input argument, minus 1. | ✓ | ✓ | | double fabs ( double x )
Calculate the absolute value of the input argument. | ✓ | ✓ | -| double fdim ( double x, double y )
Compute the positive difference between `x` and `y`. | ✓ | ✓ | +| double fdim ( double x, double y )
Compute the positive difference between `x` and `y`. | ✓ | ✓ | | double floor ( double x )
Calculate the largest integer less than or equal to `x`. | ✓ | ✓ | | double fma ( double x, double y, double z )
Compute `x × y + z` as a single operation. | ✓ | ✓ | | double fmax ( double , double )
Determine the maximum numeric value of the arguments. | ✓ | ✓ | @@ -794,5 +794,7 @@ The file format for binary is `.co` which means Code Object. The following comma [OUTPUT FILE] = Name of the generated code object file ``` +Note: When using binary code objects is that the number of arguments to the kernel is different on HIP-Clang and NVCC path. Refer to the sample in samples/0_Intro/module_api for differences in the arguments to be passed to the kernel. + ## gfx-arch-specific-kernel Clang defined '__gfx*__' macros can be used to execute gfx arch specific codes inside the kernel. Refer to the sample 14_gpu_arch in samples/2_Cookbook. diff --git a/projects/hip/docs/markdown/hip_porting_guide.md b/projects/hip/docs/markdown/hip_porting_guide.md index 6518cb1a4a..74c7fb77d0 100644 --- a/projects/hip/docs/markdown/hip_porting_guide.md +++ b/projects/hip/docs/markdown/hip_porting_guide.md @@ -163,7 +163,7 @@ Often, it's useful to know whether the underlying compiler is HIP-Clang or nvcc. ``` ``` -#ifdef __NVCC__ +#ifdef __HIP_PLATFORM_NVIDIA__ // Compiled with nvcc // Could be compiling with CUDA language extensions enabled (for example, a ".cu file) // Could be in pass-through mode to an underlying host compile OR (for example, a .cpp file) @@ -283,17 +283,15 @@ HIP_PATH ?= $(shell hipconfig --path) ## Identifying HIP Runtime -HIP can depend on ROCclr, or NVCC as runtime +HIP can depend on rocclr, or cuda as runtime - AMD platform On AMD platform, HIP uses Radeon Open Compute Common Language Runtime, called ROCclr. ROCclr is a virtual device interface that HIP runtimes interact with different backends which allows runtimes to work on Linux , as well as Windows without much efforts. -Note, `HIP_ROCclr` needed to be defined on AMD platform that HIP used ROCclr, only when platform was defined as deprecated `HIP_PLATFORM_HCC`. - - NVIDIA platform On Nvidia platform, HIP is just a thin layer on top of CUDA. -On non-AMD platform, HIP runtime determines if nvcc is available and can be used. If available, HIP_PLATFORM is set to nvcc and underneath CUDA path is used. +On non-AMD platform, HIP runtime determines if cuda is available and can be used. If available, HIP_PLATFORM is set to nvidia and underneath CUDA path is used. ## hipLaunchKernel