Table of Contents generated with DocToc
- Introduction
- Function-Type Qualifiers
- Calling
__global__Functions - Kernel-Launch Example
- Variable-Type Qualifiers
- Built-In Variables
- Vector Types
- Memory-Fence Instructions
- Synchronization Functions
- Math Functions
- Texture Functions
- Surface Functions
- Timer Functions
- Atomic Functions
- Warp Cross-Lane Functions
- Profiler Counter Function
- Assert
- Printf
- Device-Side Dynamic Global Memory Allocation
__launch_bounds__- Register Keyword
- Pragma Unroll
- In-Line Assembly
- C++ Support
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
Supported __device__ functions are
- Executed on the device
- Called from the device only
The __device__ keyword can combine with the host keyword (see host).
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.
HIP lacks dynamic-parallelism support, so __global__ functions cannot be called from the device.
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.
__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)
- hipStream_t: stream where the kernel should execute. A value of 0 corresponds to the NULL stream (see 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. 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.
// 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);
}
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.
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 memory, including the __managed__ keyword, are not supported in HIP.
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.
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 |
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.
Note that these types are defined in hip_runtime.h and are not automatically provided by the compiler.
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 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) {};
};
HIP support for __threadfence(), __threadfence_block() and __threadfence_system() is under development.
The __syncthreads() built-in function is supported in HIP. The __syncthreads_count(int), __syncthreads_and(int) and __syncthreads_or(int) functions are under development.
hcc supports a set of math operations callable from the device.
Following is the list of supported single precision mathematical functions.
| Function | Supported on Host | Supported on Device |
|---|---|---|
| float acosf ( float x ) Calculate the arc cosine of the input argument. |
✓ | ✓ |
| float acoshf ( float x ) Calculate the nonnegative arc hyperbolic cosine of the input argument. |
✓ | ✓ |
| float asinf ( float x ) Calculate the arc sine of the input argument. |
✓ | ✓ |
| float asinhf ( float x ) Calculate the arc hyperbolic sine of the input argument. |
✓ | ✓ |
| float atan2f ( float y, float x ) Calculate the arc tangent of the ratio of first and second input arguments. |
✓ | ✓ |
| float atanf ( float x ) Calculate the arc tangent of the input argument. |
✓ | ✓ |
| float atanhf ( float x ) Calculate the arc hyperbolic tangent of the input argument. |
✓ | ✓ |
| float cbrtf ( float x ) Calculate the cube root of the input argument. |
✓ | ✓ |
| float ceilf ( float x ) Calculate ceiling of the input argument. |
✓ | ✓ |
| float copysignf ( float x, float y ) Create value with given magnitude, copying sign of second value. |
✓ | ✓ |
| float cosf ( float x ) Calculate the cosine of the input argument. |
✓ | ✓ |
| float coshf ( float x ) Calculate the hyperbolic cosine of the input argument. |
✓ | ✓ |
| float erfcf ( float x ) Calculate the complementary error function of the input argument. |
✓ | ✓ |
| float erff ( float x ) Calculate the error function of the input argument. |
✓ | ✓ |
| float exp10f ( float x ) Calculate the base 10 exponential of the input argument. |
✓ | ✓ |
| float exp2f ( float x ) Calculate the base 2 exponential of the input argument. |
✓ | ✓ |
| float expf ( float x ) Calculate the base e exponential of the input argument. |
✓ | ✓ |
| float expm1f ( float x ) Calculate the base e exponential of the input argument, minus 1. |
✓ | ✓ |
| float fabsf ( float x ) Calculate the absolute value of its argument. |
✓ | ✓ |
| float fdimf ( float x, float y ) Compute the positive difference between x and y. |
✓ | ✓ |
| float floorf ( float x ) Calculate the largest integer less than or equal to x. |
✓ | ✓ |
| float fmaf ( float x, float y, float z ) Compute x × y + z as a single operation. |
✓ | ✓ |
| float fmaxf ( float x, float y ) Determine the maximum numeric value of the arguments. |
✓ | ✓ |
| float fminf ( float x, float y ) Determine the minimum numeric value of the arguments. |
✓ | ✓ |
| float fmodf ( float x, float y ) Calculate the floating-point remainder of x / y. |
✓ | ✓ |
| float frexpf ( float x, int* nptr ) Extract mantissa and exponent of a floating-point value. |
✓ | ✗ |
| float hypotf ( float x, float y ) Calculate the square root of the sum of squares of two arguments. |
✓ | ✓ |
| int ilogbf ( float x ) Compute the unbiased integer exponent of the argument. |
✓ | ✓ |
| __RETURN_TYPE1 isfinite ( float a ) Determine whether argument is finite. |
✓ | ✓ |
| __RETURN_TYPE1 isinf ( float a ) Determine whether argument is infinite. |
✓ | ✓ |
| __RETURN_TYPE1 isnan ( float a ) Determine whether argument is a NaN. |
✓ | ✓ |
| float ldexpf ( float x, int exp ) Calculate the value of x ⋅ 2exp. |
✓ | ✓ |
| float log10f ( float x ) Calculate the base 10 logarithm of the input argument. |
✓ | ✓ |
| float log1pf ( float x ) Calculate the value of loge( 1 + x ). |
✓ | ✓ |
| float logbf ( float x ) Calculate the floating point representation of the exponent of the input argument. |
✓ | ✓ |
| float log2f ( float x ) Calculate the base 2 logarithm of the input argument. |
✓ | ✓ |
| float logf ( float x ) Calculate the natural logarithm of the input argument. |
✓ | ✓ |
| float modff ( float x, float* iptr ) Break down the input argument into fractional and integral parts. |
✓ | ✗ |
| float nanf ( const char* tagp ) Returns "Not a Number" value. |
✗ | ✓ |
| float nearbyintf ( float x ) Round the input argument to the nearest integer. |
✓ | ✓ |
| float powf ( float x, float y ) Calculate the value of first argument to the power of second argument. |
✓ | ✓ |
| float remainderf ( float x, float y ) Compute single-precision floating-point remainder. |
✓ | ✓ |
| float remquof ( float x, float y, int* quo ) Compute single-precision floating-point remainder and part of quotient. |
✓ | ✗ |
| float roundf ( float x ) Round to nearest integer value in floating-point. |
✓ | ✓ |
| float scalbnf ( float x, int n ) Scale floating-point input by integer power of two. |
✓ | ✓ |
| __RETURN_TYPE1 signbit ( float a ) Return the sign bit of the input. |
✓ | ✓ |
| void sincosf ( float x, float* sptr, float* cptr ) Calculate the sine and cosine of the first input argument. |
✓ | ✗ |
| float sinf ( float x ) Calculate the sine of the input argument. |
✓ | ✓ |
| float sinhf ( float x ) Calculate the hyperbolic sine of the input argument. |
✓ | ✓ |
| float sqrtf ( float x ) Calculate the square root of the input argument. |
✓ | ✓ |
| float tanf ( float x ) Calculate the tangent of the input argument. |
✓ | ✓ |
| float tanhf ( float x ) Calculate the hyperbolic tangent of the input argument. |
✓ | ✓ |
| float truncf ( float x ) Truncate input argument to the integral part. |
✓ | ✓ |
| float tgammaf ( float x ) Calculate the gamma function of the input argument. |
✓ | ✓ |
| [1] __RETURN_TYPE is dependent on compiler. It is usually 'int' for C compilers and 'bool' for C++ compilers. ↩ |
Following is the list of supported double precision mathematical functions.
| Function | Supported on Host | Supported on Device |
|---|---|---|
| double acos ( double x ) Calculate the arc cosine of the input argument. |
✓ | ✓ |
| double acosh ( double x ) Calculate the nonnegative arc hyperbolic cosine of the input argument. |
✓ | ✓ |
| double asin ( double x ) Calculate the arc sine of the input argument. |
✓ | ✓ |
| double asinh ( double x ) Calculate the arc hyperbolic sine of the input argument. |
✓ | ✓ |
| double atan ( double x ) Calculate the arc tangent of the input argument. |
✓ | ✓ |
| double atan2 ( double y, double x ) Calculate the arc tangent of the ratio of first and second input arguments. |
✓ | ✓ |
| double atanh ( double x ) Calculate the arc hyperbolic tangent of the input argument. |
✓ | ✓ |
| double cbrt ( double x ) Calculate the cube root of the input argument. |
✓ | ✓ |
| double ceil ( double x ) Calculate ceiling of the input argument. |
✓ | ✓ |
| double copysign ( double x, double y ) Create value with given magnitude, copying sign of second value. |
✓ | ✓ |
| double cos ( double x ) Calculate the cosine of the input argument. |
✓ | ✓ |
| double cosh ( double x ) Calculate the hyperbolic cosine of the input argument. |
✓ | ✓ |
| double erf ( double x ) Calculate the error function of the input argument. |
✓ | ✓ |
| double erfc ( double x ) Calculate the complementary error function of the input argument. |
✓ | ✓ |
| double exp ( double x ) Calculate the base e exponential of the input argument. |
✓ | ✓ |
| double exp10 ( double x ) Calculate the base 10 exponential of the input argument. |
✓ | ✓ |
| double exp2 ( double x ) Calculate the base 2 exponential of the input argument. |
✓ | ✓ |
| double expm1 ( double x ) Calculate the base e exponential of the input argument, minus 1. |
✓ | ✓ |
| double fabs ( double x ) Calculate the absolute value of the input argument. |
✓ | ✓ |
| double fdim ( double x, double y ) Compute the positive difference between x and y. |
✓ | ✓ |
| double floor ( double x ) Calculate the largest integer less than or equal to x. |
✓ | ✓ |
| double fma ( double x, double y, double z ) Compute x × y + z as a single operation. |
✓ | ✓ |
| double fmax ( double , double ) Determine the maximum numeric value of the arguments. |
✓ | ✓ |
| double fmin ( double x, double y ) Determine the minimum numeric value of the arguments. |
✓ | ✓ |
| double fmod ( double x, double y ) Calculate the floating-point remainder of x / y. |
✓ | ✓ |
| double frexp ( double x, int* nptr ) Extract mantissa and exponent of a floating-point value. |
✓ | ✗ |
| double hypot ( double x, double y ) Calculate the square root of the sum of squares of two arguments. |
✓ | ✓ |
| int ilogb ( double x ) Compute the unbiased integer exponent of the argument. |
✓ | ✓ |
| __RETURN_TYPE1 isfinite ( double a ) Determine whether argument is finite. |
✓ | ✓ |
| __RETURN_TYPE1 isinf ( double a ) Determine whether argument is infinite. |
✓ | ✓ |
| __RETURN_TYPE1 isnan ( double a ) Determine whether argument is a NaN. |
✓ | ✓ |
| double ldexp ( double x, int exp ) Calculate the value of x ⋅ 2exp. |
✓ | ✓ |
| double log ( double x ) Calculate the base e logarithm of the input argument. |
✓ | ✓ |
| double log10 ( double x ) Calculate the base 10 logarithm of the input argument. |
✓ | ✓ |
| double log1p ( double x ) Calculate the value of loge( 1 + x ). |
✓ | ✓ |
| double log2 ( double x ) Calculate the base 2 logarithm of the input argument. |
✓ | ✓ |
| double logb ( double x ) Calculate the floating point representation of the exponent of the input argument. |
✓ | ✓ |
| double modf ( double x, double* iptr ) Break down the input argument into fractional and integral parts. |
✓ | ✗ |
| double nan ( const char* tagp ) Returns "Not a Number" value. |
✗ | ✓ |
| double nearbyint ( double x ) Round the input argument to the nearest integer. |
✓ | ✓ |
| double pow ( double x, double y ) Calculate the value of first argument to the power of second argument. |
✓ | ✓ |
| double remainder ( double x, double y ) Compute double-precision floating-point remainder. |
✓ | ✓ |
| double remquo ( double x, double y, int* quo ) Compute double-precision floating-point remainder and part of quotient. |
✓ | ✗ |
| double round ( double x ) Round to nearest integer value in floating-point. |
✓ | ✓ |
| double scalbn ( double x, int n ) Scale floating-point input by integer power of two. |
✓ | ✓ |
| __RETURN_TYPE1 signbit ( double a ) Return the sign bit of the input. |
✓ | ✓ |
| double sin ( double x ) Calculate the sine of the input argument. |
✓ | ✓ |
| void sincos ( double x, double* sptr, double* cptr ) Calculate the sine and cosine of the first input argument. |
✓ | ✗ |
| double sinh ( double x ) Calculate the hyperbolic sine of the input argument. |
✓ | ✓ |
| double sqrt ( double x ) Calculate the square root of the input argument. |
✓ | ✓ |
| double tan ( double x ) Calculate the tangent of the input argument. |
✓ | ✓ |
| double tanh ( double x ) Calculate the hyperbolic tangent of the input argument. |
✓ | ✓ |
| double tgamma ( double x ) Calculate the gamma function of the input argument. |
✓ | ✓ |
| double trunc ( double x ) Truncate input argument to the integral part. |
✓ | ✓ |
| [1] __RETURN_TYPE is dependent on compiler. It is usually 'int' for C compilers and 'bool' for C++ compilers. ↩ |
Following is the list of supported integer intrinsics. Note that intrinsics are supported on device only.
| Function |
|---|
| unsigned int __brev ( unsigned int x ) Reverse the bit order of a 32 bit unsigned integer. |
| unsigned long long int __brevll ( unsigned long long int x ) Reverse the bit order of a 64 bit unsigned integer. |
| int __clz ( int x ) Return the number of consecutive high-order zero bits in a 32 bit integer. |
| unsigned int __clz(unsigned int x) Return the number of consecutive high-order zero bits in 32 bit unsigned integer. |
| int __clzll ( long long int x ) Count the number of consecutive high-order zero bits in a 64 bit integer. |
| unsigned int __clzll(long long int x) Return the number of consecutive high-order zero bits in 64 bit signed integer. |
| unsigned int __ffs(unsigned int x) Find the position of least signigicant bit set to 1 in a 32 bit unsigned integer.1 |
| unsigned int __ffs(int x) Find the position of least signigicant bit set to 1 in a 32 bit signed integer. |
| unsigned int __ffsll(unsigned long long int x) Find the position of least signigicant bit set to 1 in a 64 bit unsigned integer.1 |
| unsigned int __ffsll(long long int x) Find the position of least signigicant bit set to 1 in a 64 bit signed integer. |
| unsigned int __popc ( unsigned int x ) Count the number of bits that are set to 1 in a 32 bit integer. |
| int __popcll ( unsigned long long int x ) Count the number of bits that are set to 1 in a 64 bit integer. |
| [1] |
| 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 -1, while for ffs the index starts at 0. |
Following is the list of supported floating-point intrinsics. Note that intrinsics are supported on device only.
| Function |
|---|
| float __cosf ( float x ) Calculate the fast approximate cosine of the input argument. |
| float __expf ( float x ) Calculate the fast approximate base e exponential of the input argument. |
| float __frsqrt_rn ( float x ) Compute 1 / √x in round-to-nearest-even mode. |
| float __fsqrt_rd ( float x ) Compute √x in round-down mode. |
| float __fsqrt_rn ( float x ) Compute √x in round-to-nearest-even mode. |
| float __fsqrt_ru ( float x ) Compute √x in round-up mode. |
| float __fsqrt_rz ( float x ) Compute √x in round-towards-zero mode. |
| float __log10f ( float x ) Calculate the fast approximate base 10 logarithm of the input argument. |
| float __log2f ( float x ) Calculate the fast approximate base 2 logarithm of the input argument. |
| float __logf ( float x ) Calculate the fast approximate base e logarithm of the input argument. |
| float __powf ( float x, float y ) Calculate the fast approximate of xy. |
| float __sinf ( float x ) Calculate the fast approximate sine of the input argument. |
| float __tanf ( float x ) Calculate the fast approximate tangent of the input argument. |
| double __dsqrt_rd ( double x ) Compute √x in round-down mode. |
| double __dsqrt_rn ( double x ) Compute √x in round-to-nearest-even mode. |
| double __dsqrt_ru ( double x ) Compute √x in round-up mode. |
| double __dsqrt_rz ( double x ) Compute √x in round-towards-zero mode. |
Texture functions are not supported.
Surface functions are not supported.
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 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) | ✗ | ✓ |
| unsigned int atomicDec(unsigned int* address) | ✗ | ✓ |
| 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)) | ✓ | ✓ |
- 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 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)
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.
The following warp shuffle instructions are under development.
Half-float shuffles are not supported. The default width is warpSize---see 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);
The Cuda __prof_trigger() instruction is not supported.
The assert function is under development.
The printf function is under development.
Device-side dynamic global memory allocation is not supported.
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.
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.
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, 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.
The following C++ features are not supported:
- Run-time-type information (RTTI)
- Virtual functions
- Try/catch