Skip to content

Bfloat16 #652

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
leekillough merged 3 commits into from
Sep 4, 2019
Merged

Bfloat16 #652

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
46a50e8
Use a similar method as in rocBLAS to determine whether to use C++ an…
Jul 17, 2019
61bb391
Make bfloat16 changes similar to rocBLAS PR #678
Sep 1, 2019
5663a50
Merge branch 'bfloat16' of github.com:ROCmSoftwarePlatform/Tensile in…
Sep 1, 2019
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
39 changes: 18 additions & 21 deletions Tensile/Source/MathTemplates.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -86,15 +86,15 @@ template<> TensileComplexDouble tensileGetOne<TensileComplexDouble>() {
#ifdef Tensile_ENABLE_HALF
template<> TensileHalf tensileGetRandom<TensileHalf>() { return static_cast<TensileHalf>((rand()%7) - 3); }
#endif
template<> uint32_t tensileGetRandom<uint32_t>() {
int8_t t0 = static_cast<int8_t>((rand()%7) - 3);
int8_t t1 = static_cast<int8_t>((rand()%7) - 3);
int8_t t2 = static_cast<int8_t>((rand()%7) - 3);
int8_t t3 = static_cast<int8_t>((rand()%7) - 3);
template<> uint32_t tensileGetRandom<uint32_t>() {
int8_t t0 = static_cast<int8_t>((rand()%7) - 3);
int8_t t1 = static_cast<int8_t>((rand()%7) - 3);
int8_t t2 = static_cast<int8_t>((rand()%7) - 3);
int8_t t3 = static_cast<int8_t>((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<int32_t>() { return static_cast<int32_t>((rand()%7) - 3); }
template<> float tensileGetRandom<float>() { return static_cast<float>((rand()%201) - 100); }
Expand Down Expand Up @@ -142,20 +142,20 @@ template<> TensileComplexDouble tensileGetTypeForInt<TensileComplexDouble>( size
#ifdef Tensile_ENABLE_HALF
template<> TensileHalf tensileGetTrig<TensileHalf>(int i) { return static_cast<TensileHalf>(sin(i)); }
#endif
template<> uint32_t tensileGetTrig<uint32_t>(int i) {
int8_t t0 = static_cast<int8_t>((rand()%7) - 3);
int8_t t1 = static_cast<int8_t>((rand()%7) - 3);
int8_t t2 = static_cast<int8_t>((rand()%7) - 3);
int8_t t3 = static_cast<int8_t>((rand()%7) - 3);
template<> uint32_t tensileGetTrig<uint32_t>(int i) {
int8_t t0 = static_cast<int8_t>((rand()%7) - 3);
int8_t t1 = static_cast<int8_t>((rand()%7) - 3);
int8_t t2 = static_cast<int8_t>((rand()%7) - 3);
int8_t t3 = static_cast<int8_t>((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<int32_t>(int i) { return static_cast<int32_t>((rand()%7) - 3); }
template<> float tensileGetTrig<float>(int i) { return static_cast<float>(sin(i)); }
template<> tensile_bfloat16 tensileGetTrig<tensile_bfloat16>(int i) { return sin(static_cast<tensile_bfloat16>(i)); }
template<> double tensileGetTrig<double>(int i) { return static_cast<double>(sin(i)); }
template<> int32_t tensileGetTrig<int32_t>(int i) { return rand() % 7 - 3; }
template<> float tensileGetTrig<float>(int i) { return sin(i); }
template<> tensile_bfloat16 tensileGetTrig<tensile_bfloat16>(int i) { return tensile_bfloat16(sinf(i)); }
template<> double tensileGetTrig<double>(int i) { return sin(i); }
template<> TensileComplexFloat tensileGetTrig<TensileComplexFloat>(int i) {
TensileComplexFloat r;
TENSILEREAL(r) = tensileGetTrig<float>(i);
Expand Down Expand Up @@ -531,6 +531,3 @@ template<> std::string tensileToString(TensileHalf v){
#endif
template<> std::string tensileToString(tensile_bfloat16 v){
return tensileToString(static_cast<float>(v)); }



7 changes: 1 addition & 6 deletions Tensile/Source/tensile_bfloat16.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
{
Expand Down Expand Up @@ -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))); }

Expand Down