HIP: Heterogenous-computing Interface for Portability
Loading...
Searching...
No Matches
amd_hip_bf16.h File Reference

hip_bf16.h provides struct for __hip_bfloat16 types More...

#include <hip/amd_detail/amd_hip_common.h>
#include "amd_hip_vector_types.h"
#include "device_library_decls.h"
#include "math_fwd.h"
#include <algorithm>
#include <climits>
#include <cmath>

Go to the source code of this file.

Classes

struct  __hip_bfloat16
 Struct to represent a 16 bit brain floating point number. More...
 
struct  __hip_bfloat162
 Struct to represent two 16 bit brain floating point numbers. More...
 

Macros

#define __HOST_DEVICE__   __host__ __device__ static inline
 
#define HIPRT_ONE_BF16   __float2bfloat16(1.0f)
 
#define HIPRT_ZERO_BF16   __float2bfloat16(0.0f)
 
#define HIPRT_INF_BF16   __ushort_as_bfloat16((unsigned short)0x7F80U)
 
#define HIPRT_MAX_NORMAL_BF16   __ushort_as_bfloat16((unsigned short)0x7F7FU)
 
#define HIPRT_MIN_DENORM_BF16   __ushort_as_bfloat16((unsigned short)0x0001U)
 
#define HIPRT_NAN_BF16   __ushort_as_bfloat16((unsigned short)0x7FFFU)
 
#define HIPRT_NEG_ZERO_BF16   __ushort_as_bfloat16((unsigned short)0x8000U)
 

Functions

__HOST_DEVICE__ float __bfloat162float (__hip_bfloat16 a)
 Converts bfloat16 to float.
 
__HOST_DEVICE__ __hip_bfloat16 __float2bfloat16 (float f)
 Converts float to bfloat16.
 
__HOST_DEVICE__ float2 __bfloat1622float2 (const __hip_bfloat162 a)
 Converts and moves bfloat162 to float2.
 
__HOST_DEVICE__ __hip_bfloat162 __bfloat162bfloat162 (const __hip_bfloat16 a)
 Moves bfloat16 value to bfloat162.
 
__HOST_DEVICE__ short int __bfloat16_as_short (const __hip_bfloat16 h)
 Reinterprets bits in a __hip_bfloat16 as a signed short integer.
 
__HOST_DEVICE__ unsigned short int __bfloat16_as_ushort (const __hip_bfloat16 h)
 Reinterprets bits in a __hip_bfloat16 as an unsigned signed short integer.
 
__HOST_DEVICE__ __hip_bfloat16 __double2bfloat16 (const double a)
 Convert double to __hip_bfloat16.
 
__HOST_DEVICE__ __hip_bfloat162 __float22bfloat162_rn (const float2 a)
 Convert float2 to __hip_bfloat162.
 
__HOST_DEVICE__ __hip_bfloat162 __halves2bfloat162 (const __hip_bfloat16 a, const __hip_bfloat16 b)
 Combine two __hip_bfloat16 to __hip_bfloat162.
 
__HOST_DEVICE__ __hip_bfloat16 __high2bfloat16 (const __hip_bfloat162 a)
 Returns high 16 bits of __hip_bfloat162.
 
__HOST_DEVICE__ __hip_bfloat162 __high2bfloat162 (const __hip_bfloat162 a)
 Returns high 16 bits of __hip_bfloat162.
 
__HOST_DEVICE__ float __high2float (const __hip_bfloat162 a)
 Converts high 16 bits of __hip_bfloat162 to float and returns the result.
 
__HOST_DEVICE__ __hip_bfloat162 __highs2bfloat162 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Extracts high 16 bits from each and combines them.
 
__HOST_DEVICE__ __hip_bfloat16 __low2bfloat16 (const __hip_bfloat162 a)
 Returns low 16 bits of __hip_bfloat162.
 
__HOST_DEVICE__ __hip_bfloat162 __low2bfloat162 (const __hip_bfloat162 a)
 Returns low 16 bits of __hip_bfloat162.
 
