diff --git a/Tensile/Source/MathTemplates.cpp b/Tensile/Source/MathTemplates.cpp index d24960fe1..5268c2bb1 100644 --- a/Tensile/Source/MathTemplates.cpp +++ b/Tensile/Source/MathTemplates.cpp @@ -86,15 +86,15 @@ template<> TensileComplexDouble tensileGetOne() { #ifdef Tensile_ENABLE_HALF template<> TensileHalf tensileGetRandom() { return static_cast((rand()%7) - 3); } #endif -template<> uint32_t tensileGetRandom() { - int8_t t0 = static_cast((rand()%7) - 3); - int8_t t1 = static_cast((rand()%7) - 3); - int8_t t2 = static_cast((rand()%7) - 3); - int8_t t3 = static_cast((rand()%7) - 3); +template<> uint32_t tensileGetRandom() { + int8_t t0 = static_cast((rand()%7) - 3); + int8_t t1 = static_cast((rand()%7) - 3); + int8_t t2 = static_cast((rand()%7) - 3); + int8_t t3 = static_cast((rand()%7) - 3); int8_t t1x4[4] = {t0, t1, t2, t3}; - uint32_t tmp; + uint32_t tmp; memcpy(&tmp, t1x4, sizeof(uint32_t)); - return tmp; + return tmp; } template<> int32_t tensileGetRandom() { return static_cast((rand()%7) - 3); } template<> float tensileGetRandom() { return static_cast((rand()%201) - 100); } @@ -142,20 +142,20 @@ template<> TensileComplexDouble tensileGetTypeForInt( size #ifdef Tensile_ENABLE_HALF template<> TensileHalf tensileGetTrig(int i) { return static_cast(sin(i)); } #endif -template<> uint32_t tensileGetTrig(int i) { - int8_t t0 = static_cast((rand()%7) - 3); - int8_t t1 = static_cast((rand()%7) - 3); - int8_t t2 = static_cast((rand()%7) - 3); - int8_t t3 = static_cast((rand()%7) - 3); +template<> uint32_t tensileGetTrig(int i) { + int8_t t0 = static_cast((rand()%7) - 3); + int8_t t1 = static_cast((rand()%7) - 3); + int8_t t2 = static_cast((rand()%7) - 3); + int8_t t3 = static_cast((rand()%7) - 3); int8_t t1x4[4] = {t0, t1, t2, t3}; - uint32_t tmp; + uint32_t tmp; memcpy(&tmp, t1x4, sizeof(uint32_t)); - return tmp; + return tmp; } -template<> int32_t tensileGetTrig(int i) { return static_cast((rand()%7) - 3); } -template<> float tensileGetTrig(int i) { return static_cast(sin(i)); } -template<> tensile_bfloat16 tensileGetTrig(int i) { return sin(static_cast(i)); } -template<> double tensileGetTrig(int i) { return static_cast(sin(i)); } +template<> int32_t tensileGetTrig(int i) { return rand() % 7 - 3; } +template<> float tensileGetTrig(int i) { return sin(i); } +template<> tensile_bfloat16 tensileGetTrig(int i) { return tensile_bfloat16(sinf(i)); } +template<> double tensileGetTrig(int i) { return sin(i); } template<> TensileComplexFloat tensileGetTrig(int i) { TensileComplexFloat r; TENSILEREAL(r) = tensileGetTrig(i); @@ -531,6 +531,3 @@ template<> std::string tensileToString(TensileHalf v){ #endif template<> std::string tensileToString(tensile_bfloat16 v){ return tensileToString(static_cast(v)); } - - - diff --git a/Tensile/Source/tensile_bfloat16.h b/Tensile/Source/tensile_bfloat16.h index c806b6884..3bef81a58 100644 --- a/Tensile/Source/tensile_bfloat16.h +++ b/Tensile/Source/tensile_bfloat16.h @@ -61,7 +61,7 @@ struct tensile_bfloat16 explicit __host__ __device__ tensile_bfloat16(float f) : data(float_to_bfloat16(f)) { } // zero extend lower 16 bits of bfloat16 to convert to IEEE float - explicit __host__ __device__ operator float() const + __host__ __device__ operator float() const { union { @@ -184,11 +184,6 @@ inline __host__ __device__ tensile_bfloat16 operator--(tensile_bfloat16& a, int) inline __host__ __device__ bool isinf(tensile_bfloat16 a) { return !(~a.data & 0x7f80) && !(a.data & 0x7f); } inline __host__ __device__ bool isnan(tensile_bfloat16 a) { return !(~a.data & 0x7f80) && +(a.data & 0x7f); } inline __host__ __device__ bool iszero(tensile_bfloat16 a) { return !(a.data & 0x7fff); } -inline __host__ __device__ tensile_bfloat16 abs(tensile_bfloat16 a) -{ - a.data &= 0x7fff; - return a; -} inline tensile_bfloat16 sin(tensile_bfloat16 a) { return tensile_bfloat16(sinf(float(a))); } inline tensile_bfloat16 cos(tensile_bfloat16 a) { return tensile_bfloat16(cosf(float(a))); }