diff --git a/projects/clr/hipamd/README.md b/projects/clr/hipamd/README.md index 38d4d147ed..4b69485ba3 100644 --- a/projects/clr/hipamd/README.md +++ b/projects/clr/hipamd/README.md @@ -12,16 +12,15 @@ Key features include: 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. ## More Info: -- [Installation](INSTALL.md) and [clang-hipify](clang-hipify/README.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://gpuopen-professionalcompute-tools.github.io/HIP) - [HIP Porting Guide](docs/markdown/hip_porting_guide.md) - [HIP Terminology](docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) +- [clang-hipify](clang-hipify/README.md) - [Developer/CONTRIBUTING Info](CONTRIBUTING.md) - [Release Notes](RELEASE.md) - - ## How do I get set up? See the [Installation](INSTALL.md) notes. diff --git a/projects/clr/hipamd/RELEASE.md b/projects/clr/hipamd/RELEASE.md index dfec3285ae..f60dafd4b5 100644 --- a/projects/clr/hipamd/RELEASE.md +++ b/projects/clr/hipamd/RELEASE.md @@ -1,11 +1,10 @@ # Release notes -We have attempted to document known bugs and limitations - in particular the [HIP Kernel Language](docs/markdown/hip_kernel_language.md) document uses the phrase "Under Development", and the [HIP Runtime API bug list](http://gpuopen-professionalcompute-tools.github.io/HIP/bug.html) lists known bugs. Some of the key items we are working on: -- Tuning built-in functions, including shfl. -- Performance optimization. +We have attempted to document known bugs and limitations - in particular the [HIP Kernel Language](docs/markdown/hip_kernel_language.md) document uses the phrase "Under Development", and the [HIP Runtime API bug list](http://gpuopen-professionalcompute-tools.github.io/HIP/bug.html) lists known bugs. - -Stay tuned - the work for many of these features is already in-flight. +=================================================================================================== +Upcoming: +- hipLaunchKernel supports one-dimensional grid and/or block dims, without explicit cast to dim3 type. =================================================================================================== Release:0.90.00 diff --git a/projects/clr/hipamd/docs/doxygen-input/mainpage.txt b/projects/clr/hipamd/docs/doxygen-input/mainpage.txt index 5424464f5c..48177198cc 100644 --- a/projects/clr/hipamd/docs/doxygen-input/mainpage.txt +++ b/projects/clr/hipamd/docs/doxygen-input/mainpage.txt @@ -7,11 +7,11 @@ * The HIP interface makes it very easy to port existing CUDA apps to run on AMD GPUs, * or to develop new apps that can run on either CUDA or AMD GPUs from a common source base. * - * - HIP is very thin and has little or no performance impact over coding directly in CUDA NVCC or Kalmar HC mode. + * - HIP is very thin and has little or no performance impact over coding directly in CUDA NVCC or HCC HC mode. * - HIP allows developers to use the "best" development environment and tools on each target platform. * - HIP allows coding in a single-source C++ programming language including features such as templates, C++11 lambdas,and more. - * - "hipify" script automatically converts source from CUDA to HIP. - * - Developers can specialize for CUDA or Kalmar to tune for performance or handle tricky cases with #ifdef. + * - "hipify" tool automatically converts source from CUDA to HIP. + * - Developers can specialize for CUDA or HCC to tune for performance or handle tricky cases with #ifdef. * - See the @ref API. diff --git a/projects/clr/hipamd/docs/markdown/hip_faq.md b/projects/clr/hipamd/docs/markdown/hip_faq.md index 5e722383df..e71dcdfb36 100644 --- a/projects/clr/hipamd/docs/markdown/hip_faq.md +++ b/projects/clr/hipamd/docs/markdown/hip_faq.md @@ -1,13 +1,12 @@ # FAQ - - -**Table of Contents** *generated with [DocToc](https://github.com/thlorenz/doctoc)* - -- [What APIs does HIP support ?](#what-apis-does-hip-support-) -- [What is not supported?](#what-is-not-supported) +- [What APIs and features does HIP support ?](#Q1) +- [What is not supported?](#Q2) - [Run-time features](#run-time-features) - [Kernel language features](#kernel-language-features) +- [Is HIP a drop-in replacement for CUDA?](#Q3) +- [What version of CUDA is supported?](#Q4) +- [What libraries does HIP support?](#Q5) - [How does HIP compare with OpenCL?](#how-does-hip-compare-with-opencl) - [What hardware does HIP support?](#what-hardware-does-hip-support) - [Does Hipify automatically convert all source code?](#does-hipify-automatically-convert-all-source-code) @@ -20,38 +19,79 @@ - [What's the difference between HIP and hc?](#whats-the-difference-between-hip-and-hc) - [HIP detected my platform (hcc vs nvcc) incorrectly - what should I do?](#hip-detected-my-platform-hcc-vs-nvcc-incorrectly---what-should-i-do) - -================= -### What APIs does HIP support ? +### What APIs and features does HIP support? HIP provides the following: -- Devices (hipSetDevice(), hipGetDeviceProperties(), etc) -- Memory management (hipMalloc(), hipMemcpy(), hipFree()) -- Streams (hipStreamCreate(), etc.)---under development +- Devices (hipSetDevice(), hipGetDeviceProperties(), etc.) +- Memory management (hipMalloc(), hipMemcpy(), hipFree(), etc.) +- Streams (hipStreamCreate(), etc.) - Events (hipEventRecord(), hipEventElapsedTime(), etc.) - Kernel launching (hipLaunchKernel is a standard C/C++ function that replaces <<< >>>) -- CUDA-style kernel indexing -- Device-side math built-ins +- CUDA-style kernel coordinate functions (threadIdx, blockIdx, blockDim, gridDim) +- Most device-side math built-ins - Error reporting (hipGetLastError(), hipGetErrorString()) -The HIP documentation describes each API and its limitations, if any, compared with the equivalent CUDA API. +The HIP API documentation describes each API and its limitations, if any, compared with the equivalent CUDA API. -### What is not supported? +### What is not supported? #### Run-time features - Textures -- Dynamic parallelism -- Managed memory +- MemcpyToSymbol functions +- Dynamic parallelism (CUDA 5.0) +- Managed memory (CUDA 6.5) - Graphics interoperation with OpenGL or Direct3D - CUDA array, mipmappedArray and pitched memory - CUDA Driver API - + #### Kernel language features -- Device-side dynamic memory allocations (malloc, free, new, delete) -- Virtual functions, indirect functions and try/catch +- Device-side dynamic memory allocations (malloc, free, new, delete) (CUDA 4.0) +- Virtual functions, indirect functions and try/catch (CUDA 4.0) - `__prof_trigger` -- PTX assembly -- See the [HIP Kernel Language](hip_kernel_language.md) for more information. +- PTX assembly (CUDA 4.0) +- 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? +No. HIP provides porting tools which do most of the work do convert CUDA code into portable C++ code that uses the HIP APIs. +Most developers will port their code from CUDA to HIP and then maintain the HIP version. +HIP code provides the same performance as coding in native CUDA, plus the benefit that the code can also run on AMD platforms. + +### What version of CUDA is supported? +HIP APIs and features do not map to a specific CUDA version. HIP provides a strong subset of functionality provided in CUDA, and the hipify tools can +scan code to identify any unsupported CUDA functions - this is very useful for identifying the specific features required by a given application. + +However, we can provide a rough summary of the features included in each CUDA SDK and the support level in HIP: + +- CUDA 4.0 and earlier : + - HIP supports CUDA 4.0 except for the limitations described [above](#Q2). +- CUDA 5.0 : + - Dynamic Parallelism (not supported) + - cuIpc functions (under development). +- CUDA 5.5 : + - CUPTI (not directly supported, [AMD GPUPerfAPI](http://developer.amd.com/tools-and-sdks/graphics-development/gpuperfapi/) can be used as an alternative in some cases) +- CUDA 6.0 + - Managed memory (under development) +- CUDA 6.5 + - __shfl instriniscs (supported) +- CUDA 7.0 + - Per-thread-streams (under development) + - C++11 (HCC supports all of C++11, all of C++14 and some C++17 features) +- CUDA 7.5 + - float16 (under development) +- CUDA 8.0 + - No new language features. + +### What libraries does HIP support? +HIP includes growing support for the 4 key math libraries using hcBlas, hcFft, hcrng, and hcsparse). +These offer pointer-based memory interfaces (as opposed to opaque buffers) and can be easily interfaces with other HCC code. Developers should use conditional compliation if portability to nvcc systems is desired - using calls to cu* routines on one path and hc* routines on the other. + +- [hcblas](https://bitbucket.org/multicoreware/hcblas) +- [hcfft](https://bitbucket.org/multicoreware/hcfft) +- [hcsparse](https://bitbucket.org/multicoreware/hcsparse) +- [hcrng](https://bitbucket.org/multicoreware/hcrng) + +Additionally, some of the cublas routines are automatically converted to hipblas equivalents by the clang-hipify tool. These APIs use cublas or hcblas depending on the platform, and replace the need +to use conditional compilation. ### How does HIP compare with OpenCL? Both AMD and Nvidia support OpenCL 1.2 on their devices, so developers can write portable code. @@ -65,7 +105,7 @@ HIP offers several benefits over OpenCL: - HIP offers an offline compilation model. ### What hardware does HIP support? -- For AMD platforms, HIP runs on the same hardware that the HCC "hc" mode supports - specifically AMD Kaveri, Carrizo and Fiji. +- For AMD platforms, HIP runs on the same hardware that the HCC "hc" mode supports. See the ROCM documentation for the list of supported platforms. - For Nvidia platforms, HIP requires Unified Memory and should run on a device which runs the CUDA SDK 6.0 or newer. We have tested the Nvidia Titan and K40. ### Does Hipify automatically convert all source code? @@ -79,7 +119,7 @@ Developers should always expect to perform some platform-specific tuning and opt NVCC is Nvidia's compiler driver for compiling "CUDA C++" code into PTX or device code for Nvidia GPUs. It's a closed-source binary product that comes with CUDA SDKs. ### What is HCC? -HCC is AMD's compiler driver which compiles "heterogenous C++" code into HSAIL or GCN device code for AMD GPUs. HCC is an open-source compiler based on recent versions of CLANG/LLVM. +HCC is AMD's compiler driver which compiles "heterogenous C++" code into HSAIL or GCN device code for AMD GPUs. It's an open-source compiler based on recent versions of CLANG/LLVM. ### Why use HIP rather than supporting CUDA directly? While HIP is a strong subset of the CUDA, it is a subset. The HIP layer allows that subset to be clearly defined and documented. diff --git a/projects/clr/hipamd/docs/markdown/hip_kernel_language.md b/projects/clr/hipamd/docs/markdown/hip_kernel_language.md index 9ac1a48a25..478de7cda6 100644 --- a/projects/clr/hipamd/docs/markdown/hip_kernel_language.md +++ b/projects/clr/hipamd/docs/markdown/hip_kernel_language.md @@ -1,6 +1,4 @@ - - -**Table of Contents** *generated with [DocToc](https://github.com/thlorenz/doctoc)* +**Table of Contents** - [Introduction](#introduction) - [Function-Type Qualifiers](#function-type-qualifiers) @@ -45,8 +43,6 @@ - [In-Line Assembly](#in-line-assembly) - [C++ Support](#c-support) - - ## Introduction @@ -98,9 +94,9 @@ HIP parses the `__noinline__` and `__forceinline__` keywords and converts them t `__global__` functions are often referred to as *kernels,* and calling one is termed *launching the kernel.* These functions require the caller to specify an "execution configuration" that includes the grid and block dimensions. The execution configuration can also include other information for the launch, such as the amount of additional shared memory to allocate and the stream where the kernel should execute. HIP introduces a standard C++ calling convention to pass the execution configuration to the kernel (this convention replaces the Cuda <<< >>> syntax). In HIP, - Kernels launch with the "hipLaunchKernel" function - The first five parameters to hipLaunchKernel are the following: - - **symbol kernelName**: the name of the kernel to launch - - **dim3 gridDim**: 3D-grid dimensions - - **dim3 blockDim**: 3D-block dimensions + - **symbol kernelName**: the name of the kernel to launch. To support template kernels which contains "," use the HIP_KERNEL_NAME macro. The hipify tools insert this automatically. + - **dim3 gridDim**: 3D-grid dimensions specifying the number of blocks to launch. + - **dim3 blockDim**: 3D-block dimensions specifying the number of threads in each block. - **size_t dynamicShared**: amount of additional shared memory to allocate when launching the kernel (see [__shared__](#__shared__)) - **hipStream_t**: stream where the kernel should execute. A value of 0 corresponds to the NULL stream (see [Synchronization Functions](#synchronization-functions)). - Kernel arguments follow these first five parameters @@ -629,6 +625,9 @@ The compiler ensures that the kernel uses fewer registers than both allowed maxi HIP/hcc will parse the `launch_bounds` attribute but silently ignores the performance hint. Full support is under development. +The hcc compiler does not support the "--maxregcount" option like nvcc. Instead, users are encouraged to use the hip_launch_bounds directive since the parameters are more intuitive and portable than +micro-architecture details like registers, and also the directive allows per-kernel control rather than an entire file. hip_launch_bounds works on both hcc and nvcc targets. + ## Register Keyword The register keyword affects code generation in neither nvcc nor hcc. It’s deprecated in standard C++, so hcc will generate a warning. (nvcc silently ignores use of this keyword.) To disable the warning, you can pass the option `-Wno-deprecated-register` to hcc. @@ -636,23 +635,25 @@ The register keyword affects code generation in neither nvcc nor hcc. It’s de ## Pragma Unroll -hcc support for the unroll pragma is under development and is slated to arrive with the Lightning Compiler. - +Unroll with a bounds that is known at compile-time is supported. For example: ``` #pragma unroll 16 /* hint to compiler to unroll next loop by 16 */ for (int i=0; i<16; i++) ... ``` +``` +#pragma unroll 1 /* tell compiler to never unroll the loop */ +for (int i=0; i<16; i++) ... +``` + + +Unbounded loop unroll is under development on HCC compiler. ``` #pragma unroll /* hint to compiler to completely unroll next loop. */ for (int i=0; i<16; i++) ... ``` -``` -#pragma unroll 1 /* tell compiler to never unroll the loop */ -for (int i=0; i<16; i++) ... -``` ## In-Line Assembly diff --git a/projects/clr/hipamd/docs/markdown/hip_porting_guide.md b/projects/clr/hipamd/docs/markdown/hip_porting_guide.md index 50282c3eff..76cf86aa49 100644 --- a/projects/clr/hipamd/docs/markdown/hip_porting_guide.md +++ b/projects/clr/hipamd/docs/markdown/hip_porting_guide.md @@ -1,95 +1,50 @@ - - -**Table of Contents** *generated with [DocToc](https://github.com/thlorenz/doctoc)* - -- [HIP Porting Guide](#hip-porting-guide) -- [###Table of Contents](#table-of-contents) - - [Porting a New Cuda Project](#porting-a-new-cuda-project) - - [General Tips](#general-tips) - - [Scanning existing CUDA code to scope the porting effort](#scanning-existing-cuda-code-to-scope-the-porting-effort) - - [Converting a project "in-place"](#converting-a-project-in-place) - - [Distinguishing Compiler Modes](#distinguishing-compiler-modes) - - [Identifying HIP Target Platform](#identifying-hip-target-platform) - - [Identifying the Compiler: hcc or nvcc](#identifying-the-compiler-hcc-or-nvcc) - - [Identifying Current Compilation Pass: Host or Device](#identifying-current-compilation-pass-host-or-device) - - [Compiler Defines: Summary](#compiler-defines-summary) - - [Identifying Architecture Features](#identifying-architecture-features) - - [HIP_ARCH Defines](#hip_arch-defines) - - [Device-Architecture Properties](#device-architecture-properties) - - [Table of Architecture Properties](#table-of-architecture-properties) - - [Finding HIP](#finding-hip) - - [Compiler Options](#compiler-options) - - [Linking Issues](#linking-issues) - - [Linking With hipcc](#linking-with-hipcc) - - [-lm Option](#-lm-option) - - [Linking Code With Other Compilers](#linking-code-with-other-compilers) - - [libc++ and libstdc++](#libc-and-libstdc) - - [HIP Headers (hip_runtime.h, hip_runtime_api.h)](#hip-headers-hip_runtimeh-hip_runtime_apih) - - [Using a Standard C++ Compiler](#using-a-standard-c-compiler) - - [cuda.h](#cudah) - - [Choosing HIP File Extensions](#choosing-hip-file-extensions) - - [Workarounds](#workarounds) - - [warpSize](#warpsize) - - [Textures and Cache Control](#textures-and-cache-control) - - [More Tips](#more-tips) - - [hcc CPU Mode](#hcc-cpu-mode) - - [HIPTRACE Mode](#hiptrace-mode) - - [Environment Variables](#environment-variables) - - [Debugging hipcc](#debugging-hipcc) - - [What Does This Error Mean?](#what-does-this-error-mean) - - [/usr/include/c++/v1/memory:5172:15: error: call to implicitly deleted default constructor of 'std::__1::bad_weak_ptr' throw bad_weak_ptr();](#usrincludecv1memory517215-error-call-to-implicitly-deleted-default-constructor-of-std__1bad_weak_ptr-throw-bad_weak_ptr) - - [grid_launch kernel dispatch - fallback](#grid_launch-kernel-dispatch---fallback) - - [HIP Environment Variables](#hip-environment-variables) - - [Editor Highlighting](#editor-highlighting) - - - # HIP Porting Guide In addition to providing a portable C++ programmming environement for GPUs, HIP is designed to ease the porting of existing CUDA code into the HIP environment. This section describes the available tools and provides practical suggestions on how to port CUDA code and work through common issues. -###Table of Contents +## Table of Contents ================= - * [HIP Porting Guide](#hip-porting-guide) - * [Table of Contents](#table-of-contents) - * [Porting a New Cuda Project To HIP](#porting-a-new-cuda-project) - * [General Tips](#general-tips" aria-hidden="true"> ## Porting a New Cuda Project @@ -463,17 +418,13 @@ void myFunc () ``` +Additionally, many of the Rodinia benchmarks demonstrate how to modify hipified programs so that textures are not required - search for USE_TEXTURES define in the rodinia source directory. +For example, [here + Cuda programs that employ sampler hardware must either wait for hcc texture support or use more-sophisticated workarounds. ## More Tips -### hcc CPU Mode -Recent hcc versions support CPU accelerator targets. This feature enables some interesting possibilities for HIP porting: - -- hcc can run on any machine, including perhaps a cross-compiling environment on a machine also running nvcc -- Standard CPU debuggers can debug CPU code -- A single code path can run on an AMD or Nvidia GPU or CPU, but the CPU accelerator is a low-performance target---it’s just a single core and lacks SIMD acceleration - ### HIPTRACE Mode On an hcc/AMD platform, set the HIP_TRACE_API environment variable to see a textural API trace. Use the following bit mask: @@ -490,7 +441,7 @@ On hcc/AMD platforms, set the HIP_PRINT_ENV environment variable to 1 and run an - HIP_TRACE_API = 1: trace each HIP API call. Print the function name and return code to stderr as the program executes. - HIP_LAUNCH_BLOCKING = 0: make HIP APIs “host-synchronous” so they are blocked until any kernel launches or data-copy commands are complete (an alias is CUDA_LAUNCH_BLOCKING) -- KMDUMPISA = 1 : Will dump the GCN ISA for all kernels into the local directory. +- KMDUMPISA = 1 : Will dump the GCN ISA for all kernels into the local directory. (This flag is provided by HCC). ### Debugging hipcc