__HOST_DEVICE__ float __low2float (const __hip_bfloat162 a)
 Converts low 16 bits of __hip_bfloat162 to float and returns the result.
 
__HOST_DEVICE__ __hip_bfloat162 __lowhigh2highlow (const __hip_bfloat162 a)
 Swaps both halves.
 
__HOST_DEVICE__ __hip_bfloat162 __lows2bfloat162 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Extracts low 16 bits from each and combines them.
 
__HOST_DEVICE__ __hip_bfloat16 __short_as_bfloat16 (const short int a)
 Reinterprets short int into a bfloat16.
 
__HOST_DEVICE__ __hip_bfloat16 __ushort_as_bfloat16 (const unsigned short int a)
 Reinterprets unsigned short int into a bfloat16.
 
__HOST_DEVICE__ __hip_bfloat16 __hadd (const __hip_bfloat16 a, const __hip_bfloat16 b)
 Adds two bfloat16 values.
 
__HOST_DEVICE__ __hip_bfloat16 __hsub (const __hip_bfloat16 a, const __hip_bfloat16 b)
 Subtracts two bfloat16 values.
 
__HOST_DEVICE__ __hip_bfloat16 __hdiv (const __hip_bfloat16 a, const __hip_bfloat16 b)
 Divides two bfloat16 values.
 
__device__ __hip_bfloat16 __hfma (const __hip_bfloat16 a, const __hip_bfloat16 b, const __hip_bfloat16 c)
 Performs FMA of given bfloat16 values.
 
__HOST_DEVICE__ __hip_bfloat16 __hmul (const __hip_bfloat16 a, const __hip_bfloat16 b)
 Multiplies two bfloat16 values.
 
__HOST_DEVICE__ __hip_bfloat16 __hneg (const __hip_bfloat16 a)
 Negate a bfloat16 value.
 
__HOST_DEVICE__ __hip_bfloat16 __habs (const __hip_bfloat16 a)
 Returns absolute of a bfloat16.
 
