Initial commit for GPUOpen Launch

2016-01-26 20:14:33 -06:00
@@ -0,0 +1,81 @@
+# FAQ
+
+
+### What APIs does HIP support?
+HIP provides the following:
+- Devices (hipSetDevice(), hipGetDeviceProperties(), etc)
+- Memory management (hipMalloc(), hipMemcpy(), hipFree())
+- Streams (hipStreamCreate(), etc.)---under development
+- Events (hipEventRecord(), hipEventElapsedTime(), etc.)
+- Kernel launching (hipLaunchKernel is a standard C/C++ function that replaces <<< >>>)
+- Cuda-style kernel indexing
+- 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.
+
+### What is not supported?
+#### Run-time features:
+- Textures 
+- Dynamic parallelism
+- Managed memory
+- 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
+- `__prof_trigger` 
+- PTX assembly
+- See the [HIP Kernel Language](hip_kernel_language.md) for more information.
+
+### How does HIP compare with OpenCL?
+Both AMD and Nvidia support OpenCL 1.2 on their devices, so developers can write portable code.
+HIP offers several benefits over OpenCL:
+- Developers can code in C++ as well as mix host and device C++ code in their source files. HIP C++ code can use templates, lambdas, classes and so on.
+- The HIP API is less verbose than OpenCL and is familiar to Cuda developers.
+- Because both Cuda and HIP are C++ languages, porting from Cuda to HIP is significantly easier than porting from Cuda to OpenCL.
+- HIP uses the best available development tools on each platform: on Nvidia GPUs, HIP code compiles using NVCC and can employ the nSight profiler and debugger (unlike OpenCL on Nvidia GPUs).
+- HIP provides pointers and host-side pointer arithmetic.
+- HIP provides device-level control over memory allocation and placement.
+- 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 Nvidia platforms, it should run on a device that uses the Cuda SDK 5.0 or newer. We have tested the Nvidia Titan and K40.
+- For NVIDIA platform, HIP should run on an device which runs the CUDA SDK 5.0 or newer.  We have tested nvidia Titan and K40.
+
+### Does Hipify automatically convert all source code?
+Typically, Hipify can automatically convert almost all run-time code, and the coordinate indexing device code. 
+Most device code needs no additional conversion, since HIP and Cuda have similar names for math and built-in functions. 
+HIP currently requires manual addition of one more arguments to the kernel so that the host can communicate the execution configuration to the device. 
+Additional porting may be required to deal with architecture feature queries or with Cuda capabilities that HIP doesnt support. 
+Developers should always expect to perform some platform-specific tuning and optimization.
+
+### What is NVCC?
+NVCC is Nvidia's compiler driver for compiling "Cuda C++" code into PTX or device code for Nvidia GPUs. Its 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.
+
+### Why use HIP rather than supporting Cuda run time 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.
+Developers who code to the HIP API can be assured there code will remain portable across NVIDIA and AMD platforms.
+
+### Can I develop HIP code on an NVIDIA CUDA platform?
+Yes!  HIP's CUDA path only exposes the APIs and functionality that work on both NVCC and HCC back-ends.
+"Extra" APIs, parameters, and features which exist in CUDA but not in HCC will typically result in compile-time or run-time errors.
+Developers need to use the HIP API for most accelerator code, and bracket any CUDA-specific code with appropriate ifdefs.
+Developers concerned about portability should of course run on both platforms, and should expect to tune for performance.
+In some cases CUDA has a richer set of modes for some APIs, and some C++ capabilities such as virtual functions - see the HIP @API documentation for more details.
+
+### Can I develop HIP code on an AMD HCC platform?
+Yes! HIP's HCC path only exposes the APIs and functions that work on both NVCC and HCC back ends. "Extra" APIs, parameters and features that appear in HCC but not Cuda will typically cause compile- or run-time errors. Developers must use the HIP API for most accelerator code and bracket any HCC-specific code with appropriate ifdefs. Those concerned about portability should, of course, test their code on both platforms and should tune it for performance. Typically, HCC supports a more modern set of C++11/C++14/C++17 features, so HIP developers who want portability should be careful when using advanced C++ features on the hc path.
+
+### Can a HIP binary run on both AMD and Nvidia platforms?
+HIP is a source-portable language that can be compiled to run on either the HCC or NVCC platform. HIP tools dont create a "fat binary" that can run on either platform, however.
+
+### What's the difference between HIP and hc?
+HIP is a portable C++ language that supports a strong subset of the Cuda run-time APIs and device-kernel language. Its designed to simplify Cuda conversion to portable C++. HIP provides a C-compatible run-time API, C-compatible kernel-launch mechanism, C++ kernel language and pointer-based memory management.
+
+A C++ dialect, hc is supported by the AMD HCC compiler. It provides C++ run time, C++ kernel-launch APIs (parallel_for_each), C++ kernel language, and several memory-management options, including pointers, arrays and array_view (with implicit data synchronization). Its intended to be a leading indicator of the ISO C++ standard.
@@ -0,0 +1,583 @@
+# HIP Kernel Language
+
+
+###Table of Contents
+=================
+
+  * [HIP Kernel Language](#hip-kernel-language" aria-hidden="true"><span aria-hidden="true)
+    * [Table of Contents](#table-of-contents" aria-hidden="true"><span aria-hidden="true)
+    * [Function-Type Qualifiers](#function-type-qualifiers" aria-hidden="true"><span aria-hidden="true)
+      * [__device__ ](#__device__" aria-hidden="true"><span aria-hidden="true)
+      * [__global__ ](#__global__" aria-hidden="true"><span aria-hidden="true)
+      * [__host__ ](#__host__" aria-hidden="true"><span aria-hidden="true)
+    * [Calling __global__ Functions](#calling-__global__-functions" aria-hidden="true"><span aria-hidden="true)
+    * [Kernel-Launch Example](#kernel-launch-example" aria-hidden="true"><span aria-hidden="true)
+    * [Variable-Type Qualifiers](#variable-type-qualifiers" aria-hidden="true"><span aria-hidden="true)
+      * [__constant__ ](#__constant__" aria-hidden="true"><span aria-hidden="true)
+      * [__shared__ ](#__shared__" aria-hidden="true"><span aria-hidden="true)
+      * [__managed__ ](#__managed__" aria-hidden="true"><span aria-hidden="true)
+      * [__restrict__ ](#__restrict__" aria-hidden="true"><span aria-hidden="true)
+    * [Built-In Variables](#built-in-variables" aria-hidden="true"><span aria-hidden="true)
+      * [Coordinate Built-Ins](#coordinate-built-ins" aria-hidden="true"><span aria-hidden="true)
+      * [warpSize](#warpsize" aria-hidden="true"><span aria-hidden="true)
+    * [Vector Types](#vector-types" aria-hidden="true"><span aria-hidden="true)
+      * [Short Vector Types](#short-vector-types" aria-hidden="true"><span aria-hidden="true)
+      * [dim3](#dim3" aria-hidden="true"><span aria-hidden="true)
+    * [Memory-Fence Instructions](#memory-fence-instructions" aria-hidden="true"><span aria-hidden="true)
+    * [Synchronization Functions](#synchronization-functions" aria-hidden="true"><span aria-hidden="true)
+    * [Math Functions](#math-functions" aria-hidden="true"><span aria-hidden="true)
+      * [Single Precision Mathematical Functions](#single-precision-mathematical-functions" aria-hidden="true"><span aria-hidden="true)
+      * [Double Precision Mathematical Functions](#double-precision-mathematical-functions" aria-hidden="true"><span aria-hidden="true)
+      * [Integer Intrinsics](#integer-intrinsics" aria-hidden="true"><span aria-hidden="true)
+    * [Texture Functions](#texture-functions" aria-hidden="true"><span aria-hidden="true)
+    * [Surface Functions](#surface-functions" aria-hidden="true"><span aria-hidden="true)
+    * [Timer Functions](#timer-functions" aria-hidden="true"><span aria-hidden="true)
+    * [Atomic Functions](#atomic-functions" aria-hidden="true"><span aria-hidden="true)
+      * [Caveats and Features Under-Development:](#caveats-and-features-under-development" aria-hidden="true"><span aria-hidden="true)
+    * [Warp Cross-Lane Functions](#warp-cross-lane-functions" aria-hidden="true"><span aria-hidden="true)
+      * [Warp Vote and Ballot Functions](#warp-vote-and-ballot-functions" aria-hidden="true"><span aria-hidden="true)
+      * [Warp Shuffle Functions](#warp-shuffle-functions" aria-hidden="true"><span aria-hidden="true)
+    * [Profiler Counter Function](#profiler-counter-function" aria-hidden="true"><span aria-hidden="true)
+    * [Assert](#assert" aria-hidden="true"><span aria-hidden="true)
+    * [Printf](#printf" aria-hidden="true"><span aria-hidden="true)
+    * [Device-Side Dynamic Global Memory Allocation](#device-side-dynamic-global-memory-allocation" aria-hidden="true"><span aria-hidden="true)
+    * [__launch_bounds__ ](#__launch_bounds__" aria-hidden="true"><span aria-hidden="true)
+    * [Register Keyword](#register-keyword" aria-hidden="true"><span aria-hidden="true)
+    * [Pragma Unroll](#pragma-unroll" aria-hidden="true"><span aria-hidden="true)
+    * [In-Line Assembly](#in-line-assembly" aria-hidden="true"><span aria-hidden="true)
+    * [C   Support](#c-support" aria-hidden="true"><span aria-hidden="true)
+
+
+
+HIP provides a C++ syntax that is suitable for compiling most code that commonly appears in compute kernels, including classes, namespaces, operator overloading, templates and more. Additionally, it defines other language features designed specifically to target accelerators, such as the following:
+- A kernel-launch syntax that uses standard C++, resembles a function call and is portable to all HIP targets
+- Short-vector headers that can serve on a host or a device
+- Math functions resembling those in the "math.h" header included with standard C++ compilers
+- Built-in functions for accessing specific GPU hardware capabilities
+
+This section describes the built-in variables and functions accessible from the HIP kernel. It’s intended for readers who are familiar with Cuda kernel syntax and want to understand how HIP is different.
+
+Features are marked with one of the following keywords:
+- **Supported**---HIP supports the feature with a Cuda-equivalent function
+- **Not supported**---HIP does not support the feature
+- **Under development**---the feature is under development but not yet available
+
+## Function-Type Qualifiers
+### `__device__`
+Supported  `__device__` functions are
+  - Executed on the device
+  - Called from the device only
+
+The `__device__` keyword can combine with the host keyword (see [__host__](#host)).
+
+### `__global__`
+Supported `__global__` functions are
+  - Executed on the device
+  - Called ("launched") from the host
+
+HIP `__global__` functions must have a `void` return type, and the first parameter to a HIP `__global__` function must have the type `hipLaunchParm`. See [Kernel-Launch Example](#kernel-launch-example). 
+
+HIP lacks dynamic-parallelism support, so `__global__ ` functions cannot be called from the device.
+
+### `__host__`
+Supported `__host__` functions are
+  - Executed on the host
+  - Called from the host
+
+`__host__` can combine with `__device__`, in which case the function compiles for both the host and device. These functions cannot use the HIP grid coordinate functions (for example, "hipThreadIdx_x"). A possible workaround is to pass the necessary coordinate info as an argument to the function.
+
+`__host__` cannot combine with `__global__`.
+
+HIP parses the `__noinline__` and `__forceinline__` keywords and converts them to the appropriate Clang attributes. The hcc compiler, however, currently in-lines all device functions, so they are effectively ignored.
+
+
+## Calling `__global__` Functions
+
+`__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
+   - **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
+
+
+```
+// Example psuedocode introducing hipLaunchKernel:
+__global__ MyKernel(hipLaunchParm lp, float *A, float *B, float *C, size_t N)
+{
+...
+}
+
+// Replace MyKernel<<<dim3(gridDim), dim3(gridDim), 0, 0>>> (a,b,c,n);
+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. The Hipify script automatically converts 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).
+
+
+## Kernel-Launch Example
+```
+// Example showing device function, __device__ __host__   
+// <- compile for both device and host 
+float PlusOne(float x) 
+{
+    return x + 1.0;
+}
+
+__global__ 
+void 
+MyKernel (hipLaunchParm lp, /*lp parm for execution configuration */
+          const float *a, const float *b, float *c, unsigned N)
+{
+    unsigned gid = hipThreadIdx_x; // <- coordinate index function
+    if (gid < N) {
+        c[gid] = a[gid] + PlusOne(b[gid]);
+    }
+}
+void callMyKernel()
+{
+    float *a, *b, *c; // initialization not shown...
+    unsigned N = 1000000;
+    const unsigned blockSize = 256;
+
+    hipLaunchKernel(MyKernel, dim3(N/blockSize), dim3(blockSize), 0, 0,  a,b,c,N);
+}
+```
+
+## Variable-Type Qualifiers
+
+### `__constant__`
+The `__constant__` keyword is supported. The host writes constant memory before launching the kernel; from the GPU, this memory is read-only during kernel execution. The functions for accessing constant memory (hipGetSymbolAddress(), hipGetSymbolSize(), hipMemcpyToSymbol(), hipMemcpyToSymbolAsync, hipMemcpyFromSymbol, hipMemcpyFromSymbolAsync) are under development.
+
+### `__shared__` 
+The `__shared__` keyword is supported.
+
+`extern __shared__` allows the host to dynamically allocate shared memory and is specified as a launch parameter. This feature is under development.
+
+### `__managed__`
+Managed memory, including the `__managed__` keyword, are not supported in HIP.
+
+### `__restrict__`
+The `__restrict__` keyword tells the compiler that the associated memory pointer will not alias with any other pointer in the kernel or function.  This feature can help the compiler generate better code. In most cases, all pointer arguments must use this keyword to realize the benefit. hcc support for the `__restrict__` qualifier on kernel arguments is under development.  
+
+
+## Built-In Variables
+
+### Coordinate Built-Ins
+These built-ins determine the coordinate of the active work item in the execution grid. They are defined in hip_runtime.h (rather than being implicitly defined by the compiler).   
+
+| **HIP Syntax** | **Cuda Syntax** |
+| --- | --- |
+| hipThreadIdx_x | threadIdx.x |
+| hipThreadIdx_y | threadIdx.y |
+| hipThreadIdx_z | threadIdx.z |
+|                |             |
+| hipBlockIdx_x  | blockIdx.x  |
+| hipBlockIdx_y  | blockIdx.y  |
+| hipBlockIdx_z  | blockIdx.z  |
+|                |             |
+| hipBlockDim_x  | blockDim.x  |
+| hipBlockDim_y  | blockDim.y  |
+| hipBlockDim_z  | blockDim.z  |
+|                |             |
+| hipGridDim_x   | gridDim.x   |
+| hipGridDim_y   | gridDim.y   |
+| hipGridDim_z   | gridDim.z   |
+
+### 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.
+
+
+## Vector Types
+
+Note that these types are defined in hip_runtime.h and are not automatically provided by the compiler. 
+
+
+### Short Vector Types
+Short vector types derive from the basic integer and floating-point types. They are structures defined in hip_vector_types.h. The first, second, third and fourth components of the vector are accessible through the ```x```, ```y```, ```z``` and ```w``` fields, respectively. All the short vector types support a constructor function of the form ```make_<type_name>()```. For example, ```float4 make_float4(float x, float y, float z, float w)``` creates a vector of type ```float4``` and value ```(x,y,z,w)```.
+
+HIP supports the following short vector formats:
+- Signed Integers:
+    - char1, char2, char3, char4
+    - short1, short2, short3, short4
+    - int1, int2, int3, int4
+    - long1, long2, long3, long4
+    - longlong1, longlong2, longlong3, longlong4
+- Unsigned Integers:
+    - uchar1, uchar2, uchar3, uchar4
+    - ushort1, ushort2, ushort3, ushort4
+    - uint1, uint2, uint3, uint4
+    - ulong1, ulong2, ulong3, ulong4
+    - ulonglong1, ulonglong2, ulonglong3, ulonglong4
+- Floating Points
+    - float1, float2, float3, float4
+    - double1, double2, double3, double4
+
+### dim3
+dim3 is a three-dimensional integer vector type commonly used to specify grid and group dimensions. Unspecified dimensions are initialized to 1. 
+```
+typedef struct dim3 {
+  uint32_t x; 
+  uint32_t y; 
+  uint32_t z; 
+
+  dim3(uint32_t _x=1, uint32_t _y=1, uint32_t _z=1) : x(_x), y(_y), z(_z) {};
+};
+
+```
+
+## Memory-Fence Instructions
+HIP support for __threadfence(), __threadfence_block() and __threadfence_system() is under development.
+
+## Synchronization Functions
+The __syncthreads() built-in function is supported in HIP. The __syncthreads_count(int), __syncthreads_and(int) and __syncthreads_or(int) functions are under development.  
+
+## Math Functions
+hcc supports a set of math operations callable from the device.
+
+### Single Precision Mathematical Functions
+Following is the list of supported single precision mathematical functions.
+
+| **Function** | **Supported on Host** | **Supported on Device** |
+| --- | --- | --- |
+| float acosf ( float  x ) <br><sub>Calculate the arc cosine of the input argument.</sub> | ✓ | ✓ |
+| float acoshf ( float  x ) <br><sub>Calculate the nonnegative arc hyperbolic cosine of the input argument.</sub> | ✓ | ✓ |
+| float asinf ( float  x ) <br><sub>Calculate the arc sine of the input argument.</sub> | ✓ | ✓ |
+| float asinhf ( float  x ) <br><sub>Calculate the arc hyperbolic sine of the input argument.</sub> | ✓ | ✓ |
+| float atan2f ( float  y, float  x ) <br><sub>Calculate the arc tangent of the ratio of first and second input arguments.</sub> | ✓ | ✓ |
+| float atanf ( float  x ) <br><sub>Calculate the arc tangent of the input argument.</sub> | ✓ | ✓ |
+| float atanhf ( float  x ) <br><sub>Calculate the arc hyperbolic tangent of the input argument.</sub> | ✓ | ✓ |
+| float cbrtf ( float  x ) <br><sub>Calculate the cube root of the input argument.</sub> | ✓ | ✓ |
+| float ceilf ( float  x ) <br><sub>Calculate ceiling of the input argument.</sub> | ✓ | ✓ |
+| float copysignf ( float  x, float  y ) <br><sub>Create value with given magnitude, copying sign of second value.</sub> | ✓ | ✓ |
+| float cosf ( float  x ) <br><sub>Calculate the cosine of the input argument.</sub> | ✓ | ✓ |
+| float coshf ( float  x ) <br><sub>Calculate the hyperbolic cosine of the input argument.</sub> | ✓ | ✓ |
+| float erfcf ( float  x ) <br><sub>Calculate the complementary error function of the input argument.</sub> | ✓ | ✓ |
+| float erff ( float  x ) <br><sub>Calculate the error function of the input argument.</sub> | ✓ | ✓ |
+| float exp10f ( float  x ) <br><sub>Calculate the base 10 exponential of the input argument.</sub> | ✓ | ✓ |
+| float exp2f ( float  x ) <br><sub>Calculate the base 2 exponential of the input argument.</sub> | ✓ | ✓ |
+| float expf ( float  x ) <br><sub>Calculate the base e exponential of the input argument.</sub> | ✓ | ✓ |
+| float expm1f ( float  x ) <br><sub>Calculate the base e exponential of the input argument, minus 1.</sub> | ✓ | ✓ |
+| float fabsf ( float  x ) <br><sub>Calculate the absolute value of its argument.</sub> | ✓ | ✓ |
+| float fdimf ( float  x, float  y ) <br><sub>Compute the positive difference between `x` and `y`.</sub> | ✓ | ✓ |
+| float floorf ( float  x ) <br><sub>Calculate the largest integer less than or equal to `x`.</sub> | ✓ | ✓ |
+| float fmaf ( float  x, float  y, float  z ) <br><sub>Compute `x × y + z` as a single operation.</sub> | ✓ | ✓ |
+| float fmaxf ( float  x, float  y ) <br><sub>Determine the maximum numeric value of the arguments.</sub> | ✓ | ✓ |
+| float fminf ( float  x, float  y ) <br><sub>Determine the minimum numeric value of the arguments.</sub> | ✓ | ✓ |
+| float fmodf ( float  x, float  y ) <br><sub>Calculate the floating-point remainder of `x / y`.</sub> | ✓ | ✓ |
+| float frexpf ( float  x, int* nptr ) <br><sub>Extract mantissa and exponent of a floating-point value.</sub> | ✓ | ✗ |
+| float hypotf ( float  x, float  y ) <br><sub>Calculate the square root of the sum of squares of two arguments.</sub> | ✓ | ✓ |
+| int ilogbf ( float  x ) <br><sub>Compute the unbiased integer exponent of the argument.</sub> | ✓ | ✓ |
+| __RETURN_TYPE<sup id="a1">[1](#f1)</sup> isfinite ( float  a ) <br><sub>Determine whether argument is finite.</sub> | ✓ | ✓ |
+| __RETURN_TYPE<sup>[1](#f1)</sup> isinf ( float  a ) <br><sub>Determine whether argument is infinite.</sub> | ✓ | ✓ |
+| __RETURN_TYPE<sup>[1](#f1)</sup> isnan ( float  a ) <br><sub>Determine whether argument is a NaN.</sub> | ✓ | ✓ |
+| float ldexpf ( float  x, int  exp ) <br><sub>Calculate the value of x ⋅ 2<sup>exp</sup>.</sub> | ✓ | ✓ |
+| float log10f ( float  x ) <br><sub>Calculate the base 10 logarithm of the input argument.</sub> | ✓ | ✓ |
+| float log1pf ( float  x ) <br><sub>Calculate the value of log<sub>e</sub>( 1 + x ).</sub> | ✓ | ✓ |
+| float logbf ( float  x ) <br><sub>Calculate the floating point representation of the exponent of the input argument.</sub> | ✓ | ✓ |
+| float log2f ( float  x ) <br><sub>Calculate the base 2 logarithm of the input argument.</sub> | ✓ | ✓ | 
+| float logf ( float  x ) <br><sub>Calculate the natural logarithm of the input argument.</sub> | ✓ | ✓ |
+| float modff ( float  x, float* iptr ) <br><sub>Break down the input argument into fractional and integral parts.</sub> | ✓ | ✗ |
+| float nanf ( const char* tagp ) <br><sub>Returns "Not a Number" value.</sub> | ✗ | ✓ |
+| float nearbyintf ( float  x ) <br><sub>Round the input argument to the nearest integer.</sub> | ✓ | ✓ |
+| float powf ( float  x, float  y ) <br><sub>Calculate the value of first argument to the power of second argument.</sub> | ✓ | ✓ |
+| float remainderf ( float  x, float  y ) <br><sub>Compute single-precision floating-point remainder.</sub> | ✓ | ✓ |
+| float remquof ( float  x, float  y, int* quo ) <br><sub>Compute single-precision floating-point remainder and part of quotient.</sub> | ✓ | ✗ |
+| float roundf ( float  x ) <br><sub>Round to nearest integer value in floating-point.</sub> | ✓ | ✓ |
+| float scalbnf ( float  x, int  n ) <br><sub>Scale floating-point input by integer power of two.</sub> | ✓ | ✓ |
+| __RETURN_TYPE<sup>[1](#f1)</sup> signbit ( float  a ) <br><sub>Return the sign bit of the input.</sub> | ✓ | ✓ |
+| void sincosf ( float  x, float* sptr, float* cptr ) <br><sub>Calculate the sine and cosine of the first input argument.</sub> | ✓ | ✗ |
+| float sinf ( float  x ) <br><sub>Calculate the sine of the input argument.</sub> | ✓ | ✓ |
+| float sinhf ( float  x ) <br><sub>Calculate the hyperbolic sine of the input argument.</sub> | ✓ | ✓ |
+| float sqrtf ( float  x ) <br><sub>Calculate the square root of the input argument.</sub> | ✓ | ✓ |
+| float tanf ( float  x ) <br><sub>Calculate the tangent of the input argument.</sub> | ✓ | ✓ |
+| float tanhf ( float  x ) <br><sub>Calculate the hyperbolic tangent of the input argument.</sub> | ✓ | ✓ |
+| float truncf ( float  x ) <br><sub>Truncate input argument to the integral part.</sub> | ✓ | ✓ |
+| float tgammaf ( float  x ) <br><sub>Calculate the gamma function of the input argument.</sub> | ✓ | ✓ |
+<sub><b id="f1"><sup>[1]</sup></b> __RETURN_TYPE is dependent on compiler. It is usually 'int' for C compilers and 'bool' for C++ compilers.</sub> [↩](#a1)
+
+### Double Precision Mathematical Functions
+Following is the list of supported double precision mathematical functions.
+
+| **Function** | **Supported on Host** | **Supported on Device** |
+| --- | --- | --- |
+| double acos ( double  x ) <br><sub>Calculate the arc cosine of the input argument.</sub> | ✓ | ✓ |
+| double acosh ( double  x ) <br><sub>Calculate the nonnegative arc hyperbolic cosine of the input argument.</sub> | ✓ | ✓ |
+| double asin ( double  x ) <br><sub>Calculate the arc sine of the input argument.</sub> | ✓ | ✓ |
+| double asinh ( double  x ) <br><sub> Calculate the arc hyperbolic sine of the input argument.</sub> | ✓ | ✓ |
+| double atan ( double  x ) <br><sub>Calculate the arc tangent of the input argument.</sub> | ✓ | ✓ |
+| double atan2 ( double  y, double  x ) <br><sub>Calculate the arc tangent of the ratio of first and second input arguments.</sub> | ✓ | ✓ |
+| double atanh ( double  x ) <br><sub>Calculate the arc hyperbolic tangent of the input argument.</sub> | ✓ | ✓ |
+| double cbrt ( double  x ) <br><sub>Calculate the cube root of the input argument.</sub> | ✓ | ✓ |
+| double ceil ( double  x ) <br><sub>Calculate ceiling of the input argument.</sub> | ✓ | ✓ |
+| double copysign ( double  x, double  y ) <br><sub>Create value with given magnitude, copying sign of second value.</sub> | ✓ | ✓ |
+| double cos ( double  x ) <br><sub>Calculate the cosine of the input argument.</sub> | ✓ | ✓ |
+| double cosh ( double  x ) <br><sub>Calculate the hyperbolic cosine of the input argument.</sub> | ✓ | ✓ |
+| double erf ( double  x ) <br><sub>Calculate the error function of the input argument.</sub> | ✓ | ✓ |
+| double erfc ( double  x ) <br><sub>Calculate the complementary error function of the input argument.</sub> | ✓ | ✓ |
+| double exp ( double  x ) <br><sub>Calculate the base e exponential of the input argument.</sub> | ✓ | ✓ |
+| double exp10 ( double  x ) <br><sub>Calculate the base 10 exponential of the input argument.</sub> | ✓ | ✓ |
+| double exp2 ( double  x ) <br><sub>Calculate the base 2 exponential of the input argument.</sub> | ✓ | ✓ |
+| double expm1 ( double  x ) <br><sub>Calculate the base e exponential of the input argument, minus 1.</sub> | ✓ | ✓ |
+| double fabs ( double  x ) <br><sub>Calculate the absolute value of the input argument.</sub> | ✓ | ✓ |
+| double fdim ( double  x, double  y ) <br><sub>Compute the positive difference between `x` and `y`.</sub> | ✓ | ✓ | 
+| double floor ( double  x ) <br><sub>Calculate the largest integer less than or equal to `x`.</sub> | ✓ | ✓ |
+| double fma ( double  x, double  y, double  z ) <br><sub>Compute `x × y + z` as a single operation.</sub> | ✓ | ✓ |
+| double fmax ( double , double ) <br><sub>Determine the maximum numeric value of the arguments.</sub> | ✓ | ✓ |
+| double fmin ( double  x, double  y ) <br><sub>Determine the minimum numeric value of the arguments.</sub> | ✓ | ✓ |
+| double fmod ( double  x, double  y ) <br><sub>Calculate the floating-point remainder of `x / y`.</sub> | ✓ | ✓ |
+| double frexp ( double  x, int* nptr ) <br><sub>Extract mantissa and exponent of a floating-point value.</sub> | ✓ | ✗ |
+| double hypot ( double  x, double  y ) <br><sub>Calculate the square root of the sum of squares of two arguments.</sub> | ✓ | ✓ |
+| int ilogb ( double  x ) <br><sub>Compute the unbiased integer exponent of the argument.</sub> | ✓ | ✓ |
+| __RETURN_TYPE<sup id="a2">[1](#f2)</sup> isfinite ( double  a ) <br><sub>Determine whether argument is finite.</sub> | ✓ | ✓ |
+| __RETURN_TYPE<sup>[1](#f2)</sup> isinf ( double  a ) <br><sub>Determine whether argument is infinite.</sub> | ✓ | ✓ |
+| __RETURN_TYPE<sup>[1](#f2)</sup> isnan ( double  a ) <br><sub>Determine whether argument is a NaN.</sub> | ✓ | ✓ |
+| double ldexp ( double  x, int  exp ) <br><sub>Calculate the value of x ⋅ 2<sup>exp</sup>.</sub> | ✓ | ✓ |
+| double log ( double  x ) <br><sub>Calculate the base e logarithm of the input argument.</sub> | ✓ | ✓ |
+| double log10 ( double  x ) <br><sub>Calculate the base 10 logarithm of the input argument.</sub> | ✓ | ✓ |
+| double log1p ( double  x ) <br><sub>Calculate the value of log<sub>e</sub>( 1 + x ).</sub> | ✓ | ✓ |
+| double log2 ( double  x ) <br><sub>Calculate the base 2 logarithm of the input argument.</sub> | ✓ | ✓ |
+| double logb ( double  x ) <br><sub>Calculate the floating point representation of the exponent of the input argument.</sub> | ✓ | ✓ |
+| double modf ( double  x, double* iptr ) <br><sub>Break down the input argument into fractional and integral parts.</sub> | ✓ | ✗ |
+| double nan ( const char* tagp ) <br><sub>Returns "Not a Number" value.</sub> | ✗ | ✓ |
+| double nearbyint ( double  x ) <br><sub>Round the input argument to the nearest integer.</sub> | ✓ | ✓ |
+| double pow ( double  x, double  y ) <br><sub>Calculate the value of first argument to the power of second argument.</sub> | ✓ | ✓ |
+| double remainder ( double  x, double  y ) <br><sub>Compute double-precision floating-point remainder.</sub> | ✓ | ✓ |
+| double remquo ( double  x, double  y, int* quo ) <br><sub>Compute double-precision floating-point remainder and part of quotient.</sub> | ✓ | ✗ |
+| double round ( double  x ) <br><sub>Round to nearest integer value in floating-point.</sub> | ✓ | ✓ |
+| double scalbn ( double  x, int  n ) <br><sub>Scale floating-point input by integer power of two.</sub> | ✓ | ✓ |
+| __RETURN_TYPE<sup>[1](#f2)</sup> signbit ( double  a ) <br><sub>Return the sign bit of the input.</sub> | ✓ | ✓ |
+| double sin ( double  x ) <br><sub>Calculate the sine of the input argument.</sub> | ✓ | ✓ |
+| void sincos ( double  x, double* sptr, double* cptr ) <br><sub>Calculate the sine and cosine of the first input argument.</sub> | ✓ | ✗ |
+| double sinh ( double  x ) <br><sub>Calculate the hyperbolic sine of the input argument.</sub> | ✓ | ✓ |
+| double sqrt ( double  x ) <br><sub>Calculate the square root of the input argument.</sub> | ✓ | ✓ |
+| double tan ( double  x ) <br><sub>Calculate the tangent of the input argument.</sub> | ✓ | ✓ |
+| double tanh ( double  x ) <br><sub>Calculate the hyperbolic tangent of the input argument.</sub> | ✓ | ✓ |
+| double tgamma ( double  x ) <br><sub>Calculate the gamma function of the input argument.</sub> | ✓ | ✓ |
+| double trunc ( double  x ) <br><sub>Truncate input argument to the integral part.</sub> | ✓ | ✓ |
+<sub><b id="f2"><sup>[1]</sup></b> __RETURN_TYPE is dependent on compiler. It is usually 'int' for C compilers and 'bool' for C++ compilers.</sub> [↩](#a2)
+
+### Integer Intrinsics
+Following is the list of supported integer intrinsics.
+
+| **Function** | **Supported on Host** | **Supported on Device** |
+| --- | --- | --- |
+| unsigned int __brev ( unsigned int x ) <br><sub>Reverse the bit order of a 32 bit unsigned integer.</sub> | ✓ | ✓ |
+| unsigned long long int __brevll ( unsigned long long int x ) <br><sub>Reverse the bit order of a 64 bit unsigned integer. </sub> | ✓ | ✓ |
+| int __clz ( int  x ) <br><sub>Return the number of consecutive high-order zero bits in a 32 bit integer.</sub> | ✓ | ✓ |
+| unsigned int __clz(unsigned int x) <br><sub>Return the number of consecutive high-order zero bits in 32 bit unsigned integer.</sub> | ✓ | ✗ |  
+| int __clzll ( long long int x ) <br><sub>Count the number of consecutive high-order zero bits in a 64 bit integer.</sub> | ✓ | ✓ |
+| unsigned int __clzll(long long int x) <br><sub>Return the number of consecutive high-order zero bits in 64 bit signed integer.</sub> | ✓ | ✗ |
+| unsigned int __ffs(unsigned int x) <br><sub>Find the position of least signigicant bit set to 1 in a 32 bit unsigned integer.<sup id="a3">[1](#f3)</sup></sub> | ✓ | ✓|
+| unsigned int __ffs(int x) <br><sub>Find the position of least signigicant bit set to 1 in a 32 bit signed integer.</sub> | ✗ | ✓ |
+| unsigned int __ffsll(unsigned long long int x) <br><sub>Find the position of least signigicant bit set to 1 in a 64 bit unsigned integer.<sup>[1](#f3)</sup></sub> | ✓ | ✓ |
+| unsigned int __ffsll(long long int x) <br><sub>Find the position of least signigicant bit set to 1 in a 64 bit signed integer.</sub> | ✗ | ✓ |
+| unsigned int __popc ( unsigned int x ) <br><sub>Count the number of bits that are set to 1 in a 32 bit integer.</sub> | ✓ | ✓ |
+| int __popcll ( unsigned long long int x )<br><sub>Count the number of bits that are set to 1 in a 64 bit integer.</sub> | ✓ | ✓ |
+<sub><b id="f3"><sup>[1]</sup></b> 
+The hcc implementation of __ffs() and __ffsll() contains code to add a constant +1 to produce the ffs result format.
+For the cases where this overhead is not acceptable and programmer is willing to specialize for the platform, 
+hcc provides hc::__lastbit_u32_u32(unsigned int input) and  hc::__lastbit_u32_u64(unsigned long long int input).
+The index returned by __lastbit_ instructions starts at 0, while for ffs the index starts at 1.  
+
+
+## Texture Functions
+Texture functions are not supported.
+
+## Surface Functions
+Surface functions are not supported.
+
+## Timer Functions
+HIP provides the following built-in functions for reading a high-resolution timer from the device.
+```
+clock_t clock()
+long long int clock64()
+```
+
+AMD devices employ a per-GPU timer that increments at a constant time interval regardless of any dynamic frequency changes. All compute units in the system share the timer. 
+Nvidia devices implement the timer as a per-compute-unit clock that increments on every clock cycle.  
+
+To obtain the clock frequency, use the hipDeviceProp_t.clockInstructionRate field:
+
+```
+hipDeviceGetProperties(&deviceProps, deviceId);
+// Compute time in ms--device_ticks is based on values reported from clock() device function
+float time = device_ticks / (float)deviceProps.clockInstructionRate;
+```
+
+## Atomic Functions
+
+Atomic functions execute as read-modify-write operations residing in global or shared memory. No other device or thread can observe or modify the memory location during an atomic operation. If multiple instructions from different devices or threads target the same memory location, the 
+instructions are serialized in an undefined order.  
+
+HIP supports the following atomic operations.
+
+| **Function** | **Supported in HIP** | **Supported in CUDA** |
+| --- | --- | --- |
+| int atomicAdd(int* address, int val) | ✓ | ✓ |
+| unsigned int atomicAdd(unsigned int* address,unsigned int val) | ✓ | ✓ |
+| unsigned long long int atomicAdd(unsigned long long int* address,unsigned long long int val) | ✓ | ✓ |
+| float atomicAdd(float* address, float val) | ✓ | ✓ |
+| int atomicSub(int* address, int val) | ✓ | ✓ |
+| unsigned int atomicSub(unsigned int* address,unsigned int val) | ✓ | ✓ |
+| int atomicExch(int* address, int val) | ✓ | ✓ |
+| unsigned int atomicExch(unsigned int* address,unsigned int val) | ✓ | ✓ |
+| unsigned long long int atomicExch(unsigned long long int* address,unsigned long long int val) | ✓ | ✓ |
+| float atomicExch(float* address, float val) | ✓ | ✓ |
+| int atomicMin(int* address, int val) | ✓ | ✓ |
+| unsigned int atomicMin(unsigned int* address,unsigned int val) | ✓ | ✓ |
+| unsigned long long int atomicMin(unsigned long long int* address,unsigned long long int val) | ✓ | ✓ |
+| int atomicMax(int* address, int val) | ✓ | ✓ |
+| unsigned int atomicMax(unsigned int* address,unsigned int val) | ✓ | ✓ |
+| unsigned long long int atomicMax(unsigned long long int* address,unsigned long long int val) | ✓ | ✓ |
+| unsigned int atomicInc(unsigned int* address)| ✓ <br><sub>Takes one argument.</sub> | ✓ <br><sub> Wrapping increment,takes two arguments.</sub> |
+| unsigned int atomicDec(unsigned int* address)| ✓ <br><sub>Takes one argument.</sub> | ✓ <br><sub> Wrapping decrement,takes two arguments.</sub> |
+| int atomicCAS(int* address, int compare, int val) | ✓ | ✓ |
+| unsigned int atomicCAS(unsigned int* address,unsigned int compare,unsigned int val) | ✓ | ✓ |
+| unsigned long long int atomicCAS(unsigned long long int* address,unsigned long long int compare,unsigned long long int val) | ✓ | ✓ |
+| int atomicAnd(int* address, int val) | ✓ | ✓ |
+| unsigned int atomicAnd(unsigned int* address,unsigned int val) | ✓ | ✓ |
+| unsigned long long int atomicAnd(unsigned long long int* address,unsigned long long int val) | ✓ | ✓ |
+| int atomicOr(int* address, int val) | ✓ | ✓ |
+| unsigned int atomicOr(unsigned int* address,unsigned int val) | ✓ | ✓ |
+| unsigned long long int atomicOr(unsigned long long int* address,unsigned long long int val) | ✓ | ✓ |
+| int atomicXor(int* address, int val) | ✓ | ✓ |
+| unsigned int atomicXor(unsigned int* address,unsigned int val) | ✓ | ✓ |
+| unsigned long long int atomicXor(unsigned long long int* address,unsigned long long int val)) | ✓ | ✓ |
+
+### Caveats and Features Under-Development:
+
+- HIP enables atomic operations on 32-bit integers. Additionally, it supports an atomic float add. AMD hardware, however, implements the float add using a CAS loop, so this function may not perform efficiently.
+- hcc currently maps `__shared__` atomics to `__device__` atomics. Optimal support is under development.
+- wrapping increment and decrement are under development.
+
+## Warp Cross-Lane Functions
+
+Warp cross-lane functions operate across all lanes in a warp. The hardware guarantees that all warp lanes will execute in lockstep, so additional synchronization is unnecessary, and the instructions use no shared memory.
+
+Note that Nvidia and AMD devices have different warp sizes, so portable code should use the warpSize built-ins to query the warp size. Hipified code from the Cuda path requires careful review to ensure it doesn’t assume a waveSize of 32. "Wave-aware" code that assumes a waveSize of 32 will run on a wave-64 machine, but it will utilize only half of the machine resources. In addition to the warpSize device function, host code can obtain the warpSize from the device properties:
+
+```
+	cudaDeviceProp props;
+	cudaGetDeviceProperties(&props, deviceID);
+    int w = props.warpSize;  
+    // implement portable algorithm based on w (rather than assume 32 or 64)
+```
+
+
+### Warp Vote and Ballot Functions
+
+```
+int __all(int predicate)
+int __any(int predicate)
+uint64_t __ballot(int predicate)
+```
+
+Threads in a warp are referred to as *lanes* and are numbered from 0 to warpSize -- 1. For these functions, each warp lane contributes 1 -- the bit value (the predicate), which is efficiently broadcast to all lanes in the warp. The 32-bit int predicate from each lane reduces to a 1-bit value: 0 (predicate = 0) or 1 (predicate != 0). `__any` and `__all` provide a summary view of the predicates that the other warp lanes contribute:
+
+- `__any()` returns 1 if any warp lane contributes a nonzero predicate, or 0 otherwise
+- `__all()` returns 1 if all other warp lanes contribute nonzero predicates, or 0 otherwise
+
+Applications can test whether the target platform supports the any/all instruction using the `hasWarpVote` device property or the HIP_ARCH_HAS_WARP_VOTE compiler define.
+
+`__ballot` provides a bit mask containing the 1-bit predicate value from each lane. The nth bit of the result contains the 1 bit contributed by the nth warp lane. Note that HIP's `__ballot` function supports a 64-bit return value (compared with Cuda’s 32 bits). Code ported from Cuda should support the larger warp sizes that the HIP version of this instruction supports. Applications can test whether the target platform supports the ballot instruction using the `hasWarpBallot` device property or the HIP_ARCH_HAS_WARP_BALLOT compiler define.
+
+
+### Warp Shuffle Functions 
+
+The following warp shuffle instructions are under development.
+
+Half-float shuffles are not supported. The default width is warpSize---see [Warp Cross-Lane Functions](#warp-cross-lane-functions). Applications should not assume the warpSize is 32 or 64.
+
+```
+int   __shfl      (int var,   int srcLane, int width=warpSize);
+float __shfl      (float var, int srcLane, int width=warpSize);
+int   __shfl_up   (int var,   unsigned int delta, int width=warpSize);
+float __shfl_up   (float var, unsigned int delta, int width=warpSize);
+int   __shfl_down (int var,   unsigned int delta, int width=warpSize);
+float __shfl_down (float var, unsigned int delta, int width=warpSize) ;
+int   __shfl_xor  (int var,   int laneMask, int width=warpSize) 
+float __shfl_xor  (float var, int laneMask, int width=warpSize);
+
+```
+
+## Profiler Counter Function
+
+The Cuda `__prof_trigger()` instruction is not supported.
+
+## Assert
+
+The assert function is under development.
+
+## Printf
+
+The printf function is under development.
+
+## Device-Side Dynamic Global Memory Allocation
+
+Device-side dynamic global memory allocation is not supported.
+
+## `__launch_bounds__`
+GPU multiprocessors have a fixed pool of resources (primarily registers and shared memory) that are shared among the active warps. Using more resources can increase the kernel’s IPC, but it reduces the resources available for other warps and limits the number of warps that can run simultaneously. Thus, GPUs exhibit a complex relationship between resource usage and performance. `__launch_bounds__` allows the application to provide usage hints that influence the resources (primarily registers) employed by the generated code. It’s a function attribute that must be attached to a `__global__` function:
+
+```
+__global__ void
+`__launch_bounds__`(requiredMaxThreadsPerBlock, minBlocksPerMultiprocessor)  
+MyKernel(hipGridLaunch lp, ...) 
+...
+```
+
+`__launch_bounds__` supports two parameters:
+
+- **requiredMaxThreadsPerBlock**---the programmer guarantees that the kernel will launch with threadsPerBlock less than requiredMaxThreadsPerBlock. (In nvcc, this parameter maps to the _.maxntid_ PTX directive; in hcc, it maps to the HSAIL _requiredworkgroupsize_ directive.) If launch_bounds is unspecified, requiredMaxThreadsPerBlock is the maximum block size that the device supports (typically 1,024 or larger). Specifying requiredMaxThreadsPerBlock less than the maximum effectively allows the compiler to use more resources than a default unconstrained compilation supporting all possible block sizes at launch time. The threadsPerBlock value is the product hipBlockDim_x * hipBlockDim_y * hipBlockDim_z.
+- **minBlocksPerMultiprocessor**---directs the compiler to minimize resource usage so that the requested number of blocks can be simultaneously active on a multiprocessor. Because active blocks compete for the same fixed resource pool, the compiler must reduce the resource requirements of each block (primarily registers). minBlocksPerMultiprocessor is optional and defaults to 1 if unspecified. Selecting a minBlocksPerMultiprocessor value greater than 1 effectively constrains the compiler's resource usage. 
+
+The compiler uses these two parameters as follows:
+- It employs the hints only to manage register usage and does not automatically reduce shared memory or other resources.
+- Compilation fails if the compiler cannot generate a kernel that meets the requirements of the specified launch bounds.
+- From requiredMaxThreadsPerBlock, the compiler derives the maximum number of warps per block that are usable at launch time. Values less than the default allow the compiler to use a larger register pool: each warp uses registers, and this hint constrains the launch to a warps-per-block size less than maximum. 
+- From minBlocksPerMultiprocessor, the compiler derives a maximum number of registers that the kernel can use (to meet the required number of simultaneously active blocks). If the value is 1, the kernel can use all registers supported by the multiprocessor.  
+
+The compiler ensures that the kernel uses fewer registers than both allowed maxima specify, typically by spilling to shared memory or using more instructions. It may use heuristics to increase register usage or may simply be able to avoid spilling. The requiredMaxThreadsPerBlock parameter is particularly useful in this case, since it allows the compiler to use more registers---avoiding situations where the compiler constrains the register usage (potentially spilling) to meet the requirements of a large block size never sees use at launch time.  
+
+HIP/hcc will parse the `launch_bounds` attribute but silently ignores the performance hint. Full support is under development.
+
+
+## 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.
+
+
+## Pragma Unroll
+
+hcc support for the unroll pragma is under development and is slated to arrive with the Lightning Compiler.
+
+
+```
+#pragma unroll 16 /* hint to compiler to unroll next loop by 16 */
+for (int i=0; i<16; i++) ...
+```
+
+```
+#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
+
+In-line assembly, including in-line PTX, in-line HSAIL and in-line GCN ISA, is not supported. Users who need these features should employ conditional compilation to provide different functionally equivalent implementations on each target platform.
+
+
+## C++ Support
+The following C++ features are not supported:
+- Run-time-type information (RTTI)
+- Virtual functions
+- Try/catch
@@ -0,0 +1,466 @@
+# 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
+=================
+
+ * [HIP Porting Guide](#hip-porting-guide" aria-hidden="true"><span aria-hidden="true)
+      * [Table of Contents](#table-of-contents" aria-hidden="true"><span aria-hidden="true)
+    * [Porting a New Cuda Project](#porting-a-new-cuda-project" aria-hidden="true"><span aria-hidden="true)
+      * [General Tips](#general-tips" aria-hidden="true"><span aria-hidden="true)
+      * [Scanning existing CUDA code to scope the porting effort](#scanning-existing-cuda-code-to-scope-the-porting-effort" aria-hidden="true"><span aria-hidden="true)
+    * [Distinguishing Compiler Modes](#distinguishing-compiler-modes" aria-hidden="true"><span aria-hidden="true)
+      * [Identifying HIP Target Platform](#identifying-hip-target-platform" aria-hidden="true"><span aria-hidden="true)
+      * [Identifying the Compiler: hcc or nvcc](#identifying-the-compiler-hcc-or-nvcc" aria-hidden="true"><span aria-hidden="true)
+      * [Identifying Current Compilation Pass: Host or Device](#identifying-current-compilation-pass-host-or-device" aria-hidden="true"><span aria-hidden="true)
+      * [Compiler Defines: Summary](#compiler-defines-summary" aria-hidden="true"><span aria-hidden="true)
+    * [Identifying Architecture Features](#identifying-architecture-features" aria-hidden="true"><span aria-hidden="true)
+      * [HIP_ARCH Defines](#hip_arch-defines" aria-hidden="true"><span aria-hidden="true)
+      * [Device-Architecture Properties](#device-architecture-properties" aria-hidden="true"><span aria-hidden="true)
+      * [Table of Architecture Properties](#table-of-architecture-properties" aria-hidden="true"><span aria-hidden="true)
+    * [Finding HIP](#finding-hip" aria-hidden="true"><span aria-hidden="true)
+    * [Compiler Options](#compiler-options" aria-hidden="true"><span aria-hidden="true)
+    * [Linking Issues](#linking-issues" aria-hidden="true"><span aria-hidden="true)
+      * [Linking With hipcc](#linking-with-hipcc" aria-hidden="true"><span aria-hidden="true)
+      * [-lm Option](#-lm-option" aria-hidden="true"><span aria-hidden="true)
+    * [Linking Code With Other Compilers](#linking-code-with-other-compilers" aria-hidden="true"><span aria-hidden="true)
+      * [libc   and libstdc  ](#libc-and-libstdc" aria-hidden="true"><span aria-hidden="true)
+      * [HIP Headers (hip_runtime.h, hip_runtime_api.h)](#hip-headers-hip_runtimeh-hip_runtime_apih" aria-hidden="true"><span aria-hidden="true)
+      * [Using a Standard C   Compiler](#using-a-standard-c-compiler" aria-hidden="true"><span aria-hidden="true)
+        * [cuda.h](#cudah" aria-hidden="true"><span aria-hidden="true)
+      * [Choosing HIP File Extensions](#choosing-hip-file-extensions" aria-hidden="true"><span aria-hidden="true)
+      * [Workarounds](#workarounds" aria-hidden="true"><span aria-hidden="true)
+        * [warpSize](#warpsize" aria-hidden="true"><span aria-hidden="true)
+        * [Textures and Cache Control](#textures-and-cache-control" aria-hidden="true"><span aria-hidden="true)
+    * [More Tips](#more-tips" aria-hidden="true"><span aria-hidden="true)
+      * [hcc CPU Mode](#hcc-cpu-mode" aria-hidden="true"><span aria-hidden="true)
+      * [HIPTRACE Mode](#hiptrace-mode" aria-hidden="true"><span aria-hidden="true)
+      * [Environment Variables](#environment-variables" aria-hidden="true"><span aria-hidden="true)
+      * [Debugging hipcc](#debugging-hipcc" aria-hidden="true"><span aria-hidden="true)
+      * [What Does This Error Mean?](#what-does-this-error-mean" aria-hidden="true"><span aria-hidden="true)
+        * [/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" aria-hidden="true"><span aria-hidden="true)
+        * [grid_launch kernel dispatch - fallback](#grid_launch-kernel-dispatch---fallback" aria-hidden="true"><span aria-hidden="true)
+        * [Editor Highlighting](#editor-highlighting" aria-hidden="true"><span aria-hidden="true)
+
+
+## Porting a New Cuda Project
+
+### General Tips
+- Starting the port on a Cuda machine is often the easiest approach, since you can incrementally port pieces of the code to HIP while leaving the rest in Cuda. (Recall that on Cuda machines HIP is just a thin layer over Cuda, so the two code types can interoperate on nvcc platforms.) Also, the HIP port can be compared with the original Cuda code for function and performance.
+- Once the Cuda code is ported to HIP and is running on the Cuda machine, compile the HIP code using hcc on an AMD machine.
+- HIP ports can replace Cuda versions---HIP can deliver the same performance as a native Cuda implementation, with the benefit of portability to both Nvidia and AMD architectures as well as a path to future C++ standard support. You can handle platform-specific features through conditional compilation or by adding them to the open-source HIP infrastructure.
+- Use **bin/hipconvertinplace.sh** to hipify all code files in the Cuda source directory.
+
+### Scanning existing CUDA code to scope the porting effort
+The hipexamine.sh tool will scan a source directory to determine which files contain CUDA code and how much of that code can be automatically hipified, 
+```
+> cd examples/rodinia_3.0/cuda/kmeans
+> $HIP_DIR/bin/hipexamine.sh .
+info: hipify ./kmeans.h =====>
+info: hipify ./unistd.h =====>
+info: hipify ./kmeans.c =====>
+info: hipify ./kmeans_cuda_kernel.cu =====>
+  info: converted 40 CUDA->HIP refs( dev:0 mem:0 kern:0 builtin:37 math:0 stream:0 event:0 err:0 def:0 tex:3 other:0 ) warn:0 LOC:185
+info: hipify ./getopt.h =====>
+info: hipify ./kmeans_cuda.cu =====>
+  info: converted 49 CUDA->HIP refs( dev:3 mem:32 kern:2 builtin:0 math:0 stream:0 event:0 err:0 def:0 tex:12 other:0 ) warn:0 LOC:311
+info: hipify ./rmse.c =====>
+info: hipify ./cluster.c =====>
+info: hipify ./getopt.c =====>
+info: hipify ./kmeans_clustering.c =====>
+info: TOTAL-converted 89 CUDA->HIP refs( dev:3 mem:32 kern:2 builtin:37 math:0 stream:0 event:0 err:0 def:0 tex:15 other:0 ) warn:0 LOC:3607
+  kernels (1 total) :   kmeansPoint(1)
+```
+
+hipexamine scans each code file (cpp, c, h, hpp, etc) found in the specified directory:
+
+   * Files with no CUDA code (ie kmeans.h) print one line summary just listing the source file name.
+   * Files with CUDA code print a summary of what was found - for example the kmeans_cuda_kernel.cu file:
+```
+info: hipify ./kmeans_cuda_kernel.cu =====>
+  info: converted 40 CUDA->HIP refs( dev:0 mem:0 kern:0 builtin:37 math:0 stream:0 event:0 
+```
+* Some of the most interesting information in kmeans_cuda_kernel.cu :
+       * How many CUDA calls were converted to HIP (40)
+       * Breakdown of the different CUDA functionality used (dev:0 mem:0 etc).  This file uses many CUDA builtins (37) and texture functions (3).
+       * Warning for code that looks like CUDA API but was not converted (0 in this file).
+       * Count Lines-of-Code (LOC) - 185 for this file. 
+
+* hipexamine also presents a summary at the end of the process for the statistics collected across all files. This has similar format to the per-file reporting, and also includes a list of all kernels which have been called.  An example from above:
+
+```shell
+info: TOTAL-converted 89 CUDA->HIP refs( dev:3 mem:32 kern:2 builtin:37 math:0 stream:0 event:0 err:0 def:0 tex:15 other:0 ) warn:0 LOC:3607
+  kernels (1 total) :   kmeansPoint(1)
+```
+ 
+ 
+## Distinguishing Compiler Modes
+ 
+ 
+### Identifying HIP Target Platform
+All HIP projects target either the hcc or nvcc platform. The platform affects which headers are included and which libraries are used for linking. 
+ 
+- `HIPCC_PLATFORM_HCC` is defined if the HIP platform targets hcc
+- `HIPCC_PLATFORM_NVCC` is defined if the HIP platform targets nvcc
+ 
+Many projects use a mixture of an accelerator compiler (hcc or nvcc) and a standard compiler (e.g., g++). These defines are set for both accelerator and standard compilers and thus are often the best option when writing code that uses conditional compilation.
+ 
+ 
+### Identifying the Compiler: hcc or nvcc
+Often, its useful to know whether the underlying compiler is hcc or nvcc. This knowledge can guard platform-specific code (features that only work on the nvcc or hcc path but not both) or aid in platform-specific performance tuning.   
+ 
+```
+#ifdef __HCC__
+// Compiled with hcc 
+ 
+```
+ 
+```
+#ifdef __NVCC__
+// 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)
+ 
+```
+ 
+```
+#ifdef __CUDACC__
+// Compiled with nvcc (Cuda language extensions enabled) 
+```
+ 
+hcc directly generates the host code (using the Clang x86 target) and passes the code to another host compiler. Thus, it lacks the equivalent of the \__CUDA_ACC define.
+ 
+The macro `__HIPCC__` is set if either `__HCC__` or `__CUDACC__` is defined. This configuration is useful in determining when code is being compiled using an accelerator-enabled compiler (hcc or nvcc) as opposed to a standard host compiler (GCC, ICC, Clang, etc.).
+ 
+### Identifying Current Compilation Pass: Host or Device
+ 
+Both nvcc and hcc make two passes over the code: one for host code and one for device code. `__HIP_DEVICE_COMPILE__` is set to a nonzero value when the compiler (hcc or nvcc) is compiling code for a device inside a `__global__` kernel or for a device function. `__HIP_DEVICE_COMPILE__` can replace #ifdef checks on the `__CUDA_ARCH__` define.
+ 
+```
+// #ifdef __CUDA_ARCH__  
+#ifdef __HIP_DEVICE_COMPILE__ 
+```
+ 
+Unlike `__CUDA_ARCH__`, the `__HIP_DEVICE_COMPILE__` value is 0 or 1, and it doesnt represent the feature capability of the target device.  
+
+
+### Compiler Defines: Summary
+|Define  		|  hcc      | nvcc 		|  Other (GCC, ICC, Clang, etc.) 
+|--- | --- | --- |---|
+|HIP-related defines:|
+|`__HIP_PLATFORM_HCC___`| Defined | Undefined |  Defined if targeting hcc platform; undefined otherwise |
+|`__HIP_PLATFORM_NVCC___`| Undefined | Defined |  Defined if targeting nvcc platform; undefined otherwise |
+|`__HIP_DEVICE_COMPILE__`     | 1 if compiling for device; 0 if compiling for host  |1 if compiling for device; 0 if compiling for host  | Undefined 
+|`__HIPCC__`		| Defined   | Defined 		|  Undefined
+|`__HIP_ARCH_*` | 0 or 1 depending on feature support (see below) | 0 or 1 depending on feature support (see below) | 0 
+|nvcc-related defines:|
+|`__CUDACC__` 		| Undefined | Defined if compiling for Cuda device; undefined otherwise 		|  Undefined
+|`__NVCC__` 		| Undefined | Defined 		|  Undefined
+|`__CUDA_ARCH__`		| Undefined | Unsigned representing compute capability (e.g., "130") if in device code; 0 if in host code  | Undefined 
+|hcc-related defines:|
+|`__HCC__`  		| Defined   | Undefined   	|  Undefined
+|`__HCC_ACCELERATOR__`  	| Nonzero if in device code; otherwise undefined | Undefined | Undefined 
+|`__clang__`		| Defined   | Undefined 	|  Defined if using Clang; otherwise undefined
+
+
+## Identifying Architecture Features
+
+### HIP_ARCH Defines
+
+Some Cuda code tests `__CUDA_ARCH__` for a specific value to determine whether the machine supports a certain architectural feature. For instance,
+
+```
+#if (__CUDA_ARCH__ >= 130) 
+// doubles are supported
+```
+This type of code requires special attention, since hcc/AMD and nvcc/Cuda devices have different architectural capabilities. Moreover, you cant determine the presence of a feature using a simple comparison against an architectures version number. HIP provides a set of defines and device properties to query whether a specific architectural feature is supported.
+
+The `__HIP_ARCH_*` defines can replace comparisons of `__CUDA_ARCH__` values: 
+```
+//#if (__CUDA_ARCH__ >= 130)   // non-portable
+if __HIP_ARCH_HAS_DOUBLES__ {  // portable HIP feature query
+   // doubles are supported
+}
+```
+
+For host code, the `__HIP_ARCH__*` defines are set to 0. You should only use the __HIP_ARCH__ fields in device code.
+
+### Device-Architecture Properties
+
+Host code should query the architecture feature flags in the device properties that hipDeviceGetProperties returns, rather than testing the "major" and "minor" fields directly:
+
+```
+hipDeviceGetProperties(&deviceProp, device);
+//if ((deviceProp.major == 1 && deviceProp.minor < 2))  // non-portable
+if (deviceProp.arch.hasSharedInt32Atomics) {            // portable HIP feature query
+    // has shared int32 atomic operations ...
+}
+```
+
+### Table of Architecture Properties
+The table below shows the full set of architectural properties that HIP supports.
+
+|Define (use only in device code) | Device Property (run-time query) | Comment |
+|------- | ---------   | -----   |
+|32-bit atomics:||
+|`__HIP_ARCH_HAS_GLOBAL_INT32_ATOMICS__`        |    hasGlobalInt32Atomics      |32-bit integer atomics for global memory
+|`__HIP_ARCH_HAS_GLOBAL_FLOAT_ATOMIC_EXCH__`    |    hasGlobalFloatAtomicExch   |32-bit float atomic exchange for global memory
+|`__HIP_ARCH_HAS_SHARED_INT32_ATOMICS__`        |    hasSharedInt32Atomics      |32-bit integer atomics for shared memory
+|`__HIP_ARCH_HAS_SHARED_FLOAT_ATOMIC_EXCH__`    |    hasSharedFloatAtomicExch   |32-bit float atomic exchange for shared memory
+|`__HIP_ARCH_HAS_FLOAT_ATOMIC_ADD__`            |     hasFloatAtomicAdd         |32-bit float atomic add in global and shared memory
+|64-bit atomics:                           |                                     |                                                     
+|`__HIP_ARCH_HAS_GLOBAL_INT64_ATOMICS__`        |    hasGlobalInt64Atomics      |64-bit integer atomics for global memory                                                 
+|`__HIP_ARCH_HAS_SHARED_INT64_ATOMICS__`        |    hasSharedInt64Atomics      |64-bit integer atomics for shared memory
+|Doubles:                                  |                                     |
+|`__HIP_ARCH_HAS_DOUBLES__`                    |     hasDoubles                 |Double-precision floating point                                                     
+|Warp cross-lane operations:              |                                      |                                                     
+|`__HIP_ARCH_HAS_WARP_VOTE__`                  |     hasWarpVote                |Warp vote instructions (any, all)                                                      
+|`__HIP_ARCH_HAS_WARP_BALLOT__`                |     hasWarpBallot              |Warp ballot instructions 
+|`__HIP_ARCH_HAS_WARP_SHUFFLE__`               |     hasWarpShuffle             |Warp shuffle operations (shfl\_\*)                                                     
+|`__HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__`          |     hasFunnelShift             |Funnel shift two input words into one
+|Sync:                                    |                                      |
+|`__HIP_ARCH_HAS_THREAD_FENCE_SYSTEM__`        |     hasThreadFenceSystem       |threadfence\_system
+|`__HIP_ARCH_HAS_SYNC_THREAD_EXT__`            |     hasSyncThreadsExt         |syncthreads\_count, syncthreads\_and, syncthreads\_or 
+|                                                       
+|Miscellaneous:                                     |                                     |                                                     
+|`__HIP_ARCH_HAS_SURFACE_FUNCS__`              |   hasSurfaceFuncs              |                                                     
+|`__HIP_ARCH_HAS_3DGRID__`                     |   has3dGrid                    | Grids and groups are 3D                                                     
+|`__HIP_ARCH_HAS_DYNAMIC_PARALLEL__`           |   hasDynamicParallelism        | 
+                                                                                 
+
+## Finding HIP
+
+Makefiles can use the following syntax to conditionally provide a default HIP_PATH if one does not exist:
+
+```
+HIP_PATH ?= $(shell hipconfig --path)
+```
+
+## Compiler Options
+
+hipcc is a portable compiler driver that will call nvcc or hcc (depending on the target system) and attach all required include and library options. It passes options through to the target compiler. Tools that call hipcc must ensure the compiler options are appropriate for the target compiler. The `hipconfig` script may helpful in making
+infrastructure that identifies the target platform and sets options appropriately. It returns either "nvcc" or "hcc." The following sample shows the script in a makefile:
+
+```
+HIP_PLATFORM=$(shell hipconfig --compiler)
+
+ifeq (${HIP_PLATFORM}, nvcc)
+	HIPCC_FLAGS = -gencode=arch=compute_20,code=sm_20 
+endif
+ifeq (${HIP_PLATFORM}, hcc)
+	HIPCC_FLAGS = -Wno-deprecated-register
+endif
+
+```
+
+
+## Linking Issues
+
+### Linking With hipcc
+
+hipcc adds the necessary libraries for HIP as well as for the accelerator compiler (nvcc or hcc). We recommend linking with hipcc.
+
+### -lm Option
+ 
+hcc does not add -lm by default. If you see errors about missing math functions at link time (e.g., "sqrt@@GLIBC_2.2.5"), ensure that -lm is in the link options.
+
+
+## Linking Code With Other Compilers
+
+Cuda code often uses nvcc for accelerator code (defining and launching kernels, typically defined in .cu or .cuh files). It also uses a standard compiler (G++) for the rest of the application. nvcc is a preprocessor that employs a standard host compiler (e.g., GCC) to generate the host code. Code compiled using this tool can employ the intersection of language features supported by both nvcc and the host compiler. In some cases, you must take care to ensure the data types and alignment of the host compiler are identical to those of the device compiler. Only some host compilers are supported---for example, recent nvcc versions lack Clang host-compiler capability.  
+
+hcc generates both device and host code using the same Clang-based compiler. The code uses the same API as GCC, which allows code generated by different GCC-compatible compilers to be linked together. For example, code compiled using hcc can link with code compiled using "standard" compilers (such as GCC, ICC and Clang). You must take care to ensure all compilers use the same standard C++ header and library formats.
+
+
+### libc++ and libstdc++
+
+hcc uses the LLVM "libc++" standard library, whereas GCC uses the "libstdc++" standard library. Generally, libc++ provides a broader set of C++ features; libstdc++ is the standard for more compilers.
+
+When cross-linking C++ code, any C++ functions that use types from the C++ standard library (including std::string, std::vector and other containers) must use the same standard-library implementation. They include the following:  
+
+- Functions or kernels defined in hcc that are called from a standard compiler
+- Functions defined in a standard compiler that are called from hcc.
+
+Note that C++ code that doesnt use the standard library, or pure-C code, can be linked across compiler boundaries. The following suggestions may help you to meet these requirements:
+
+- Use the libc++ header and library for standard compilers. For GCC you can do so by passing the following:
+    - `g++ -std=c++11 -I/usr/include/c++/v1 `; note that the host code will compile as C++ 11 code.
+- Use hipcc to compile the entire application 
+- Partition the code into separate files so that standard C++ types dont apply across compiler boundaries. For example, placing the kernels in a separate file compiled using hcc will resolve the issue (provided they avoid using types from the C++ library).
+
+A future hcc version may support libstdc++ in addition to libc++.
+
+
+### HIP Headers (hip_runtime.h, hip_runtime_api.h)
+
+The hip_runtime.h and hip_runtime_api.h files define the types, functions and enumerations needed to compile a HIP program:
+
+- hip_runtime_api.h: defines all the HIP runtime APIs (e.g., hipMalloc) and the types required to call them. A source file that is only calling HIP APIs but neither defines nor launches any kernels can include hip_runtime_api.h. hip_runtime_api.h uses no custom C++ features and can be compiled using a standard C++ compiler.
+- hip_runtime.h: included in hip_runtime_api.h. It additionally provides the types and defines required to create and launch kernels. hip_runtime.h does use custom C++ features, but they are guarded by ifdef checks. It can be compiled using a standard C++ compiler but will expose a subset of the available functions.
+
+Cuda has slightly different contents for these two files. In some cases you may need to convert hipified code to include the richer hip_runtime.h instead of hip_runtime_api.h.
+
+### Using a Standard C++ Compiler
+You can compile hip\_runtime\_api.h using a standard C or C++ compiler (e.g., GCC or ICC). The HIP include paths and defines (`__HIP_PLATFORM_HCC__` or `__HIP_PLATFORM_NVCC__`) must pass to the standard compiler; hipconfig then returns the necessary options:
+```
+> hipconfig --cxx_config 
+ -D__HIP_PLATFORM_HCC__ -I/home/user1/hip/include
+```
+
+You can capture the hipconfig output and passed it to the standard compiler; below is a sample makefile syntax:
+
+```
+CPPFLAGS += $(shell $(HIP_PATH)/bin/hipconfig --cpp_config)
+```
+
+nvcc includes some headers by default. Files that call HIP run-time APIs or define HIP kernels must explicitly include HIP headers. If the compilation process reports that it cannot find necessary APIs (for example, "error: identifier hipSetDevice is undefined"),
+ensure that the file includes hip_runtime.h (or hip_runtime_api.h, if appropriate). The hipify script automatically converts "cuda_runtime.h" to "hip_runtime.h," and it converts "cuda_runtime_api.h" to "hip_runtime_api.h", but it may miss nested headers or macros.  
+
+#### cuda.h
+
+The hcc path provides an empty cuda.h file. Some existing Cuda programs include this file but don't require any of the functions.
+
+### Choosing HIP File Extensions
+
+Many existing Cuda projects use the ".cu" and ".cuh" file extensions to indicate code that should be run through the nvcc compiler. 
+For quick HIP ports, leaving these file extensions unchanged is often easier, as it minimizes the work required to change file names in the directory and #include statements in the files.
+
+For new projects or ports which can be re-factored, we recommend the use of the extension ".hip.cpp" for header files, and
+".hip.h" or ".hip.hpp" for header files.
+This indicates that the code is standard C++ code, but also provides a unique indication for make tools to
+run hipcc when appropriate.
+
+### Workarounds
+
+#### warpSize
+Code should not assume a warp size of 32 or 64.  See [Warp Cross-Lane Functions](hip_kernel_language.md#warp-cross-lane-functions) for information on how to write portable wave-aware code.
+
+
+#### Textures and Cache Control
+
+Compute programs sometimes use textures either to access dedicated texture caches or to use the texture-sampling hardware for interpolation and clamping. The former approach uses simple point samplers with linear interpolation, essentially only reading a single point. The latter approach uses the sampler hardware to interpolate and combine multiple
+point samples. AMD hardware, as well as recent competing hardware,
+has a unified texture/L1 cache, so it no longer has a dedicated texture cache. But the nvcc path often caches global loads in the L2 cache, and some programs may benefit from explicit control of the L1 cache contents.  We recommend the __ldg instruction for this purpose.
+
+HIP currently lacks texture support; a future revision will add this capability. Also, AMD compilers currently load all data into both the L1 and L2 caches, so __ldg is treated as a no-op. 
+
+We recommend the following for functional portability:
+
+- For programs that use textures only to benefit from improved caching, use the __ldg instruction
+- Alternatively, use conditional compilation (see [Identify HIP Target Platform](#identify-hip-target-platform)) 
+   - For the `__HIP_PLATFORM_NVCC__` path, use the full texture path
+   - For the `__HIP_PLATFORM_HCC__` path, pass an additional pointer to the kernel and reference it using regular device memory-load instructions rather than texture loads. Some applications may already take this step, since it allows experimentation with caching behavior.
+
+``` 
+texture<float, 1, cudaReadModeElementType> t_features;
+
+void __global__ MyKernel(float *d_features /* pass pointer parameter, if not already available */...) 
+{
+    // ... 
+
+#ifdef __HIP_PLATFORM_NVCC__
+    float tval = tex1Dfetch(t_features,addr);
+#else
+    float tval = d_features[addr];
+#endif
+        
+}
+
+// Host code:
+void myFunc () 
+{
+    // ...
+
+#ifdef __HIP_PLATFORM_NVCC__
+    cudaChannelFormatDesc chDesc0 = cudaCreateChannelDesc<float>();
+    t_features.filterMode = cudaFilterModePoint;   
+    t_features.normalized = false;
+    t_features.channelDesc = chDesc0;
+
+	cudaBindTexture(NULL, &t_features, d_features, &chDesc0, npoints*nfeatures*sizeof(float));
+#endif
+
+``` 
+
+
+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---its 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:
+    
+- 0x1 = trace APIs
+- 0x2 = trace synchronization operations
+- 0x4 = trace memory allocation / deallocation
+
+### Environment Variables
+
+On hcc/AMD platforms, set the HIP_PRINT_ENV environment variable to 1 and run an application that calls a HIP API to see all HIP-supported environment variables and their current values:
+
+- HIP_PRINT_ENV = 1: print HIP environment variables
+- HIP_TRACE_API = 0: 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)
+
+
+
+### Debugging hipcc
+To see the detailed commands that hipcc issues, set the environment variable HIPCC_VERBOSE to 1. Doing so will print to stderr the hcc (or nvcc) commands that hipcc generates. 
+
+```
+export HIPCC_VERBOSE=1
+make
+...
+hipcc-cmd: /opt/hcc/bin/hcc  -hc -I/opt/hcc/include -stdlib=libc++ -I../../../../hc/include -I../../../../include/hcc_detail/cuda -I../../../../include -x c++ -I../../common -O3 -c backprop_cuda.cu
+```
+
+### 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();
+
+If you pass a ".cu" file, hcc will attempt to compile it as a Cuda language file. You must tell hcc that its in fact a C++ file: use the "-x c++" option.
+
+
+#### grid_launch kernel dispatch - fallback
+HIP uses an hcc language feature called "grid_launch". The [[hc_grid_launch]] attribute that can be attached to a function definition, and the first parameter is of type grid_launch_parm.
+When a [[hc_grid_launch]] function is called, hcc runtime uses the grid_launch_parm to control the execution configuration of the kernel 
+(including the grid and group dimensions, the queue, and dynamic group memory allocations).   By default, the hipLaunchKernel macro creates a grid_launch_parm structure and launches a
+[[hc_grid_launch]] kernel.  grid_launch is a relatively new addition to hcc so this section describes how to fall back to a traditional calling sequence which invokes a standard host function
+which calls a hc::parallel_for_each to launch the kernel.  
+
+First, set DISABLE_GRID_LAUNCH:
+include/hip_common.h
+```
+// Set this define to disable GRID_LAUNCH
+#define DISABLE_GRID_LAUNCH
+```
+
+Inside any kernel use the KERNELBEGIN as the first line in the kernel function, and KERNELEND as the last line.  For example:
+```
+__global__ void
+MyKernel(hipLaunchParm lp, float *C, const float *A, size_t N)
+{
+    KERNELBEGIN; // Required if hc_grid_launch is disabled
+
+	int tid = hipBlockIdx_x*MAX_THREADS_PER_BLOCK + hipThreadIdx_x;
+
+    if (tid < N) {
+        C[tid] = A[tid];
+    }
+
+    KERNELEND; // Required if hc_grid_launch is disabled
+}
+```
+
+
+#### Editor Highlighting
+See the utils/vim or utils/gedit directories to add handy highlighting to hip files.
@@ -0,0 +1,42 @@
+# Table Comparing Syntax for Different Compute APIs
+
+|Term|CUDA|HIP|HC|C++AMP|OpenCL|
+|---|---|---|---|---|---|
+|Device|`int deviceId`|`int deviceId`|`hc::accelerator`|`concurrency::`<br>`accelerator`|`cl_device`
+|Queue|`cudaStream_t`|`hipStream_t`|`hc::`<br>`accelerator_view`|`concurrency::`<br>`accelerator_view`|`cl_command_queue`
+|Event|`cudaEvent_t`|`hipEvent_t`|`hc::`<br>`completion_future`|`concurrency::`<br>`completion_future`|`cl_event`
+|Memory|`void *`|`void *`|`void *`; `hc::array`; `hc::array_view`|`concurrency::array`;<br>`concurrency::array_view`|`cl_mem`
+|||||
+| |grid|grid|extent|extent|NDRange
+| |block|block|tile|tile|work-group
+| |thread|thread|thread|thread|work-item
+| |warp|warp|wavefront|N/A|sub-group
+|||||
+|Thread-<br>index | threadIdx.x | hipThreadIdx_x | t_idx.local[0] | t_idx.local[0] | get_local_id(0) |
+|Block-<br>index  | blockIdx.x  | hipBlockIdx_x  | t_idx.tile[0]  | t_idx.tile[0]  | get_group_id(0) |
+|Block-<br>dim    | blockDim.x  | hipBlockDim_x  | t_ext.tile_dim[0]| t_idx.tile_dim0 | get_local_size(0) |
+|Grid-dim     | gridDim.x   | hipGridDim_x   | t_ext[0]| t_ext[0] | get_global_size(0) |
+|||||
+|Device Kernel|`__global__`|`__global__`|lambda inside `hc::`<br>`parallel_for_each` or [[hc]]|`restrict(amp)`|`__kernel`
+|Device Function|`__device__`|`__device__`|`[[hc]]` (detected automatically in many case)|`restrict(amp)`|Implied in device compilation
+|Host Function|`__host_` (default)|`__host_` (default)|`[[cpu]]`  (default)|`restrict(cpu)` (default)|Implied in host compilation.
+|Host + Device Function|`__host__` `__device__`|`__host__` `__device__`|  `[[hc]]` `[[cpu]]`|`restrict(amp,cpu)`|No equivalent
+|Kernel Launch|`<<< >>>`|`hipLaunchKernel`|`hc::`<br>`parallel_for_each`|`concurrency::`<br>`parallel_for_each`|`clEnqueueNDRangeKernel`
+||||||
+|Global Memory|`__global__`|`__global__`|Unnecessary / Implied|Unnecessary / Implied|`__global`
+|Group Memory|`__shared__`|`__shared__`|`tile_static`|`tile_static`|`__local`
+|Constant|`__constant__`|`__constant__`|Unnecessary / Implied|Unnecessary / Implied|`__constant`
+||||||
+||`__syncthreads`|`__syncthreads`|`tile_static.barrier()`|`t_idx.barrier()`|`barrier(CLK_LOCAL_MEMFENCE)`
+|Atomic Builtins|`atomicAdd`|`atomicAdd`|`hc::atomic_fetch_add`|`concurrency::`<br>`atomic_fetch_add`|`atomic_add`
+|Precise Math|`cos(f)`|`cos(f)`|`hc::`<br>`precise_math::cos(f)`|`concurrency::`<br>`precise_math::cos(f)`|`cos(f)`
+|Fast Math|`__cos(f)`|`__cos(f)`|`hc::`<br>`fast_math::cos(f)`|`concurrency::`<br>`fast_math::cos(f)`|`native_cos(f)`
+|Vector|`float4`|`float4`|`hc::`<br>`short_vector::float4`|`concurrency::`<br>`graphics::float_4`|`float4`
+
+###Notes
+1. For HC and C++AMP, assume captured _tiled_ext_ t_ext and captured _extent_ ext.  These languages use captured variables to pass information to the kernel rather than special built-in functions so variable name may vary.
+2. The indexig functions (starting with thread-index) show the terminology for a 1D grid.  Some APIs use reverse order of xyz / 012 indexing for 3D grids.
+3. HC allos  tile dimensions to be specified at runtime while C++AMP requires that tile dimensions are specified at compile-time.  Thus hc syntax for tile dims is `t_ext.tile_dim[0]` while C++AMP is t_ext.tile_dim0.
+
+
+
@@ -0,0 +1,21 @@
+
+# Terms used in HIP Documentation
+
+- host,  host cpu : Executes the HIP runtime API and is capable of initiating kernel launches to one or more devices.
+- default device : Each host thread maintains a "default device".  
+Most HIP runtime APIs (including memory allocation, copy commands, kernel launches) do not use accept an explicit device
+argument but instead implicitly use the default device.
+The default device can be set with hipSetDevice.
+
+- "active host thread" - the thread which is running the HIP APIs.   
+
+- completion_future becomes ready.  "Completes"
+
+- hcc = Heterogeneous Compute Compiler (https://bitbucket.org/multicoreware/hcc/wiki/Home).  
+
+- hipify - tool to convert CUDA(R) code to portable C++ code.
+- hipconfig - tool to report various confoguration properties of the target platform.
+
+
+
+