__HOST_DEVICE__ __hip_bfloat162 __h2div (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Divides bfloat162 values.
 
__HOST_DEVICE__ __hip_bfloat162 __habs2 (const __hip_bfloat162 a)
 Returns absolute of a bfloat162.
 
__HOST_DEVICE__ __hip_bfloat162 __hadd2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Adds two bfloat162 values.
 
__device__ __hip_bfloat162 __hfma2 (const __hip_bfloat162 a, const __hip_bfloat162 b, const __hip_bfloat162 c)
 Performs FMA of given bfloat162 values.
 
__HOST_DEVICE__ __hip_bfloat162 __hmul2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Multiplies two bfloat162 values.
 
__HOST_DEVICE__ __hip_bfloat162 __hneg2 (const __hip_bfloat162 a)
 Converts a bfloat162 into negative.
 
__HOST_DEVICE__ __hip_bfloat162 __hsub2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Subtracts two bfloat162 values.
 
__HOST_DEVICE__ __hip_bfloat16 operator* (const __hip_bfloat16 &l, const __hip_bfloat16 &r)
 Operator to multiply two __hip_bfloat16 numbers.
 
__HOST_DEVICE__ __hip_bfloat16operator*= (__hip_bfloat16 &l, const __hip_bfloat16 &r)
 Operator to multiply-assign two __hip_bfloat16 numbers.
 
__HOST_DEVICE__ __hip_bfloat16 operator+ (const __hip_bfloat16 &l)
 Operator to unary+ on a __hip_bfloat16 number.
 
__HOST_DEVICE__ __hip_bfloat16 operator+ (const __hip_bfloat16 &l, const __hip_bfloat16 &r)
 Operator to add two __hip_bfloat16 numbers.
 
__HOST_DEVICE__ __hip_bfloat16 operator- (const __hip_bfloat16 &l)
 Operator to negate a __hip_bfloat16 number.
 
__HOST_DEVICE__ __hip_bfloat16 operator- (const __hip_bfloat16 &l, const __hip_bfloat16 &r)
 Operator to subtract two __hip_bfloat16 numbers.
 
__HOST_DEVICE__ __hip_bfloat16 operator++ (__hip_bfloat16 &l, const int)
 Operator to post increment a __hip_bfloat16 number.
 
__HOST_DEVICE__ __hip_bfloat16operator++ (__hip_bfloat16 &l)
 Operator to pre increment a __hip_bfloat16 number.
 
__HOST_DEVICE__ __hip_bfloat16 operator-- (__hip_bfloat16 &l, const int)
 Operator to post decrement a __hip_bfloat16 number.
 
__HOST_DEVICE__ __hip_bfloat16operator-- (__hip_bfloat16 &l)
 Operator to pre decrement a __hip_bfloat16 number.
 
__HOST_DEVICE__ __hip_bfloat16operator+= (__hip_bfloat16 &l, const __hip_bfloat16 &r)
 Operator to add-assign two __hip_bfloat16 numbers.
 
__HOST_DEVICE__ __hip_bfloat16operator-= (__hip_bfloat16 &l, const __hip_bfloat16 &r)
 Operator to subtract-assign two __hip_bfloat16 numbers.
 
__HOST_DEVICE__ __hip_bfloat16 operator/ (const __hip_bfloat16 &l, const __hip_bfloat16 &r)
 Operator to divide two __hip_bfloat16 numbers.
 
__HOST_DEVICE__ __hip_bfloat16operator/= (__hip_bfloat16 &l, const __hip_bfloat16 &r)
 Operator to divide-assign two __hip_bfloat16 numbers.
 
__HOST_DEVICE__ __hip_bfloat162 operator* (const __hip_bfloat162 &l, const __hip_bfloat162 &r)
 Operator to multiply two __hip_bfloat162 numbers.
 
__HOST_DEVICE__ __hip_bfloat162operator*= (__hip_bfloat162 &l, const __hip_bfloat162 &r)
 Operator to multiply-assign two __hip_bfloat162 numbers.
 
__HOST_DEVICE__ __hip_bfloat162 operator+ (const __hip_bfloat162 &l)
 Operator to unary+ on a __hip_bfloat162 number.
 
__HOST_DEVICE__ __hip_bfloat162 operator+ (const __hip_bfloat162 &l, const __hip_bfloat162 &r)
 Operator to add two __hip_bfloat162 numbers.
 
__HOST_DEVICE__ __hip_bfloat162 operator- (const __hip_bfloat162 &l)
 Operator to negate a __hip_bfloat162 number.
 
__HOST_DEVICE__ __hip_bfloat162 operator- (const __hip_bfloat162 &l, const __hip_bfloat162 &r)
 Operator to subtract two __hip_bfloat162 numbers.
 
__HOST_DEVICE__ __hip_bfloat162 operator++ (__hip_bfloat162 &l, const int)
 Operator to post increment a __hip_bfloat162 number.
 
__HOST_DEVICE__ __hip_bfloat162operator++ (__hip_bfloat162 &l)
 Operator to pre increment a __hip_bfloat162 number.
 
__HOST_DEVICE__ __hip_bfloat162 operator-- (__hip_bfloat162 &l, const int)
 Operator to post decrement a __hip_bfloat162 number.
 
__HOST_DEVICE__ __hip_bfloat162operator-- (__hip_bfloat162 &l)
 Operator to pre decrement a __hip_bfloat162 number.
 
__HOST_DEVICE__ __hip_bfloat162operator+= (__hip_bfloat162 &l, const __hip_bfloat162 &r)
 Operator to add-assign two __hip_bfloat162 numbers.
 
__HOST_DEVICE__ __hip_bfloat162operator-= (__hip_bfloat162 &l, const __hip_bfloat162 &r)
 Operator to subtract-assign two __hip_bfloat162 numbers.
 
__HOST_DEVICE__ __hip_bfloat162 operator/ (const __hip_bfloat162 &l, const __hip_bfloat162 &r)
 Operator to divide two __hip_bfloat162 numbers.
 
__HOST_DEVICE__ __hip_bfloat162operator/= (__hip_bfloat162 &l, const __hip_bfloat162 &r)
 Operator to divide-assign two __hip_bfloat162 numbers.
 
__HOST_DEVICE__ bool __heq (const __hip_bfloat16 a, const __hip_bfloat16 b)
 Compare two bfloat162 values.
 
__HOST_DEVICE__ bool __hequ (const __hip_bfloat16 a, const __hip_bfloat16 b)
 Compare two bfloat162 values - unordered equal.
 
__HOST_DEVICE__ bool __hgt (const __hip_bfloat16 a, const __hip_bfloat16 b)
 Compare two bfloat162 values - greater than.
 
__HOST_DEVICE__ bool __hgtu (const __hip_bfloat16 a, const __hip_bfloat16 b)
 Compare two bfloat162 values - unordered greater than.
 
__HOST_DEVICE__ bool __hge (const __hip_bfloat16 a, const __hip_bfloat16 b)
 Compare two bfloat162 values - greater than equal.
 
__HOST_DEVICE__ bool __hgeu (const __hip_bfloat16 a, const __hip_bfloat16 b)
 Compare two bfloat162 values - unordered greater than equal.
 
__HOST_DEVICE__ bool __hne (const __hip_bfloat16 a, const __hip_bfloat16 b)
 Compare two bfloat162 values - not equal.
 
__HOST_DEVICE__ bool __hneu (const __hip_bfloat16 a, const __hip_bfloat16 b)
 Compare two bfloat162 values - unordered not equal.
 
__HOST_DEVICE__ __hip_bfloat16 __hmax (const __hip_bfloat16 a, const __hip_bfloat16 b)
 Compare two bfloat162 values - return max.
 
__HOST_DEVICE__ __hip_bfloat16 __hmin (const __hip_bfloat16 a, const __hip_bfloat16 b)
 Compare two bfloat162 values - return min.
 
__HOST_DEVICE__ bool __hlt (const __hip_bfloat16 a, const __hip_bfloat16 b)
 Compare two bfloat162 values - less than operator.
 
__HOST_DEVICE__ bool __hltu (const __hip_bfloat16 a, const __hip_bfloat16 b)
 Compare two bfloat162 values - unordered less than.
 
__HOST_DEVICE__ bool __hle (const __hip_bfloat16 a, const __hip_bfloat16 b)
 Compare two bfloat162 values - less than equal.
 
__HOST_DEVICE__ bool __hleu (const __hip_bfloat16 a, const __hip_bfloat16 b)
 Compare two bfloat162 values - unordered less than equal.
 
__HOST_DEVICE__ int __hisinf (const __hip_bfloat16 a)
 Checks if number is inf.
 
__HOST_DEVICE__ bool __hisnan (const __hip_bfloat16 a)
 Checks if number is nan.
 
__HOST_DEVICE__ bool __hbeq2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Checks if two numbers are equal.
 
__HOST_DEVICE__ bool __hbequ2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Checks if two numbers are equal - unordered.
 
__HOST_DEVICE__ bool __hbge2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Check for a >= b.
 
__HOST_DEVICE__ bool __hbgeu2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Check for a >= b - unordered.
 
__HOST_DEVICE__ bool __hbgt2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Check for a > b.
 
__HOST_DEVICE__ bool __hbgtu2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Check for a > b - unordered.
 
__HOST_DEVICE__ bool __hble2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Check for a <= b.
 
__HOST_DEVICE__ bool __hbleu2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Check for a <= b - unordered.
 
__HOST_DEVICE__ bool __hblt2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Check for a < b.
 
__HOST_DEVICE__ bool __hbltu2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Check for a < b - unordered.
 
__HOST_DEVICE__ bool __hbne2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Check for a != b.
 
__HOST_DEVICE__ bool __hbneu2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Check for a != b.
 
__HOST_DEVICE__ __hip_bfloat162 __heq2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Check for a != b, returns 1.0 if equal, otherwise 0.0.
 
__HOST_DEVICE__ __hip_bfloat162 __hge2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Check for a >= b, returns 1.0 if greater than equal, otherwise 0.0.
 
__HOST_DEVICE__ __hip_bfloat162 __hgt2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Check for a > b, returns 1.0 if greater than equal, otherwise 0.0.
 
__HOST_DEVICE__ __hip_bfloat162 __hisnan2 (const __hip_bfloat162 a)
 Check for a is NaN, returns 1.0 if NaN, otherwise 0.0.
 
__HOST_DEVICE__ __hip_bfloat162 __hle2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Check for a <= b, returns 1.0 if greater than equal, otherwise 0.0.
 
__HOST_DEVICE__ __hip_bfloat162 __hlt2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Check for a < b, returns 1.0 if greater than equal, otherwise 0.0.
 
__HOST_DEVICE__ __hip_bfloat162 __hmax2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Returns max of two elements.
 
__HOST_DEVICE__ __hip_bfloat162 __hmin2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Returns min of two elements.
 
__HOST_DEVICE__ __hip_bfloat162 __hne2 (const __hip_bfloat162 a, const __hip_bfloat162 b)
 Checks for not equal to.
 
__HOST_DEVICE__ bool operator== (const __hip_bfloat16 &l, const __hip_bfloat16 &r)
 Operator to perform an equal compare on two __hip_bfloat16 numbers.
 
__HOST_DEVICE__ bool operator!= (const __hip_bfloat16 &l, const __hip_bfloat16 &r)
 Operator to perform a not equal on two __hip_bfloat16 numbers.
 
__HOST_DEVICE__ bool operator< (const __hip_bfloat16 &l, const __hip_bfloat16 &r)
 Operator to perform a less than on two __hip_bfloat16 numbers.
 
__HOST_DEVICE__ bool operator<= (const __hip_bfloat16 &l, const __hip_bfloat16 &r)
 Operator to perform a less than equal on two __hip_bfloat16 numbers.
 
__HOST_DEVICE__ bool operator> (const __hip_bfloat16 &l, const __hip_bfloat16 &r)
 Operator to perform a greater than on two __hip_bfloat16 numbers.
 
__HOST_DEVICE__ bool operator>= (const __hip_bfloat16 &l, const __hip_bfloat16 &r)
 Operator to perform a greater than equal on two __hip_bfloat16 numbers.
 
__HOST_DEVICE__ bool operator== (const __hip_bfloat162 &l, const __hip_bfloat162 &r)
 Operator to perform an equal compare on two __hip_bfloat16 numbers.
 
__HOST_DEVICE__ bool operator!= (const __hip_bfloat162 &l, const __hip_bfloat162 &r)
 Operator to perform a not equal on two __hip_bfloat16 numbers.
 
__HOST_DEVICE__ bool operator< (const __hip_bfloat162 &l, const __hip_bfloat162 &r)
 Operator to perform a less than on two __hip_bfloat16 numbers.
 
__HOST_DEVICE__ bool operator<= (const __hip_bfloat162 &l, const __hip_bfloat162 &r)
 Operator to perform a less than equal on two __hip_bfloat16 numbers.
 
__HOST_DEVICE__ bool operator> (const __hip_bfloat162 &l, const __hip_bfloat162 &r)
 Operator to perform a greater than on two __hip_bfloat16 numbers.
 
__HOST_DEVICE__ bool operator>= (const __hip_bfloat162 &l, const __hip_bfloat162 &r)
 Operator to perform a greater than equal on two __hip_bfloat16 numbers.
 
__device__ __hip_bfloat16 hceil (const __hip_bfloat16 h)
 Calculate ceil of bfloat16.
 
__device__ __hip_bfloat16 hcos (const __hip_bfloat16 h)
 Calculate cosine of bfloat16.
 
__device__ __hip_bfloat16 hexp (const __hip_bfloat16 h)
 Calculate exponential of bfloat16.
 
__device__ __hip_bfloat16 hexp10 (const __hip_bfloat16 h)
 Calculate exponential 10 of bfloat16.
 
__device__ __hip_bfloat16 hexp2 (const __hip_bfloat16 h)
 Calculate exponential 2 of bfloat16.
 
__device__ __hip_bfloat16 hfloor (const __hip_bfloat16 h)
 Calculate floor of bfloat16.
 
__device__ __hip_bfloat16 hlog (const __hip_bfloat16 h)
 Calculate natural log of bfloat16.
 
__device__ __hip_bfloat16 hlog10 (const __hip_bfloat16 h)
 Calculate log 10 of bfloat16.
 
__device__ __hip_bfloat16 hlog2 (const __hip_bfloat16 h)
 Calculate log 2 of bfloat16.
 
__device__ __hip_bfloat16 hrcp (const __hip_bfloat16 h)
 Calculate reciprocal.
 
__device__ __hip_bfloat16 hrint (const __hip_bfloat16 h)
 Round to nearest int.
 
__device__ __hip_bfloat16 hrsqrt (const __hip_bfloat16 h)
 Reciprocal square root.
 
__device__ __hip_bfloat16 hsin (const __hip_bfloat16 h)
 Calculate sin of bfloat16.
 
__device__ __hip_bfloat16 hsqrt (const __hip_bfloat16 h)
 Calculate sqrt of bfloat16.
 
__device__ __hip_bfloat16 htrunc (const __hip_bfloat16 h)
 Calculate truncate of bfloat16.
 
__device__ __hip_bfloat162 h2ceil (const __hip_bfloat162 h)
 Calculate ceil of bfloat162.
 
__device__ __hip_bfloat162 h2cos (const __hip_bfloat162 h)
 Calculate cosine of bfloat162.
 
__device__ __hip_bfloat162 h2exp (const __hip_bfloat162 h)
 Calculate exponential of bfloat162.
 
__device__ __hip_bfloat162 h2exp10 (const __hip_bfloat162 h)
 Calculate exponential 10 of bfloat162.
 
__device__ __hip_bfloat162 h2exp2 (const __hip_bfloat162 h)
 Calculate exponential 2 of bfloat162.
 
__device__ __hip_bfloat162 h2floor (const __hip_bfloat162 h)
 Calculate floor of bfloat162.
 
__device__ __hip_bfloat162 h2log (const __hip_bfloat162 h)
 Calculate natural log of bfloat162.
 
__device__ __hip_bfloat162 h2log10 (const __hip_bfloat162 h)
 Calculate log 10 of bfloat162.
 
__device__ __hip_bfloat162 h2log2 (const __hip_bfloat162 h)
 Calculate log 2 of bfloat162.
 
__device__ __hip_bfloat162 h2rcp (const __hip_bfloat162 h)
 Calculate vector reciprocal.
 
__device__ __hip_bfloat162 h2rint (const __hip_bfloat162 h)
 Calculate vector round to nearest int.
 
__device__ __hip_bfloat162 h2rsqrt (const __hip_bfloat162 h)
 Calculate vector reciprocal square root.
 
__device__ __hip_bfloat162 h2sin (const __hip_bfloat162 h)
 Calculate sin of bfloat162.
 
__device__ __hip_bfloat162 h2sqrt (const __hip_bfloat162 h)
 Calculate sqrt of bfloat162.
 
__device__ __hip_bfloat162 h2trunc (const __hip_bfloat162 h)
 Calculate truncate of bfloat162.
 

Detailed Description

hip_bf16.h provides struct for __hip_bfloat16 types

MIT License

Copyright (c) 2019 - 2023 Advanced Micro Devices, Inc. All rights reserved.

Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of the Software, and to permit persons to whom the Software is furnished to do so, subject to the following conditions:

The above copyright notice and this permission notice shall be included in all copies or substantial portions of the Software.

THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.